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

Merged

JoeyBF merged 28 commits into

SpectralSequences:masterfrom

JoeyBF:fp_cuda_hopper

Aug 15, 2026

Commits on Jul 17, 2026

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

JoeyBF
and
claude
committed
c25c43f

Commits on Jul 18, 2026

fp-cuda: address review feedback

JoeyBF
and
claude
committed
70f2e02

Commits on Jul 29, 2026

fp-cuda: per-thread CUDA streams — lock-free concurrent GEMM dispatch

JoeyBF
and
claude
committed
7c9fece

Commits on Jul 30, 2026

fix(fp-cuda): scope the per-thread stream cache to its GpuContext

JoeyBF
and
claude
committed
6b4e62f
fp-cuda: green CI — GPU-availability test gate, fix private doc link, rustfmt

JoeyBF
and
claude
committed
10ca2ae

Commits on Aug 15, 2026

fp-cuda: trim comments before merge

JoeyBF
and
claude
committed
9f59320
fp-cuda: single-source the tuning knobs in params.h

JoeyBF
and
claude
committed
992a0f6
fp-cuda: return anyhow::Result instead of boxed errors

JoeyBF
and
claude
committed
eceb5b7
fp-cuda: remove the thread-block cluster layer from the GEMM

JoeyBF
and
claude
committed
5916e86
fp-cuda: re-clean comments after the cluster removal

JoeyBF
and
claude
committed
c81d2d8
fp-cuda: address review comments on the stub PTX and doc structure

JoeyBF
and
claude
committed
473021a
fp-cuda: fence each writer before the epilogue barrier

JoeyBF
and
claude
committed
516893a
fp-cuda: cover a fully non-aligned shape in the demo

JoeyBF
and
claude
committed
4caad21
fp-cuda: label bench_shapes throughput as kernel-only

JoeyBF
and
claude
committed
d2839e8