Skip to content

Add ObjectCode.get_module() for legacy driver API integration #2339

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
lijinf2 wants to merge 2 commits into
base: main
Choose a base branch
from
+57 −2
Open

Add ObjectCode.get_module() for legacy driver API integration #2339

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
600176e
add an api to return module in ObjectCode to support calling legacy d…
lijinf2 Jul 10, 2026
b56e37e
cuda.core: clarify ObjectCode.handle vs get_module() docstrings
lijinf2 Jul 18, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 18 additions & 1 deletion cuda_core/cuda/core/_module.pyi
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -458,6 +458,20 @@ class ObjectCode:

"""

def get_module(self) -> object:
"""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
-------
: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."""
Expand All @@ -476,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::

Expand Down
25 changes: 24 additions & 1 deletion cuda_core/cuda/core/_module.pyx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -796,6 +797,25 @@ cdef class ObjectCode:
HANDLE_RETURN(get_last_error())
return Kernel._from_handle(h_kernel)

def get_module(self) -> object:
"""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
-------
: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(<intptr_t>mod)

@leofang leofang Jul 14, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Q: Shouldn't we use as_py() to do the return value conversion?

Q: How do we teach the difference between .get_module() and .handle?

@lijinf2 lijinf2 Jul 18, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry for not getting back sooner. I shifted focus to issue #2179, for which I need to set up cuda bindings development environment for the first time.

It looks as_py() does not provide support for CUmodule type yet. I am not sure if we should add the support there per CUmodule is deemed as legacy API. But let me know if you find the support favorable. Happy to add that.

For .handle vs .get_module(): I'll clarify in the docstring.

@lijinf2 lijinf2 Jul 18, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done


@property
def code(self) -> CodeTypeT:
"""Return the underlying code object."""
Expand All @@ -818,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::

Expand Down
15 changes: 15 additions & 0 deletions cuda_core/tests/test_module.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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",
[
Expand Down
Loading