Add __weakref__ support to core API classes (#1533)

* Add __weakref__ support to core API classes

Enable weak referencing for Cython cdef classes:
- Stream, Event, Context (cdef classes via pxd)
- Buffer, LaunchConfig (cdef classes via pxd)

Enable weak referencing for Python classes with __slots__:
- Device, ObjectCode (added __weakref__ to __slots__)

Note: Kernel already had __weakref__ in its __slots__.
Memory resource classes inherit __weakref__ from _MemPool.

Add test_weakref.py to verify all core API classes are
weak-referenceable.

* Fix test_weakref kernel fixture PTX version incompatibility

Change object_code fixture to compile to cubin instead of ptx to avoid
CUDA_ERROR_UNSUPPORTED_PTX_VERSION when the toolkit version is newer
than the driver version on test machines.

Author: Andy-Jost

cuda_core/cuda/core/_context.pxd

@@ -14,6 +14,7 @@ cdef class Context:
     cdef:
         ContextHandle _h_context
         int _device_id
+        object __weakref__
 
     @staticmethod
     cdef Context _from_handle(type cls, ContextHandle h_context, int device_id)

cuda_core/cuda/core/_device.pyx

@@ -955,7 +955,7 @@ class Device:
         Default value of `None` return the currently used device.
 
     """
-    __slots__ = ("_device_id", "_memory_resource", "_has_inited", "_properties", "_uuid", "_context")
+    __slots__ = ("_device_id", "_memory_resource", "_has_inited", "_properties", "_uuid", "_context", "__weakref__")
 
     def __new__(cls, device_id: Device | int | None = None):
         if isinstance(device_id, Device):

cuda_core/cuda/core/_event.pxd

@@ -16,6 +16,7 @@ cdef class Event:
         bint _ipc_enabled
         object _ipc_descriptor
         int _device_id
+        object __weakref__
 
     @staticmethod
     cdef Event _init(type cls, int device_id, ContextHandle h_context, options, bint is_free)

cuda_core/cuda/core/_launch_config.pxd

@@ -17,6 +17,7 @@ cdef class LaunchConfig:
         public bint cooperative_launch
 
         vector[cydriver.CUlaunchAttribute] _attrs
+        object __weakref__
 
     cdef cydriver.CUlaunchConfig _to_native_launch_config(self)
 

cuda_core/cuda/core/_memory/_buffer.pxd

@@ -23,6 +23,7 @@ cdef class Buffer:
         object          _owner
         _MemAttrs       _mem_attrs
         bint            _mem_attrs_inited
+        object          __weakref__
 
 
 cdef class MemoryResource:

cuda_core/cuda/core/_module.py

@@ -546,7 +546,7 @@ class ObjectCode:
     :class:`~cuda.core.Program`
     """
 
-    __slots__ = ("_handle", "_code_type", "_module", "_loader", "_sym_map", "_name")
+    __slots__ = ("_handle", "_code_type", "_module", "_loader", "_sym_map", "_name", "__weakref__")
     _supported_code_type = ("cubin", "ptx", "ltoir", "fatbin", "object", "library")
 
     def __new__(self, *args, **kwargs):

cuda_core/cuda/core/_stream.pxd

@@ -13,6 +13,7 @@ cdef class Stream:
         int _device_id
         int _nonblocking
         int _priority
+        object __weakref__
 
     @staticmethod
     cdef Stream _from_handle(type cls, StreamHandle h_stream)

cuda_core/tests/test_weakref.py

@@ -0,0 +1,75 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import weakref
+
+import pytest
+from cuda.core import Device
+
+
+@pytest.fixture(scope="module")
+def device():
+    dev = Device()
+    dev.set_current()
+    return dev
+
+
+@pytest.fixture
+def stream(device):
+    return device.create_stream()
+
+
+@pytest.fixture
+def event(device):
+    return device.create_event()
+
+
+@pytest.fixture
+def context(device):
+    return device.context
+
+
+@pytest.fixture
+def buffer(device):
+    return device.allocate(1024)
+
+
+@pytest.fixture
+def launch_config():
+    from cuda.core import LaunchConfig
+
+    return LaunchConfig(grid=(1,), block=(1,))
+
+
+@pytest.fixture
+def object_code():
+    from cuda.core import Program
+
+    prog = Program('extern "C" __global__ void test_kernel() {}', "c++")
+    return prog.compile("cubin")
+
+
+@pytest.fixture
+def kernel(object_code):
+    return object_code.get_kernel("test_kernel")
+
+
+WEAK_REFERENCEABLE = [
+    "device",
+    "stream",
+    "event",
+    "context",
+    "buffer",
+    "launch_config",
+    "object_code",
+    "kernel",
+]
+
+
+@pytest.mark.parametrize("fixture_name", WEAK_REFERENCEABLE)
+def test_weakref(fixture_name, request):
+    """Core API classes should be weak-referenceable."""
+    obj = request.getfixturevalue(fixture_name)
+    ref = weakref.ref(obj)
+    assert ref() is obj