perf(ops): choose the predicated W8 GEMM cache policy instead of inheriting it - #201
perf(ops): choose the predicated W8 GEMM cache policy instead of inheriting it#201MichaelDementii wants to merge 1 commit into
Conversation
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
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[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 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".
| } else { | ||
| const int valid = (nn < n && kk < k) ? min(8, k - kk) * 2 : 0; | ||
| ninfer::ops::cp_async_zfill<16>( | ||
| ninfer::ops::cp_async_zfill<16, ninfer::ops::Cache::cg>( |
There was a problem hiding this comment.
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 👍 / 👎.
66d2eee to
ed15090
Compare
MichaelDementii
changed the title
|
The Codex finding is correct, it understated the cost, and the change is now a different change: What the measurement found. The bot said the blanket 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 So the policy became a schedule property — Result. The regression is gone ( The residual is not zero and is not rounded away. Two errors of mine, corrected in the body rather than dropped. The headline cell at Not measured. The end-to-end effect in a 27B round; the cause of the 🤖 Generated with Claude Code |
|
The branch was replaced after the finding above; the review on this PR still points at the previous head. Current head is @codex review |
|
Codex Review: Didn't find any major issues. Delightful! Reviewed commit: ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
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". |
|
Local review of There are useful gains, and numerical validation passed. However, for the feature Linear shape N=5120, K=25600, two opposite-order paired comparisons gave these candidate Op-time changes:
Positive values mean slower. T=57 was essentially neutral. Conditions: RTX 5090, CUDA 13.1, Release sm_120a, g++ 13.3, Ubuntu 24.04/WSL2; cold-cache timing, 5 warmups and 50 samples. These small effects have not been causally established with the same-process/null-controlled experiment used for #200. They are an unresolved acceptance concern, not a confirmed regression of that strength. The positive Op results are clear in these runs: feature T=129 took approximately 11.4% less time; generic Linear [9216,2048] at T=8191 took approximately 2.8% less time; DFlash2 attention T=49/56/63 improved by 0.44–3.24%. The implementation makes activation policy explicit per schedule, with no arithmetic, allocation, workspace, or API changes. Targeted independent FP64 checks passed on both arms: 28 Linear and 42 Q/K/V comparisons across feature/generic/attention route boundaries, including graph replay with updated inputs. Error statistics agreed exactly. The temporary attention test driver settled the phase-1 upload on both arms to avoid the existing #196 test ordering issue; production synchronization was unchanged. Whole-Engine prefill measurements did not resolve the tradeoff: Qwen3.6-35B-A3B, int8 KV, prompts 8191/8192, prefill chunk 12288, max context 12288, 5 repeats after 2 warmups. Ragged pp8191 throughput improved +0.881% / +2.163%, but the unchanged aligned pp8192 control also improved +0.778% / +2.543%. Relative gains after that control were only +0.102% / -0.371%, so these runs do not establish an attributable end-to-end gain. Please qualify the BM=16 feature-route weight/scale policy with a focused controlled comparison around T=57/60/64, retaining the positive T=129 result as a check. If the residual cost is real, adjust the per-schedule policy or demonstrate the workload-level benefit that justifies accepting it. Lack of an end-to-end gain alone is not the reason for deferral; the unresolved cost on the changed feature route is. No broad profiling campaign is needed to address this question. These measurements belong to the baseline above, not today's master. |
…riting it
w8_rowsplit_gemm_mma_kernel loads activations, codes and scales one way on full tiles and another
on predicated ones:
if constexpr (Full) { cp_async<16, Cache::cg>(...); // chosen
} else { cp_async_zfill<16>(...); // Cache::ca, the memory.cuh default
}
The predicated side never chose ca; it inherited cp_async_zfill's default. Which side runs is
decided per launch, not per tile - full needs (w.n % BM) == 0 && (x.ne[1] % BN) == 0 && ... - so any
token count that is not a multiple of BN sends the whole GEMM down the predicated path. This makes
the policy explicit on the schedule and sets it where it has been measured.
One parameter, defaulting to Cache::ca, so adding it changes no instantiation. One schedule opts in:
MmaR64C128, the wide route of the linear family.
WHY ONLY THAT ONE
Because the sign is not a property of the schedule. Sweeping an arm with every predicated
instantiation on cg, against master, on the shapes the dispatch actually sends to each:
schedule shape cg flip null p95
mma_r128_c64 34816x5120 -10.59 % 0.025 %
mma_r64x32_c64_k128_a1 248320x5120 -3.69 % 0.188 %
mma_r64_c96 2048x16384 -3.22 % 0.000 %
mma_r48_c96 / c64 / c128 2048x16384 -2.09..-1.21 % <=0.03 %
mma_r32_c64 / c96 / c112 2048x16384 -0.94..-0.56 % <=0.10 %
mma_r32_c128 2048x16384 -2.92 % 0.149 %
mma_r32_c128 9216x2048 0.00 % 0.000 % unresolved
mma_r64x16_c48_k128_a1 34816x5120 -3.27 % 0.000 %
mma_r64x16_c48_k128_a1 248320x5120 +1.07 % 0.205 % LOSES
Most schedules gain, several by a lot. But MmaR64x16C48K128A1 - a BM = 64 tile, not one of the
BM = 16 ones - gains 3.3 % on one shape it serves and loses 1.1 % on another, 6 of its 8 token
counts positive there. A per-schedule constant cannot express that, so it cannot be set from one
shape's reading. The two BM = 16 schedules are worse still: forced onto cg they cost +16.8 % on the
feature route and +21.0 % on the dflash2 attn route.
MmaR64C128 is the one schedule measured across enough shapes to show a consistent sign, so it is
the one that moves. The knob exists so the next schedule can follow with its own measurement.
EFFECT
One RTX 5090, one session, SM clock recorded per cell (2663 to 2872 MHz, mean 2770), base b88c0f6,
ninfer_linear_bench --execution graph --warmup 5 --repeat 40 --flush-mib 256, three passes with the
arm order rotated; ratios paired inside a pass, median over passes per T, then median over T; CI
from a paired bootstrap over the per-T medians; the null is a second build of master carried through
the same rotation. No number below comes from any other capture.
shape T range grid this patch CI95 null p95
feature 5120x25600 50..56 small_t +0.045 % [+0.00, +1.00] 0.034 %
feature 5120x25600 57..64 BM=16 +0.007 % [+0.00, +0.38] 0.014 %
feature 5120x25600 65..128 BM=32 +0.000 % [+0.00, +0.00] 0.013 %
feature 5120x25600 129..256 160 CTAs -11.27 % [-11.59, -11.22] 0.045 %
feature 5120x25600 257..384 240 CTAs -2.02 % [-2.26, -1.57] 0.163 %
feature 5120x25600 385..400 320 CTAs -0.86 % 0.000 %
generic 9216x2048 4090..4130 144 x 33 -2.76 % [-2.77, -2.76] 0.250 %
mid 34816x5120 129..320 544 x 2..3 -3.77 % [-3.94, -3.52] 0.084 %
lm head 248320x5120 129..320 3880 x 2..3 -3.79 % [-4.07, -3.45] 0.752 %
sparse 2048x4096 896..1100 32 x 7..9 -1.39 % [-1.60, -0.08] 0.158 %
The first three rows are the routes the patch deliberately does not touch, and they read zero to
within their nulls. That is the check that it changes only what it claims.
The -11.27 % is a partial-wave effect and is not the headline: 160 CTAs is 0.94 of a wave on a
170-SM part, and the gain falls to -2.02 % at 1.41 waves and -0.86 % at 1.88. Nothing past T=400 was
measured on that shape. The number to carry is the -2.8 to -3.8 % on the three large shapes, whose
grids run many waves.
Against the machine, from the same rows: 8 to 25 % of the 1674.5 GB/s sustained read, so nothing
here is bandwidth-bound, and 48 to 91 % of the 209.5 TFLOP/s dense bf16 tensor peak. The generic
shape sits at 90.6 % of that peak on master and still gains 2.76 %, which is where the change is
least expected and most worth having. This is not a roofline wall moving; it is L1 tag and
allocation work removed from a path that gets nothing back from L1.
A control the sweep already contains: full needs (T % BN) == 0, so T=4096 is the one point in the
generic sweep the patch cannot act on. It reads +0.0000 % while T=4090 reads -2.71 %, and the spread
over the 40 predicated T is 0.410 pp, so this is not an alignment-phase artefact.
WHY ca IS RIGHT WHERE IT STAYS
Counters from the same card (ncu; hit rates and byte counts only - no timing from that capture is
used here, because ncu serialises launches and controls caches between replays, which removes the
L2 reuse the policy depends on):
capture L1 sector hit rate L2 bytes 1/(1-hit)
BM=16 feature route T=57..64 35.61 % -> 0.64 % x1.549 1.553
BM=32 feature route T=65..128 0.30 % -> 1.07 % x0.992 1.003
MmaR64C128 n=5120 T=129..256 0.29 % -> 0.37 % x1.000 1.003
MmaR64C128 n=248320 T=129..256 1.59 % -> 1.53 % x1.017 1.016
Only the BM=16 schedule gets reuse out of L1, and bypassing it turns exactly those hits into L2
traffic: x1.549 measured against 1.553 predicted from the hit rate alone. At BM=16 the schedule is
128 threads and 22.8 KB of shared memory, so four CTAs sit on an SM and m/BM = 320 of them read the
same 16 KB tile. MmaR64C128 is 46 KB and two CTAs, and at n=5120 its grid is 160 CTAs on 170 SMs, so
no two share an SM. This does not separate "L1 has reuse" from "the tile is shared" - one captured
kernel has a non-trivial hit rate and it is also the only BM=16 one - and the r64x16_c48 result
above shows the rule is not simply about BM either.
VERIFICATION
A cache policy selects where a load is cached, not what it returns, so the output should be
bit-identical. Witnessed rather than argued: the six op tests covering the eight translation units
that include this header were run on master and on this commit with NINFER_OP_REPORT_STATS=1, which
prints max_abs, max_rel and rel_l2 per case at %.17g. That is 4145 numeric records, byte-identical
on both sides:
ninfer_linear_w8_a16_test 600 records
ninfer_linear_add_w8_a16_test 123
ninfer_linear_swiglu_w8_a16_test 192
ninfer_linear_pair_w8_a16_test 288
ninfer_attn_input_proj_test 2816
ninfer_gdn_input_proj_test 126
ctest is 114/114 on both sides. clang-format adds no replacement to either file.
The two sides are separate builds from sources whose hashes differ, and the gate compares the
resulting binaries and fails if they turn out equal. That check is there because the previous
acceptance run for this branch did not have it: it rebuilt one tree twice and compared a build
against itself.
"Changes no instantiation but the one that opts in" is checked rather than asserted. Comparing the
SASS of ninfer_ops between the two builds: 30 instantiations of the kernel and 454 LDGSTS on both
sides, and exactly one instantiation differs - the predicated BM=64 BN=128 one, whose 15 LDGSTS all
gain the BYPASS modifier (0 -> 15 of 15). The other 29 are unchanged.
WHAT IS NOT HERE
No end-to-end number. The one offered in the first version did not survive its own control: ragged
pp8191 improved, but so did the unchanged aligned pp8192, leaving +0.102 % and -0.371 %. This claim
is operator-level.
Five of the nineteen dispatch cells that reach MmaR64C128 are measured, covering five of the fifteen
distinct (n, k) shapes; the other fourteen cells are not.
Every other predicated instantiation keeps the inherited ca, including the ones measured above to
gain from cg. Moving them is a separate change with its own per-shape measurement, because
r64x16_c48 shows a schedule can want opposite things on two shapes it serves.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
ed15090 to
d621ba1
Compare
MichaelDementii
changed the title
|
Replacing the branch. Everything below is one card, one session, SM clock recorded per cell (2663 to 2872 MHz). Nothing The finding that decided the shape of this
So this version moves the one schedule measured across enough shapes to show a consistent sign. What this version isOne parameter, Also here and missing from Effect
The first three rows are the routes this deliberately leaves alone, and they read zero to within The -11.27 % is a partial-wave effect and is not the headline. 160 CTAs is 0.94 of a wave on a Against the machine, from the same rows: 8 to 25 % of the 1674.5 GB/s sustained read, so nothing A control the sweep already contains: Why
|
| capture | L1 sector hit rate | L2 bytes | 1/(1-hit) |
|---|---|---|---|
| BM=16 feature route, T=57..64 | 35.61 % → 0.64 % | x1.549 | 1.553 |
| BM=32 feature route, T=65..128 | 0.30 % → 1.07 % | x0.992 | 1.003 |
| MmaR64C128, n=5120 | 0.29 % → 0.37 % | x1.000 | 1.003 |
| MmaR64C128, n=248320 | 1.59 % → 1.53 % | x1.017 | 1.016 |
One kernel gets reuse out of L1, and bypassing it converts exactly those hits into L2 traffic —
x1.549 measured against 1.553 predicted from the hit rate alone. Forced onto cg the two BM=16
schedules cost +16.8 % (feature) and +21.0 % (dflash2 attn) in this session. But as the
r64x16_c48 row above shows, the rule is not simply about BM, and this capture does not separate
"L1 has reuse" from "the tile is shared".
Verification
A cache policy selects where a load is cached, not what it returns, so the output should be
bit-identical — witnessed, not argued. The six op tests covering the eight translation units that
include this header were run on master and on this commit with NINFER_OP_REPORT_STATS=1, which
prints max_abs, max_rel and rel_l2 per case at %.17g: 4145 numeric records,
byte-identical on both sides (600 / 123 / 192 / 288 / 2816 / 126). ctest 114/114 on both.
clang-format adds no replacement.
"Changes no instantiation but the one that opts in" is checked, not asserted: comparing the SASS
of ninfer_ops between the two builds gives 30 instantiations and 454 LDGSTS on both sides, with
exactly one instantiation differing — the predicated BM=64 BN=128 one, whose 15 LDGSTS all
gain the BYPASS modifier (0 → 15 of 15). The other 29 are unchanged.
The two sides are separate builds from sources whose hashes differ, and the gate compares the
resulting binaries and fails if they are equal. My previous acceptance run on this branch did not do
that, and was worthless because of it: it rebuilt one tree twice and compared a build against
itself.
Not here
- No end-to-end number. The one I offered first did not survive its own control: ragged pp8191
improved, but so did the unchanged aligned pp8192, leaving +0.102 % and -0.371 %. This is an
operator-level claim. - Five of the nineteen dispatch cells reaching
MmaR64C128are measured, covering five of the
fifteen distinct (n, k) shapes. - Every other predicated instantiation keeps
ca, including the ones measured above to gain from
cg. Moving them is a separate change with its own per-shape measurement, because
r64x16_c48shows one schedule can want opposite things on two shapes it serves.
🤖 Generated with Claude Code
|
Codex Review: Something went wrong. Try again later by commenting “@codex review”. ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
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". |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: d621ba1e87
ℹ️ 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".
| template <int BM_, int BN_, int WM_, int WN_, int MIN_BLOCKS_, int STAGES_ = 2, int BK_ = 64, | ||
| int ACTIVATION_STAGES_ = STAGES_> | ||
| int ACTIVATION_STAGES_ = STAGES_, Cache PredicatedCache_ = Cache::ca> |
There was a problem hiding this comment.
When ragged inputs hit consumers outside linear/w8/w8_rowsplit_gemm_mma.cu—including attention-input, GDN-input, linear-add, linear-pair, and SwiGLU—their schedule instantiations omit the new ninth argument, so this Cache::ca default leaves their predicated kernels on the parent's cache policy. Only the local alias in w8_rowsplit_gemm_mma.cu:92-93 opts into cg, meaning most of the seven configurations cited by the performance report do not receive the measured change and the carried-forward operator/end-to-end results do not describe this commit; opt the measured routes into cg or remeasure and narrow the claims.
AGENTS.md reference: AGENTS.md:L91-L94
Useful? React with 👍 / 👎.
| ninfer::ops::cp_async_zfill<16, Cfg::kPredicatedCache>( | ||
| dst, &codes[gi * 32 + chunk * 16], valid_row ? 16 : 0); |
There was a problem hiding this comment.
For the BM=16 feature and DFlash2 routes, the default ca policy is needed for activation loads, but this changed call also applies it to weight codes (and the scale call below). The reported revision specifically attributes the DFlash2 improvement to retaining activation on ca while moving all weight-side copies to cg; the emitted kernel instead keeps both sides on ca, so it cannot deliver that measured behavior. Split the activation and weight policies, or keep the weight copies explicitly on cg.
AGENTS.md reference: AGENTS.md:L91-L94
Useful? React with 👍 / 👎.
|
@codex review |
|
Codex Review: Didn't find any major issues. Delightful! Reviewed commit: ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
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". |
3 participants
w8_rowsplit_gemm_mma_kernelbranches on the template parameterFull, so a ragged token countdoes 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 isCache::ca, while the alignedbody staged the same bytes with
cp_async<16, Cache::cg>. This revision makes the activationpolicy a property of the schedule instead of a constant of the kernel:
Cache::cgby default,Cache::caon the two schedules with aBM = 16tile. The weight side - codes and scales - goes toCache::cgunconditionally. 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 sameFullflag.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
cgalso movesw8_feature, a route that hard-wiresFull = falseeven on full tiles (src/ops/linear/w8/w8_feature.cu:45), and the branch's sweep onlycovered
Tnear 8192. I built the arms and measured it. The bot was right and it understated thecost: 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 decoderoute, not a prefill corner.
Tfeature_r16_c64,[5120, 25600]r16_c64_k128The 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 = 100andT = 128on thefeature_r32_c64rows. Nothing in the arms changed - the regression arm reads 270.336 us at T = 64in 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 = 57the 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 withgit ls-remote https://github.com/Neroued/ninfer.git refs/heads/master.a16b6442is one commitpast
487f8977, the base of the previous edition of this report, and that commit isdocumentation only -
git diff --stat 487f8977 a16b6442 -- src include tests apps benchisempty. Every
file:linein this report was re-read ona16b6442.Provenance, said once. Three groups of numbers, on three bases, and they are labelled where they
appear.
a16b6442, stand B, 2026-09-07. Everything under The regression, Mechanism,Effect of this revision, Machine code, Bit-exactness gate,
ctest. 984 measured cells inthe second campaign, 1194 in the first, zero rejected in either.
ad0f3d38, host R. The seven-family operator table and theend-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 = 16schedules in the tree).a140e7ae, host C. The instruction and register counts under Change and the nulling probeunder 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
a16b6442with a plaingit apply(git diff --numstat: 5/1, 8/1, 13/7). One new templateparameter 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_projin two formats,linear_addat twok), mapping onto the six registered tests named under Coverage; "family"means a configuration. The two routes this revision sets to
caare not among those seven.Environment
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); underload 2542-2662 MHz, median 2557, falling to 195 MHz between cells. All numbers on
a16b6442.Four arms per point:
A= master,B= the previous head66d2eee5,N= this revision,Z= a second independent build of master.AandZare the same tree built twice and theirbinaries differ in md5, so
Zis 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; oneexclusive 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.
ad0f3d38. Only the tables underWhat the branch claims; 60 clean cells, none rejected.
a140e7ae. Instruction and register counts, nullingprobe.
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 atT = 49(57.3 us) it is3.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 max0.887%, attention 0.000%, attention windows 0.000%, tooth +0.214 pp; median 0.000
everywhere. The single outlying cell is
T = 64in the widest sweep window and is exactly one timerstep. 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.Amaster,Bthe previous head,X= master plus the activationpolicy only,
W= master plus the weight policy only,Zthe zero control. The source ofXwiththe
Wpatch applied was checked to be byte-identical toBbefore building.Where it lost
bench/ops/linear_bench.cu --qtype w8 --n 5120 --k 25600, medians of five passes, us:Tw8_dispatch.cpp:61-66)feature_r16_c64feature_r16_c64feature_r16_c64feature_r32_c64mma_r64_c128,Full=falsemma_r64_c128,Full=trueRe-measured in three further sweep windows (
T +/- 4,+/- 8,+/- 12, each its own process): fourindependent measurements agree on sign and magnitude at
T = 57..64and atT = 129.T <= 56uses a different kernel entirely (w8_feature.cu:19->w8_small_t_mma.cuh) and readsexactly +0.000% at
T = 16 / 32 / 48 / 56.bench/ops/attn_input_proj_bench.cu --format w8-dflash2-qkv, medians of five passes, us:Tsmall_t, other kernelr16_c64_k128r16_c64_k128r16_c64_k128r32_c32_k128,Full=truer32_c64*The whole regression is one line, and it is also most of the gain
Tfeature_r16_c64feature_r16_c64feature_r16_c64mma_r64_c128r16_c64_k128r16_c64_k128Arm
Xalone reproducesBto the last digit at every regressing point. And the same line carriesmost of what the branch is worth: on the alignment tooth (
t(8191)/t(8192) - 1, both points insideone process) of two of this branch's own families,
attn_input_projw8-qgkvattn_input_projw8-qkvSo 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
Fullis chosenAt
BM = 16each CTA owns half as many rows, so twice as many CTAs read the same 64-columnactivation tile:
5120 / 16 = 320against 160 on the feature route,6144 / 16 = 384against 192on the attention route.
Cache::caleaves that tile in L1, where CTAs co-resident on one SM hit itagain;
Cache::cgbypasses L1 and sends every one of those reads back to L2. AtBM >= 32thebalance is the other way and
cgis neutral or better.This is falsifiable and it was tested. The obvious alternative rule - "
cawhereFullishard-wired,
cgwhere it is chosen by alignment" - is refuted inside one file:w8_feature.cu:45hard-wiresFull = falsefor both of its instantiations,launch<16>at:59and
launch<32>at:63, and only theBM = 16one suffers. TheBM = 32one (T = 72..128)gains -0.70 to -0.80% from
cgand is left oncghere.How many schedules this rule touches.
git grep -h -o "W8RowSplitMmaGemmSchedule<[^>]*" -- src/on
a16b6442gives 87 instantiations in eight translation units. By first template argument(
BM):BMRowsThe last two are the two this commit touches:
w8_feature.cu:42withRows == 16, andw8_dflash2_attn_input.cu:127. The other 85 keep thecgdefault. Selecting bygit grep "W8RowSplitMmaGemmSchedule<16"would be wrong twice over - it matches16in theWM/WNpositions (e.g.
<32, 64, 16, 16, 3, 2, 128, 1>atw8_dflash2_attn_input.cu:141) and missesw8_feature.cu, whoseBMarrives as a template parameter.Change
Three files. The policy becomes a schedule parameter with a
cgdefault, which is the idiom fivesibling headers in this tree already use:
Both branches of
stage_xthen read it - theFullcopy at:132and the predicatedcp_async_zfillat:135- exactly asbf16_gemm_mma.cuh:208,212does. TheFullbranch wasalready on
cg, so for it this is a rename, not a policy change. Weight codes (:159) and weightscales (
:176) takeninfer::ops::Cache::cgoutright: no schedule wantscathere, and onr16_c64_k128the weight side is a gain.The member is spelled
kActivationCachewhile the rest of the struct is upper-case. That isdeliberate:
git grep kActivationCachethen 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:65and this one.Then the two opt-ins. On the feature route
BMarrives as the template parameterRows, so the ruleis an expression of it, kept on one line at the point where
RowsbecomesBM:Named weakness of that form: the rule is written on
Rowswhile the thing that decides isBM.Here they are the same -
Rowsis the schedule's first argument - but the link is a convention, nota compiler check, and
Rows == 16would keep compiling and silently pick the wrong policy if thatever stopped holding. Checking
Schedule::BMis not possible at that point, the schedule not yetbeing built. A stricter shape would be a
constexprhelper next to the schedule itself, called fromboth 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::cagoes straight into the schedule atw8_dflash2_attn_input.cu:127. Both sites carry a comment giving the CTA counts above, so the nextreader sees the reason rather than the policy.
From the ISA.
cp.async.ca.shared.globalandcp.async.cg.shared.globalboth transfersrc-sizebytes and both zero-fill up tocp-size = 16; the difference is the cache level the copysettles in. A cache does not change values. All the sites are already
<16>, satisfying thestatic_assertatmemory.cuh:39and:56that thecgform requires a 16-byte copy.From the binary (Host C,
a140e7ae, the previous head).LDGSTStotals were unchanged at 2372in both arms across 158 instantiations of this kernel,
0 of 158differing in issue count, andcuobjdump -res-usagediffered in0of all 3135 bodies onREG/STACK/SHARED/LOCAL.Within one arm the 79 pairs of instantiations differing only in
Fullagree on issue count79 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
IDENTICALcolumn under Machine code.Effect of this revision
Feature route,
[5120, 25600]Medians of five passes, us. Stand B,
a16b6442.Tfeature_r16_c64feature_r16_c64feature_r16_c64feature_r32_c64feature_r32_c64feature_r32_c64feature_r32_c64feature_r32_c64feature_r32_c64feature_r32_c64mma_r64_c128,Full=falsemma_r64_c128,Full=falsemma_r64_c128,Full=trueT = 256is the real control here:256 % 128 == 0, soFull = trueand neither form is on theexecuted path. It reads +0.000%.
T = 129andT = 192are not controls -Fullisfalsethere too, and they are the points with the largest gain.
One difference between the two campaigns, both readings kept.
B - AatT = 57was+12.84% in the first campaign (
B = 251.904) and +13.75% here (B = 253.920); the per-passvalues 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 readB - A = -0.70%and this onereads
+0.00%. Both are right for their own run, andNis compared against theBof the samerun throughout.
The residual at
T = 60is not zeroN - AatT = 60reads +0.909% in the single point and in all three sweep windows - fourindependent measurements out of four. In the main run all five per-pass medians of
Nread227.328 us against a master median of 225.280; across the three windows exactly one pass of fifteen
falls below that.
TZ-A = +0.887in the same cell)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
Wread atthat point in the decomposition above, so the most likely reading is that it is the price of moving
the weight side to
cgatBM = 16. I did not establish that: noncuprofile was taken and noL2 traffic was measured. At
T = 57the single point shows one step and all three windows showzero; at
T = 64three of four measurements show zero and the fourth shows one step against a zerocontrol 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 gainTsmall_t, other kernelr16_c64_k128r16_c64_k128r16_c64_k128r32_c32_k128,Full=truer32_c64_k128r32_c64_k128,Full=truer32_c64,Full=trueThree 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
Wof the first campaign to three decimals. Thisgain does not exist on the previous head; it appears only because the weight side goes to
cgwhilethe 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) - 1inside one process:attn_input_projw8-qgkvattn_input_projw8-qkvIdentical 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 allBMin{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;
Nis compared with theBof 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 atbench/README.md:255and
:258, and 209.5 TFLOP/s for bf16 with an f32 accumulator,kRtx5090Bf16Fp32AccumulateTFLOPsatbench/ops/linear_bench.cu:48.us,GB/sandTFLOP/sarethe bench's own columns; the two percentage columns and the MMA column are those divided by those
constants.
T = 57T = 57T = 57T = 64T = 64T = 64T = 129T = 129T = 129The previous head took 5.6 points of read-ceiling share off this route at
T = 64; this revisiontakes 0.3 and adds 2.6 at
T = 129. On the attention route atT = 49, GB/s derived from thebench's
logical_bytes = 34527232and its median: A 602.1 GB/s = 33.6% DRAM / 36.0% of the readceiling; B 510.9 = 28.5% / 30.5%; N 624.4 = 34.8% / 37.3%.
Machine code: which routes moved
cuobjdump -sassover all eight translation units instantiating this kernel, normalised for thebuild path and for the extra template argument in the mangled names.
ops/linear/w8/w8_feature.cuops/attn_input_proj/w8/w8_dflash2_attn_input.cuops/linear/w8/w8_rowsplit_gemm_mma.cuops/attn_input_proj/w8/w8_attn_input_gemm_mma.cuops/gdn_input_proj/w8/w8_gdn_input_gemm_mma.cuops/linear_add/w8/w8_linear_add_gemm_mma.cuops/linear_pair/w8/w8_pair_gemm_concat.cuops/linear_swiglu/w8/w8_linear_swiglu_gemm_mma.cuSix of eight are byte-identical to the previous head, and the two that differ are exactly the two
holding a
BM = 16schedule. This is the direct answer to "are the branch's other numbers still thesame numbers": on six of eight units the instructions are literally the same bytes. The
A vs Bcolumn is the discriminating control - the comparison does see a cache-policy change when there is
one.
Copy census inside the kernels (
LDGSTS.E.BYPASSiscg, withoutBYPASSisca;ZFILLis thepredicated copy):
w8_feature,BM = 32caZFILL + 6cacgZFILL + 6cgcgZFILL + 6cgw8_feature,BM = 16caZFILL + 6cacgZFILL + 6cgcaZFILL + 4cgZFILL + 6cgBM = 32(three kernels)cacgcgBM = 16caZFILL + 6cacgZFILL + 6cgcaZFILL + 4cgZFILL + 6cgBM = 64cacgcgSix copies stay on
cain bothBM = 16kernels - the activation feed - and all ten weight-sidecopies 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 headof 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 = 16schedule:w8_attn_input_gemm_mma.cuholdsBMin {32, 64, 128},gdn_input_projonly 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 thefigure is the median of three.
attn_input_projw8-qgkvattn_input_projw8-qkvgdn_input_projw8linear_swigluw8linear_addw8, k=4096linear_addw8, k=6144linear_pairw8Two zero controls: at
T = 8192the untouchedFull == truebody runs in every arm and absolutet(8192)spans 0.00 to 0.42% across all arms and families; a second independent build of theidentical 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 = 8191andT = 8096, and within 0.10 pp of zero at the alignedlength. A mechanism proportional to the remainder cannot produce that. The
spancolumn refutes theproportionality 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 referenceT = 8192is re-read between everytwo ragged points inside a pass; three passes per arm, arms alternated, one exclusive card grab and
one witness verdict per cell.
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,384is aligned (2048 mod 128 == 0), so this change is not on theexecuted 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_rounds678 and acceptance0.7994100295 with the product default
ProposalHead::Full(include/ninfer/types.h:82), 690 and0.7819767442 with
--lm-head-draft- which is what a bit-exact change must do.Observation: the omission this branch started from
Fullis a template parameter (w8_rowsplit_gemm_mma.cuh:65) consumed byif constexpr, so aragged call enters a different kernel body, compiled separately. Three consumers feed it:
stage_xactivations (131-138),stage_wweight codes (153-161) and weight scales (168-177)used
cp_async<16, Cache::cg>underFulland a barecp_async_zfill<16>otherwise. The policy isa 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
a16b6442withgit grep -n "cp_async_zfill<" -- src include apps bench tests, excludingthe header that defines it: 39 sites in 22 files; eighteen files pass a policy at every site,
three at none, one in part.
sparse_moe/prefill/*,softmax_attention/*(3),gdn_gating_proj/bf16/*,dynamic_grouped_conv/bf16/*Cache::cglinear/bf16/*(2),linear/fp8/*(3),linear/w8/w8_small_t_mma.cuh(2),linear_topk/w8.cuSchedule::kActivationCachecontext_kv_materialize/*,linear_topk/*(the other four)Cache::ca, chosenlinear/nvfp4/nvfp4_w4a4_mma.cuh:96yes; activation scales:111and:130, the two branches of oneif constexpr, nolinear/q4/q4_small_t_mma.cuh:114)linear/w8/w8_rowsplit_gemm_mma.cuhlinear_pair/w8/w8_pair_gemm_mma.cuhcp_asyncwithout zfill is called withCache::cgin 50 further places (51 matches, one of themthe
gemm_cp_asyncwrapper atrowsplit_mma.cuh:52). This commit fixes the first of the threeno-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::cgout.The default itself must not be changed. Flipping it in
memory.cuh:35/:50does not compile:cgrequires a 16-byte copy (static_assertat:39,:56) andcp_asyncis instantiated at 4 and8 bytes in
rowsplit_grouped_mma.cuh:163,199. The correct shape is the one adopted here, where thepolicy 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 freeVRAM, 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 afresh single-shot CLI process, and the CLI disables the context cache unconditionally
(
apps/cli/main.cpp:281-283).REFLEX, master against a copy of its own binary in another directory - two separate runs, not a file compared with itselfMAIN, master against this commitDISCRIM, master against the saboteurINVALIDin all three blocksThe 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) * 2becomes
* 2 - 2in the activation feed, so every predicated copy arrives two bytes short. It isinert while
Full == trueand corrupts numbers as soon as a column is ragged - a literal model ofhow this change could have gone wrong if the
cgform treatedsrc-sizedifferently fromca. Itcomes 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
calive atT = 49..63andT = 57..64, while--prefill-chunkmust 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 twoBM = 16routes theevidence is the copy census under Machine code (their activation feed carries the same
cacopiesas master) and
ctestbelow, which checks them against an oracle at everyTfrom 1 to 128.ctest-DBUILD_TESTING=ON(off by default atCMakeLists.txt:19, andcteston an empty set returnszero after printing "No tests were found", so everything below is a count and not a return code).
The seven skips are marked
***Skippedbyctestitself - 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:16runsfor (int t = 1; t <= 128; ++t)on geometry{5120, 25600}(:18,:26), so the whole offeature_r16_c64'sT = 57..64is compared againstan oracle; that test passes.
tests/ops/test_attn_input_proj.cpp:509runsfor (int tokens = 1; tokens <= 128; ++tokens)on the DFlash2 shape (:514), covering the whole ofr16_c64_k128'sT = 49..63, and its eager path passes.The single failure is a defect of master, and here are the numbers.
ninfer_attn_input_proj_testfails 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:
a16b6442The 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. thereplay = truecase attests/ops/test_attn_input_proj.cpp:519; the eager variant passes on both. The token counts on whichit 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
:518forsixteen 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-formatTwo versions, both on
a16b6442and ona16b6442plus this change:src/ops/linear/w8/w8_rowsplit_gemm_mma.cuhsrc/ops/linear/w8/w8_feature.cusrc/ops/attn_input_proj/w8/w8_dflash2_attn_input.cuw8_dflash2_attn_input.cuis 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
kActivationCacheshifted an alignment group andSCALE_CACHE_BYTESneeded 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 tozero - 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_testMatch on names: the ordinals from
a140e7aedo not carry to this base. The probe branch is notsubmitted, 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 aproperty of the schedule, not of the ragged path. Three files, +26 / -9.
2. Whether
cacould ever be right here. Previously, under Tradeoffs:"It moves data off L1 for a body that previously used it. Could
caever be right here? For thiskernel 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 = 16is 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.cppbetween the bases is guarded byt <= 192ortighter ... and at
T = 8096...8192every geometry still returnslaunch_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 = 192mattered, and that was wrong: the changereached
t <= 64andt <= 128onw8_dispatch.cpp:64-65and 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.syncbf16 with an f32 accumulator", givinggdn_input_proj w8at 83.6% of the MMA ceiling on 7.96% of the read ceiling atT = 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 thefirst of those, the
gdn_input_projaligned point reads 8.03% of the read ceiling; its211.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
FullintoFullCols/FullRowsso a ragged token count stopspredicating 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 numbersstand; only its account of the mechanism is withdrawn.
Checkable inside this package: if what the split leaves is exactly the activation term, then
activation-onlymust equalbothminusweights-only. Five of seven families agree within0.38 pp;
linear_add k=6144(1.13) andlinear_pair(0.56) do not, so the two policy terms are notstrictly additive there and the whole exceeds the sum of its parts.
For merge order: this change and the
Fullsplit are expected not to be additive. The splitmoves both weight branches to
if constexpr (FullRows), so when only the columns are ragged - thecase this change was written for - two of the three sites already take the fast
cgpath and onlythe 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
do not opt in produce the same instantiations, and six of the eight translation units compile to
the same bytes as the previous head.
measured, and the falsifier is in the tree: at
BM = 32in the same file, under the samehard-wired
Full = false,cggains 0.70 to 0.80%.kActivationCachedefaults tocg; a futureBM = 16schedulewould get
cgand would have to be switched by hand. There is no thirdBM = 16schedule in thetree to check the rule against.
cgform treatedsrc-sizedifferently fromcathe 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_featureis aQwen3.6-27Broute -dflash2/feature_projection, bound atsrc/targets/qwen3_6_27b/impl/load/bindings.cpp:400and materialised at5120 x 25600at:654,which is exactly the dispatch case at
w8_dispatch.cpp:61-66. So the two routes this revisionprotects 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
Every cell file opens with a state block: md5 of every binary taking part,
HEADand dirty-filecount of its tree, count of
NINFER_*markers in the environment (0), the card's occupancy, clockand 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 andstops 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
BM = 16routes is at the operator level; the share of those calls in a 27B or 35B round was nottaken. So "+18.9% on an operator" is not converted into a round-level number in either direction.
T = 60is unexplained. It is one timer step, it reproduces fourtimes 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 = 16was not establishedncuprofile, no L2 traffic measurement.Ton this revision. The argument thatthey 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.
BM = 16wantsca" rests on two schedules against 85. There is no thirdBM = 16schedule 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
cgdefault.cawould also win somewhere atBM = 32was not swept. The change keepscgtherebecause
cgmeasures faster on everyBM >= 32route tested, but a full "policy x tile shape"sweep was not run.
caroutes - see the end of Bit-exactness gate.ctesthas one failure and it is master's. Shown by three runs on each side; the cause of thegraph phase=1failure was not investigated, and the retained logs keep only the first 20 failurelines per run.
--cache cold. One card, one driver, onehost, clocks not pinned (2542-2662 MHz under load).
linear_pair/w8/w8_pair_gemm_mma.cuh(four sites, shortT, selected atw8_pair_plan.cpp:46-66) andlinear/q4/q4_small_t_mma.cuh:114. Neither is measured here.linear/nvfp4/nvfp4_w4a4_mma.cuh:111and:130stage activation scales without a policy whilethe codes at
:96have one; also not measured.Fullsplit and with the W8 A8 series is read off the code, never off abinary carrying both.
🤖 Generated with Claude Code