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)

**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

Start with the measurements and raw verdict files named in the report, then compare the observations with the shape-specialization discussion in #385. Determine whether the cost model selects schedules as well as tile sizes and whether JIT-per-shape affects bit-exact determinism; document the supported conclusion or remaining uncertainty.

Written by the indexing model from the issue text.