cuBLAS 60% MatMul Bug on RTX 5090

🔑 Enhanced Key Takeaways

• •The issue stems from a heuristic failure in the cuBLAS dispatch engine that incorrectly prioritizes legacy register-file-based tiling over the Blackwell architecture's Tensor Memory Accelerator (TMA) pipeline for specific batched FP32 workloads.
• •NVIDIA's internal engineering response suggests the bug is limited to the 'consumer' firmware branch of the RTX 5090, as the same cuBLAS library version correctly dispatches optimized TMA kernels on workstation-class Blackwell silicon.
• •Community-developed workarounds involve forcing a custom CUDA graph that bypasses the cuBLAS heuristic layer, effectively 'pinning' the execution to the high-throughput TMA path that the driver currently ignores.

🛠️ Technical Deep Dive

• •The bottleneck is identified as a 'stall-on-issue' state in the SM90 scheduler, where the warp scheduler fails to hide latency because the cuBLAS kernel does not utilize the asynchronous copy (cp.async) instructions required for the Blackwell TMA unit.
• •The custom kernel implementation utilizes a double-buffering strategy in Shared Memory (SRAM), allowing the TMA to pre-fetch the next tile of the matrix while the current tile is being processed by the Tensor Cores.
• •NCU (NVIDIA Nsight Compute) profiling reveals that the cuBLAS-dispatched kernel exhibits a high 'L1/TEX Cache Hit Rate' but suffers from 'Instruction Issue Stall: Barrier' due to the lack of overlap between memory movement and compute cycles.

🔮 Future ImplicationsAI analysis grounded in cited sources

⏳ Timeline