From 600176e9de036d4c8685fd53f7e5082e21cb731c Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Fri, 10 Jul 2026 23:31:00 +0000 Subject: [PATCH 1/2] add an api to return module in ObjectCode to support calling legacy driver APIs --- cuda_core/cuda/core/_module.pyi | 10 ++++++++++ cuda_core/cuda/core/_module.pyx | 16 ++++++++++++++++ cuda_core/tests/test_module.py | 15 +++++++++++++++ 3 files changed, 41 insertions(+) diff --git a/cuda_core/cuda/core/_module.pyi b/cuda_core/cuda/core/_module.pyi index 5125b99131a..6c0cc85b50a 100644 --- a/cuda_core/cuda/core/_module.pyi +++ b/cuda_core/cuda/core/_module.pyi @@ -458,6 +458,16 @@ class ObjectCode: """ + def get_module(self) -> object: + """Return the :obj:`~driver.CUmodule` associated with this object code. + + Returns + ------- + :obj:`~driver.CUmodule` + Module handle for the current CUDA context, suitable for legacy + driver APIs that accept ``CUmodule``. + """ + @property def code(self) -> CodeTypeT: """Return the underlying code object.""" diff --git a/cuda_core/cuda/core/_module.pyx b/cuda_core/cuda/core/_module.pyx index 91c8ad43895..990a2c8445c 100644 --- a/cuda_core/cuda/core/_module.pyx +++ b/cuda_core/cuda/core/_module.pyx @@ -6,6 +6,7 @@ from __future__ import annotations cimport cython from libc.stddef cimport size_t +from libc.stdint cimport intptr_t from collections import namedtuple from os import fsencode, fspath, PathLike @@ -796,6 +797,21 @@ cdef class ObjectCode: HANDLE_RETURN(get_last_error()) return Kernel._from_handle(h_kernel) + def get_module(self) -> object: + """Return the :obj:`~driver.CUmodule` associated with this object code. + + Returns + ------- + :obj:`~driver.CUmodule` + Module handle for the current CUDA context, suitable for legacy + driver APIs that accept ``CUmodule``. + """ + self._lazy_load_module() + cdef cydriver.CUmodule mod + with nogil: + HANDLE_RETURN(cydriver.cuLibraryGetModule(&mod, as_cu(self._h_library))) + return driver.CUmodule(mod) + @property def code(self) -> CodeTypeT: """Return the underlying code object.""" diff --git a/cuda_core/tests/test_module.py b/cuda_core/tests/test_module.py index a0600966619..b5c9b4bee7a 100644 --- a/cuda_core/tests/test_module.py +++ b/cuda_core/tests/test_module.py @@ -226,6 +226,21 @@ def test_get_kernel(init_cuda): assert object_code.get_kernel(b"ABC").handle is not None +def test_object_code_get_module_for_legacy_integration(init_cuda): + src = """ + extern "C" __global__ void ABC() { } + extern "C" __global__ void DEF() { } + """ + object_code = Program(src, "c++").compile("cubin") + + # Bridge: CUlibrary (new) → CUmodule (legacy) + module = object_code.get_module() + + # Legacy module-only API consumes it directly + count = handle_return(driver.cuModuleGetFunctionCount(module)) + assert count == 2 + + @pytest.mark.parametrize( "attr, expected_type", [ From b56e37ebc2a8c7e512eb7737119941ef9ff100f9 Mon Sep 17 00:00:00 2001 From: Jinfeng Date: Sat, 18 Jul 2026 00:38:47 +0000 Subject: [PATCH 2/2] cuda.core: clarify ObjectCode.handle vs get_module() docstrings Document that .handle returns the native context-independent CUlibrary for cuda.core/newer APIs, while get_module() is a legacy-interop bridge returning a context-dependent CUmodule via cuLibraryGetModule. --- cuda_core/cuda/core/_module.pyi | 11 +++++++++-- cuda_core/cuda/core/_module.pyx | 11 +++++++++-- 2 files changed, 18 insertions(+), 4 deletions(-) diff --git a/cuda_core/cuda/core/_module.pyi b/cuda_core/cuda/core/_module.pyi index 6c0cc85b50a..f51b4cb2817 100644 --- a/cuda_core/cuda/core/_module.pyi +++ b/cuda_core/cuda/core/_module.pyi @@ -459,7 +459,11 @@ class ObjectCode: """ def get_module(self) -> object: - """Return the :obj:`~driver.CUmodule` associated with this object code. + """Return a context-dependent :obj:`~driver.CUmodule` for legacy interop. + + Bridges the native :obj:`~driver.CUlibrary` (see :attr:`handle`) to a + ``CUmodule`` via ``cuLibraryGetModule``, for use with legacy driver APIs + that only accept ``CUmodule``. Returns ------- @@ -486,7 +490,10 @@ class ObjectCode: @property def handle(self) -> object: - """Return the underlying handle object. + """Return the native, context-independent :obj:`~driver.CUlibrary` handle. + + Used by ``cuda.core`` and newer driver library APIs. For legacy APIs + that only accept a ``CUmodule``, use :meth:`get_module` instead. .. caution:: diff --git a/cuda_core/cuda/core/_module.pyx b/cuda_core/cuda/core/_module.pyx index 990a2c8445c..5734e31414b 100644 --- a/cuda_core/cuda/core/_module.pyx +++ b/cuda_core/cuda/core/_module.pyx @@ -798,7 +798,11 @@ cdef class ObjectCode: return Kernel._from_handle(h_kernel) def get_module(self) -> object: - """Return the :obj:`~driver.CUmodule` associated with this object code. + """Return a context-dependent :obj:`~driver.CUmodule` for legacy interop. + + Bridges the native :obj:`~driver.CUlibrary` (see :attr:`handle`) to a + ``CUmodule`` via ``cuLibraryGetModule``, for use with legacy driver APIs + that only accept ``CUmodule``. Returns ------- @@ -834,7 +838,10 @@ cdef class ObjectCode: @property def handle(self) -> object: - """Return the underlying handle object. + """Return the native, context-independent :obj:`~driver.CUlibrary` handle. + + Used by ``cuda.core`` and newer driver library APIs. For legacy APIs + that only accept a ``CUmodule``, use :meth:`get_module` instead. .. caution::