[Bug] [llvm] Fix FP<->UInt castings (#3560)

* [Bug] [llvm] Fix FP<->UInt castings

* Auto Format

Co-authored-by: Taichi Gardener <taichigardener@gmail.com>

taichi/codegen/codegen_llvm.cpp

@@ -373,9 +373,11 @@ void CodeGenLLVM::visit(UnaryOpStmt *stmt) {
       llvm_val[stmt] = llvm_val[stmt->operand];
     } else if (is_real(from) != is_real(to)) {
       if (is_real(from) && is_integral(to)) {
-        cast_op = llvm::Instruction::CastOps::FPToSI;
+        cast_op = is_signed(to) ? llvm::Instruction::CastOps::FPToSI
+                                : llvm::Instruction::CastOps::FPToUI;
       } else if (is_integral(from) && is_real(to)) {
-        cast_op = llvm::Instruction::CastOps::SIToFP;
+        cast_op = is_signed(from) ? llvm::Instruction::CastOps::SIToFP
+                                  : llvm::Instruction::CastOps::UIToFP;
       } else {
         TI_P(data_type_name(from));
         TI_P(data_type_name(to));

tests/python/test_cast.py

@@ -3,6 +3,26 @@
 import taichi as ti
 
 
+@pytest.mark.parametrize('dtype', [ti.u8, ti.u16, ti.u32])
+@ti.test(exclude=ti.opengl)
+def test_cast_uint_to_float(dtype):
+    @ti.kernel
+    def func(a: dtype) -> ti.f32:
+        return ti.cast(a, ti.f32)
+
+    assert func(255) == 255
+
+
+@pytest.mark.parametrize('dtype', [ti.u8, ti.u16, ti.u32])
+@ti.test(exclude=ti.opengl)
+def test_cast_float_to_uint(dtype):
+    @ti.kernel
+    def func(a: ti.f32) -> dtype:
+        return ti.cast(a, dtype)
+
+    assert func(255) == 255
+
+
 @ti.test()
 def test_cast_f32():
     z = ti.field(ti.i32, shape=())