[opengl] [test] "test_bit_operations::test_bit_shr" fails on opengl

[xumingkuan]

Describe the bug
ti.bit_shr is demoted to bit_sar on opengl, but the data type becomes u32, which is not supported on opengl...

To Reproduce
ti test bit_operations -v -t 1.

@ti.test()
def test_bit_shr():
    @ti.kernel
    def shr(a: ti.i32, b: ti.i32) -> ti.i32:
        return ti.bit_shr(a, b)

    n = 8
    test_num = 2**n
    neg_test_num = -test_num
    for i in range(n):
        assert shr(test_num, i) == 2**(n - i)
    for i in range(n):
        offset = 0x100000000 if i > 0 else 0
        assert shr(neg_test_num, i) == (neg_test_num + offset) >> i

Log/Screenshots

..\tests\python\test_bit_operations.py::test_bit_shr[opengl] [Taichi] Starting
 on arch=opengl
[Taichi] materializing...
[I 10/08/20 18:27:05.793] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Initial IR:
kernel {
  <i32 x1> $0 : kernel return (cast_value<int32> (arg[0] (dt=int32) shr arg[1]
 (dt=int32)))
}
[I 10/08/20 18:27:05.794] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Lowered:
kernel {
  <i32 x1> $0 = arg[0]
  <i32 x1> $1 = arg[1]
  $2 = bit_shr $0 $1
  $3 = cast_value<i32> $2
  <i32 x1> $4 : kernel return $3
}
[I 10/08/20 18:27:05.794] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Typechecked:
kernel {
  <i32 x1> $0 = arg[0]
  <i32 x1> $1 = arg[1]
  <i32 x1> $2 = bit_shr $0 $1
  <i32 x1> $3 = cast_value<i32> $2
  <i32 x1> $4 : kernel return $3
}
[I 10/08/20 18:27:05.795] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Simplified I:
kernel {
  <i32 x1> $0 = arg[0]
  <i32 x1> $1 = arg[1]
  <i32 x1> $2 = bit_shr $0 $1
  <i32 x1> $3 : kernel return $2
}
[I 10/08/20 18:27:05.796] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Access flagged I:
kernel {
  <i32 x1> $0 = arg[0]
  <i32 x1> $1 = arg[1]
  <i32 x1> $2 = bit_shr $0 $1
  <i32 x1> $3 : kernel return $2
}
[I 10/08/20 18:27:05.797] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Simplified II:
kernel {
  <i32 x1> $0 = arg[0]
  <i32 x1> $1 = arg[1]
  <i32 x1> $2 = bit_shr $0 $1
  <i32 x1> $3 : kernel return $2
}
[I 10/08/20 18:27:05.797] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Offloaded:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.798] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Optimized by CFG:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.798] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Access flagged II:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.799] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Simplified III:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.800] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Start offload_to_executable:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.800] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Dense struct-for demoted:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.801] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Make thread local:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.801] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Remove range assumption:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.802] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Access lowered:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.803] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] DIE:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.804] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Access flagged III:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.804] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Atomics demoted:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <i32 x1> $3 = bit_shr $1 $2
    <i32 x1> $4 : kernel return $3
  }
}
[I 10/08/20 18:27:05.805] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Operations demoted:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <u32 x1> $3 = reinterpret_cast_bits<u32> $1
    <u32 x1> $4 = cast_value<u32> $2
    <u32 x1> $5 = bit_sar $3 $4
    <i32 x1> $6 = reinterpret_cast_bits<i32> $5
    <i32 x1> $7 : kernel return $6
  }
}
[I 10/08/20 18:27:05.806] [compile_to_offloads.cpp:taichi::lang::irpass::`anon
ymous-namespace'::make_pass_printer::<lambda_fe1d620add3df83d4ee306f9c0ab10ca>
::operator ()@18] [shr_c14_0] Simplified IV:
kernel {
  $0 = offloaded
  body {
    <i32 x1> $1 = arg[0]
    <i32 x1> $2 = arg[1]
    <u32 x1> $3 = reinterpret_cast_bits<u32> $1
    <u32 x1> $4 = cast_value<u32> $2
    <u32 x1> $5 = bit_sar $3 $4
    <i32 x1> $6 = reinterpret_cast_bits<i32> $5
    <i32 x1> $7 : kernel return $6
  }
}
[E 10/08/20 18:27:05.807] [taichi/backends/opengl/opengl_data_types.h:taichi::
lang::opengl::opengl_data_type_name@20] Not supported.