Problem
- Symptom: On Blackwell (B200 CC10.0 / B300 CC10.3) under CUDA MPS multi-tenancy, when the cuDNN fused SDPA decode kernel (q_len=1) runs alongside ≥2 large D2D copies every step, the GPU enters a permanent, silent hang at ≥4 tenants. No CUDA error, no exception, no Xid. GPU utilization pinned at 100%, never recovers.
- Kernel:
cudnn_generated_fort_native_sdpa_sm100_flash_fprop_f16_knob_1_1x128x128_4x1x1_cga1x1x1_kernel0_0(grid 1×NH×B, block 384, shmem 232448; shared by bf16 and fp16). - Three necessary conditions:
- cuDNN SDPA with q_len=1
- ≥2 large D2D copies per step — it is the count, not the total volume (670MB × 1 survives; 335MB × 2 wedges)
- ≥4 MPS clients (2 tenants survive; a single process never hangs)
Environment
- B200, driver 580.126.20
- cuDNN 9.19 / 9.23 / 9.24
- PyTorch 2.11 (cu130) and 2.13 (NGC 26.06)