Module m16_tc 

The TENSOR-CORE 5..=32-row decode tier (ATLAS_FFN_M16_TC, #927) — same child-module reason as the two above: it reads this layer’s private kernel handles, and its rule, lever and the numerics seam it opens need room. The 5..=32-row native-FP8 dense-FFN decode tier on TENSOR CORES — w8a16_gemm_m16, behind ATLAS_FFN_M16_TC (#927).

WHY. Measured on 1xH100, 2026-09-11, Qwen/Qwen3.8-27B-FP8, tip 2962cfed7: at a decode batch of 16 the step is 86.7 ms, of which the 48 SSM layers are 63.3 ms and the dense FFN inside them is 63% (~833 us/layer). The tier that serves those widths today, w8a16_gemv_batch16 (rungs 2-3 of dense_ffn_batch16_decode.rs), is bit-exact but FP32-FMA-bound at M=16, not bandwidth-bound:

shape	batch16 GEMV @ M=16	HBM3
gate/up N=17408 K=5120	0.260 ms / 342 GB/s	~3,000 GB/s
down N=5120 K=17408	0.330 ms / 270 GB/s	~3,000 GB/s

An 89 MB FP8 weight matrix should stream in ~30 us. The GEMV spends ~37 ALU ops per weight BYTE (16 scalar FFMA across the 16 rows, 16 BF16->FP32 converts, a LUT lookup, a scale multiply), which caps it near 350 GB/s no matter how fast the DRAM is. w8a16_gemm_m16 replaces those 16 FFMA with one mma.sync.m16n8k16 lane-slot — the M tile IS 16 rows, so nothing is padded, which is the whole difference from the tile GEMMs that pad M to 128 and waste 7/8 of every tile — and cuts the dequant to ~2 instructions per byte. Target: >= 1,500 GB/s at M=16, >= 1,000 GB/s at M=8.

🔴 NUMERICS — THIS ARM REASSOCIATES; THE BATCH16 ARM DOES NOT. w8a16_gemv_batch16 reduces each output in ONE FP32 accumulator walked in strict K order, which makes it bit-identical to the scalar w8a16_gemv that M=1 decode runs. An MMA reduces 16 K-products in the tensor core’s own order first, so THIS arm is not. Its contract is <= 2 BF16 ULP per element, with the 128-K block scale still folded once per block onto an FP32 outer accumulator (the two-level fold, preserved exactly). That is a seam, and it is why the lever exists and defaults OFF.

It is not a NEW seam, though: the arm the FFN reached at these widths BEFORE #927 was w8a16_gemm_n128_m128 / w8a16_gemm_pipelined, both m16n8k16 MMA kernels with exactly this reassociation. Turning the lever on returns 5..=32 to MMA numerics while keeping the ONE-weight-pass property #927 bought.

ARM ORDER in w8_gemm! (dense_ffn.rs) with the lever ON:

1. m <= 4 -> w8a16_gemv_batch4 (bit-exact)
2. m 5..=16 -> w8a16_gemm_m16 (here, MMA)
3. m 17..=32 -> w8a16_gemm_m16 x2 halves (here, MMA)
4. m 5..=16 -> w8a16_gemv_batch16 (bit-exact)
5. m 17..=32 -> w8a16_gemv_batch16 x2 halves (bit-exact)
6. W8A8 block-scaled prefill (#917/#928)
7. transposed / pipelined / base W8A16 tile GEMMs With the lever OFF (the default) rungs 2-3 vanish and the ladder is exactly what #927 shipped.

17..=32 runs the kernel TWICE on contiguous row halves, for the same reason batch16_decode.rs does: the FFN activations and outputs are contiguous [m, k] / [m, n], so a half is a plain byte offset, and two weight passes still beat one M-padded MMA tile at these widths.

── THE LEVER IS SPLIT PER PROJECTION FAMILY (round 6) ───────────────────── Round 6’s serving A/B on 1xH100 (2026-09-11, bs16, ATLAS_MS_PROFILE) with the single old lever turned the WHOLE route on at once and measured two opposite results in one number:

phase	tier	Δ step time
attention	QKV + o_proj (w8a16_gemm_m16{,_strided})	−21.7%
SSM layers	dense FFN arm (w8a16_gemm_m16)	+13.7%
net		+5.2%

One lever could only ship both or neither, so the win was unbuyable. The grammar is now three presence-based variables, all default OFF:

variable	turns on
ATLAS_ATTN_M16_TC	the QKV and o_proj tiers
ATLAS_FFN_M16_TC	the dense-FFN arm (rungs 2-3 above)
ATLAS_M16_TC	BOTH — the umbrella, i.e. round 6’s behaviour

⚠ ATLAS_FFN_M16_TC=1 MEANS SOMETHING NARROWER THAN IT DID IN ROUND 6. Before this commit it was the only lever and it reached all three tiers; round 6’s serve J and its +5.2% were measured with it. The recipe that reproduces round 6 is now ATLAS_M16_TC=1. The recipe that buys the attention win WITHOUT the FFN loss — the point of the split — is ATLAS_ATTN_M16_TC=1 alone.

── WHY THE FFN ARM LOSES WHERE THE ATTENTION TIERS WIN (HYPOTHESIS) ─────── Same kernel, same M, same weight format; the one thing that differs is N, and therefore the grid:

tier	N	CTAs at N_TILE=32
o_proj	5120	160
QKV (K, V)	1024	32
QKV (Q)	6144	192
dense FFN gate/up	17408	544
dense FFN down	5120	160

The kernel is 4 warps at 19,456 B of smem under __launch_bounds__(128, 4), so an H100 SM holds 4 CTAs and the machine holds 132 × 4 = 528. Every attention tier fits inside one partial wave and runs at full occupancy from the first instruction. gate/up at 544 is one full wave plus a 16-CTA tail: 3% of the work costs a second wave’s worth of launch, prologue and HBM-latency ramp, none of which is overlapped with anything, because by then 116 SMs are idle. That is the leading explanation for a tier that beats w8a16_gemv_batch16 3.71× in the microtest (which times ONE shape in isolation, with no tail to pay) and still loses 13.7% in the serve.

A second, non-exclusive explanation: at N_TILE=32 each 16-row A tile is read by twice as many CTAs as at 64, and gate/up’s 89 MB weight evicts A from a 50 MB L2 between passes, so the “A stays L2-resident” claim in the kernel header — which holds comfortably at N=1024 — may not hold at N=17408.

Both hypotheses predict the same fix, which is why ATLAS_FFN_M16_TC_NTILE exists: =64 selects w8a16_gemm_m16_n64, taking gate/up to 272 CTAs (inside one wave) and doubling A reuse. Default stays 32 — the tile with the receipt. NEITHER hypothesis has been measured; the A/B that settles it is ATLAS_FFN_M16_TC=1 ATLAS_FFN_M16_TC_NTILE=64 against ATLAS_FFN_M16_TC=1 on the same serve.

── THE ROUND-6 M=32 RED CELL WAS THE ORACLE, NOT THE SPLIT ──────────────── Round 6’s microtest reported gate/up M=32 at max_ulp 28, 5 of 557,056 elements over the 2-ULP budget, sign_flips 0, rel_rms 4.2e-5, while down M=32 and every M ≤ 16 cell was green. It was read as a possible row/pitch defect in the two-halves rung. It is not: a host simulation of the exact geometry (dense_ffn_m16_tc_m32_tests.rs) reproduces the signature — 5 over-budget elements, none in rows 0..15 — with NO offset arithmetic at all. Every one of them is an output that cancelled to |ref| between 5.7e-6 and 1.6e-4 against a reference RMS of 39.1, i.e. to ~1e-7..4e-6 of the matrix scale, where one FP32 accumulation rounding spans hundreds of ordinal BF16 ULP. M=32 trips it and M=16 does not because M=32 samples twice the outputs; gate/up trips it and down does not because gate/up has 3.4× the columns. The fix is in the oracle’s comparison (a mixed absolute/relative criterion), not here — see examples/native_fp8_ffn_m16_tc_microtest.rs.

Re-exports

pub use oracle::M16_TC_ACC_FLOOR_MARGIN;

pub use oracle::M16_TC_MAX_ULP;

pub use oracle::bf16_ord;

pub use oracle::m16_tc_acc_floor;

pub use oracle::within_m16_tc_budget;

Modules

oracle

The tier’s NUMERICS CONTRACT — the one comparison the GPU oracle (examples/native_fp8_ffn_m16_tc_microtest.rs) and the host simulation both evaluate, so a receipt and a unit test cannot be grading different things. The numerics contract for the tensor-core decode tiers (w8a16_gemm_m16, #927; dense_gemm_m16_bf16, #927/#928) — ONE comparison, evaluated by the GPU oracles (examples/native_fp8_ffn_m16_tc_microtest.rs, examples/native_bf16_lm_head_m16_microtest.rs) and by the host simulations (dense_ffn_m16_tc_m32_tests.rs, ops/dense_gemm_m16_bf16_tests.rs, ops/dense_gemm_m16_bf16_floor_tests.rs), so a receipt and a unit test cannot drift into grading different things.

Structs

M16TcLevers

Which projection families the tensor-core tier serves, and at what CTA width. SSOT for the whole lever grammar; every call site resolves it ONCE at construction into a field.

Functions

m16_tc_levers

The resolved levers for this process.