Add tests for library/kernel lifetime vs context lifetime

Experiment to detect segfaults on CI platforms where cuLibraryUnload
after context destruction may crash. LibraryBox does not currently
store a ContextHandle, so nothing prevents the context from being
destroyed before the library.

Co-authored-by: Cursor <cursoragent@cursor.com>

Author: Andy-Jost

cuda_core/tests/test_library_context_lifetime.py

@@ -0,0 +1,149 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: LicenseRef-NVIDIA-SOFTWARE-LICENSE
+
+"""Tests for library/kernel lifetime vs context lifetime.
+
+These tests exercise scenarios where a CUDA library (CUlibrary) or kernel
+(CUkernel) outlives the context in which it was created. This can happen
+when:
+  - The primary context is reset (e.g. numba-cuda's test teardown)
+  - A non-primary context is explicitly destroyed
+
+Currently, LibraryBox in resource_handles.cpp does NOT store a
+ContextHandle, so nothing prevents the context from being destroyed
+before the library. The library's RAII deleter then calls cuLibraryUnload
+on a library whose owning context is gone, which may segfault on some
+driver versions.
+"""
+
+import gc
+
+import pytest
+from cuda.bindings import driver
+from cuda.core import Device, Program
+from cuda.core._utils.cuda_utils import handle_return
+
+KERNEL_SOURCE = 'extern "C" __global__ void test_kernel() {}'
+
+
+def _compile_and_get_kernel():
+    """Compile a trivial kernel and return (ObjectCode, Kernel)."""
+    prog = Program(KERNEL_SOURCE, "c++")
+    obj = prog.compile("cubin")
+    kernel = obj.get_kernel("test_kernel")
+    return obj, kernel
+
+
+def _create_nonprimary_context(dev=0):
+    """Create a non-primary context, handling CUDA 12.x vs 13.x API differences."""
+    try:
+        return handle_return(driver.cuCtxCreate(None, 0, dev))
+    except TypeError:
+        return handle_return(driver.cuCtxCreate(0, dev))
+
+
+@pytest.fixture(autouse=True)
+def _restore_primary_context():
+    """Re-establish a valid primary context after each test."""
+    yield
+    ctx = handle_return(driver.cuDevicePrimaryCtxRetain(0))
+    handle_return(driver.cuCtxSetCurrent(ctx))
+
+
+class TestPrimaryContextReset:
+    """Library/kernel destroyed after primary context reset."""
+
+    def test_objectcode_outlives_primary_context_reset(self):
+        dev = Device(0)
+        dev.set_current()
+
+        obj, kernel = _compile_and_get_kernel()
+        del kernel
+
+        handle_return(driver.cuDevicePrimaryCtxReset(0))
+
+        del obj
+        gc.collect()
+
+    def test_kernel_outlives_primary_context_reset(self):
+        dev = Device(0)
+        dev.set_current()
+
+        obj, kernel = _compile_and_get_kernel()
+        del obj
+
+        handle_return(driver.cuDevicePrimaryCtxReset(0))
+
+        del kernel
+        gc.collect()
+
+    def test_objectcode_outlives_primary_context_reset_and_reretain(self):
+        """The numba-cuda pattern: reset, then re-init for the next test."""
+        dev = Device(0)
+        dev.set_current()
+
+        obj, kernel = _compile_and_get_kernel()
+        del kernel
+
+        handle_return(driver.cuDevicePrimaryCtxReset(0))
+
+        ctx = handle_return(driver.cuDevicePrimaryCtxRetain(0))
+        handle_return(driver.cuCtxSetCurrent(ctx))
+
+        del obj
+        gc.collect()
+
+
+class TestNonPrimaryContextDestroy:
+    """Library/kernel destroyed after non-primary context is destroyed."""
+
+    def test_objectcode_outlives_nonprimary_context(self):
+        ctx = _create_nonprimary_context()
+
+        obj, kernel = _compile_and_get_kernel()
+        del kernel
+
+        handle_return(driver.cuCtxDestroy(ctx))
+
+        del obj
+        gc.collect()
+
+    def test_kernel_outlives_nonprimary_context(self):
+        ctx = _create_nonprimary_context()
+
+        obj, kernel = _compile_and_get_kernel()
+        del obj
+
+        handle_return(driver.cuCtxDestroy(ctx))
+
+        del kernel
+        gc.collect()
+
+
+class TestKernelOutlivesObjectCode:
+    """Kernel transitively holds library; both outlive context."""
+
+    def test_kernel_outlives_objectcode_and_primary_context(self):
+        dev = Device(0)
+        dev.set_current()
+
+        obj, kernel = _compile_and_get_kernel()
+        del obj
+        gc.collect()
+
+        handle_return(driver.cuDevicePrimaryCtxReset(0))
+
+        del kernel
+        gc.collect()
+
+    def test_kernel_outlives_objectcode_and_nonprimary_context(self):
+        ctx = _create_nonprimary_context()
+
+        obj, kernel = _compile_and_get_kernel()
+        del obj
+        gc.collect()
+
+        handle_return(driver.cuCtxDestroy(ctx))
+
+        del kernel
+        gc.collect()