Skip to content

tests: add coverage tests for cuda core#1923

Merged
rwgk merged 3 commits intoNVIDIA:mainfrom
rluo8:coverage-tests
Apr 17, 2026
Merged

tests: add coverage tests for cuda core#1923
rwgk merged 3 commits intoNVIDIA:mainfrom
rluo8:coverage-tests

Conversation

@rluo8
Copy link
Copy Markdown
Contributor

@rluo8 rluo8 commented Apr 16, 2026

Add coverage tests for cuda core about files _event.pyx, _program.pyx, _memoryview.pyx,_linker.pyx, _kernel_arg_handler.pyx.
Co-worked with cursor-agent, target is mainly for error paths.

As the test couldn't be run using the runners before commit, the data got from local run on system with A10.
_event.pyx 85.59% -> 90.68%
_program.pyx 81.48% -> 84.93%
_memoryview.pyx 60.55% -> 65.14%
_linker.pyx 63.61% -> 64.44%
_kernel_arg_handler.pyx 80.28% -> 85.45%

All the tests passed on Linux/Windows:


Results: 39 passed, 0 failed, 0 skipped, 39 total
PASSED:
  cuda_core/tests/test_event.py::test_event_isub_not_implemented
  cuda_core/tests/test_event.py::test_event_rsub_not_implemented
  cuda_core/tests/test_event.py::test_event_get_ipc_descriptor_non_ipc
  cuda_core/tests/test_event.py::test_event_is_done_false
  cuda_core/tests/test_event.py::test_ipc_event_descriptor_direct_init
  cuda_core/tests/test_launcher.py::test_kernel_arg_ctypes_subclass_isinstance_fallback
  cuda_core/tests/test_launcher.py::test_kernel_arg_numpy_subclass_isinstance_fallback
  cuda_core/tests/test_launcher.py::test_kernel_arg_python_isinstance_fallbacks
  cuda_core/tests/test_linker.py::test_linker_handle
  cuda_core/tests/test_linker.py::test_linker_options_nvjitlink_options_as_str
  cuda_core/tests/test_program.py::test_program_options_repr
  cuda_core/tests/test_program.py::test_program_options_bad_define_macro_short_tuple
  cuda_core/tests/test_program.py::test_program_options_bad_define_macro_non_str_value
  cuda_core/tests/test_program.py::test_program_options_bad_define_macro_list_non_str
  cuda_core/tests/test_program.py::test_program_options_bad_define_macro_list_bad_tuple
  cuda_core/tests/test_program.py::test_ptx_program_extra_sources_unsupported
  cuda_core/tests/test_program.py::test_ptx_program_handle_is_linker_handle
  cuda_core/tests/test_program.py::test_nvvm_program_wrong_code_type
  cuda_core/tests/test_program.py::test_extra_sources_not_sequence
  cuda_core/tests/test_program.py::test_extra_sources_bad_module_not_tuple
  cuda_core/tests/test_program.py::test_extra_sources_bad_module_name_not_str
  cuda_core/tests/test_program.py::test_extra_sources_bad_module_source_wrong_type
  cuda_core/tests/test_program.py::test_extra_sources_empty_source
  cuda_core/tests/test_program.py::test_nvrtc_compile_with_logs_capture
  cuda_core/tests/test_utils.py::test_strided_memory_view_repr
  cuda_core/tests/test_utils.py::test_strided_memory_view_copy_to_raises
  cuda_core/tests/test_utils.py::test_strided_memory_view_get_layout_error
  cuda_core/tests/test_utils.py::test_strided_memory_view_deprecated_cai_init
  cuda_core/tests/test_utils.py::test_from_any_interface_cai_fallback
  cuda_core/tests/test_utils.py::test_strided_memory_view_copy_from_raises
  cuda_core/tests/test_utils.py::test_strided_memory_view_view_no_args_returns_self
  cuda_core/tests/test_utils.py::test_strided_memory_view_view_with_dtype_only
  cuda_core/tests/test_utils.py::test_dlpack_export_structured_dtype_raises
  cuda_core/tests/test_utils.py::test_dlpack_export_unsupported_dtype_raises
  cuda_core/tests/test_utils.py::test_cai_v2_rejected
  cuda_core/tests/test_utils.py::test_cai_mask_rejected
  cuda_core/tests/test_utils.py::test_cai_stream_none_rejected
  cuda_core/tests/test_utils.py::test_array_interface_v2_rejected
  cuda_core/tests/test_utils.py::test_array_interface_mask_rejected
============================================================
ALL PATCH TESTS PASSED

@rluo8
Copy link
Copy Markdown
Contributor Author

rluo8 commented Apr 16, 2026

Hi @mdboom , could you please help review it? Thanks!

@github-actions
Copy link
Copy Markdown

@leofang leofang added CI/CD CI/CD infrastructure cuda.core Everything related to the cuda.core module labels Apr 16, 2026
@leofang leofang added this to the cuda.core v1.0.0 milestone Apr 16, 2026
@leofang leofang requested review from leofang and mdboom April 16, 2026 14:48
@rwgk
Copy link
Copy Markdown
Contributor

rwgk commented Apr 16, 2026

I spot-checked the CI failures: all infrastructure flakes. I'll update the branch so that the tests run fresh.

@rwgk
Copy link
Copy Markdown
Contributor

rwgk commented Apr 16, 2026

Generated with Cursor GPT-5.4 Extra High Fast

__

Nice coverage addition. One way to make the launcher fallback tests stronger would be to keep the existing ParamHolder(...) smoke tests, but add one end-to-end launch case for a subclassed scalar too. Right now the new tests show that the fallback path accepts the object shape, but they do not prove that a launched kernel receives the correct ABI/value.

Something like this would exercise the same prepare_ctypes_arg isinstance(...) fallback through a real kernel launch:

diff --git a/cuda_core/tests/test_launcher.py b/cuda_core/tests/test_launcher.py
@@
 def test_launch_scalar_argument(python_type, cpp_type, init_value):
     ...
 
+@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

Agent hints if you want to apply the same pattern elsewhere in this PR:

  • test_kernel_arg_numpy_subclass_isinstance_fallback: good candidate for the same treatment with a subclassed np.int32 or np.float32, since those scalar paths are unambiguous and already have nearby end-to-end launch coverage.
  • test_kernel_arg_python_isinstance_fallbacks: probably only worth doing for float or complex subclasses. I would be careful with int/bool subclasses here, because ParamHolder intentionally treats Python int values as pointer-like, so an end-to-end launch test may not exercise the scalar ABI you want.
  • I would keep the existing ParamHolder(...)-only tests even if you add one real launch test; they still provide cheap branch coverage across multiple accepted scalar families.

@copy-pr-bot
Copy link
Copy Markdown
Contributor

copy-pr-bot Bot commented Apr 17, 2026

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@rluo8
Copy link
Copy Markdown
Contributor Author

rluo8 commented Apr 17, 2026

Updated according to Ralf's comment.
Tests passed:
cuda_core/tests/test_launcher.py::test_launch_scalar_argument_ctypes_subclass_fallback
cuda_core/tests/test_launcher.py::test_launch_scalar_argument_numpy_subclass_fallback

@rwgk
Copy link
Copy Markdown
Contributor

rwgk commented Apr 17, 2026

LGTM, thanks @rluo8

We're having many infrastructure flakes at the moment. I'll try reruns tomorrow (PT) to get this merged.

@rwgk rwgk merged commit 330b30e into NVIDIA:main Apr 17, 2026
171 of 175 checks passed
github-actions Bot pushed a commit that referenced this pull request Apr 18, 2026
Removed preview folders for the following PRs:
- PR #1593
- PR #1908
- PR #1923
- PR #1925
- PR #1933
- PR #1934
- PR #1938
- PR #1939
mdboom pushed a commit to mdboom/cuda-python that referenced this pull request Apr 20, 2026
* tests: add coverage tests for cuda core

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

---------

Co-authored-by: Ralf W. Grosse-Kunstleve <rwgkio@gmail.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CI/CD CI/CD infrastructure cuda.core Everything related to the cuda.core module

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants