feat: add sm120 support for DeepGEMM #324

Open

leavelet wants to merge 55 commits into

deepseek-ai:nv_devfrom

Commits on May 26, 2026

Add SM120a (RTX PRO 6000) dense FP8 GEMM support

leavelet
and
claude
committed
Add SM120a dense FP4 GEMM support and consolidate kernel to warp-spec only

leavelet
and
claude
committed
Add SM120a K-grouped contiguous FP8 GEMM and use custom MMA wrappers

leavelet
and
claude
committed
Add SM120a Einsum (batched GEMM), M-grouped BF16, and HC prenorm TF32

leavelet
and
claude
committed
Add SM120a FP8 MQA logits (dense + paged) for attention

leavelet
and
claude
committed
feat: bf16 mn gemm
leavelet
committed
Add SM120a skip_head_mid, MN-major B for M-grouped, and FP8 stride_d fix
leavelet
committed
Optimize MN-major B with ldmatrix.trans.x2, replacing 4 scalar loads
leavelet
committed
Enable BLOCK_N=128 for MN-major B via multi-atom ldmatrix.trans
leavelet
committed
Add SM120a Einsum: BF16 bmk,bnk->mn kernel, FP8 bhd,hdr->bhr and bhd,bhr->hdr
leavelet
committed
Add SM120a K-grouped FP8 TN layout and Paged MQA varlen support
leavelet
committed
Add SM120a FP4 MQA logits kernels (dense + paged)
leavelet
committed
Add SM120a FP8×FP4 mixed precision GEMM via native .b4x16_p64 TMA + ldmatrix
leavelet
committed
Enable FP8×FP4 mixed precision for K-grouped, M-grouped, and batched GEMM
leavelet
committed
Fix BF16 TMA store epilogue missing apply_index_n and add CUDA_ARCH guards to MQA kernels
leavelet
committed
SM120 FP4/FP8 perf: sub-tile TMA store epilogue + ldmatrix.x4 B load
leavelet
committed
SM120: per-N-tile ldmatrix.x4 B load — zero MOV inner loop
leavelet
committed
SM120: cooperative warp layout (4M×2N) — halve per-warp B loads
leavelet
committed
SM120: tune heuristic to prefer BN=128 for cooperative warp layout
leavelet
committed
Remove K-grouped FP4 support (SM120-only, unused in production)
leavelet
committed
SM120: MN-major support for FP8/FP4 GEMM (align with SM100 API)
leavelet
committed
SM120: enable block_kv=32 for FP4 paged MQA, guard FP8 to block_kv=64
leavelet
committed
SM120: SF-major loop with compile-time byte extraction and SwizzleContext hoist
leavelet
committed
SM120: BLOCK_K=64 with 4 pipeline stages for large-M FP8/FP4 GEMM
leavelet
committed
SM120: enable BLOCK_M=64 for grouped FP8/FP4 GEMM at small expected_m

lucifer1004
authored and
leavelet
committed
SM120: enable BLOCK_M=32 for grouped FP8/FP4 GEMM via 2×4 warp layout

lucifer1004
authored and
leavelet
committed
Revert "SM120: enable BLOCK_M=32 for grouped FP8/FP4 GEMM via 2×4 warp layout"

lucifer1004
authored and
leavelet
committed
SM120: AB-swap path for small-M FP8/FP4 BMM

authored and
leavelet
committed
SM120: extend AB-swap to plain FP8/FP4 GEMM + asymmetric recipe support

authored and
leavelet
committed
SM120: wave-packing heuristic + BLOCK_M=64 + AB-swap for dense GEMM
leavelet
committed
SM120: AB-swap with cp.async B-tile loader for M=1..16 dense GEMM
leavelet
committed
Revert "SM120: AB-swap with cp.async B-tile loader for M=1..16 dense GEMM"
leavelet
committed
SM121: reuse SM120 family cubin on NVCC/NVRTC >= 12.9
leavelet
committed
SM120: AB-swap via TMA for M=1..16 dense GEMM + per-element epilogue
leavelet
committed
SM120: prefer BM=64 for AB-swap path (2x SM utilization)
leavelet
committed
SM120: fix BMM swap stride, num_groups alignment, and BF16 heuristic bugs
leavelet
committed
SM120: split-K for small-dimension FP8 GEMM
leavelet
committed
SM120: disable split-K for AB-swap path
leavelet
committed
SM120: BF16 kernel kNWarps=2 for BM=64 + runtime_align guard
leavelet
committed
SM120: fix rebase conflicts with nv_dev, guard paged MQA next_n <= 2
leavelet
committed
SM120: remove sm120-specific test files and fix SM90 test regression
leavelet
committed

Commits on May 27, 2026

Commits on May 29, 2026

Commits on May 31, 2026

SM120: complete odd/large next_n (kPadOddN) port for paged MQA logits
leavelet
committed

Commits on Jun 4, 2026

Commits on Jun 5, 2026

SM120: fix SM90/SM100 MQA-logits SMEM regressions from shared-launcher changes
leavelet
committed

Commits on Jun 7, 2026

SM120: clean up branch for upstream (drop dev artifacts, trim comments)
leavelet
committed

Commits on Jun 12, 2026

Commits on Jun 16, 2026