tests: add launch-level coverage for ctypes/numpy subclass fallback

Author: rluo8

cuda_core/tests/test_launcher.py

@@ -407,6 +407,46 @@ class MyBool(ctypes.c_bool):
     assert holder.ptr != 0
 
 
+@requires_module(np, "2.1")
+def test_launch_scalar_argument_ctypes_subclass_fallback():
+    """Subclassed ctypes scalars survive the launch path and reach the kernel correctly."""
+
+    class MyInt32(ctypes.c_int32):
+        pass
+
+    dev = Device()
+    dev.set_current()
+
+    mr = LegacyPinnedMemoryResource()
+    b = mr.allocate(np.dtype(np.int32).itemsize)
+    arr = np.from_dlpack(b).view(np.int32)
+    arr[:] = 0
+
+    scalar = MyInt32(-123456)
+
+    code = r"""
+    template <typename T>
+    __global__ void write_scalar(T* arr, T val) {
+        arr[0] = val;
+    }
+    """
+
+    arch = "".join(f"{i}" for i in dev.compute_capability)
+    pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}")
+    prog = Program(code, code_type="c++", options=pro_opts)
+    ker_name = "write_scalar<signed int>"
+    mod = prog.compile("cubin", name_expressions=(ker_name,))
+    ker = mod.get_kernel(ker_name)
+
+    # This exercises the prepare_ctypes_arg isinstance fallback through a real launch.
+    stream = dev.default_stream
+    config = LaunchConfig(grid=1, block=1)
+    launch(stream, config, ker, arr.ctypes.data, scalar)
+    stream.sync()
+
+    assert arr[0] == scalar.value
+
+
 def test_kernel_arg_numpy_subclass_isinstance_fallback():
     """Subclassed numpy scalars hit the isinstance fallback in prepare_numpy_arg."""
     from cuda.core._kernel_arg_handler import ParamHolder
@@ -421,6 +461,46 @@ class MyFloat32(np.float32):
     assert holder.ptr != 0
 
 
+@requires_module(np, "2.1")
+def test_launch_scalar_argument_numpy_subclass_fallback():
+    """Subclassed numpy scalars survive the launch path and reach the kernel correctly."""
+
+    class MyFloat32(np.float32):
+        pass
+
+    dev = Device()
+    dev.set_current()
+
+    mr = LegacyPinnedMemoryResource()
+    b = mr.allocate(np.dtype(np.float32).itemsize)
+    arr = np.from_dlpack(b).view(np.float32)
+    arr[:] = 0.0
+
+    scalar = MyFloat32(3.14)
+
+    code = r"""
+    template <typename T>
+    __global__ void write_scalar(T* arr, T val) {
+        arr[0] = val;
+    }
+    """
+
+    arch = "".join(f"{i}" for i in dev.compute_capability)
+    pro_opts = ProgramOptions(std="c++17", arch=f"sm_{arch}")
+    prog = Program(code, code_type="c++", options=pro_opts)
+    ker_name = "write_scalar<float>"
+    mod = prog.compile("cubin", name_expressions=(ker_name,))
+    ker = mod.get_kernel(ker_name)
+
+    # This exercises the prepare_numpy_arg isinstance fallback through a real launch.
+    stream = dev.default_stream
+    config = LaunchConfig(grid=1, block=1)
+    launch(stream, config, ker, arr.ctypes.data, scalar)
+    stream.sync()
+
+    assert arr[0] == scalar
+
+
 def test_kernel_arg_python_isinstance_fallbacks():
     """Subclassed Python builtins hit the isinstance fallback in ParamHolder."""
     from cuda.core._kernel_arg_handler import ParamHolder