[GH-ISSUE #402] Measured data point on fixed-tile vs shape-adaptive GEMM: where a bit-exact fixed-tile kernel stops (and why we couldn't fix it) #184

New issue

Open

opened 2026-10-26 21:14:24 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 21:14:24 +00:00

Owner

Originally created by @dancinlife on GitHub (Aug 7, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/402

Title: Measured data point on fixed-tile vs shape-adaptive GEMM: where a bit-exact fixed-tile kernel stops (and why we couldn't fix it)

Hi — not a bug report. We ran a multi-month campaign building a cuBLAS-free GEMM inside our own compiler, hit a wall that DeepGEMM's JIT/shape-specialization design appears built to avoid, and exhausted the lever family trying to close it. The project is now shelved, so I'd rather the measurements be useful to someone than sit in a dead repo. Close this if it's off-topic for your tracker.

What we measured

A from-scratch TF32 GEMM, no cuBLAS call, device-resident (callable inline from a persistent megakernel, which a host API can't be):

shape	own TFLOP/s	cuBLAS-TF32	ratio	rel_rms
D=2048 (H100)	~315	~342	1.08× (~93% of roofline)	0.000e+00
D=4096 (H100)	~284	~427	~1.50×	0.000e+00
D=768 (RTX 5070, sm_120)	—	—	0.95× (own edges cuBLAS)	~1.3e-5

rel_rms = 0.000e+00 is an equality, not a tolerance — bit-exact at every config.

The wall, and why we think it's structural

Our kernel is one fixed 128×128 tile at every D. At D=4096 cuBLAS scales up +24.6% while ours scales down −9.9% (2× K-loop drain). We statically excluded register spill, occupancy drop, and a ptxas ceiling as causes, then burned two GPU builds on the remaining lever:

CTA swizzle in isolation (non-persistent): regresses −1.6% — 280.5 vs 285.1 TFLOP/s
A 2-CTA/SM-preserving 128×256 tile (90 regs, same 2 CTA/SM as the baseline): regresses −7.1% — 263.3 vs 283.5. The sequential-halves schedule serializes the wgmma pipeline (ptxas C7515) and doubles the K-drain.

Our conclusion was that no bit-exact 256-N schedule on sm_90a is both 2 CTA/SM and non-serialized — i.e. for us the large-D gap was bit-exactness-bound, not a tuning backlog.

The generalizable version, which is the part I think is worth sharing: what a vendor BLAS actually sells is not one excellent kernel — it's shape adaptivity. A single good tile reaches parity in a band and loses outside it. We budgeted for a kernel and needed a dispatcher.

One more result that surprised us

The parity is dtype-scoped. Porting the same kernel to 16-bit operands did not transfer:

dtype	own @D=4096	cuBLAS roofline	ratio
TF32	(315 @D=2048)	342	1.08×
FP16	71.6	827.2	11.55× off
BF16	71.1	816.1	11.48× off

Our kernel ran at the same absolute throughput in FP16 as TF32 (~71–76 TFLOP/s — same occupancy-bound design), but the cuBLAS FP16 roofline doubled. The precision change moved the target away without lifting us. Worth flagging because "just go lower precision" is common advice and it can widen the gap rather than close it.

The actual question

DeepGEMM's JIT + runtime shape specialization looks like the direct answer to what stopped us. Two things I'd genuinely like to know, if anyone has them handy:

Does tile selection recover large-D scaling, or does the tile schedule have to change with shape? Our −7.1% result says the schedule was the binding term, not the tile size — a larger tile that kept 2 CTA/SM still lost, because the schedule serialized the pipeline. We only tested two points, so this is a weak inference on our side. I noticed from #385 that you calibrate cost-model constants per arch; I'd be curious whether that model selects among schedules or only among sizes.
Does JIT-per-shape cost you determinism, and did you care? Our binding constraint was bit-exactness — rel_rms = 0 as a gate, not a tolerance — which is what removed split-K and most swizzles from our option set before we ever benchmarked them. If DeepGEMM doesn't carry that constraint, then a much wider option space is a cleaner explanation of our wall than "sm_90a can't do it," and our conclusion above is overstated. I'd rather find that out than leave a wrong claim in a dead repo.

Happy to point at the raw verdict files if any of this is useful. Sources for the numbers above: F-GPU-ROUTEA-KEEPBAND-MEASURE, F-OP45GPU-OCCUPANCY-SWEEP, F-OP52-TF32-GAP-CLOSE, F-OP55-NEWTILE-D4096, F-FUSION-SM90-WGMMA-W14-FP16, F-OP54-SUMMER-OWNGEMM-TF32.

Full post-mortem (MIT, includes three headline numbers we got wrong and retracted): https://github.com/dancinlab/hexa-lang/issues/5035

Originally created by @dancinlife on GitHub (Aug 7, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/402 **Title:** Measured data point on fixed-tile vs shape-adaptive GEMM: where a bit-exact fixed-tile kernel stops (and why we couldn't fix it) Hi — not a bug report. We ran a multi-month campaign building a cuBLAS-free GEMM inside our own compiler, hit a wall that DeepGEMM's JIT/shape-specialization design appears built to avoid, and exhausted the lever family trying to close it. The project is now shelved, so I'd rather the measurements be useful to someone than sit in a dead repo. Close this if it's off-topic for your tracker. ## What we measured A from-scratch TF32 GEMM, no cuBLAS call, device-resident (callable inline from a persistent megakernel, which a host API can't be): | shape | own TFLOP/s | cuBLAS-TF32 | ratio | rel_rms | |---|---|---|---|---| | D=2048 (H100) | ~315 | ~342 | **1.08×** (~93% of roofline) | **0.000e+00** | | D=4096 (H100) | ~284 | ~427 | ~1.50× | 0.000e+00 | | D=768 (RTX 5070, sm_120) | — | — | **0.95×** (own edges cuBLAS) | ~1.3e-5 | `rel_rms = 0.000e+00` is an equality, not a tolerance — bit-exact at every config. ## The wall, and why we think it's structural Our kernel is **one fixed 128×128 tile at every D**. At D=4096 cuBLAS scales *up* +24.6% while ours scales *down* −9.9% (2× K-loop drain). We statically excluded register spill, occupancy drop, and a ptxas ceiling as causes, then burned two GPU builds on the remaining lever: - **CTA swizzle in isolation** (non-persistent): **regresses** −1.6% — 280.5 vs 285.1 TFLOP/s - **A 2-CTA/SM-preserving 128×256 tile** (90 regs, same 2 CTA/SM as the baseline): **regresses −7.1%** — 263.3 vs 283.5. The sequential-halves schedule serializes the wgmma pipeline (ptxas C7515) and doubles the K-drain. Our conclusion was that no bit-exact 256-N schedule on sm_90a is *both* 2 CTA/SM *and* non-serialized — i.e. for us the large-D gap was **bit-exactness-bound**, not a tuning backlog. The generalizable version, which is the part I think is worth sharing: **what a vendor BLAS actually sells is not one excellent kernel — it's shape adaptivity.** A single good tile reaches parity in a band and loses outside it. We budgeted for a kernel and needed a dispatcher. ## One more result that surprised us The parity is **dtype-scoped**. Porting the same kernel to 16-bit operands did not transfer: | dtype | own @D=4096 | cuBLAS roofline | ratio | |---|---|---|---| | TF32 | (315 @D=2048) | 342 | 1.08× | | FP16 | 71.6 | **827.2** | **11.55× off** | | BF16 | 71.1 | 816.1 | 11.48× off | Our kernel ran at the *same absolute throughput* in FP16 as TF32 (~71–76 TFLOP/s — same occupancy-bound design), but the cuBLAS FP16 roofline **doubled**. The precision change moved the target away without lifting us. Worth flagging because "just go lower precision" is common advice and it can widen the gap rather than close it. ## The actual question DeepGEMM's JIT + runtime shape specialization looks like the direct answer to what stopped us. Two things I'd genuinely like to know, if anyone has them handy: 1. **Does tile *selection* recover large-D scaling, or does the tile *schedule* have to change with shape?** Our −7.1% result says the schedule was the binding term, not the tile size — a larger tile that kept 2 CTA/SM still lost, because the schedule serialized the pipeline. We only tested two points, so this is a weak inference on our side. I noticed from #385 that you calibrate cost-model constants per arch; I'd be curious whether that model selects among schedules or only among sizes. 2. **Does JIT-per-shape cost you determinism, and did you care?** Our binding constraint was bit-exactness — `rel_rms = 0` as a *gate*, not a tolerance — which is what removed split-K and most swizzles from our option set before we ever benchmarked them. If DeepGEMM doesn't carry that constraint, then a much wider option space is a cleaner explanation of our wall than "sm_90a can't do it," and our conclusion above is overstated. I'd rather find that out than leave a wrong claim in a dead repo. Happy to point at the raw verdict files if any of this is useful. Sources for the numbers above: `F-GPU-ROUTEA-KEEPBAND-MEASURE`, `F-OP45GPU-OCCUPANCY-SWEEP`, `F-OP52-TF32-GAP-CLOSE`, `F-OP55-NEWTILE-D4096`, `F-FUSION-SM90-WGMMA-W14-FP16`, `F-OP54-SUMMER-OWNGEMM-TF32`. Full post-mortem (MIT, includes three headline numbers we got wrong and retracted): https://github.com/dancinlab/hexa-lang/issues/5035

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

deepseek-ai/DeepGEMM#184

No description provided.

Rows
Columns