test(cuda.bindings): capture thread-locally in tests run on parallel threads

With pooled contexts the worker threads of a test start together; before,
serialized context creation kept them apart. Three tests begin a stream
capture in global mode, and in that mode a call such as cuGraphInstantiate
in one thread fails while another thread is capturing. The host-function
graph tests failed this way on the free-threaded rows. Capture in
thread-local mode instead; what the tests check does not depend on it.

Drop the synchronize at the end of each pooled test for the same reason:
it could break a global capture running in another thread. Pooled contexts
now behave like the module-scoped ctx fixture, which does not synchronize
between tests either.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Author: Andy-Jost

cuda_bindings/tests/conftest.py

@@ -49,8 +49,11 @@ class _ThreadContextPool:
     The pool gives each running thread a context of its own, as before, and
     keeps it for the next test. The plugin empties the pool at the end of each
     test module, so a context lives no longer than the module-scoped ``ctx``
-    fixture does in a run without threads. A context that reports an error at
-    the end of a test is destroyed rather than reused.
+    fixture does in a run without threads.
+
+    The worker threads of a test now run at the same time. Do not add a
+    synchronizing call between tests here: while another thread captures a
+    stream in global mode, such a call fails and breaks that capture.
     """
 
     def __init__(self):
@@ -81,18 +84,13 @@ def context(self):
         try:
             yield device, ctx
         finally:
-            (sync_err,) = cuda.cuCtxSynchronize()
             # Pop rather than unset, so that the context that was current
             # before the test (on the main thread, for tests that are not run
             # in parallel) is current again.
             err, _ = cuda.cuCtxPopCurrent()
             assert err == cuda.CUresult.CUDA_SUCCESS
-            if sync_err == cuda.CUresult.CUDA_SUCCESS:
-                with self._lock:
-                    self._idle.append(entry)
-            else:
-                (err,) = cuda.cuCtxDestroy(ctx)
-                assert err == cuda.CUresult.CUDA_SUCCESS
+            with self._lock:
+                self._idle.append(entry)
 
     def clear(self):
         """Destroy the idle contexts. Call only when no test is running."""

cuda_bindings/tests/test_cuda.py

@@ -706,7 +706,10 @@ def host_fn(userData):
     err, stream = cuda.cuStreamCreate(0)
     assert err == cuda.CUresult.CUDA_SUCCESS
     try:
-        (err,) = cuda.cuStreamBeginCapture(stream, cuda.CUstreamCaptureMode.CU_STREAM_CAPTURE_MODE_GLOBAL)
+        # Thread-local mode: pytest-run-parallel runs this test in several threads
+        # at once, and in global mode a call such as cuGraphInstantiate in one
+        # thread fails while another thread is capturing.
+        (err,) = cuda.cuStreamBeginCapture(stream, cuda.CUstreamCaptureMode.CU_STREAM_CAPTURE_MODE_THREAD_LOCAL)
         assert err == cuda.CUresult.CUDA_SUCCESS
         (err,) = cuda.cuLaunchHostFunc(stream, callback, ctypes.addressof(counter))
         assert err == cuda.CUresult.CUDA_SUCCESS

cuda_bindings/tests/test_cudart.py

@@ -616,7 +616,10 @@ def test_cudart_cudaStreamGetCaptureInfo():
     assert status == cudart.cudaStreamCaptureStatus.cudaStreamCaptureStatusNone
 
     # start capture
-    (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeGlobal)
+    # Thread-local mode: pytest-run-parallel runs this test in several threads
+    # at once, and in global mode a call such as cudaGraphInstantiate in one
+    # thread fails while another thread is capturing.
+    (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeThreadLocal)
     assertSuccess(err)
 
     # validate that stream is capturing now
@@ -1469,7 +1472,10 @@ def host_fn(userData):
     err, stream = cudart.cudaStreamCreate()
     assertSuccess(err)
     try:
-        (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeGlobal)
+        # Thread-local mode: pytest-run-parallel runs this test in several threads
+        # at once, and in global mode a call such as cudaGraphInstantiate in one
+        # thread fails while another thread is capturing.
+        (err,) = cudart.cudaStreamBeginCapture(stream, cudart.cudaStreamCaptureMode.cudaStreamCaptureModeThreadLocal)
         assertSuccess(err)
         (err,) = cudart.cudaLaunchHostFunc(stream, callback, ctypes.addressof(counter))
         assertSuccess(err)