Repository navigation
[BUG]: Device methods operate on the current context, not the device they're called on #2311
Description
Type of Bug
Silent Failure / Wrong Result
Component
cuda.core
Describe the bug
Several Device methods read as if they create a resource on that Device, but they actually create it on whatever CUDA context is current on the calling thread. The device's own context (self._context) is passed through only as metadata (for lifetime tracking and device_id labeling); it is never made current before the driver call, and the methods only guard with _check_context_initialized(), which checks "has set_current() ever been called on this object" — not "is this device current right now."
The result: device1.create_stream() can silently create a stream on device 0, and label it device_id == 1. The returned object's metadata disagrees with where the resource physically lives.
Affected methods
Device.create_stream(cuda_core/cuda/core/_device.pyx) →cuStreamCreateWithPriority/cuGreenCtxStreamCreate(no context arg; current context). AlsocuCtxGetStreamPriorityRangein_stream.pyxreads the current context.Device.create_event(cuda_core/cuda/core/_device.pyx) →cuEventCreate(no context arg; current context).- The four texture allocation methods currently in flight in cuda.core.texture: API-style follow-up before v1.1.0 (#2292) #2307 —
create_opaque_array,create_mipmapped_array,create_texture_object,create_surface_object— which delegate to helpers that use_get_current_device_id()and current-context driver calls. Their docstrings already document the footgun ("created in the current CUDA context, so make this device current withset_currentbefore calling").
How to Reproduce
from cuda.core import Device
dev0 = Device(0); dev0.set_current()
dev1 = Device(1); dev1.set_current() # dev1 marked initialized
dev0.set_current() # device 0 is now current
s = dev1.create_stream() # passes the init check (dev1 was inited once)
# but the stream is created on device 0!
# s._device_id reports 1, contradicting where the stream actually livesExpected behavior
deviceN.create_*(...) should create the resource on device N's context regardless of which context is currently active on the thread, and the returned object's device_id should be accurate. At minimum, if we intend to keep the "current context" semantics, the methods should assert that this device is current and raise rather than silently doing the wrong thing.
Proposed fix
Once #1586 lands and Device is a context manager, each method can wrap its body in a with self: block, which activates this device's context and restores the previous context on exit (exception-safe). That neatly fixes all six methods with a one-line change each, e.g.:
def create_stream(self, obj=None, options=None):
self._check_context_initialized()
with self:
return Stream._init(obj=obj, options=options, device_id=self._device_id, ctx=self._context)This depends on #1586 (Device context manager for temporary device switching), which already proposes stateless, driver-level save/restore that interoperates with other libraries.
Related
- Feature Request: Device context manager for temporary device switching #1586 — Device context manager (enables the
with self:fix) - cuda.core.texture: API-style follow-up before v1.1.0 (#2292) #2307 — texture API follow-up, which introduces the four
create_*texture methods in scope here - Deprecate
Device.set_current()? #548 — broader discussion of movingcuda.coreaway from implicit global device state (this bug is a concrete symptom)
Update (2026-08-24): Device.sync() has the same current-device behavior
Device.sync() is also affected. Its implementation checks only that self has been initialized and then calls cudaDeviceSynchronize(), which takes no device argument and synchronizes the device currently active on the calling thread.
Consequently:
dev0 = Device(0)
dev0.set_current()
dev1 = Device(1)
dev1.set_current()
dev0.set_current()
dev1.sync() # Synchronizes device 0, not device 1.This can silently fail to wait for work submitted to dev1. Under the current contract, callers must call dev1.set_current() immediately before dev1.sync().
This should be addressed consistently with the other affected Device methods: either make self current temporarily, assert that self is already current, or clearly document the current-device requirement.
Activity
Andy-Jost commentedon Jul 7, 2026
Assigning to @leofang for triage. This seems like a significant bug and perhaps a reason to revisit #1586.
Device.set_current()? #548leofang commentedon Jul 7, 2026
It is strictly speaking not a bug, but a design compromise that we deliberately made.
By doing anything smarter internally in cuda.core (ex: switching CUDA contexts), we risk performance regression that some of our users are extremely sensitive about, because calling any CUDA APIs is not cheap, and certain things like current device IDs must not be cached to avoid interoperability issues. So, the current contract (following cudart, which e.g. PyTorch and CuPy rely on) is that
dev.set_current()dev.do_somthing()with the same devicedevAny deviation is not considered compatible with the programming model.
dev.do_something()should not be read as "do something on this devicedev," but "do something on the devicedevthat is set to current."IIRC we wanted to have a Tips & Tricks (like the one for cuda-bindings) or Known Pitfalls to document things like this, but haven't had time to get to it...
I labeled #548 as the parent issue of this issue. If we move away from the cudart model as per #548, we need to bump the major version. A minimal, nonbreaking thing that we can do today at the cost of a small perf penalty is to do the suggested assertion:
Andy-Jost commentedon Jul 7, 2026
We currently have a situation where
s = dev1.create_stream()can create a stream on a device other than 1 while settings._device_idequal to 1. I think you're saying that's acceptable because that call represents undefined behavior (UB).I don't see how the performance or compatibility concerns constrain us. (1) Keeping the current context cached outside the driver means
with devicecan be a no-op at the same cost as the proposed assert -- when the device is already current, just check (without a driver call) and do nothing. The performance cost would be on the order of ns. (2) extending the semantics to permitwith devicedoes not break existing programs that follow the two rules you provided, so no major version bump is needed. What was previously UB becomes defined, and existing programs that follow the old rules continue to work as expected.leofang commentedon Jul 7, 2026
How would you achieve this without at least calling
cudaGetDeviceorcuCtxGetCurrent...?Andy-Jost commentedon Jul 7, 2026
I will answer more fully later, but in the meantime I got curious and measured the overhead of these calls so we can reason from real numbers.
Results (NVIDIA A10, driver 595.58.03, CUDA 13.2 driver API)
cuCtxGetCurrent(the call in question)cuCtxGetDevice(device-granularity check)cuCtxSetCurrent(same)(redundant switch)get + set + setwith-cycle (same device)Each is a tight-loop average over 20M iterations x 15 repeats; "best" and "median" agreed to within ~1%.
What this establishes
cuCtxGetCurrentis ~9 ns — solidly in the "ns-level, trivial" bucket.with self:/create_*can query live (~9 ns) and, when the device is already current, skip thecuCtxSetCurrententirely.Andy-Jost commentedon Jul 7, 2026
I agree caching is a non-starter. I had initially thought that a cache in cuda.core outside the driver could work, but I'm convinced now that's the wrong approach. Keeping track of the current device/context on our side will fail the moment a dependent like CuPy or PyTorch switches the device out from under us.
Fortunately, the performance numbers above suggest that checking the current context is not expensive. A call to
cuCtxGetCurrentdoes nothing more than read a value from TLS, and in particular does not make a syscall or acquire any lock. It looks to me like we can safely call it without even releasing the GIL.Given that, it seems to me that the best overall outcome is to resurrect #1597 and then update the
create_*methods to something like:This costs the same as the proposed assertion, but with a better result:
deviceN.create_*()does the least surprising and most useful thing by creating the resource on device N rather than raising when it notices a mismatch. And this is a purely additive change, so existing programs will continue to work and no version bump is needed.[-][BUG]: Device.create_* methods operate on the current context, not the device they're called on[/-][+][BUG]: Device methods operate on the current context, not the device they're called on[/+]cuda.core: create stream-ordering events in the recorded stream's con…
cuda.core: make Device methods use their bound context (#2750)