Cythonize _module.py (#1520)

* Cythonize _module.py (Phase 2a: cdef classes)

Convert Kernel, ObjectCode, and KernelOccupancy to cdef classes with
proper .pxd declarations. This phase establishes the Cython structure
while maintaining Python driver module usage.

Changes:
- Rename _module.py to _module.pyx
- Create _module.pxd with cdef class declarations
- Convert Kernel, ObjectCode, KernelOccupancy to cdef class
- Remove _backend dict in favor of direct driver calls
- Add _init_py() Python-accessible factory for ObjectCode
- Update _program.py and _linker.py to use _init_py()
- Fix test to handle cdef class property descriptors

Phase 2b will convert driver calls to cydriver with nogil blocks.
Phase 2c will add RAII handles to resource_handles.

* Phase 2a refinements: hide private attrs, add public properties

- Use strong types in .pxd (ObjectCode, KernelOccupancy)
- Remove cdef public - attributes now private to C level
- Add Kernel.handle property for external access
- Add ObjectCode.symbol_mapping property (symmetric with input)
- Update _launcher.pyx, _linker.py, tests to use public APIs

* Convert module-level functions and Kernel._get_arguments_info to cdef

- Module globals: _inited, _py_major_ver, _py_minor_ver, _driver_ver,
  _kernel_ctypes, _paraminfo_supported -> cdef typed
- Module functions: _lazy_init, _get_py_major_ver, _get_py_minor_ver,
  _get_driver_ver, _get_kernel_ctypes, _is_paraminfo_supported,
  _make_dummy_library_handle -> cdef inline with exception specs
- Module constant: _supported_code_type -> cdef tuple
- Kernel._get_arguments_info -> cdef tuple

Note: KernelAttributes remains a regular Python class due to
segfaults when converted to cdef class (likely due to weakref
interaction with cdef class properties).

* Convert KernelAttributes to cdef class

Follow the _MemPoolAttributes pattern:
- cdef class with inline cdef attributes (_kernel_weakref, _cache)
- _init as @classmethod (not @staticmethod cdef)
- _get_cached_attribute and _resolve_device_id use except? -1
- Explicit cast when dereferencing weakref

* Add LibraryHandle and KernelHandle to resource_handles infrastructure

Extends the RAII handle system to support CUlibrary and CUkernel driver
objects used in _module.pyx. This provides automatic lifetime management
and proper cleanup for library and kernel handles.

Changes:
- Add LibraryHandle/KernelHandle types with factory functions
- Update Kernel, ObjectCode, KernelOccupancy to use typed handles
- Move KernelAttributes cdef block to .pxd for strong typing
- Update _launcher.pyx to access kernel handle directly via cdef

* Convert _module.pyx driver calls to cydriver with nogil

Replaces Python-level driver API calls with low-level cydriver calls
wrapped in nogil blocks for improved performance. This allows the GIL
to be released during CUDA driver operations.

Changes:
- cuDriverGetVersion, cuKernelGetAttribute, cuKernelGetParamInfo
- cuOccupancy* functions (with appropriate GIL handling for callbacks)
- cuKernelGetLibrary
- Update KernelAttributes._get_cached_attribute to use cydriver types

* Fix SEGV in Kernel.from_handle with non-int types

Remove type annotation from handle parameter to prevent Cython's
automatic float-to-int coercion, which caused a segmentation fault.
The manual isinstance check properly validates all non-int types.

* Refactor ObjectCode._init and add kernel lifetime test

- Change ObjectCode._init from cdef to @classmethod def, matching the
  pattern used by Buffer, Stream, Graph, and other classes
- Remove _init_py wrapper (no longer needed)
- Update callers in _program.py and _linker.py
- Add test_kernel_keeps_library_alive to verify that a Kernel keeps its
  underlying library alive after ObjectCode goes out of scope

* Simplify resource handle patterns and clean up tests

- Remove Kernel._module (ObjectCode reference no longer needed since
  KernelHandle keeps library alive via LibraryHandle dependency)
- Simplify Kernel._from_obj signature (remove unused ObjectCode param)
- KernelAttributes: store KernelHandle instead of weakref to Kernel
- Rename get_kernel_from_library to create_kernel_handle for consistency
- Remove fragile annotation introspection from test_saxpy_arguments

* Simplify _MemPoolAttributes to use direct MemoryPoolHandle

Replace weakref pattern with direct MemoryPoolHandle storage in
_MemPoolAttributes. The handle's shared_ptr keeps the underlying
pool alive, so attributes remain accessible after the MR is deleted.

Note: _MemPool retains __weakref__ because the IPC subsystem uses
WeakValueDictionary to track memory resources across processes.

* Fix access violation in occupancy queries with uninitialized hStream

Zero-initialize CUlaunchConfig struct to prevent garbage values in
hStream field when no stream is provided. The driver dereferences
hStream even when querying occupancy, causing access violations on
some platforms (observed on Windows with RTX Pro 6000).

Author: Andy-Jost

cuda_core/cuda/core/_cpp/resource_handles.cpp

@@ -51,6 +51,11 @@ decltype(&cuMemFreeHost) p_cuMemFreeHost = nullptr;
 
 decltype(&cuMemPoolImportPointer) p_cuMemPoolImportPointer = nullptr;
 
+decltype(&cuLibraryLoadFromFile) p_cuLibraryLoadFromFile = nullptr;
+decltype(&cuLibraryLoadData) p_cuLibraryLoadData = nullptr;
+decltype(&cuLibraryUnload) p_cuLibraryUnload = nullptr;
+decltype(&cuLibraryGetKernel) p_cuLibraryGetKernel = nullptr;
+
 // ============================================================================
 // GIL management helpers
 // ============================================================================
@@ -682,4 +687,81 @@ DevicePtrHandle deviceptr_import_ipc(const MemoryPoolHandle& h_pool, const void*
     }
 }
 
+// ============================================================================
+// Library Handles
+// ============================================================================
+
+namespace {
+struct LibraryBox {
+    CUlibrary resource;
+};
+}  // namespace
+
+LibraryHandle create_library_handle_from_file(const char* path) {
+    GILReleaseGuard gil;
+    CUlibrary library;
+    if (CUDA_SUCCESS != (err = p_cuLibraryLoadFromFile(&library, path, nullptr, nullptr, 0, nullptr, nullptr, 0))) {
+        return {};
+    }
+
+    auto box = std::shared_ptr<const LibraryBox>(
+        new LibraryBox{library},
+        [](const LibraryBox* b) {
+            GILReleaseGuard gil;
+            p_cuLibraryUnload(b->resource);
+            delete b;
+        }
+    );
+    return LibraryHandle(box, &box->resource);
+}
+
+LibraryHandle create_library_handle_from_data(const void* data) {
+    GILReleaseGuard gil;
+    CUlibrary library;
+    if (CUDA_SUCCESS != (err = p_cuLibraryLoadData(&library, data, nullptr, nullptr, 0, nullptr, nullptr, 0))) {
+        return {};
+    }
+
+    auto box = std::shared_ptr<const LibraryBox>(
+        new LibraryBox{library},
+        [](const LibraryBox* b) {
+            GILReleaseGuard gil;
+            p_cuLibraryUnload(b->resource);
+            delete b;
+        }
+    );
+    return LibraryHandle(box, &box->resource);
+}
+
+LibraryHandle create_library_handle_ref(CUlibrary library) {
+    auto box = std::make_shared<const LibraryBox>(LibraryBox{library});
+    return LibraryHandle(box, &box->resource);
+}
+
+// ============================================================================
+// Kernel Handles
+// ============================================================================
+
+namespace {
+struct KernelBox {
+    CUkernel resource;
+    LibraryHandle h_library;  // Keeps library alive
+};
+}  // namespace
+
+KernelHandle create_kernel_handle(const LibraryHandle& h_library, const char* name) {
+    GILReleaseGuard gil;
+    CUkernel kernel;
+    if (CUDA_SUCCESS != (err = p_cuLibraryGetKernel(&kernel, *h_library, name))) {
+        return {};
+    }
+
+    return create_kernel_handle_ref(kernel, h_library);
+}
+
+KernelHandle create_kernel_handle_ref(CUkernel kernel, const LibraryHandle& h_library) {
+    auto box = std::make_shared<const KernelBox>(KernelBox{kernel, h_library});
+    return KernelHandle(box, &box->resource);
+}
+
 }  // namespace cuda_core

cuda_core/cuda/core/_cpp/resource_handles.hpp

@@ -61,6 +61,12 @@ extern decltype(&cuMemFreeHost) p_cuMemFreeHost;
 
 extern decltype(&cuMemPoolImportPointer) p_cuMemPoolImportPointer;
 
+// Library
+extern decltype(&cuLibraryLoadFromFile) p_cuLibraryLoadFromFile;
+extern decltype(&cuLibraryLoadData) p_cuLibraryLoadData;
+extern decltype(&cuLibraryUnload) p_cuLibraryUnload;
+extern decltype(&cuLibraryGetKernel) p_cuLibraryGetKernel;
+
 // ============================================================================
 // Handle type aliases - expose only the raw CUDA resource
 // ============================================================================
@@ -69,6 +75,8 @@ using ContextHandle = std::shared_ptr<const CUcontext>;
 using StreamHandle = std::shared_ptr<const CUstream>;
 using EventHandle = std::shared_ptr<const CUevent>;
 using MemoryPoolHandle = std::shared_ptr<const CUmemoryPool>;
+using LibraryHandle = std::shared_ptr<const CUlibrary>;
+using KernelHandle = std::shared_ptr<const CUkernel>;
 
 // ============================================================================
 // Context handle functions
@@ -218,6 +226,40 @@ StreamHandle deallocation_stream(const DevicePtrHandle& h) noexcept;
 // Set the deallocation stream for a device pointer handle.
 void set_deallocation_stream(const DevicePtrHandle& h, const StreamHandle& h_stream) noexcept;
 
+// ============================================================================
+// Library handle functions
+// ============================================================================
+
+// Create an owning library handle by loading from a file path.
+// When the last reference is released, cuLibraryUnload is called automatically.
+// Returns empty handle on error (caller must check).
+LibraryHandle create_library_handle_from_file(const char* path);
+
+// Create an owning library handle by loading from memory data.
+// The driver makes an internal copy of the data; caller can free it after return.
+// When the last reference is released, cuLibraryUnload is called automatically.
+// Returns empty handle on error (caller must check).
+LibraryHandle create_library_handle_from_data(const void* data);
+
+// Create a non-owning library handle (references existing library).
+// Use for borrowed libraries (e.g., from foreign code).
+// The library will NOT be unloaded when the handle is released.
+LibraryHandle create_library_handle_ref(CUlibrary library);
+
+// ============================================================================
+// Kernel handle functions
+// ============================================================================
+
+// Get a kernel from a library by name.
+// The kernel structurally depends on the provided library handle.
+// Kernels have no explicit destroy - their lifetime is tied to the library.
+// Returns empty handle on error (caller must check).
+KernelHandle create_kernel_handle(const LibraryHandle& h_library, const char* name);
+
+// Create a non-owning kernel handle with library dependency.
+// Use for borrowed kernels. The library handle keeps the library alive.
+KernelHandle create_kernel_handle_ref(CUkernel kernel, const LibraryHandle& h_library);
+
 // ============================================================================
 // Overloaded helper functions to extract raw resources from handles
 // ============================================================================
@@ -243,6 +285,14 @@ inline CUdeviceptr as_cu(const DevicePtrHandle& h) noexcept {
     return h ? *h : 0;
 }
 
+inline CUlibrary as_cu(const LibraryHandle& h) noexcept {
+    return h ? *h : nullptr;
+}
+
+inline CUkernel as_cu(const KernelHandle& h) noexcept {
+    return h ? *h : nullptr;
+}
+
 // as_intptr() - extract handle as intptr_t for Python interop
 // Using signed intptr_t per C standard convention and issue #1342
 inline std::intptr_t as_intptr(const ContextHandle& h) noexcept {
@@ -265,6 +315,14 @@ inline std::intptr_t as_intptr(const DevicePtrHandle& h) noexcept {
     return static_cast<std::intptr_t>(as_cu(h));
 }
 
+inline std::intptr_t as_intptr(const LibraryHandle& h) noexcept {
+    return reinterpret_cast<std::intptr_t>(as_cu(h));
+}
+
+inline std::intptr_t as_intptr(const KernelHandle& h) noexcept {
+    return reinterpret_cast<std::intptr_t>(as_cu(h));
+}
+
 // as_py() - convert handle to Python driver wrapper object (returns new reference)
 namespace detail {
 // n.b. class lookup is not cached to avoid deadlock hazard, see DESIGN.md
@@ -300,4 +358,12 @@ inline PyObject* as_py(const DevicePtrHandle& h) noexcept {
     return detail::make_py("CUdeviceptr", as_intptr(h));
 }
 
+inline PyObject* as_py(const LibraryHandle& h) noexcept {
+    return detail::make_py("CUlibrary", as_intptr(h));
+}
+
+inline PyObject* as_py(const KernelHandle& h) noexcept {
+    return detail::make_py("CUkernel", as_intptr(h));
+}
+
 }  // namespace cuda_core

cuda_core/cuda/core/_launch_config.pyx

@@ -2,6 +2,8 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
+from libc.string cimport memset
+
 from cuda.core._utils.cuda_utils cimport (
     HANDLE_RETURN,
 )
@@ -152,9 +154,9 @@ cdef class LaunchConfig:
 
     cdef cydriver.CUlaunchConfig _to_native_launch_config(self):
         _lazy_init()
-        # TODO: memset to zero?
         cdef cydriver.CUlaunchConfig drv_cfg
         cdef cydriver.CUlaunchAttribute attr
+        memset(&drv_cfg, 0, sizeof(drv_cfg))
         self._attrs.resize(0)
 
         # Handle grid dimensions and cluster configuration

cuda_core/cuda/core/_launcher.pyx

@@ -8,6 +8,7 @@ from cuda.bindings cimport cydriver
 
 from cuda.core._launch_config cimport LaunchConfig
 from cuda.core._kernel_arg_handler cimport ParamHolder
+from cuda.core._module cimport Kernel
 from cuda.core._resource_handles cimport as_cu
 from cuda.core._stream cimport Stream_accept, Stream
 from cuda.core._utils.cuda_utils cimport (
@@ -77,11 +78,11 @@ def launch(stream: Stream | GraphBuilder | IsStreamT, config: LaunchConfig, kern
     cdef ParamHolder ker_args = ParamHolder(kernel_args)
     cdef void** args_ptr = <void**><uintptr_t>(ker_args.ptr)
 
-    # TODO: cythonize Module/Kernel/...
     # Note: We now use CUkernel handles exclusively (CUDA 12+), but they can be cast to
     # CUfunction for use with cuLaunchKernel, as both handle types are interchangeable
     # for kernel launch purposes.
-    cdef cydriver.CUfunction func_handle = <cydriver.CUfunction>(<uintptr_t>(kernel._handle))
+    cdef Kernel ker = <Kernel>kernel
+    cdef cydriver.CUfunction func_handle = <cydriver.CUfunction>as_cu(ker._h_kernel)
 
     # Note: CUkernel can still be launched via cuLaunchKernel (not just cuLaunchKernelEx).
     # We check both binding & driver versions here mainly to see if the "Ex" API is

cuda_core/cuda/core/_linker.py

@@ -444,29 +444,29 @@ def __init__(self, *object_codes: ObjectCode, options: LinkerOptions = None):
             self._add_code_object(code)
 
     def _add_code_object(self, object_code: ObjectCode):
-        data = object_code._module
+        data = object_code.code
         with _exception_manager(self):
             name_str = f"{object_code.name}"
             if _nvjitlink and isinstance(data, bytes):
                 _nvjitlink.add_data(
                     self._mnff.handle,
-                    self._input_type_from_code_type(object_code._code_type),
+                    self._input_type_from_code_type(object_code.code_type),
                     data,
                     len(data),
                     name_str,
                 )
             elif _nvjitlink and isinstance(data, str):
                 _nvjitlink.add_file(
                     self._mnff.handle,
-                    self._input_type_from_code_type(object_code._code_type),
+                    self._input_type_from_code_type(object_code.code_type),
                     data,
                 )
             elif (not _nvjitlink) and isinstance(data, bytes):
                 name_bytes = name_str.encode()
                 handle_return(
                     _driver.cuLinkAddData(
                         self._mnff.handle,
-                        self._input_type_from_code_type(object_code._code_type),
+                        self._input_type_from_code_type(object_code.code_type),
                         data,
                         len(data),
                         name_bytes,
@@ -481,7 +481,7 @@ def _add_code_object(self, object_code: ObjectCode):
                 handle_return(
                     _driver.cuLinkAddFile(
                         self._mnff.handle,
-                        self._input_type_from_code_type(object_code._code_type),
+                        self._input_type_from_code_type(object_code.code_type),
                         data.encode(),
                         0,
                         None,

cuda_core/cuda/core/_memory/_memory_pool.pxd

@@ -19,6 +19,16 @@ cdef class _MemPool(MemoryResource):
         object                __weakref__
 
 
+cdef class _MemPoolAttributes:
+    cdef:
+        MemoryPoolHandle _h_pool
+
+    @staticmethod
+    cdef _MemPoolAttributes _init(MemoryPoolHandle h_pool)
+
+    cdef int _getattribute(self, cydriver.CUmemPool_attribute attr_enum, void* value) except? -1
+
+
 cdef class _MemPoolOptions:
 
     cdef:

cuda_core/cuda/core/_memory/_memory_pool.pyx

@@ -29,7 +29,6 @@ from cuda.core._utils.cuda_utils cimport (
 )
 
 import platform  # no-cython-lint
-import weakref
 
 from cuda.core._utils.cuda_utils import driver
 
@@ -45,16 +44,15 @@ cdef class _MemPoolOptions:
 
 
 cdef class _MemPoolAttributes:
-    cdef:
-        object _mr_weakref
+    """Provides access to memory pool attributes."""
 
     def __init__(self, *args, **kwargs):
         raise RuntimeError("_MemPoolAttributes cannot be instantiated directly. Please use MemoryResource APIs.")
 
-    @classmethod
-    def _init(cls, mr):
-        cdef _MemPoolAttributes self = _MemPoolAttributes.__new__(cls)
-        self._mr_weakref = mr
+    @staticmethod
+    cdef _MemPoolAttributes _init(MemoryPoolHandle h_pool):
+        cdef _MemPoolAttributes self = _MemPoolAttributes.__new__(_MemPoolAttributes)
+        self._h_pool = h_pool
         return self
 
     def __repr__(self):
@@ -64,12 +62,8 @@ cdef class _MemPoolAttributes:
         )
 
     cdef int _getattribute(self, cydriver.CUmemPool_attribute attr_enum, void* value) except?-1:
-        cdef _MemPool mr = <_MemPool>(self._mr_weakref())
-        if mr is None:
-            raise RuntimeError("_MemPool is expired")
-        cdef cydriver.CUmemoryPool pool_handle = as_cu(mr._h_pool)
         with nogil:
-            HANDLE_RETURN(cydriver.cuMemPoolGetAttribute(pool_handle, attr_enum, value))
+            HANDLE_RETURN(cydriver.cuMemPoolGetAttribute(as_cu(self._h_pool), attr_enum, value))
         return 0
 
     @property
@@ -197,8 +191,7 @@ cdef class _MemPool(MemoryResource):
     def attributes(self) -> _MemPoolAttributes:
         """Memory pool attributes."""
         if self._attributes is None:
-            ref = weakref.ref(self)
-            self._attributes = _MemPoolAttributes._init(ref)
+            self._attributes = _MemPoolAttributes._init(self._h_pool)
         return self._attributes
 
     @property