Skip to content

[CUDA] Set current device before allocating memory #3110

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
zcbenz merged 1 commit into from
Feb 8, 2026
Merged

[CUDA] Set current device before allocating memory #3110

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 1 addition & 4 deletions mlx/backend/cuda/allocator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -168,9 +168,6 @@ CudaAllocator::CudaAllocator()
free_limit_ = total_memory_ - memory_limit_;
max_pool_size_ = memory_limit_;

int curr;
CHECK_CUDA_ERROR(cudaGetDevice(&curr));

int device_count = gpu::device_count();
free_streams_.resize(device_count);
mem_pools_.resize(device_count);
Expand All @@ -181,7 +178,6 @@ CudaAllocator::CudaAllocator()
CHECK_CUDA_ERROR(cudaDeviceGetDefaultMemPool(&mem_pools_[i], i));
}
}
CHECK_CUDA_ERROR(cudaSetDevice(curr));
}

Buffer
Expand Down Expand Up @@ -223,6 +219,7 @@ CudaAllocator::malloc_async(size_t size, int device, cudaStream_t stream) {
if (device == -1) {
data = unified_malloc(size);
} else {
cu::device(device).make_current();
if (mem_pools_[device]) { // supports memory pools
CHECK_CUDA_ERROR(cudaMallocAsync(&data, size, stream));
} else {
Expand Down
8 changes: 5 additions & 3 deletions mlx/backend/cuda/event.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -213,7 +213,7 @@ auto check_gpu_coherency() {
return coherency;
}

AtomicEvent::AtomicEvent() {
AtomicEvent::AtomicEvent(Device& d) {
void* buf;
cudaError_t (*cuda_free)(void*);
// There are 3 kinds of systems we are implementing for:
Expand All @@ -223,6 +223,7 @@ AtomicEvent::AtomicEvent() {
// => use cuda::atom_ref on pinned host memory
// 2. no hardware cpu/gpu coherency
// => use cuda::atom_ref on device memory
d.make_current();
auto [concurrent_managed_access, host_native_atomic] = check_gpu_coherency();
if (concurrent_managed_access) {
CHECK_CUDA_ERROR(cudaMallocManaged(&buf, sizeof(uint32_t)));
Expand Down Expand Up @@ -347,11 +348,12 @@ struct EventImpl {
if (is_created()) {
return;
}
auto& d = cu::device(s.device);
if (s.device == mlx::core::Device::cpu || signal_value > 1) {
nvtx3::mark("Using slow AtomicEvent");
atomic = std::make_unique<cu::AtomicEvent>();
atomic = std::make_unique<cu::AtomicEvent>(d);
} else {
cuda = std::make_unique<cu::CopyableCudaEvent>(cu::device(s.device));
cuda = std::make_unique<cu::CopyableCudaEvent>(d);
}
}
};
Expand Down
2 changes: 1 addition & 1 deletion mlx/backend/cuda/event.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -54,7 +54,7 @@ class CudaEvent {
// CudaEvent so the latter should always be preferred when possible.
class AtomicEvent {
public:
AtomicEvent();
AtomicEvent(Device& d);

void wait(uint32_t value);
void wait(cudaStream_t stream, uint32_t value);
Expand Down
3 changes: 2 additions & 1 deletion mlx/backend/cuda/fence.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,8 @@ struct FenceImpl {

Fence::Fence(Stream s) {
fence_ = std::shared_ptr<void>(
new FenceImpl{0}, [](void* ptr) { delete static_cast<FenceImpl*>(ptr); });
new FenceImpl{0, cu::device(s.device)},
[](void* ptr) { delete static_cast<FenceImpl*>(ptr); });
}

void Fence::wait(Stream s, const array&) {
Expand Down
Loading