cuMemSetAccess Failed to create GPU mapping

Accelerated Computing CUDA CUDA on Windows Subsystem for Linux
,
1

Sorry for cross posting, but I might have stumbled onto a WSL2 specific problem and want some clarifications.

Symptom 1: cuMemSetAccess() returns CUDA_ERROR_NOT_READY

While growing a virtual-memory-management (VMM) backed memory pool toward a device’s memory limit, cuMemSetAccess() intermittently returns CUDA_ERROR_NOT_READY
(600). Nothing in the driver documentation indicates this call should ever return that code under these conditions.

The workflow being exercised is the standard VMM allocation pattern:

  1. cuMemCreate() physical handles
  2. cuMemMap() it into a reserved virtual address range
  3. cuMemSetAccess() to grant read/write permission on the mapped range

Step 3 is the one that fails — and only after step 2 has already succeeded on the same handle, from the same thread, with no other thread touching that memory.

What does CUDA_ERROR_NOT_READY mean when returned by cuMemSetAccess?

From local experiments, it seems retrying a few times fixes the issue. Is small retry the recommended approach when receiving CUDA_ERROR_NOT_READY from cuMemSetAccess?

The repro

A minimal repro is here: random_stuff/device_not_ready at main · thisisjimmyfb/random_stuff · GitHub
Simply run build_and_run.sh, succesfully reproducing the issue looks like this:

=== device_not_ready ===
VMM min granularity: 2097152 bytes, using handle size: 33554432 bytes
Device memory: 14953 / 16302 MB free
Done: 257 allocations created, 0MB free remaining.
cuMemSetAccess returned CUDA_ERROR_NOT_READY 1 time(s)

Compute Sanitizer:

=== error_mem_alloc ===
========= COMPUTE-SANITIZER
Sizes (MB): 7 14 28 57 115
Primed near-OOM: filler_chunks=12 headroom_target=5635MB
[start] nvidia_free=6085MB / 16302MB
========= CUDA API Error: Failed to create GPU mapping
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: cuMemSetAccess [0x3562c9] in libcuda.so.1.1
=========         Host Frame: VMMAllocator::try_allocate_block(unsigned long) [0x8f5c] in error_mem_alloc
=========         Host Frame: VMMAllocator::allocate(unsigned long) [0x830e] in error_mem_alloc
=========         Host Frame: main [0x5dc9] in error_mem_alloc

========= Program hit CUDA_ERROR_NOT_READY (error 600) due to "device not ready" on CUDA API call to cuMemSetAccess.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: VMMAllocator::try_allocate_block(unsigned long) [0x8f5c] in error_mem_alloc
=========         Host Frame: VMMAllocator::allocate(unsigned long) [0x830e] in error_mem_alloc
=========         Host Frame: main [0x5dc9] in error_mem_alloc

========= Program hit CUDA_ERROR_OUT_OF_MEMORY (error 2) due to "out of memory" on CUDA API call to cuMemSetAccess.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: VMMAllocator::try_allocate_block(unsigned long) [0x8f5c] in error_mem_alloc
=========         Host Frame: VMMAllocator::allocate(unsigned long) [0x830e] in error_mem_alloc
=========         Host Frame: main [0x5dc9] in error_mem_alloc

Observed on Driver: 610.43.02 (CUDA UMD 13.3) · CUDA Toolkit (nvcc): 13.3.33 · GPU: RTX 5080 (SM120), WSL2

Symptom 2: cudaStreamSynchronize() / cudaEventSynchronize() returning cudaErrorMemoryAllocation

Setup (one single foreground thread, multiple background threads):

Foreground thread, once per round for a total of 20 rounds, running concurrently with 8 background threads:

  1. allocate memory via cuMemCreate() / cuMemMap() / cuMemSetAccess()
  2. if allocation fails: cudaEventSynchronize() on previous cudaEventRecord, and then cuMemUnmap() / cuMemRelease (potential failure point, much more rare compared to other failure points)
  3. launch a kernel using the allocated memory on the foreground stream
  4. cudaGetLastError() to check the launch
  5. cudaStreamSynchronize() on the foreground stream (high chance failure point here)

Each background thread (8 of them, running until the foreground loop finishes), in a tight loop:

  1. allocate memory via cuMemCreate() / cuMemMap() / cuMemSetAccess()
  2. if allocation fails: cudaEventSynchronize() on previous cudaEventRecord, and then cuMemUnmap() / cuMemRelease
  3. cudaStreamSynchronize() on the background stream (potential failure point)

Assuming cudaStreamSynchronize() / cudaEventSynchronize() returning cudaErrorMemoryAllocation is simply a deferred asynchronous error code, is there a good way to find out which asynchronous call caused the error? I have tried checking the return code of all prior calls and they all succeeded. What is the good action to take after receiving cudaErrorMemoryAllocation from synchronization calls?

The repro

A minimal repro is here: random_stuff/error_mem_alloc at main · thisisjimmyfb/random_stuff · GitHub
Simply run build_and_run.sh, successfully reproducing the problem looks like this:

=== error_mem_alloc ===
Sizes (MB): 7 14 29 58 116
Primed near-OOM: filler_chunks=12 headroom_target=5712MB
  [start] nvidia_free=6313MB / 16302MB
  [fg round 0] nvidia_free=0MB / 16302MB
  [fg round 1] nvidia_free=0MB / 16302MB
  [end] nvidia_free=0MB / 0MB
  [bg0] iters=787
  [bg1] iters=1305
  [bg2] iters=978
  [bg3] iters=1119
  [bg4] iters=1580
  [bg5] iters=1729
  [bg6] iters=1205
  [bg7] iters=1181
TOTALS: sync_oom=1 meminfo_oom=0 bg_iters=9884 fg_oom=0
REPRODUCED: a sync/meminfo call reported OOM even though it does not itself allocate memory.
FG FAILED: CUDA error: out of memory at vmm_allocator.h:128

Compute Sanitizer:

=== device_not_ready ===
========= COMPUTE-SANITIZER
VMM min granularity: 2097152 bytes, using handle size: 33554432 bytes
Device memory: 14725 / 16302 MB free
========= CUDA API Error: Failed to create GPU mapping
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: cuMemSetAccess [0x3562c9] in libcuda.so.1.1
=========         Host Frame: (anonymous namespace)::Allocation::try_create(unsigned long long, int, (anonymous namespace)::Allocation*, int*) [0xa9d8] in device_not_ready
=========         Host Frame: main [0xa37c] in device_not_ready

========= Program hit CUDA_ERROR_NOT_READY (error 600) due to "device not ready" on CUDA API call to cuMemSetAccess.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: (anonymous namespace)::Allocation::try_create(unsigned long long, int, (anonymous namespace)::Allocation*, int*) [0xa9d8] in device_not_ready
=========         Host Frame: main [0xa37c] in device_not_ready

========= Program hit CUDA_ERROR_OUT_OF_MEMORY (error 2) due to "out of memory" on CUDA API call to cuMemSetAccess.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: (anonymous namespace)::Allocation::try_create(unsigned long long, int, (anonymous namespace)::Allocation*, int*) [0xa9d8] in device_not_ready
=========         Host Frame: main [0xa37c] in device_not_ready

Done: 257 allocations created, 0MB free remaining.
cuMemSetAccess returned CUDA_ERROR_NOT_READY 1 time(s)
========= ERROR SUMMARY: 3 errors

Observed on Driver: 610.43.02 (CUDA UMD 13.3) · CUDA Toolkit (nvcc): 13.3.33 · GPU: RTX 5080 (SM120), WSL2