Prefer to read without ads? Become a member — from $10/month — and support the work. Already a member? Log in to read ad-free on this device.

17.7 The Gap to cuBLAS

Table 17-1 summarizes the journey. The best kernel here reaches about 82% of cuBLAS’s TF32 throughput – respectable for a few hundred lines of readable code, and revealing in what remains.

Table 17-1. SGEMM optimization journey (RTX 3060, M = 4096, N = 2048, K = 4096).

Stage	Precision	GFLOP/s	vs. cuBLAS TF32	Bottleneck
Naive	FP32	886	7%	DRAM bandwidth
Cache-blocked (8×8, no shared memory)	FP32	3762	31%	L1/L2 bandwidth
Register-blocked (128×64 tile)	FP32	6172	50%	FP32 compute
Tensor Cores (WMMA)	TF32	8718	71%	global load
+ `cp.async` double-buffer	TF32	9987	82%	global load (overlapped)
cuBLAS `Sgemm` (default)	FP32	8318	68%	–
cuBLAS TF32	TF32	12229	100%	–

The remaining 18% is not mysterious; it is a catalog of techniques applied by a production library that this chapter deliberately does not:

Raw mma.sync and ldmatrix PTX instead of the WMMA C++ API, avoiding the per-fragment load-and-convert overhead the API imposes.
Deeper pipelines – three and four cp.async stages in flight rather than the two used here – to hide more of the load latency.
Shared-memory swizzling so that fragment loads never collide on a bank, the bank-conflict problem solved by permuting rather than padding.
Split-K and stream-K scheduling that partitions the K dimension across thread blocks to keep every SM busy on awkward matrix shapes.

They also are architecture-specific, and the frontier keeps moving. Hopper replaces cp.async with the Tensor Memory Accelerator (TMA) and the register-free wgmma instruction, and adds thread-block clusters that share memory across SMs; Blackwell adds fifth-generation Tensor Cores and microscaling formats (FP8, FP6, FP4) that push the precision-versus-throughput trade even further. Revisiting this kernel on that hardware – and closing more of the gap to cuBLAS with TMA-fed, deeply pipelined wgmma – is a roadmap item that we’ll visit in the near future.