fp-cuda: Hopper wgmma.b1 device-resident F₂ GEMM backend#273

Open

JoeyBF wants to merge 16 commits into

SpectralSequences:masterfrom

JoeyBF:fp_cuda_hopper

Commits on Jul 17, 2026

fp-cuda: wgmma.b1 F2 GEMM kernel — host-arranged A/B operands

JoeyBF
and
claude
committed
fp-cuda: K-pipeline — TMA loads, double-buffering, warpgroup specialization

JoeyBF
and
claude
committed
Phase 3: 128B-swizzle operands + pipelined wgmmas

JoeyBF
and
claude
committed
Phase 4: widen to m64n256k256 (largest binary MMA shape)

JoeyBF
and
claude
committed
Phase 5: per-warpgroup register reallocation (setmaxnreg)

JoeyBF
and
claude
committed
Phase 6: deepen K-pipeline to STAGES=3

JoeyBF
and
claude
committed
Phase 7: write the output tile via TMA bulk store (S2G)

JoeyBF
and
claude
committed
Migrate host layer from cuda-core to cudarc (stable Rust)

JoeyBF
and
claude
committed
Phase 8: persistent kernel + grouped tile rasterization

JoeyBF
and
claude
committed
Phase 9: thread-block clusters + TMA multicast of B

JoeyBF
and
claude
committed
Phase 10: register-block the output tile to cut operand-refill bandwidth

JoeyBF
and
claude
committed
Phase 11: overlap the output-tile TMA store with the next tile's compute

JoeyBF
and
claude
committed
Phase 12: hoist wgmma.fence out of the consumer K-loop

JoeyBF
and
claude
committed
fp-cuda: occupancy-aware persistent launch; document 2-CTA/SM dead end

JoeyBF
and
claude
committed
fp-cuda: wire into fp behind a `gpu` feature; CI builds without nvcc

JoeyBF
and
claude
committed

Commits on Jul 18, 2026

fp-cuda: address review feedback

JoeyBF
and
claude
committed