ti.bit_shr is broken when the input is unsigned

[lin-hitonami]

Describe the bug
ti.bit_shr is broken when the input is unsigned because we only transform it to bit_sar(unsigned) when the input is signed, and the codegen assumes that bit_shr should not exist when the ir reaches the codegen.

taichi/taichi/transforms/demote_operations.cpp

Lines 152 to 175 in 6a0b63e

	} else if (stmt->op_type == BinaryOpType::bit_shr &&
	is_integral(lhs->element_type()) &&
	is_integral(rhs->element_type()) &&
	is_signed(lhs->element_type())) {
	// @ti.func
	// def bit_shr(a, b):
	// unsigned_a = ti.cast(a, ti.uXX)
	// shifted = ti.bit_sar(unsigned_a, b)
	// ret = ti.cast(shifted, ti.iXX)
	// return ret
	auto unsigned_cast = Stmt::make<UnaryOpStmt>(UnaryOpType::cast_bits, lhs);
	unsigned_cast->as<UnaryOpStmt>()->cast_type =
	to_unsigned(lhs->element_type());
	auto shift = Stmt::make<BinaryOpStmt>(BinaryOpType::bit_sar,
	unsigned_cast.get(), rhs);
	auto signed_cast =
	Stmt::make<UnaryOpStmt>(UnaryOpType::cast_bits, shift.get());
	signed_cast->as<UnaryOpStmt>()->cast_type = lhs->element_type();
	signed_cast->ret_type = stmt->ret_type;
	stmt->replace_usages_with(signed_cast.get());
	modifier.insert_before(stmt, std::move(unsigned_cast));
	modifier.insert_before(stmt, std::move(shift));
	modifier.insert_before(stmt, std::move(signed_cast));
	modifier.erase(stmt);

To Reproduce

import taichi as ti
ti.init()                                                                                                                        

@ti.kernel
def func(x: ti.u32, y: ti.i32)->ti.u32:
    return ti.bit_shr(x, y)

print(func(5, 2))

Log/Screenshots
Please post the full log of the program (instead of just a few lines around the error message, unless the log is > 1000 lines). This will help us diagnose what's happening. For example:

[Taichi] version 1.6.0, llvm 15.0.4, commit 2467d918, linux, python 3.10.0
[Taichi] Starting on arch=x64
[I 03/24/23 17:04:04.280 1754876] [codegen_llvm.cpp:visit@834] 
"binary_op_type_name(op)": "bit_shr"
[E 03/24/23 17:04:04.280 1754876] [codegen_llvm.cpp:visit@835] Not supported.