[BUG] [CUDA] Illegal __global__ write in copy_s under CUDA graphs on consumer Blackwell (sm_120) → corrupted output #4700

Summary

On an RTX 5070 Ti Laptop (Blackwell, sm_120, cc 12.0), MLX's CUDA backend
intermittently produces corrupted LLM output (fluent but off-task, runaway to
max tokens). compute-sanitizer memcheck pinpoints an out-of-bounds /
write-before-allocation in the scalar-broadcast copy kernel copy_s, launched
from a CUDA graph. The identical build is correct on RTX 3090 Ti (Ampere,
sm_86).

Environment

• GPU: NVIDIA GeForce RTX 5070 Ti Laptop (sm_120, cc 12.0); driver 595.79; CUDA 13.x
• MLX: <fill in your version, e.g. 0.32.x> (pip install mlx[cuda] or built with CMAKE_CUDA_ARCHITECTURES=120)
• OS: Windows 11

compute-sanitizer memcheck (repeats continuously)

Invalid __global__ write of size 16 bytes
  at void mlx::core::cu::copy_s<float, float, unsigned int, (int)4>(const T1*, T2*, T3)+0x430
  by thread (0,0,0) in block (0,0,0)
  Access to 0x... is potentially made before memory is allocated
  and is inside the nearest allocation of size 128 bytes
  Saved host backtrace up to driver entry point at kernel launch time
    Host Frame: cuGraphLaunch in nvcuda64.dll
    Host Frame: mlx::core::cu::get_command_encoder in mlx.dll
    Host Frame: mlx::core::scheduler::Scheduler::notify_task_completion in mlx.dll
    Host Frame: mlx::core::eval in mlx.dll

Eventually escalates to a sticky cudaErrorUnknown (999) on
cudaEventSynchronize / cudaPeekAtLastError.

Analysis

• copy_s (mlx/backend/cuda/copy/copy_contiguous.cu) is the scalar-broadcast
  copy; N_READS = 16/sizeof(float) = 4 → a 16-byte store_vector<4>.
• CommandEncoder::insert_graph_dependencies (mlx/backend/cuda/device.cpp)
  derives graph edges only from input-buffer data dependencies via node_map_.
  A newly cudaMallocAsync'd output has no prior writer in node_map_, so the
  copy node is added with no in-graph dependency expressing "allocation
  complete" — ordering relies solely on same-stream semantics + the pool's
  reuse events.
• get_graph_limits gives case 1200 // Consumer Blackwell a wide budget
  (ops=100), so many data-independent nodes run concurrently in one graph,
  a much wider timing window than Ampere — exposing the latent alloc→copy race
  on sm_120.

It is a timing/ordering race (not an index overrun)

All of the following make it disappear and give bit-correct output:

• CUDA_LAUNCH_BLOCKING=1
• attaching Nsight Systems
• MLX_USE_CUDA_GRAPHS=0

Repro

Run the copy-heavy inference under
compute-sanitizer --tool memcheck --kernel-name kernel_substring=copy_s on an
sm_120 device; observe the copy_s "before memory is allocated" write.

Related

• Run CUDA graph kernels one at a time on DGX Spark #4697 (graph serialization, perf, sm_121)
• [CUDA] Deterministic illegal memory access on sm_120 (RTX 5090) with glm_moe_dsa arch (GLM-5.2) — sm_121/sm_110 unaffected #3874 (sm_120 illegal access, attention/MoE path, unresolved)
• [CUDA] Buffer cache and CUDA graph #2358 (buffer cache × CUDA graph)

A fix PR follows.