[BUG]: Primary-context TLS destructor segfaults after Python finalization

[bdice]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Runtime Error

Component

cuda.core

Describe the bug

Initializing a primary CUDA context through cuda.core.Device.set_current() causes the process to segfault during interpreter shutdown. The CUDA operation itself succeeds; the crash happens after Python atexit handlers and finalization, when libc destroys CUDA Core's main-thread TLS cache.

This is reproducible without Numba or cuDF:

cuda.core.Device.set_current()
-> cuda_core::get_primary_context()
-> primary ContextHandle stored in thread_local primary_context_cache
-> CPython finalization completes
-> libc __call_tls_dtors destroys primary_context_cache
-> ContextHandle deleter calls p_cuDevicePrimaryCtxRelease
-> cuda.bindings.cydriver Cython wrapper calls PyGILState_Ensure
-> PyThreadState_New
-> SIGSEGV

The relevant CUDA Core code is:

static thread_local std::vector<ContextHandle> primary_context_cache;

// ...

[device_id](const ContextBox* b) {
    context_registry.unregister_handle(b->resource);
    GILReleaseGuard gil;
    p_cuDevicePrimaryCtxRelease(device_id);
    delete b;
}

p_cuDevicePrimaryCtxRelease is populated from cuda.bindings.cydriver.__pyx_capi__["cuDevicePrimaryCtxRelease"]. Although GILReleaseGuard recognizes that Python has finalized and avoids its own GIL operations, CUDA Core still invokes the Cython wrapper. The wrapper's generated initialization path touches Python thread state and is not safe after Py_Finalize().

This appears to be a regression introduced between CUDA Core 0.5.1 and 0.6.0. The TLS cache originated in #1368, and #1466 changed its driver function pointers from raw cuGetProcAddress entry points to Cython __pyx_capi__ wrappers.

Confirmed with the same minimal reproducer:

cuda-core	Result
0.5.0	exits 0
0.5.1	exits 0
0.6.0	exits 139
0.7.0	exits 139
1.0.0	exits 139
1.0.1	exits 139
1.1.0	exits 139
1.1.1	exits 139

Current main retains the same TLS-cache and __pyx_capi__ deleter pattern.

This was exposed by a cuDF health check in NVIDIA/cudf#21928. The downstream path is:

cuDF Series.apply
-> numba.cuda.compile_ptx_for_current_device
-> Numba-CUDA current-device discovery
-> cuda.core.Device.set_current
-> shutdown segfault

cuDF only calls Numba-CUDA's public API. Numba-CUDA currently chooses CUDA Core's Device abstraction for device/context management. The shutdown lifetime bug itself is in CUDA Core.

How to Reproduce

Run:

python -c 'from cuda.core import Device; Device().set_current(); print("initialized")'
echo $?

Observed output:

initialized
Segmentation fault (core dumped)
139

The equivalent Numba-CUDA reproducer also crashes:

python -c 'from numba import cuda; print(cuda.current_context())'

Calling cuda.close(), resetting the Numba device, deleting the context, or forcing garbage collection does not prevent the crash because none of those operations clear CUDA Core's TLS cache.

The native backtrace from catchsegv is:

python
PyThreadState_New
cuda/bindings/_bindings/cydriver.so
cuda/bindings/cydriver.so
cuda/core/_resource_handles.so
libc.so.6(__call_tls_dtors)

Additional isolation:

• A direct cuda.bindings.driver.cuDevicePrimaryCtxRetain / cuCtxSetCurrent / cuDevicePrimaryCtxRelease sequence exits 0.
• Running Device().set_current() in a worker thread and joining it exits 0 because that thread's TLS destructor runs while CPython is still alive.
• A registered Python atexit handler completes before the main-thread crash, confirming that the failure is post-atexit shutdown.

The original failing CI job is https://github.com/NVIDIA/cudf/actions/runs/33404951587/job/99539765205?pr=21928#step:12:3585.

Expected behavior

A process that successfully initializes a CUDA Core device should exit normally.

A small fix is to avoid invoking the Cython wrapper once Python is no longer usable and leave primary-context cleanup to process teardown:

if (Py_IsInitialized() && !py_is_finalizing()) {
    GILReleaseGuard gil;
    p_cuDevicePrimaryCtxRelease(device_id);
}

This preserves normal primary-context release when a thread exits while Python is running. During process shutdown, intentionally leaving the retain for CUDA/OS process teardown is safer than calling into finalized CPython.

An architectural alternative is to retain a raw cuDevicePrimaryCtxRelease entry point resolved through cuGetProcAddress, which could be invoked without touching Python during TLS destruction.

A prototype of the guarded fix plus a subprocess regression test was built against the current CUDA Python checkout. The exact shutdown probe exits 0, and the complete CUDA Core test_device.py module reports 166 passed and 2 skipped.

Operating System

Reproduced in Rocky Linux 8 CUDA CI containers. The original CI failure used Python 3.11, CUDA 12.2.2, NVIDIA driver 580.173.02, and an NVIDIA L4. It was also reproduced locally with a different GPU and driver, so it does not appear specific to the original hardware or driver.

nvidia-smi output

Original CI environment: NVIDIA L4, driver 580.173.02. Local reproduction: NVIDIA RTX 1000 Ada Generation Laptop GPU, driver 610.47.