sm_120 CUDA Kernel Expert - imp

PTX templates: references/ptx-patterns.md. Dead ends, version gotchas, load-bearing fixes: references/known-issues.md (read BEFORE proposing a lever on a mature path).

Architecture

Spec	Value
SMs / CUDA cores / TC	170 / 21,760 / 680 (4 per SM, 5th gen)
L1+SMEM per SM	128 KB; opt-in shared ~99 KB per block (cudaDeviceProp::sharedMemPerBlockOptin), not H100's 228 KB
L2	96 MB (every single Qwen3.8 weight fits: isolated benches read L2)
VRAM	32 GB GDDR7, 1792 GB/s datasheet; 1628 GB/s measured resident, ~237 spilled
MMA shapes	FP16 m16n8k16, TF32 m16n8k8, FP8 m16n8k32, FP4 block-scaled m16n8k64
Measured TC rates	FP4 mma.sync 2019 TOPS (~1/2 datasheet); f32-accumulate 1/4 rate; GeForce TF32 = 1/2 the FP16-fp32acc rate (253 TFLOPS)
Registers	255 per thread is the DEVICE attribute; a kernel "at 255" must be read from ptxas (make kernel-resources), the dense FA2 instance allocates 144

No tcgen05 / TMEM / wgmma / 2-CTA cluster MMA (SM100 only). Peak path = register mma.sync with block scaling, FA2-style. Toolkit bumps add nothing: CUDA 13.3 flipped 0 of 247 probed instructions.

The three laws

1. Batch=1 decode is at its ceiling; batched decode (M<=32) is its own regime. Order at batch=1: fast per-launch GEMM (native mxf4nvf4 smallm v2 for M<=32, gemm.nvfp4_smallm, #1766; CUTLASS at large M; never the gemm_nvfp4 dequant->cuBLAS fallback), then CUDA-graph capture (CudaGraphConditionalRunner, src/runtime/cuda_graph.h), then bytes. Qwen3.8 M=1 spec-off: 95.8% busy, no single lever >=1% left (#1789); the decode graph is strictly serial there (union == sum). A per-token host round-trip drops a path out of the conditional loop (-27..-45%). At 32 streams grid-shape and launch levers PAID: row-block RMSNorm +6.8% (#1769), shared act-quantize +4.6% (#1771), producer quantize +2.6% (#1773), BF16 GDN state +12.5% (#1776/#1778), smallm sibling pairs +1.7% (#1788), row-batched sampling +2.2% (#1790), PDL device half +0.5-1.3% (#1833); residual-accumulate LOST -0.9% (#1793: the adds overlap with the next layer's GEMMs). Measure the regime with a 32-stream two-image A/B (benchmark-cuda).
2. Occupancy is for reaching the roofline, not passing it. <=48 regs for 100% occupancy on GEMV paths. __launch_bounds__(T, N) forces N blocks by SPILLING: (128,12) on the MR GEMV went 14.5 -> 26.8 us at MR=4; (256,1) on the shipped dense FA2 instance moved ptxas 137 -> 180 regs. What shipped: (256,2) on a SEPARATE wrapper kernel (TWOSLOT 35 KB, 137 -> 128 regs, 24 B spill, FA2 -10% at pp4096, #1843) with the production instance byte-identical in SASS. Read --ptxas-options=-v per template instance; the CI kernels gate (tools/kernel_resource_baseline.txt, REG >= 240 or any local frame) ratchets it. Two CTAs/SM on a tensor-pipe-bound kernel is flat (chunkpar K2: 333 vs 331 ms) or worse under latency binding (+8%: barriers double, half the work per MMA phase).
3. Quantization type determines kernel strategy. Q8_0: bandwidth-bound, row-parallel + smem activations. Q6_K: K-parallel + warp division. NVFP4 large M: StorageTier::CUTLASS_NVFP4 (SfAtom layout); M<=32: smallm v2 on the PLAIN packed bytes (h.source_data/h.source_scales). Plain StorageTier::NVFP4 at large M = slow gemm_nvfp4.

Numerics rules for tensor-core rewrites

• A recurrent state path needs error compensation on EVERY link: plain tf32 on the GDN state GEMMs = state diff 3.4e-4, PPL +0.13%; on P@W it cancels (Qeff = D q - P W). 3xTF32 (a_hi + a_lo) costs +25 us; 3xFP16 m16n8k16 (a_hi + a_lo in fp16) runs 2x the TF32 rate and shipped on all state-feeding GEMMs (#1852; K1 -22%, K2 -18%). Output-only terms (Y_A, y) stay plain tf32/fp16.
• Judge: unit-test state diff vs the fused kernel (~1e-6 = fp32-equivalent; 1e-4 fails), deterministic PPL on Qwen3.8-27B-NVFP4-vllm (fused 4.6283), tools/analysis/layer_ab_diff.py added divergence of the changed blocks ~0. Qwen3.6-35B PPL moves +-0.2..0.5% between fp32-equivalent kernels: no verdict there.
• Online-softmax tile levers cost precision per extra rescale: HD=256 FA2 at Bkv=32 = kernel -11%, PPL +0.53%, opt-in only (attention.fa2_hd256_bkv=32, #1840). Scale inside the exp FMA is free (#1844, PPL 4.6283 both arms); exp2f forms and scale-in-Q are +0.35% PPL for the same speed.
• A kernel variant chosen by row count must produce identical per-row bits, or be gated to decode rows (cur_decode_rows_): otherwise a prompt row changes with the prefill chunk it lands in (#2152, fixed in #2171). Cases: src/compute/layernorm.cu row-block vs warp kernel, fused QK-norm+RoPE capped at n<=64 (src/exec/executor_attention.cu), mmq_q8_imma split-K at M<=32 (src/compute/mmq_q8_imma.cu).
• FP16 intermediates on out / row_scale overflow to inf on small-absmax rows (FP8 SSM_IN, 2026-09-01); keep the rescale in FP32 or fuse it. Uniform-random unit-test rows cannot reach this.

GDN chunk-parallel scan (#1847-#1852, gdn.chunkpar_scan default on)

src/compute/gdn_scan_chunkpar.cu (K1 factor kernel + launcher), gdn_scan_chunkpar_pass.cu (K2 state pass), .cuh shared. Grid (chunks x heads) for factors, (heads, 2) for the state pass; 42-126 MiB workspace (null degrades to the fused route); single-sequence prefill n >= 128, HD=SS=128. Qwen3.6-35B pp4096 e2e 12.9k -> 31.0k tok/s vs the 32-CTA fused scan (42% of the hybrid pp512 wall before).

Lever	Measured
Solve histories in GLOBAL memory	K1 379 us; in smem (aliased over dead K/Q staging) + 256 threads: 196 us
K2 scalar FMA from smem	628 us (= the sequential scan); float4 staging + column split: 242 us
Accumulator splits / 2-t unroll in K2	REGRESS 242 -> 295/309 us (registers cost the 2-CTA residency); in K1 four partial accumulators pay
Blockwise triangular solve (16-row diagonal blocks in registers, off-diagonal 3xTF32)	128 -> 8 barriers per chunk, K1 75 -> 49.5 us/CTA (#1850)
K2 8 warps + register-pipelined staging	K2 -30% (#1851)
Strip per head count under an L2 cap (gdn.chunkpar_strip=0 auto: factor set <= 2/3 L2)	strip 14-16 wins K1 -21% but K2 +11..13% (126 MB factor set > 96 MB L2 reads DRAM); net 0
XOR swizzle on stride-128 float4 tiles chunk ^ ((row & 7) << 1)	K1 bank conflicts 64% -> 25%; k16 fragment rows (2tg/2tg+1) need swz128 staging + SH = COLS + 4
K2 half staging (RS=32, 2 CTAs/SM)	flat under TF32 binding, +8% under FP16 binding: refuted twice

Traps: gdn_scan_chunkpar.cu sat at exactly 600 code LOC (kernel hard threshold) after #1850, hence the split; nvcc 13.3 segfaults on a generic lambda with a std::true_type tag inside a kernel (use a free template function); dumps for layer_ab_diff.py are 3.6 GB per arm at 300 tokens.

Paths that must stay active (check before optimizing)

Path	Evidence at init / config
CUTLASS NVFP4 weight cache	cutlass_nvfp4 weight cache log line; without it decode runs dequant->cuBLAS
CUDA-graph decode loop, graph prewarm	runtime.cuda_graphs, runtime.graph_prewarm (32/32 captures in 2.3 s)
PDL both halves	runtime.no_pdl=false; registered kernels call griddepcontrol.wait before the first global access (src/compute/pdl_device.cuh); registered = waits, blanket registrations without a wait RACED (GreedyDeterminism)
FA2 prefill family	attention.fmha_fa2, fa2_hd256, fa2_fp16qk, fa2_dense_2cta (all on); fa2_hd256_bkv 64
Stream-K on the CUTLASS prefill GEMM	gemm.nvfp4_cutlass_streamk=1: N > 2048, >= 1 wave, last wave <= half full; workspace sized as the max over every 128-tile grid (~22 MB), a 0 B workspace refuses EVERY launch (22k -> 4.2k tok/s); the SK-typed kernel in data-parallel mode is slower (109 vs 100 us) so shapes outside the rule keep the plain kernel
Batched-decode sidecars	gemm.nvfp4_smallm + _impl=2 + _pair, gdn.state_bf16 (+0.21% PPL by design), gemm.fp8_ssm_proj, gemm.nvfp4_lm_head*, gemm.fp8_attn_proj
Chunk-parallel GDN prefill	gdn.chunkpar_scan, gdn.chunkpar_strip
INT8-IMMA prefill for GGUF MoE	dense Q4_K IMMA loses to cuBLAS and is off by design

FP8 prefill stays DISABLED on sm_120 (cuBLAS NOT_SUPPORTED at non-aligned M, src/runtime/engine_init_resolver.cpp); FP8 GDN-projection prefill re-measured 2026-09-01: class is 12-13% of the steady-state sum, ceiling +5%, inside the cuBLAS band, REFUTED (record docs/plans/2026-08-31-fp8-ssm-prefill.md). Prefill levers here: FA2 family, ragged prefill batching (runtime.prefill_batch, #1780), chunk-parallel scan.

Closed classes (do not reopen without a new mechanism)

Class	Verdict
M=1 NVFP4 decode GEMV	structural ceiling; occupancy raise, KPAR->MR, v2 pipeline at M=1 (#1789) all refuted
MoE grouped GEMM (NVFP4 prefill)	~60% of weight floor; CUTLASS tile sweep, 32-row v2 grouped (-12.5% isolated, +3.5..+15% in situ), multi-tile CTA mt32/64/128 all refuted on real routing; the sm120 block-scaled builder rejects M=64 tiles (SF atom 128 rows)
FA2 occupancy at hd=128	16 warps/SM shipped (#1843); split-D warp pairing refuted; remaining path = in-CTA pipelining, and the kernel spent 5.5 scalar ALU per MMA in the softmax loop before #1844
NVFP4 decode attention traffic sharing	GQA tile -9% (#1785), smem double buffer -3%; the lever was LOAD WIDTH (#1817: 20 LDG.E.U8 per iteration -> word loads, +15.7%)
FP4-precision attention	closed 2026-07-04 (three refutations)
Batch=1 launch-class fusion	0% e2e; split-K cap -21..-35%; raising split-K -2.9..-4.7%

Shared memory budget

head_dim	Bq	SMEM
64	128	~89 KB
128	64 / 128 (TWOSLOT 2-CTA)	~81 KB / 35 KB
256	32 (Bkv=64) / Bkv=32	~68 KB / 35 KB

Padding two access patterns at once may not fit (chunkpar histories at stride 132 next to T/P padding did not fit 99 KB); swizzle instead.

Compile flags

--generate-code=arch=compute_120a,code=sm_120a (IMP_SM120_FLAGS), --expt-relaxed-constexpr --extended-lambda, release -O3 --use_fast_math. compute_120a unlocks mma.sync.aligned.kind::mxf4nvf4.block_scale, TMA-WS grouped GEMM, cluster launch, extended mbarrier, FP4 saturation. Never sm_120/sm_120f/generic compute_120 (no block-scale, no FP8 MMA). Guard with #if __CUDA_ARCH__ >= 1200.

Common mistakes

Mistake	Fix
Targeting sm_120 / sm_120f	compute_120a/sm_120a
Decode routed through gemm_nvfp4	promote to StorageTier::CUTLASS_NVFP4; check the init line
No graphs-ON re-bench after a kernel win	the tok/s win is graph-mediated
H100 SMEM assumptions	99 KB opt-in
__launch_bounds__ on regular GEMV/attention paths	-4.5..-20%; min-blocks spills to DRAM; only wrapper-kernel bounds with a SASS-identical production instance
Byte-pointer inner loops (const uint8_t* walked element-wise)	ptxas cannot merge; count LDG.* forms in cuobjdump -sass first
reinterpret_cast on Q8_0 blocks (34 B, unaligned)	memcpy()
Pointer advance without sizeof(T)	known-issues (FP8 FMHA S_tile)
PDL registration without griddepcontrol.wait	race on the producer's output; pdl::launch has no default args, register at the launch site (pdl::enable_kernel)
Isolated bench without an L2-defeating ring	reads L2; rotate >= 4 x 100 MB slabs, warm per shape
"perf-neutral" without a SASS diff	cuobjdump -sass; byte-identical = proof
tf32 on a recurrent state path checked only by MoE PPL	unit-test state diff + Qwen3.8 deterministic PPL
One CTA per row walking a context-length loop (qsa_select before #2074: 128.4 us/layer at 13.9k)	score across the grid (one warp per 256 B key, two waves), then a 1024-thread select with ballot ranks: 9.1 us
Split decision counted on the split kernel's CTAs, not the no-split kernel's	paged_attention_decode: GQA ratio > 8 without multitok falls back to paged_attention_gqa_kernel = batch x n_kv CTAs (16 rows x 24/2 heads: 32 CTAs, 1041 us vs split-K 133 us, #2074); counting n_kv at batch 1 too moved splits 15 -> 64 and cost +9 % (81.6 -> 89.0 us)

Where to look

• src/compute/: attention (attention_fmha_*, attention_paged_nvfp4.cu), GDN (gdn_scan_chunkpar*.cu, gdn_scan_tc.cu), pdl_device.cuh, sparse attention kernels.
• src/quant/: dequant, nvfp4_gemm_smallm_v2.cu, nvfp4_pack.cuh.
• src/exec/: executor_gemm_dispatch.cu, executor_gemm_smallm.cpp, executor_sampling.cu, executor_attention_decode.cu, producer-fused norm/swiglu+quantize, sparse_attn_geometry.h.
• src/runtime/: cuda_graph.h, pdl.h, engine_init_resolver.cpp.
• Gates: make kernel-resources (tools/kernel_resources.py), tools/check_launch_guards.py, tools/check_filesize.py (kernel .cu hard 600 code LOC; split kernel / launcher / instantiations).
• Docs: docs/internals/SM120.md, docs/internals/KERNELS.md, docs/internals/PROFILING.md, docs/PERF.md, docs/roadmap.md (verdict ledgers), tools/roofline/history/BASELINE.