Address review: fix int64 truncation clipping bounds and add float64 unit test

Author: shivansh023023

warp/native/builtin.h

@@ -110,20 +110,20 @@ typedef const char* str;
 template <typename F>
 CUDA_CALLABLE inline int64 safe_float_to_int64(F x) {
     if (!(x == x)) return 0;
-    constexpr F min_int64 = static_cast<F>(-9223372036854775808.0);
-    constexpr F max_int64 = static_cast<F>(9223372036854774784.0);
-    if (x <= min_int64) return -9223372036854775807LL - 1LL;
-    if (x >= max_int64) return 9223372036854775807LL;
+    constexpr F min_int64 = static_cast<F>(-9223372036854775808.0); // -2^63
+    constexpr F max_overflow = static_cast<F>(9223372036854775808.0); // 2^63
+    if (x < min_int64) return -9223372036854775807LL - 1LL;
+    if (x >= max_overflow) return 9223372036854775807LL;
     return static_cast<int64>(x);
 }
 
 template <typename F>
 CUDA_CALLABLE inline uint64 safe_float_to_uint64(F x) {
     if (!(x == x)) return 0;
     if (x <= 0.0) return static_cast<uint64>(safe_float_to_int64(x));
-    constexpr F pow2_63 = static_cast<F>(9223372036854775808.0);
-    constexpr F max_uint64 = static_cast<F>(18446744073709551615.0);
-    if (x >= max_uint64) return 18446744073709551615ULL;
+    constexpr F pow2_63 = static_cast<F>(9223372036854775808.0); // 2^63
+    constexpr F overflow_uint64 = static_cast<F>(18446744073709551616.0); // 2^64
+    if (x >= overflow_uint64) return 18446744073709551615ULL;
     if (x >= pow2_63) return static_cast<uint64>(safe_float_to_int64(x - pow2_63)) + 9223372036854775808ULL;
     return static_cast<uint64>(safe_float_to_int64(x));
 }

warp/tests/test_codegen_instancing.py

@@ -1090,6 +1090,15 @@ def k(input: float, expected: float):
 type_closure_kernel_float = create_type_closure_scalar(float)
 type_closure_kernel_uint8 = create_type_closure_scalar(wp.uint8)
 
+def create_type_closure_scalar_f64(scalar_type):
+    @wp.kernel
+    def k(input: wp.float64, expected: wp.float64):
+        x = scalar_type(input)
+        wp.expect_eq(wp.float64(x), expected)
+    return k
+
+type_closure_kernel_uint64_f64 = create_type_closure_scalar_f64(wp.uint64)
+
 
 def test_type_closure_scalar(test, device):
     with wp.ScopedDevice(device):
@@ -1101,6 +1110,9 @@ def test_type_closure_scalar(test, device):
         wp.launch(type_closure_kernel_uint8, dim=1, inputs=[255.1, 255.0])
         wp.launch(type_closure_kernel_uint8, dim=1, inputs=[128.0, 128.0])
         wp.launch(type_closure_kernel_uint8, dim=1, inputs=[-100.0, 156.0])
+        
+        # Test boundary cases for uint64 truncation safety with float64 precision
+        wp.launch(type_closure_kernel_uint64_f64, dim=1, inputs=[9223372036854774784.0, 9223372036854774784.0])
 
 
 # =======================================================================