From 8645bc4ffd65697d60722fe2a5a583961f26f53f Mon Sep 17 00:00:00 2001 From: ksimpson Date: Wed, 5 Mar 2025 16:19:02 -0800 Subject: [PATCH] add public handle to object code --- cuda_core/cuda/core/experimental/_module.py | 6 ++++++ cuda_core/tests/test_module.py | 15 +++++++++++++++ 2 files changed, 21 insertions(+) diff --git a/cuda_core/cuda/core/experimental/_module.py b/cuda_core/cuda/core/experimental/_module.py index 39a7c0f5b..4ba3cff11 100644 --- a/cuda_core/cuda/core/experimental/_module.py +++ b/cuda_core/cuda/core/experimental/_module.py @@ -341,3 +341,9 @@ def get_kernel(self, name) -> Kernel: def code(self) -> CodeTypeT: """Return the underlying code object.""" return self._module + + @property + @precondition(_lazy_load_module) + def handle(self): + """Return the underlying handle object.""" + return self._handle diff --git a/cuda_core/tests/test_module.py b/cuda_core/tests/test_module.py index c780c1a03..608ee03bb 100644 --- a/cuda_core/tests/test_module.py +++ b/cuda_core/tests/test_module.py @@ -68,6 +68,16 @@ def get_saxpy_kernel_ptx(init_cuda): return ptx, mod +@pytest.fixture(scope="function") +def get_saxpy_object_code(init_cuda): + prog = Program(SAXPY_KERNEL, code_type="c++") + mod = prog.compile( + "cubin", + name_expressions=("saxpy", "saxpy"), + ) + return mod + + def test_get_kernel(init_cuda): kernel = """extern "C" __global__ void ABC() { }""" @@ -147,3 +157,8 @@ def test_object_code_load_cubin_from_file(get_saxpy_kernel, tmp_path): mod = ObjectCode.from_cubin(str(cubin_file), symbol_mapping=sym_map) assert mod.code == str(cubin_file) mod.get_kernel("saxpy") # force loading + + +def test_object_code_handle(get_saxpy_object_code): + mod = get_saxpy_object_code + assert mod.handle is not None