Does the current active device at buffer destruction need to be the same as during buffer construction?

[fkallen]

RMM's documentation (https://github.com/rapidsai/rmm#multiple-devices) mentions the requirements for deallocation in a multi-gpu settings. In particular, RMM tracks the current device at creation time and sets it automatically in the destructor to ensure the correct device is selected during deallocation.

I could not find such statements for CCCL classes, for example for thrust::device_vector or cuda::buffer.

[Jacobfaib]

Yes, the default thrust allocators require the current device be the same during destruction. This is a limitation we should fix.

If instead you use cuda::buffer then it will remember the device similar to RMM on destruction. Or you can use the RMM allocators as the allocator in thrust and it should similarly do The Right Thing.

[pciolkosz]

cuda::buffer won't try to do anything smart related to the current device, but at the same time none of the memory resources we provide require management of the current device on deallocation, so there should be no issue

[VectorPeak]

The key distinction is the owning container vs. the allocator / memory resource it delegates to. For thrust::device_vector with the default Thrust allocator, it is safest to treat destruction as device-sensitive today: destroy it while the same CUDA device is current, or wrap the lifetime/destruction site in a device guard. The allocator path does not give the vector enough state to reliably restore the allocation device on its own. For cuda::buffer, the destructor should be understood as “call deallocate on the associated memory resource,” not as “restore the allocation device like RMM does.” Whether changing the current device matters depends on the memory resource used. CCCL-provided resources are intended not to require the user to manually restore the current device for deallocation, but a custom resource could still have that requirement if it is implemented that way. A practical rule of thumb:

{
 cudaSetDevice(device);
 thrust::device_vector<int> v(n);
 // destroy v while `device` is still current, unless using a device-aware allocator/resource
}

If you need RMM-style behavior with Thrust containers, using an allocator backed by RMM is a reasonable workaround.

---

If my answer solved your problem, you can click answered the question. I'm really here to help, and along the way I'm also collecting Galaxy Brain badges haha 😆