perf(ops): per-schedule activation cache policy in w8_rowsplit_gemm_mma, with ca on the BM=16 tiles

[MichaelDementii]

w8_rowsplit_gemm_mma_kernel branches on the template parameter Full, so a ragged token count
does not take a slow branch - it enters a separate kernel body, and that body staged every copy
it issues with a bare cp_async_zfill<16>, whose default policy is Cache::ca, while the aligned
body staged the same bytes with cp_async<16, Cache::cg>. This revision makes the activation
policy a property of the schedule instead of a constant of the kernel: Cache::cg by default,
Cache::ca on the two schedules with a BM = 16 tile. The weight side - codes and scales - goes to
Cache::cg unconditionally. Three files, +26 / -9, bit-identical output.

Claim level: operator, and end-to-end prefill on ragged prompt lengths. It returns 39 to 64%
of what a ragged token count costs on the seven W8 operator configurations that reach this kernel,
+0.95 to +1.09 pp of end-to-end prefill throughput on four ragged lengths against -0.03 pp on
the aligned control, and -3.33 to -3.57% on the DFlash2 attention input projection at
T = 49..63. It also retracts an explanation I published with the neighbouring split of the same
Full flag.

What changed in this revision, and why

The review bot filed a P1 against the previous head of this branch, and it was right. Its
finding: putting the predicated body on cg also moves w8_feature, a route that hard-wires
Full = false even on full tiles (src/ops/linear/w8/w8_feature.cu:45), and the branch's sweep only
covered T near 8192. I built the arms and measured it. The bot was right and it understated the
cost: the previous head did not change the profile there, it made the operation 18.9% slower.
And a second hard-wired route it did not name, the DFlash2 attention input projection
(src/ops/attn_input_proj/w8/w8_dflash2_attn_input.cu:127), lost 17.9% - that one is a decode
route, not a prefill corner.

route	T	previous head vs master	this revision vs master
feature_r16_c64, [5120, 25600]	57	+12.84%	+0.90%
	60	+16.36%	+0.91%
	64	+17.86% to +18.90%	+0.00%
DFlash2 attn input, r16_c64_k128	49	+17.86%	-3.57%
	56	+17.24%	-3.45%
	63	+16.67%	-3.33%

The left column is the campaign that found the regression, the right one the campaign that measured
the fix; they are separate runs on the same host and day, each compared against a master arm taken
inside its own run.
The two campaigns disagree in four cells, not one, and one of the four is the headline number.
Each took its own master arm inside its own run, and where those two medians land one 2.048 us step
apart the derived percentage moves with them: T = 57 (+12.84% then +13.75%), T = 64 (+18.90%
then +17.86%, master 227.360 us against 229.376 us), and T = 100 and T = 128 on the
feature_r32_c64 rows. Nothing in the arms changed - the regression arm reads 270.336 us at T = 64
in both - so the honest statement of the worst cell is +17.86% to +18.90%, and the tables below
print whichever number its own campaign measured rather than reconciling them. An earlier revision
of this paragraph called T = 57 the only disagreement.

The regression could not be fixed by dropping the offending line, because it is the same line that
carries most of the gain: on the alignment tooth of two of this branch's own families the full
change returns 57.6% and 59.1%, an activation-only arm returns 50.1% and 50.0%, and a weight-only arm
returns 12.5% and 13.6%. Reverting the activation site would have left roughly a fifth of the
branch. So the policy became conditional instead.

Base

Base origin/master @ a16b6442, re-read 2026-09-07 with
git ls-remote https://github.com/Neroued/ninfer.git refs/heads/master. a16b6442 is one commit
past 487f8977, the base of the previous edition of this report, and that commit is
documentation only - git diff --stat 487f8977 a16b6442 -- src include tests apps bench is
empty. Every file:line in this report was re-read on a16b6442.

Provenance, said once. Three groups of numbers, on three bases, and they are labelled where they
appear.

• This revision, a16b6442, stand B, 2026-09-07. Everything under The regression, Mechanism,
  Effect of this revision, Machine code, Bit-exactness gate, ctest. 984 measured cells in
  the second campaign, 1194 in the first, zero rejected in either.
• The previous head of this branch, ad0f3d38, host R. The seven-family operator table and the
  end-to-end tables under What the branch claims. No timing under that heading was re-taken,
  and it is not the submitted code. What carries it is Machine code below: six of the eight
  translation units that instantiate this kernel are byte-identical between this revision and
  that one, and none of the seven benchmark configurations dispatches into either of the two that
  differ (they hold the only BM = 16 schedules in the tree).
• a140e7ae, host C. The instruction and register counts under Change and the nulling probe
  under Coverage. The suite had 104 tests there; it has 114 on a16b6442.

Scope

One mechanism, three files, +26 / -9, verified by applying the patch to a fresh worktree of
a16b6442 with a plain git apply (git diff --numstat: 5/1, 8/1, 13/7). One new template
parameter with a default, so no call site outside the two that opt in changes behaviour, and no new
instantiation appears for the other 85 schedules. No signature change on the kernel, no new branch,
no new launch, no workspace.

"Seven" throughout counts operator benchmark configurations (attn_input_proj in two formats,
linear_add at two k), mapping onto the six registered tests named under Coverage; "family"
means a configuration. The two routes this revision sets to ca are not among those seven.

Environment

• Stand B - RTX 5090 sm_120a, driver 580.159.03, CUDA 13.1, 525 W, Release,
  -DCMAKE_CUDA_ARCHITECTURES=120a. Clocks not pinned (pinning is not permitted on this host); under
  load 2542-2662 MHz, median 2557, falling to 195 MHz between cells. All numbers on a16b6442.
  Four arms per point: A = master, B = the previous head 66d2eee5, N = this revision,
  Z = a second independent build of master. A and Z are the same tree built twice and their
  binaries differ in md5, so Z is another build of the same code and not a copy of a file.
  Protocol: ABAB with the arm order reversed on odd passes, 6 passes, pass 0 discarded, median of 5;
  every cell its own process; --warmup 5 --repeat 30, cold cache, 256 MiB L2 eviction buffer; one
  exclusive card grab for the campaign, with the witness sampling the card's process list started
  after the lock is taken and stopped before it is released. Zero foreign processes in any of the
  984 cells (312 main, 270 sweep windows, 192 attention, 162 attention windows, 48 tooth); nothing
  rejected. The first campaign, which found the regression, was 1194 cells under the same protocol
  with five arms, also zero rejected.
• Host R - same card, driver 580.159.03, 525 W, 256 cores, on ad0f3d38. Only the tables under
  What the branch claims; 60 clean cells, none rejected.
• Host C - same card, driver 580.173.02, on a140e7ae. Instruction and register counts, nulling
  probe.

Resolution. The operator bench reports medians in steps of 2.048 us. On the feature route at
T = 57 (223.2 us) one step is 0.92%; on the attention route at T = 49 (57.3 us) it is
3.57%. Nothing below one step is resolvable, and the report says so wherever an effect is one
step wide.

Zero-control band. |Z - A| in this campaign: main table max 0.008%, sweep windows max
0.887%, attention 0.000%, attention windows 0.000%, tooth +0.214 pp; median 0.000
everywhere. The single outlying cell is T = 64 in the widest sweep window and is exactly one timer
step. The first campaign gave 0.947%. The rule used throughout: anything under 0.95% is treated as
zero, and anything that survives it is named.

The regression the review found

Stand B, a16b6442, five arms. A master, B the previous head, X = master plus the activation
policy only, W = master plus the weight policy only, Z the zero control. The source of X with
the W patch applied was checked to be byte-identical to B before building.

Where it lost

bench/ops/linear_bench.cu --qtype w8 --n 5120 --k 25600, medians of five passes, us:

T	route (w8_dispatch.cpp:61-66)	A	B	Z	B-A, %	Z-A, %
57	feature_r16_c64	223.232	251.904	223.232	+12.84	+0.00
60	feature_r16_c64	225.280	262.144	225.280	+16.36	+0.00
64	feature_r16_c64	227.360	270.336	229.376	+18.90	+0.89
72...128	feature_r32_c64				-0.70 to -0.80	+0.00
129	mma_r64_c128, Full=false	495.616	428.032	495.616	-13.64	+0.00
256	mma_r64_c128, Full=true	417.792	417.792	415.744	+0.00	-0.49

Re-measured in three further sweep windows (T +/- 4, +/- 8, +/- 12, each its own process): four
independent measurements agree on sign and magnitude at T = 57..64 and at T = 129.
T <= 56 uses a different kernel entirely (w8_feature.cu:19 -> w8_small_t_mma.cuh) and reads
exactly +0.000% at T = 16 / 32 / 48 / 56.

bench/ops/attn_input_proj_bench.cu --format w8-dflash2-qkv, medians of five passes, us:

T	route	A	X	W	B	Z	X-A %	W-A %	B-A %
32	small_t, other kernel	43.008	43.008	43.008	43.008	43.008	+0.00	+0.00	+0.00
49	r16_c64_k128	57.344	67.584	55.296	67.584	57.344	+17.86	-3.57	+17.86
56	r16_c64_k128	59.392	69.632	57.344	69.632	59.392	+17.24	-3.45	+17.24
63	r16_c64_k128	61.440	71.680	59.392	71.680	61.440	+16.67	-3.33	+16.67
64	r32_c32_k128, Full=true	59.392	59.392	59.392	59.392	59.392	+0.00	+0.00	+0.00
96, 128, 192	r32_c64*						+0.00	+0.00	+0.00

The whole regression is one line, and it is also most of the gain

T	route	X-A % (activation only)	W-A % (weights only)	B-A % (both)
57	feature_r16_c64	+12.84	+0.90	+12.84
60	feature_r16_c64	+16.36	+0.91	+16.36
64	feature_r16_c64	+17.86	+0.00	+17.86
129	mma_r64_c128	-13.64	-12.81	-13.64
49 (attn)	r16_c64_k128	+17.86	-3.57	+17.86
63 (attn)	r16_c64_k128	+16.67	-3.33	+16.67

Arm X alone reproduces B to the last digit at every regressing point. And the same line carries
most of what the branch is worth: on the alignment tooth (t(8191)/t(8192) - 1, both points inside
one process) of two of this branch's own families,

family	tooth on master, %	B	X	W	Z
attn_input_proj w8-qgkv	4.570	-2.631 pp, 57.6% of the tooth	-2.349 pp, 50.1%	-0.693 pp, 12.5%	+0.145 pp
attn_input_proj w8-qkv	4.508	-2.664 pp, 59.1%	-2.254 pp, 50.0%	-0.614 pp, 13.6%	+0.000 pp

So reverting the activation site would have left about a fifth of the branch. The fix had to make
the policy conditional.

Mechanism: the discriminator is the tile, not how Full is chosen

At BM = 16 each CTA owns half as many rows, so twice as many CTAs read the same 64-column
activation tile: 5120 / 16 = 320 against 160 on the feature route, 6144 / 16 = 384 against 192
on the attention route. Cache::ca leaves that tile in L1, where CTAs co-resident on one SM hit it
again; Cache::cg bypasses L1 and sends every one of those reads back to L2. At BM >= 32 the
balance is the other way and cg is neutral or better.

This is falsifiable and it was tested. The obvious alternative rule - "ca where Full is
hard-wired, cg where it is chosen by alignment" - is refuted inside one file:
w8_feature.cu:45 hard-wires Full = false for both of its instantiations, launch<16> at :59
and launch<32> at :63, and only the BM = 16 one suffers. The BM = 32 one (T = 72..128)
gains -0.70 to -0.80% from cg and is left on cg here.

How many schedules this rule touches. git grep -h -o "W8RowSplitMmaGemmSchedule<[^>]*" -- src/
on a16b6442 gives 87 instantiations in eight translation units. By first template argument
(BM):

BM	32	64	48	128	96	Rows	16
count	29	28	13	10	5	1	1

The last two are the two this commit touches: w8_feature.cu:42 with Rows == 16, and
w8_dflash2_attn_input.cu:127. The other 85 keep the cg default. Selecting by git grep "W8RowSplitMmaGemmSchedule<16" would be wrong twice over - it matches 16 in the WM / WN
positions (e.g. <32, 64, 16, 16, 3, 2, 128, 1> at w8_dflash2_attn_input.cu:141) and misses
w8_feature.cu, whose BM arrives as a template parameter.

Change

Three files. The policy becomes a schedule parameter with a cg default, which is the idiom five
sibling headers in this tree already use:

// src/ops/linear/w8/w8_rowsplit_gemm_mma.cuh:29
template <int BM_, int BN_, int WM_, int WN_, int MIN_BLOCKS_, int STAGES_ = 2, int BK_ = 64,
          int ACTIVATION_STAGES_ = STAGES_, Cache ActivationCache_ = Cache::cg>
struct W8RowSplitMmaGemmSchedule {
    ...
    static constexpr Cache kActivationCache = ActivationCache_;

Both branches of stage_x then read it - the Full copy at :132 and the predicated
cp_async_zfill at :135 - exactly as bf16_gemm_mma.cuh:208,212 does. The Full branch was
already on cg, so for it this is a rename, not a policy change. Weight codes (:159) and weight
scales (:176) take ninfer::ops::Cache::cg outright: no schedule wants ca there, and on
r16_c64_k128 the weight side is a gain.

The member is spelled kActivationCache while the rest of the struct is upper-case. That is
deliberate: git grep kActivationCache then finds this schedule alongside the others carrying the same knob -
bf16_gemm_mma.cuh:49, fp8_a16_gemm_mma.cuh:34, fp8_a8_mma.cuh:49,
q4_rowsplit_gemm_mma.cuh:50, q5_rowsplit_gemm_mma.cuh:65 and this one.

Then the two opt-ins. On the feature route BM arrives as the template parameter Rows, so the rule
is an expression of it, kept on one line at the point where Rows becomes BM:

// src/ops/linear/w8/w8_feature.cu:42
constexpr Cache kActivationCache = Rows == 16 ? Cache::ca : Cache::cg;
using Schedule = W8RowSplitMmaGemmSchedule<Rows, 64, 16, 16, 1, 2, 128, 1, kActivationCache>;

Named weakness of that form: the rule is written on Rows while the thing that decides is BM.
Here they are the same - Rows is the schedule's first argument - but the link is a convention, not
a compiler check, and Rows == 16 would keep compiling and silently pick the wrong policy if that
ever stopped holding. Checking Schedule::BM is not possible at that point, the schedule not yet
being built. A stricter shape would be a constexpr helper next to the schedule itself, called from
both routes; it is not done here to keep the change minimal, and I will take it if you prefer it.
On the attention route the geometry is written out, so Cache::ca goes straight into the schedule at
w8_dflash2_attn_input.cu:127. Both sites carry a comment giving the CTA counts above, so the next
reader sees the reason rather than the policy.

From the ISA. cp.async.ca.shared.global and cp.async.cg.shared.global both transfer
src-size bytes and both zero-fill up to cp-size = 16; the difference is the cache level the copy
settles in. A cache does not change values. All the sites are already <16>, satisfying the
static_assert at memory.cuh:39 and :56 that the cg form requires a 16-byte copy.

From the binary (Host C, a140e7ae, the previous head). LDGSTS totals were unchanged at 2372
in both arms across 158 instantiations of this kernel, 0 of 158 differing in issue count, and
cuobjdump -res-usage differed in 0 of all 3135 bodies on REG / STACK / SHARED / LOCAL.
Within one arm the 79 pairs of instantiations differing only in Full agree on issue count 79 of 79, so raggedness never cost an extra feed instruction; it changed the form of every existing one.
This revision adds a template parameter with a default, which changes mangled names but instantiates
no new kernel - see the per-unit IDENTICAL column under Machine code.

Effect of this revision

Feature route, [5120, 25600]

Medians of five passes, us. Stand B, a16b6442.

T	route	A (master)	B (previous head)	N (this revision)	Z	B-A, %	N-A, %	Z-A, %
57	feature_r16_c64	223.232	253.920	225.248	223.232	+13.75	+0.90	+0.00
60	feature_r16_c64	225.280	262.144	227.328	225.280	+16.36	+0.91	+0.00
64	feature_r16_c64	229.376	270.336	229.376	229.376	+17.86	+0.00	+0.00
72	feature_r32_c64	256.000	253.952	253.952	256.000	-0.80	-0.80	+0.00
80	feature_r32_c64	256.000	253.952	253.952	256.000	-0.80	-0.80	+0.00
96	feature_r32_c64	260.096	258.048	258.048	260.096	-0.79	-0.79	+0.00
100	feature_r32_c64	266.240	264.192	264.192	266.240	-0.77	-0.77	+0.00
112	feature_r32_c64	274.432	272.384	272.384	274.432	-0.75	-0.75	+0.00
120	feature_r32_c64	280.576	278.560	278.528	280.576	-0.72	-0.73	+0.00
128	feature_r32_c64	286.720	286.720	284.704	286.720	+0.00	-0.70	+0.00
129	mma_r64_c128, Full=false	495.616	428.032	428.032	495.616	-13.64	-13.64	+0.00
192	mma_r64_c128, Full=false	497.664	430.048	430.048	497.664	-13.59	-13.59	+0.00
256	mma_r64_c128, Full=true	417.792	417.792	417.792	417.760	+0.00	+0.00	-0.01

T = 256 is the real control here: 256 % 128 == 0, so Full = true and neither form is on the
executed path. It reads +0.000%. T = 129 and T = 192 are not controls - Full is false
there too, and they are the points with the largest gain.

One difference between the two campaigns, both readings kept. B - A at T = 57 was
+12.84% in the first campaign (B = 251.904) and +13.75% here (B = 253.920); the per-pass
values in this run are 251.936, 251.936, 253.920, 253.952, 253.952, so the median crossed one timer
step. The same happened at T = 128, where the first campaign read B - A = -0.70% and this one
reads +0.00%. Both are right for their own run, and N is compared against the B of the same
run throughout.

The residual at T = 60 is not zero

N - A at T = 60 reads +0.909% in the single point and in all three sweep windows - four
independent measurements out of four. In the main run all five per-pass medians of N read
227.328 us against a master median of 225.280; across the three windows exactly one pass of fifteen
falls below that.

T	single point	window +/-4	window +/-8	window +/-12
57	+0.903	+0.000	+0.000	+0.000
60	+0.909	+0.909	+0.909	+0.909
64	+0.000	+0.000	+0.000	+0.901 (with Z-A = +0.887 in the same cell)
129	-13.636	-13.992	-13.636	-13.992
256	+0.000	+0.000	+0.000	+0.000

That value is exactly one 2.048 us step (225.280 -> 227.328), and it sits inside the zero-control
band of both campaigns (0.887% here, 0.947% before), so the formal criterion "inside the band"
passes. It is nevertheless not noise. It is the same +0.91% that the weight-only arm W read at
that point in the decomposition above, so the most likely reading is that it is the price of moving
the weight side to cg at BM = 16. I did not establish that: no ncu profile was taken and no
L2 traffic was measured. At T = 57 the single point shows one step and all three windows show
zero; at T = 64 three of four measurements show zero and the fourth shows one step against a zero
control that moved by the same step. Below the timer's step this instrument cannot separate a real
+0.9% from zero, and another window will not help - a finer timer would.

DFlash2 attention input, T = 49..63: a regression became a gain

T	route	A	B	N	Z	B-A %	N-A %	Z-A %
32	small_t, other kernel	43.008	43.008	43.008	43.008	+0.00	+0.00	+0.00
49	r16_c64_k128	57.344	67.584	55.296	57.344	+17.86	-3.57	+0.00
56	r16_c64_k128	59.392	69.632	57.344	59.392	+17.24	-3.45	+0.00
63	r16_c64_k128	61.440	71.680	59.392	61.440	+16.67	-3.33	+0.00
64	r32_c32_k128, Full=true	59.392	59.392	59.392	59.392	+0.00	+0.00	+0.00
96	r32_c64_k128	67.584	67.584	67.584	67.584	+0.00	+0.00	+0.00
128	r32_c64_k128, Full=true	83.968	83.968	83.968	83.968	+0.00	+0.00	+0.00
192	r32_c64, Full=true	108.544	108.544	108.544	108.544	+0.00	+0.00	+0.00

Three further windows (the attention bench takes a token list, not a --sweep) give the same
-3.571 / -3.448 / -3.333% in all nine cells, with the zero control at 0.000% in all nine. Each of
those is one timer step wide, but it is one step repeated in four independent measurements with zero
spread, and it reproduces the weight-only arm W of the first campaign to three decimals. This
gain does not exist on the previous head; it appears only because the weight side goes to cg while
the activation side goes back to ca.

The tooth this branch is measured on is unchanged

Same estimator as the branch's own tables, t(8191)/t(8192) - 1 inside one process:

family	tooth A, %	tooth B, %	tooth N, %	B, pp	N, pp	share, B	share, N
attn_input_proj w8-qgkv	4.570	1.939	1.939	-2.631	-2.631	56.2%	57.6%
attn_input_proj w8-qkv	4.295	1.840	1.840	-2.454	-2.454	57.1%	57.1%

Identical to three decimals in percentage points, which is what Machine code predicts: both
families dispatch through w8_attn_input_gemm_mma.cu, whose schedules are all BM in
{32, 64, 128}, and that translation unit compiles byte-identically in the two forms. The share
column is a median of per-pass shares whose per-pass spread is 54.5-60.6% and 54.6-61.8%, so the
57.6 / 59.1% quoted from the first campaign and the 56.2 / 57.1% here are the same number seen twice;
N is compared with the B of its own run.

Roofline

Ceilings taken from this repository rather than from me: 1792 GB/s DRAM and the measured
1674.5 GB/s pure-read ceiling of tools/hbm_bandwidth_probe.cu, both at bench/README.md:255
and :258, and 209.5 TFLOP/s for bf16 with an f32 accumulator,
kRtx5090Bf16Fp32AccumulateTFLOPs at bench/ops/linear_bench.cu:48. us, GB/s and TFLOP/s are
the bench's own columns; the two percentage columns and the MMA column are those divided by those
constants.

point	arm	us	GB/s	% DRAM	% read ceiling	TFLOP/s	% MMA ref
T = 57	A	223.232	639.5	35.7	38.2	66.9	31.9
T = 57	B	253.920	562.2	31.4	33.6	58.8	28.1
T = 57	N	225.280	633.7	35.4	37.8	66.3	31.7
T = 64	A	227.328	629.9	35.2	37.6	73.8	35.2
T = 64	B	270.336	529.7	29.6	31.6	62.1	29.6
T = 64	N	229.376	624.3	34.8	37.3	73.1	34.9
T = 129	A	495.616	297.0	16.6	17.7	68.2	32.6
T = 129	B	428.032	343.9	19.2	20.5	79.0	37.7
T = 129	N	428.032	343.9	19.2	20.5	79.0	37.7

The previous head took 5.6 points of read-ceiling share off this route at T = 64; this revision
takes 0.3 and adds 2.6 at T = 129. On the attention route at T = 49, GB/s derived from the
bench's logical_bytes = 34527232 and its median: A 602.1 GB/s = 33.6% DRAM / 36.0% of the read
ceiling; B 510.9 = 28.5% / 30.5%; N 624.4 = 34.8% / 37.3%.

Machine code: which routes moved

cuobjdump -sass over all eight translation units instantiating this kernel, normalised for the
build path and for the extra template argument in the mangled names.

translation unit	N vs A	N vs B	A vs B
ops/linear/w8/w8_feature.cu	DIFFERENT	DIFFERENT	DIFFERENT
ops/attn_input_proj/w8/w8_dflash2_attn_input.cu	DIFFERENT	DIFFERENT	DIFFERENT
ops/linear/w8/w8_rowsplit_gemm_mma.cu	DIFFERENT	IDENTICAL	DIFFERENT
ops/attn_input_proj/w8/w8_attn_input_gemm_mma.cu	DIFFERENT	IDENTICAL	DIFFERENT
ops/gdn_input_proj/w8/w8_gdn_input_gemm_mma.cu	DIFFERENT	IDENTICAL	DIFFERENT
ops/linear_add/w8/w8_linear_add_gemm_mma.cu	DIFFERENT	IDENTICAL	DIFFERENT
ops/linear_pair/w8/w8_pair_gemm_concat.cu	DIFFERENT	IDENTICAL	DIFFERENT
ops/linear_swiglu/w8/w8_linear_swiglu_gemm_mma.cu	DIFFERENT	IDENTICAL	DIFFERENT

Six of eight are byte-identical to the previous head, and the two that differ are exactly the two
holding a BM = 16 schedule. This is the direct answer to "are the branch's other numbers still the
same numbers": on six of eight units the instructions are literally the same bytes. The A vs B
column is the discriminating control - the comparison does see a cache-policy change when there is
one.

Copy census inside the kernels (LDGSTS.E.BYPASS is cg, without BYPASS is ca; ZFILL is the
predicated copy):

kernel	A	B	N
w8_feature, BM = 32	10 ca ZFILL + 6 ca	10 cg ZFILL + 6 cg	10 cg ZFILL + 6 cg
w8_feature, BM = 16	10 ca ZFILL + 6 ca	10 cg ZFILL + 6 cg	6 ca ZFILL + 4 cg ZFILL + 6 cg
dflash2, BM = 32 (three kernels)	all ca	all cg	all cg
dflash2, BM = 16	10 ca ZFILL + 6 ca	10 cg ZFILL + 6 cg	6 ca ZFILL + 4 cg ZFILL + 6 cg
dflash2, BM = 64	all ca	all cg	all cg

Six copies stay on ca in both BM = 16 kernels - the activation feed - and all ten weight-side
copies go to cg. That is the intent, read off the instructions rather than off the source.

What the branch claims

Provenance: every number in this section was taken on host R, base ad0f3d38, on the previous head
of this branch. None of it was re-taken. It carries to the submitted code because the seven
families dispatch only into translation units that compile byte-identically in the two forms (table
above), and none of them reaches either BM = 16 schedule: w8_attn_input_gemm_mma.cu holds
BM in {32, 64, 128}, gdn_input_proj only 64, linear_swiglu {32, 64, 128}, linear_add
{32, 48, 64, 128}, linear_pair {32, 48, 64, 96, 128}. That is an argument from the machine code,
not a measurement; it is listed again under What is not settled.

Operator

Three passes per arm, arm order rotated, point order reversed in pass 1; 30 of 30 cells clean on the
first attempt. Nothing compared in absolute time across arms: inside one pass of one binary
tooth(T) = t(T)/t(8192) - 1, an arm's effect is the per-pass difference against stock, and the
figure is the median of three.

family	tooth 8191	8160	8128	8096	span	eff 8191	8160	8128	8096	8064 aligned	share@8191
attn_input_proj w8-qgkv	4.46	4.87	4.74	4.74	0.42	+2.65	+2.65	+2.63	+2.91	0.00	59%
attn_input_proj w8-qkv	4.52	4.52	4.31	4.52	0.21	+2.87	+2.65	+2.65	+2.65	0.00	64%
gdn_input_proj w8	4.48	4.74	4.84	4.79	0.37	+2.40	+2.37	+2.37	+2.37	0.00	52%
linear_swiglu w8	5.55	5.97	6.03	5.98	0.48	+2.26	+2.36	+2.36	+2.26	+0.10	39%
linear_add w8, k=4096	4.83	5.11	4.83	4.55	0.57	+2.27	+2.84	+2.56	+2.56	0.00	47%
linear_add w8, k=6144	5.94	5.94	6.13	5.94	0.19	+3.45	+3.26	+3.45	+3.43	0.00	58%
linear_pair w8	5.65	5.65	5.65	5.09	0.56	+3.39	+3.39	+3.39	+2.83	-0.03	60%

Two zero controls: at T = 8192 the untouched Full == true body runs in every arm and absolute
t(8192) spans 0.00 to 0.42% across all arms and families; a second independent build of the
identical base source reads at worst 0.56 pp, exactly one 2.048 us tick on the shortest family.
The share column is a median of per-pass shares, not the effect divided by the tooth; dividing across
the rows would give 59.4 / 63.5 / 53.6 / 40.7 / 47.0 / 58.1 / 60.0%.

The effect is a step, not a share. It is flat on 28 cells of 28 while the number of cancelled
issues varies 96x between T = 8191 and T = 8096, and within 0.10 pp of zero at the aligned
length. A mechanism proportional to the remainder cannot produce that. The span column refutes the
proportionality hypothesis from the stock binary alone: 96x more cancelled issues moves the tooth by
0.19 to 0.57 pp, one or two timer ticks on the two widest families. The price is proportional to
the issues executed in the predicated form, all 1186 of them whenever the columns are ragged.

End to end

ninfer_bench, Qwen3.6-35B-A3B, -r 5 --warmup 2 --prefill-chunk 12288 --max-ctx 12288,
proposal_head=full, mtp_draft_tokens=0. The aligned reference T = 8192 is re-read between every
two ragged points inside a pass; three passes per arm, arms alternated, one exclusive card grab and
one witness verdict per cell.

T	T mod 128	penalty, base	penalty, branch	effect, pp	per pass	share of tooth
8191	127	+2.25%	+1.27%	+0.95	+0.93 / +0.95 / +1.03	42%
8160	96	+1.86%	+0.85%	+1.01	+0.97 / +1.01 / +1.03	54%
8128	64	+1.28%	+0.22%	+1.05	+0.91 / +1.05 / +1.05	82%
8096	32	+1.01%	-0.08%	+1.09	+0.99 / +1.09 / +1.13	108%
8064	0	-1.97%	-1.99%	-0.03	-0.04 / -0.03 / +0.04	zero control

No reference reading was discarded, out of 18 across six cells, by a rule fixed before any effect was
looked at. The effect is flat while the penalty it comes out of falls from 2.25% to 1.01% - the same
signature as the operator table - and the aligned control reads -0.03 pp with both arms running the
same code, so that is the instrument's floor rather than an effect.

Where it does nothing. Aligned token counts: both arms compile to the same instantiation. The
canonical fixture -pg 2048,384 is aligned (2048 mod 128 == 0), so this change is not on the
executed path there; measured anyway it reads -0.087% prefill against -0.041% for a second
build of the identical base whose own passes span -0.224 to +0.155%. Speculation counters are
identical between arms to the last digit under both proposal heads - spec_rounds 678 and acceptance
0.7994100295 with the product default ProposalHead::Full (include/ninfer/types.h:82), 690 and
0.7819767442 with --lm-head-draft - which is what a bit-exact change must do.

Observation: the omission this branch started from

Full is a template parameter (w8_rowsplit_gemm_mma.cuh:65) consumed by if constexpr, so a
ragged call enters a different kernel body, compiled separately. Three consumers feed it:
stage_x activations (131-138), stage_w weight codes (153-161) and weight scales (168-177)
used cp_async<16, Cache::cg> under Full and a bare cp_async_zfill<16> otherwise. The policy is
a property of the instantiated body, so every copy the kernel issues - including the 63 fully
valid tiles of 64 - changed cache level the moment the token count stopped being a multiple of BN.
That is a different cost from predication: predication is about which issues get cancelled, this is
about which form all of them execute in.

Counted on a16b6442 with git grep -n "cp_async_zfill<" -- src include apps bench tests, excluding
the header that defines it: 39 sites in 22 files; eighteen files pass a policy at every site,
three at none, one in part.

file	sites	policy passed?
sparse_moe/prefill/*, softmax_attention/* (3), gdn_gating_proj/bf16/*, dynamic_grouped_conv/bf16/*	15	all Cache::cg
linear/bf16/* (2), linear/fp8/* (3), linear/w8/w8_small_t_mma.cuh (2), linear_topk/w8.cu	8	via Schedule::kActivationCache
context_kv_materialize/*, linear_topk/* (the other four)	5	explicit Cache::ca, chosen
linear/nvfp4/nvfp4_w4a4_mma.cuh	3	activation codes :96 yes; activation scales :111 and :130, the two branches of one if constexpr, no
linear/q4/q4_small_t_mma.cuh	1	none (:114)
linear/w8/w8_rowsplit_gemm_mma.cuh	3	none
linear_pair/w8/w8_pair_gemm_mma.cuh	4	none

cp_async without zfill is called with Cache::cg in 50 further places (51 matches, one of them
the gemm_cp_async wrapper at rowsplit_mma.cuh:52). This commit fixes the first of the three
no-policy files, and moves it into the second row of that table rather than the first: the
activation reaches the policy through the schedule, the weight side spells Cache::cg out.

The default itself must not be changed. Flipping it in memory.cuh:35 / :50 does not compile:
cg requires a 16-byte copy (static_assert at :39, :56) and cp_async is instantiated at 4 and
8 bytes in rowsplit_grouped_mma.cuh:163,199. The correct shape is the one adopted here, where the
policy is a schedule property and the width is checked where used.

Correctness evidence

Bit-exactness gate

Stand B, a16b6442. Four prompts of 17907 / 21423 / 24905 / 29572 characters,
--greedy --seed 1234 --raw-output --max-new 96, both capacities pinned on every invocation
(--max-context 12288 --kv-capacity 12288) so no arm can differ through a capacity derived from free
VRAM, two chunk widths (12288 and 4096), two repeats, ABAB order. A non-zero return code or an answer
under 200 bytes counts as INVALID, never as a match. Card empty at the start
(foreign_apps=[none], 2 MiB occupied), NINFER_* variables in the environment: 0. Every cell is a
fresh single-shot CLI process, and the CLI disables the context cache unconditionally
(apps/cli/main.cpp:281-283).

block	expected	measured
REFLEX, master against a copy of its own binary in another directory - two separate runs, not a file compared with itself	IDENTICAL	16 / 16 IDENTICAL
MAIN, master against this commit	IDENTICAL	16 / 16 IDENTICAL
DISCRIM, master against the saboteur	DIFFERENT	16 / 16 DIFFERENT
INVALID in all three blocks	0	0

The strength control is the point of the exercise, and it is a separate build with a deliberate
fault on exactly the line this commit parameterises, not a mode switch: min(8, k - kk) * 2
becomes * 2 - 2 in the activation feed, so every predicated copy arrives two bytes short. It is
inert while Full == true and corrupts numbers as soon as a column is ragged - a literal model of
how this change could have gone wrong if the cg form treated src-size differently from ca. It
comes out DIFFERENT on all sixteen cells and moves answer lengths as well: 463->499, 452->457,
463->466, 463->399, 464->451, 465->399 bytes. Without that block, "16 of 16 IDENTICAL" would say
nothing.

What the gate does not cover. The two routes this commit sets to ca live at T = 49..63 and
T = 57..64, while --prefill-chunk must be a multiple of 128 (src/serve/serve_options.cpp:356),
so those windows can only be entered by a tail chunk of a prompt, and only at the right length.
Whether they were entered in these 96 runs was not checked - no route counter was added. What the
gate does prove is bit-exactness everywhere the policy actually changes: the weight side on every
route, and the activation on every schedule with BM >= 32. For the two BM = 16 routes the
evidence is the copy census under Machine code (their activation feed carries the same ca copies
as master) and ctest below, which checks them against an oracle at every T from 1 to 128.

ctest

-DBUILD_TESTING=ON (off by default at CMakeLists.txt:19, and ctest on an empty set returns
zero after printing "No tests were found", so everything below is a count and not a return code).

Total Tests: 114
106 Passed, 7 Skipped, 1 Failed        (ctest RC=8)

The seven skips are marked ***Skipped by ctest itself - the tests that need real weight artifacts
(*_real_test, *_load_plan_test).

Both changed routes are covered and both numerical tests pass.
tests/ops/linear/test_w8_a16.cpp:16 runs for (int t = 1; t <= 128; ++t) on geometry
{5120, 25600} (:18, :26), so the whole of feature_r16_c64's T = 57..64 is compared against
an oracle; that test passes. tests/ops/test_attn_input_proj.cpp:509 runs
for (int tokens = 1; tokens <= 128; ++tokens) on the DFlash2 shape (:514), covering the whole of
r16_c64_k128's T = 49..63, and its eager path passes.

The single failure is a defect of master, and here are the numbers. ninfer_attn_input_proj_test
fails on the graph-replay block. Built from an untouched master worktree and run three times
under the same card lock as this branch's binary, run three times:

arm	"criterion failed" lines in three runs
this revision	18, 33, 33
master a16b6442	36, 30, 33

The count moves from run to run on both arms - the test is flaky - and every failure on both
sides is tagged graph phase=1, i.e. the replay = true case at
tests/ops/test_attn_input_proj.cpp:519; the eager variant passes on both. The token counts on which
it fails are the same set on both arms in the retained logs (49, 53, 54, 63, 64, 65, 96, 97; the
harness keeps the first 20 failure lines of each run, and the replay case is registered at :518 for
sixteen token counts). It is filed as issue #196. The statement to check on this base is
therefore "the arm matches the base, and the one failure reproduces on bare master" - not "the suite
is green".

clang-format

Two versions, both on a16b6442 and on a16b6442 plus this change:

file	violations on master	on this branch	new
src/ops/linear/w8/w8_rowsplit_gemm_mma.cuh	0	0	0
src/ops/linear/w8/w8_feature.cu	0	0	0
src/ops/attn_input_proj/w8/w8_dflash2_attn_input.cu	7 (19.1.7) / 17 lines (21.1.8)	same set	0

w8_dflash2_attn_input.cu is already non-conforming on bare master under both versions checked,
and the violating set is identical before and after this change. One violation the change did
introduce was fixed: adding kActivationCache shifted an alignment group and SCALE_CACHE_BYTES
needed one more space. Which version upstream formats with was not established; that seven or
seventeen lines are already flagged on untouched master says it is neither of these.

Coverage: the tests do reach the changed lines

On a140e7ae, a nulling probe - the submitted tree with the payload of exactly those copies set to
zero - failed exactly six of the registered tests and nothing else:
ninfer_linear_w8_a16_test, ninfer_attn_input_proj_test, ninfer_gdn_input_proj_test,
ninfer_linear_add_w8_a16_test, ninfer_linear_pair_w8_a16_test, ninfer_linear_swiglu_w8_a16_test

• the six covering the seven benchmark configurations. That is also the statement about blast radius.
  Match on names: the ordinals from a140e7ae do not carry to this base. The probe branch is not
  submitted, and the probe was not re-run against this revision.

Corrected and superseded

Nothing that has been published on this branch is deleted silently. The previous wordings stand here
next to the corrected ones.

1. The title and the mechanism. Previously:
"perf(ops): restore the L1-bypass cache policy on the ragged feed path of w8_rowsplit_gemm_mma",
and "One ragged column therefore takes the whole feed of the kernel off the L1 bypass. This commit
adds the missing template argument at those three sites: one file, +5 / -4, bit-identical output by
construction."
Now: the L1 bypass is right for the weight side everywhere and for the activation side on 85
schedules of 87, and wrong for the activation side on the two with BM = 16. The policy is a
property of the schedule, not of the ragged path. Three files, +26 / -9.

2. Whether ca could ever be right here. Previously, under Tradeoffs:
"It moves data off L1 for a body that previously used it. Could ca ever be right here? For this
kernel on these seven families, no: 28 cells of 28 improve and the aligned control does not move. The
sweep does not cover a shape where the same weight tile is re-read by many blocks and L1 would serve
it; it is one artifact's production shapes at four ragged lengths."
Now: the answer is yes, on two schedules, at a cost of up to +18.9%. The named gap was real
and the review found it. The one detail the old paragraph had backwards: what many blocks re-read at
BM = 16 is the activation tile, not the weight tile.

3. That the branch does not move the routes below T = 192. Previously, under Base:
"the route is unmoved: every edit to w8_dispatch.cpp between the bases is guarded by t <= 192 or
tighter ... and at T = 8096...8192 every geometry still returns launch_w8_mma_r64_c128."
That statement was about the dispatch table moving between bases, and it is still true. But it
was doing duty as an argument that nothing below T = 192 mattered, and that was wrong: the change
reached t <= 64 and t <= 128 on w8_dispatch.cpp:64-65 and cost 12.8 to 18.9% there.

4. Roofline ceilings. Previously: "Ceilings are this project's own measurements on this card:
1689.4 GB/s for reads, 253.4 TFLOP/s for mma.sync bf16 with an f32 accumulator", giving
gdn_input_proj w8 at 83.6% of the MMA ceiling on 7.96% of the read ceiling at T = 8192.
Neither denominator is documented in this tree, and I cannot point at where they were measured.
The numbers in this revision use the tree's own references instead - 1674.5 GB/s
(bench/README.md:258) and 209.5 TFLOP/s (bench/ops/linear_bench.cu:48). Recomputed against the
first of those, the gdn_input_proj aligned point reads 8.03% of the read ceiling; its
211.9 TFLOP/s is above the 209.5 reference, so for that operator I quote the level and not a
share of a ceiling.

5. Host P, superseded earlier and repeated here so it is not lost. The first edition of the
operator table came from a host with a foreign process on the card and per-run inputs that had not
been kept. Everything was re-taken on host R, and the re-take moved the headline against this
change: operator share of the tooth 43-65% -> 39-64%; end to end +1.12 / +1.27 / +1.09 / +1.12 ->
+0.95 to +1.09 pp; the aligned control -0.18 pp, then reported as an unexplained systematic tail, ->
-0.03 pp, so that sign did not survive. Nothing from host P is quoted as a result.

What this retracts elsewhere

The neighbouring change - splitting Full into FullCols / FullRows so a ragged token count stops
predicating the weight feed - explained its gain as the removal of a predication link on the weight
feed. That explanation is wrong, and this measurement is what shows it. An arm restoring the
cache policy on the weight feed and nothing else - no predication removed, no flag changed, no
instruction added - is worth +0.28 to +1.34 pp on the seven families at T = 8191, against the
+0.58 to +1.78 pp the split measures on the same families at the same point by the same
estimator, from the split's own report. They coincide. The split was winning because taking the
weight feed to the fast body incidentally gave the weights back Cache::cg. The split's numbers
stand; only its account of the mechanism is withdrawn.

Checkable inside this package: if what the split leaves is exactly the activation term, then
activation-only must equal both minus weights-only. Five of seven families agree within
0.38 pp; linear_add k=6144 (1.13) and linear_pair (0.56) do not, so the two policy terms are not
strictly additive there and the whole exceeds the sum of its parts.

For merge order: this change and the Full split are expected not to be additive. The split
moves both weight branches to if constexpr (FullRows), so when only the columns are ragged - the
case this change was written for - two of the three sites already take the fast cg path and only
the activation site still runs. If the split lands first, the weight-side term is likely already
taken while the activation-side term is not, so the expected total is nearer the maximum of the two
than their sum. That is read off the two diffs; the arm that would settle it does not exist.

Tradeoffs

1. Nothing in registers or launches. One template parameter with a default; the 85 schedules that
   do not opt in produce the same instantiations, and six of the eight translation units compile to
   the same bytes as the previous head.
2. It leaves L1 out of the activation feed on 85 schedules and puts it back on two. The reason is
   measured, and the falsifier is in the tree: at BM = 32 in the same file, under the same
   hard-wired Full = false, cg gains 0.70 to 0.80%.
3. The rule is not automatic. kActivationCache defaults to cg; a future BM = 16 schedule
   would get cg and would have to be switched by hand. There is no third BM = 16 schedule in the
   tree to check the rule against.
4. Where the reasoning could be wrong. If the cg form treated src-size differently from ca
   the change would be silently wrong on ragged tiles only - exactly where no default bench looks.
   Hence a saboteur attacking the byte count on that path, and gate prompts long and ragged rather
   than the repository's 70-token fixtures.

The second model. w8_feature is a Qwen3.6-27B route - dflash2/feature_projection, bound at
src/targets/qwen3_6_27b/impl/load/bindings.cpp:400 and materialised at 5120 x 25600 at :654,
which is exactly the dispatch case at w8_dispatch.cpp:61-66. So the two routes this revision
protects are 27B routes, and the branch's measured gain is a 35B prefill effect. No end-to-end 27B
measurement was taken, on either side of this: not for the regression the review found and not for
its removal.

Reproduction

cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_ARCHITECTURES=120a \
  -DNINFER_BUILD_BENCHMARKS=ON -DBUILD_TESTING=ON && cmake --build build -j 24

# The two routes this revision protects. Four arms, six passes, pass 0 dropped, median of five,
# each cell its own process. linear_bench is cold-cache by construction and has no --cache flag.
for T in 57 60 64 72 80 96 100 112 120 128 129 192 256; do
  ninfer_linear_bench --qtype w8 --n 5120 --k 25600 --t $T --warmup 5 --repeat 30
done
ninfer_linear_bench --qtype w8 --n 5120 --k 25600 --sweep 52:68:1 --warmup 5 --repeat 30
ninfer_attn_input_proj_bench --format w8-dflash2-qkv --tokens 32,49,56,63,64,96,128,192 \
  --cache cold --warmup 5 --repeat 30

# The seven families this branch is claimed on (taken on ad0f3d38, not re-taken here)
TS=8192,8191,8160,8128,8096,8064
ninfer_attn_input_proj_bench  --format w8-qgkv --tokens $TS --cache cold
ninfer_attn_input_proj_bench  --format w8-qkv  --tokens $TS --cache cold
ninfer_gdn_input_proj_bench   --format w8      --tokens $TS --cache cold
ninfer_w8_linear_swiglu_bench --t-sweep $TS --problem companion
ninfer_w8_linear_add_bench    --k 4096 --t-sweep $TS --production-only
ninfer_w8_linear_add_bench    --k 6144 --t-sweep $TS --production-only
ninfer_linear_pair_bench      --k 2048 --tokens $TS

# End to end, ragged sweep
build/bench/ninfer_bench --weights qwen3_6_35b_a3b.ninfer \
  -p 8192,8191,8192,8160,8192,8128,8192,8096,8192,8064,8192 \
  -r 5 --warmup 2 --prefill-chunk 12288 --max-ctx 12288 -o csv

# Which schedules the ca rule touches, and the copy census in the binary
git grep -h -o "W8RowSplitMmaGemmSchedule<[^>]*" -- src/ | sed 's/.*<//' | cut -d, -f1 | sort | uniq -c
cuobjdump -sass build/bench/ninfer_linear_bench | grep -c 'LDGSTS.*BYPASS'

Every cell file opens with a state block: md5 of every binary taking part, HEAD and dirty-file
count of its tree, count of NINFER_* markers in the environment (0), the card's occupancy, clock
and power cap, and foreign compute processes. Each cell is one exclusive grab of the card under
flock, and the witness that samples the card's process list starts after the lock is taken and
stops before it is released, so a cell can only be rejected for a process present while it
measured. 984 cells in this campaign, 1194 in the one that found the regression, zero rejected in
either.

What is not settled

1. No end-to-end measurement of either the regression or its removal. Everything about the two
   BM = 16 routes is at the operator level; the share of those calls in a 27B or 35B round was not
   taken. So "+18.9% on an operator" is not converted into a round-level number in either direction.
2. The +0.909% residual at T = 60 is unexplained. It is one timer step, it reproduces four
   times out of four, it matches the weight-only arm at the same point, and it is inside the
   zero-control band. Whether it is really the weight side paying at BM = 16 was not established
   • no ncu profile, no L2 traffic measurement.
3. Five of the seven families were not re-swept at small T on this revision. The argument that
   they are unchanged is that their translation units are byte-identical to the previous head, which
   is a reading of the machine code, not a measurement. Likewise the end-to-end tables.
4. The rule "BM = 16 wants ca" rests on two schedules against 85. There is no third BM = 16
   schedule in the tree, so the rule has never been tested on a case it was not derived from, and a
   new such schedule would silently take the cg default.
5. Whether ca would also win somewhere at BM = 32 was not swept. The change keeps cg there
   because cg measures faster on every BM >= 32 route tested, but a full "policy x tile shape"
   sweep was not run.
6. The gate does not prove it entered the two ca routes - see the end of Bit-exactness gate.
7. ctest has one failure and it is master's. Shown by three runs on each side; the cause of the
   graph phase=1 failure was not investigated, and the retained logs keep only the first 20 failure
   lines per run.
8. Warm cache was not measured; every operator cell is --cache cold. One card, one driver, one
   host, clocks not pinned (2542-2662 MHz under load).
9. Two other files still pass no cache policy at any site, and neither is in this commit:
   linear_pair/w8/w8_pair_gemm_mma.cuh (four sites, short T, selected at
   w8_pair_plan.cpp:46-66) and linear/q4/q4_small_t_mma.cuh:114. Neither is measured here.
   linear/nvfp4/nvfp4_w4a4_mma.cuh:111 and :130 stage activation scales without a policy while
   the codes at :96 have one; also not measured.
10. Interaction with the Full split and with the W8 A8 series is read off the code, never off a
    binary carrying both.

🤖 Generated with Claude Code

[chatgpt-codex-connector]

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-09-08T08:08:02.937090Z	ed15090	Manual request

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 66d2eee575

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

[chatgpt-codex-connector]

Requalify the feature route before changing its cache profile

For the supported W8 [5120,25600] projection at T=57..128, w8_dispatch.cpp:61-66 selects w8_feature.cu, whose launcher deliberately instantiates Full=false even at the full T=64 tile because the Full specialization regressed there. This blanket Cache::cg change also alters that deliberately retained profile, but the reported sweep only covers the unrelated large-T routes around 8192; since this is the latency-sensitive DFlash2 feature projection, remeasure its 57..128 interval and retain a schedule-specific policy if the existing winner no longer holds.

AGENTS.md reference: AGENTS.md:L91-L94

Useful? React with 👍 / 👎.

[MichaelDementii]

The Codex finding is correct, it understated the cost, and the change is now a different change: ed150906 replaces the branch.

What the measurement found. The bot said the blanket Cache::cg also alters the deliberately-predicated w8_feature profile, and that the sweep only covered large T. Re-measured on a16b6442, that route does not merely change profile — the previous head of this branch slowed it down: +12.84% / +16.36% / +17.86% to +18.90% at T = 57 / 60 / 64 on feature_r16_c64. There is also a second hard-wired route the bot did not name, w8_dflash2_attn_input.cu:127, carrying the same <16,64,16,16,1,2,128,1> schedule: +17.86% / +17.24% / +16.67% at T = 49 / 56 / 63, and that one is a decode path. Null-control band 0.947%, witness clean in 1194 cells, four independent sweep windows.

The regression could not be fixed by dropping the offending line. Decomposing the previous head into arms showed the whole regression lives in one of its three sites — the activation load — and that same site carries about 85% of what the PR buys: on the two families the PR measures with, the tooth recovery falls from 57.6/59.1% to 12.5/13.6% without it.

The discriminator is tile shape, not how Full is chosen. My first hypothesis was "cg where Full follows alignment, ca where it is hard-wired". Measurement refuted it: w8_feature with a hard-wired false at Rows = 32 (T = 72..128) does not suffer, while Rows = 16 does. At BM = 16 twice as many CTAs read the same activation plate, so ca keeps it in L1 and they share it; cg bypasses L1 and sends all of it to L2. At BM >= 32 the balance is the other way.

So the policy became a schedule property — kActivationCache, defaulting to cg, the same knob the sibling quantised GEMM schedules already carry — the codes and scales stay on cg unconditionally, and the two BM = 16 schedules in the tree are set to ca. There are exactly two: I enumerated all 87 schedules reaching this kernel and grouped them by first argument.

Result. The regression is gone (+0.90 / +0.91 / +0.00%, inside the control band, and the residual is discussed below); the attn route now gains -3.57 / -3.45 / -3.33%, which it did not before; -13.64 / -13.59% at T = 129 / 192 is preserved; T = 256 on the Full path reads +0.00%; tooth recovery is 57.6/57.1% against the previous head's 56.2/57.1%. Bitwise gate: MAIN 16 of 16 IDENTICAL against master, DISCRIM 16 of 16 DIFFERENT on a sabotaged build, REFLEX 16 of 16. Six of the eight translation units are byte-identical in SASS to the previous head; the two that differ are exactly the ones holding a BM = 16 schedule.

The residual is not zero and is not rounded away. T = 60 reads +0.909%, reproducible 4 of 4 across the single point and three sweep windows. That is exactly one 2.048 us lattice step, formally inside the 0.887% control band, and it is the weight-side cost on BM = 16 that the earlier decomposition measured independently. I have not established its cause.

Two errors of mine, corrected in the body rather than dropped. The headline cell at T = 64 appears as +18.90% in one table and +17.86% in another: the two campaigns each took their own master arm, and those medians land one lattice step apart (227.360 us against 229.376 us). Four cells disagree that way, not one as an earlier revision claimed, and the worst cell is now stated as a range. Separately, a comment in the patch asserted that git grep kActivationCache finds every schedule carrying the policy and listed five; the same command on this base finds twelve files, including q6_rowsplit_gemm_mma.cuh. The comment no longer claims completeness.

Not measured. The end-to-end effect in a 27B round; the cause of the +0.909% residual; five of the seven operator families at small T — that they are unaffected is read off byte-identical SASS, not measured; whether the gate enters the two ca routes (no counter was placed); and the rule "BM = 16 wants ca" has no third case in the tree to test it against, so a new schedule of that shape would have to be marked by hand. ctest is 106 passed, 7 skipped, 1 failed; the failure is ninfer_attn_input_proj_test, built from master and run three times there with the same failure (18/33/33 messages against 36/30/33), i.e. #196 rather than this change.

🤖 Generated with Claude Code

[MichaelDementii]

The branch was replaced after the finding above; the review on this PR still points at the previous head. Current head is ed150906.

@codex review

[chatgpt-codex-connector]

Codex Review: Didn't find any major issues. Delightful!

Reviewed commit: ed15090610

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".