perf(ops): one split4 Q5 parent for the q4/q5 small-column input projections

[MOVIBALE]

Q4/Q5 GDN and attention input projections: one split4 Q5 parent for the small column counts

The Q5 parent of both Q4/Q5 input projections now uses the split4 shape (one CTA owns one output row,
its four warps split the K dimension and reduce their partial sums) for the low column counts instead of
three different shapes. This is a per-side projection route change: no new kernel, no change to the
complete-op structure, no change to the column catalog or to the grouped bands.

What changes, by public call

T is the column count the projection actually receives. The route is resolved per call: the two fused C/D
resolutions (T=1,2,3,5,6) and every batched call have their own resolution, and the flattened aggregate
count W x B is what selects the materialized projection route. Rows below are the public dispatch, not
the private launcher's internal range.

T	Q5 route in this PR's baseline	after this PR, GDN	after this PR, attention
1	GEMV	unchanged	unchanged
2..6	split4 (fused resolutions)	unchanged	unchanged
7..8	row-block	split4	split4
9	c4 SIMT	split4	split4
10	c4 SIMT	split4	c4 SIMT, unchanged
11..12	c4 SIMT	unchanged	unchanged
>=13	each op's own grouped dispatch	unchanged	unchanged

So this PR moves GDN's 7..10 and attention's 7..9; it does not unify every size on split4. The row-block
shape was in production for 7..8 in both projections. It was also a candidate at 9 (and at 10..12 in a
trial) during this work, but a candidate is not a production route, and 9's production route in the
baseline was the c4 SIMT tile.

Q4 dispatch is untouched: T=4 keeps its launch_t4_pdl SIMT pair, and the Q4 K-split MMA band is 7..12.

Measured result

Complete public Op, GPU-side timing of a captured graph holding one call, cold L2, --warmup 5 --repeat 50,
median, baseline and candidate alternating inside one window, three rounds. Baseline is master + #271
(9e163eee); the candidate is this branch.

Configuration	form	cache	baseline us	candidate us	faster by
B=1 W=7	Snapshot	cold	87.81	69.31	21.1%
B=1 W=8	Snapshot	cold	91.39	76.86	15.9%
B=1 W=9	Snapshot	cold	108.29	97.57	9.9%
B=1 W=10	Snapshot	cold	109.86	105.86	3.6%
B=1 W=7/8/9	Record	cold	87.81 / 91.39 / 107.81	68.86 / 75.52 / 97.54	21.6% / 17.4% / 9.5%
B=1 W=10	Record	cold	109.82	105.73	3.7%
B=2 W=4 (8 aggregate)	Snapshot / Record	cold	89.86 / 89.38	75.01 / 75.01	16.5% / 16.1%
B=8 W=1 (8 aggregate)	Snapshot	cold	89.34	74.91	16.2%
B=2 W=5 (10 aggregate)	Snapshot / Record	cold	109.54 / 109.34	104.19 / 104.19	4.9% / 4.7%
B=5 W=2 (10 aggregate)	Snapshot / Record	cold	108.16 / 107.81	103.68 / 103.71	4.1% / 3.8%
attention T=7/8/9	projection	cold	-	-	18.8% / 14.6% / 5.4%

Warm L2 ranks the same way, and the 10-aggregate organisations gain more there (about 9-11%). Every column
count from 11 up, every grouped count, B=1 W=4 and all larger organisations are unchanged, several of
them bit-identical between the two builds.

Both band ends are crossovers between two legal shapes, measured per parent, not limits of either shape:

• GDN parent (12288 rows): split4 through 10. At 10 it wins in every organisation that exposes 10
  aggregate columns, in both forms and both cache policies; at 11 and 12 it loses (118.0 and 140.5 us
  against 111.9 and 113.9 us), so the band stops at 10.
• attention parent (7168 rows): split4 through 9. At 10 it ties and at 11/12 it loses, so the band
  stays at 9.

Verification, by what each check actually compares

• Complete FP64 reference for Snapshot B=1 W=7/8/9/10: the mathematical reference is evaluated for
  every output row and every state channel of Q/K/V and Z, and the code compares them. The output counts
  printed in the log cover Q/K/V; Z and the state are compared by their own checks, so the printed count is
  not the sum of all channels.
• Sampled FP64 reference (same formula, smaller comparison set) for the batched Snapshot helper: the
  B=2/W=5 case added here (10 aggregate columns), the earlier B=2/W=4 masked, B=8/W=1 and the larger
  batched cases.
• Record run_case: the bit-level relation against the validated Snapshot outputs and record, the
  source state left read-only, the invalid tails, the guards and the workspace contract. It does not run its
  own sampled FP64 reference, so the record B=2/W=4 cases added here (dense and masked, 8 aggregate
  columns) are recorded as relation checks, not as sampled FP64.
• Graph replay: the projection helper captures a graph and replays it twice per count, re-poisoning the
  outputs at the captured addresses before each launch and changing the activation at the captured address
  before the second. T=10 is in its list. The conv record helper replays at widths 2/9/10/16 with changed
  inputs; it does the relation checks, not an independent FP64 per replay.
• memcheck 0 errors at the six aggregate-10 organisations; racecheck 0 hazards on the split4 Q5
  parent at three of them, with the instrumented medians in the same logs (328-346 ms against about
  106 us uninstrumented) showing the filtered instance really was analysed.
• ctest on the final build: 100% tests passed, 0 tests failed out of 120; 7 of the 120 are skipped
  by design because no model artifact is supplied (SKIP_RETURN_CODE 77, NINFER_TEST_ARTIFACT /
  --artifact), which is this environment's condition and not an effect of the change. The other 113 ran.
  Logs in logs/tests/.

Not in this PR: fusion and batch organisation

Three candidates were built and measured against the routed form (complete public Op, cold L2, both forms
where applicable):

candidate	routed us	candidate us
B=1 W=4, fused T=4	52.2 (Snapshot) / 50.9 (Record)	54.5 / 53.0	5 -> 4 nodes in both forms, and Snapshot needs no projected workspace (81920 -> 0 B), yet it is still slower
B=1 W=7 fused	69.4	81.2
B=1 W=8 fused	76.3	91.4
B=2 W=4 batched fused	75.0	95.5	one launch per side over the flattened (request, token) axis
B=8 W=1 batched fused	75.0	106.2	same

The candidates did not win, and those measurements only exclude those implementations. For W=7/8 and the
batched organisation, they drive the Q4 side through the row-split SIMT kernel whereas the routed form uses
the K-split MMA kernel there, and they move the convolution into the projection epilogue - an implementation
difference consistent with the size of the losses. For W=4 that explanation does not apply: the routed
form is launch_t4_pdl, whose Q4 side is also SIMT. The workspace saving belongs to the Snapshot prototype
only; the Record form writes the caller-owned conv_record and its temporary-workspace query is 0 in both
builds.

Not implemented: fusing the K-split MMA Q4 side with the Q5 split4 side and a sequence-collecting
convolution. Nothing measured here excludes it.

Each candidate's minimal patch, its per-configuration table and its trial build's test logs are in the
review bundle evidence_r7_review.zip. That bundle is a local file handed over with the review (the fork's
PR attachment facility is not used), and its data/prototypes/README.md states which parts of the patches
are a reconstruction, since the measured prototype sources were reverted before the final build.

Housekeeping

• q5_rowsplit_rowblock_small_t.cuh is removed from the production tree: it has no caller, and its
  mechanism and the __syncwarp() fix stay in git history, together with why it is not routed and the two
  parent shapes' actual mechanisms in the research notes that ship with the review bundle.
• This branch adds no markdown file. The research notes (mechanisms, boundaries and the rejected
  candidates) are delivered with the review bundle instead of living in the repository.
• Production comments were reduced to mechanism, applicable range and preconditions; the performance
  numbers they used to carry live in this body and in the round's report.
• The closing round's edits are comments and documents only, so no candidate was re-measured and the
  repository-wide gate was not re-run for them. Formatting was checked on LF copies with the style
  resolved from the repo root: as a count of lines clang-format would change, the two launchers went from
  28 to 20 and from 20 to 8 lines of pre-existing debt, and the three touched test files are unchanged
  (13/88/11 lines, worktree and reviewed revision identical).

Scope

No release or upload authorization is implied. Merge conditions are the maintainers' decision.

[MOVIBALE]

Review response — the selection evidence for this band, and three fusion attempts (all measured, none kept).

The band end was argued from the wrong candidate. The earlier round justified stopping at 9 by comparing the row-block shape against the narrow c4 SIMT tile, but the shape that is routed here is split4, and split4 vs the c4 tile had never been measured past 9. Measured now on the complete public Op:

• GDN parent: split4 wins at 10 columns in every organisation that exposes 10 aggregate columns (B=1/W=10, B=2/W=5, B=5/W=2), both forms, both cache policies: +3.7% to +11.0%, median +7.2%. At 11 and 12 it loses, so this band now ends at 10 (one extra exact instantiation, one moved boundary).
• attention parent: split4 ties at 10 and loses at 11/12, so its band stays at 9.

Fusion and batch organisation are now measured, not projected. The three candidates the round asked for were built and run against the routed form (complete public Op, cold L2):

candidate	routed	candidate
B=1 W=4, fused T=4	52.2 / 50.9 us	54.5 / 53.0 us	simpler (5→4 nodes, 81920→0 B) and still slower
B=1 W=7/8, fused	69.4 / 76.3 us	81.2 / 91.4 us
B=2 W=4 and B=8 W=1, batched fused	75.0 / 75.0 us	95.5 / 106.2 us	one launch per side over the flattened (request, token) axis

None of them is in the PR. The measured reason is which kernels the existing fused template can reach — its Q4 side is the row-split SIMT kernel whereas the routed form uses the K-split MMA kernel — plus moving the convolution into the projection epilogue. Not implemented, and not claimed either way: fusing the K-split Q4 side with the Q5 split4 side and a sequence-collecting convolution.

Housekeeping in the same commit: the row-block small-T shape is removed (no caller at any column count; its mechanism and the __syncwarp() fix stay in history, and RESEARCH_NOTES.md records why it is not routed); production comments now carry mechanism/range/preconditions only; the test report distinguishes a complete FP64 reference from a sampled one and from the Snapshot/Record bit-level relation.

Verification on the final build: complete FP64 reference at T=7/8/9/10 (the new T=10 included), sampled reference for the batched organisations including the new aggregate-10 case, graph replay at the covered counts, memcheck 0 errors and racecheck 0 hazards on the aggregate-10 organisations (with the instrumented medians showing the filtered instance ran), and ctest 120/120.

[MOVIBALE]

Closing round: the explanations are corrected, the execution code is unchanged.

The performance work ended with the previous round; this pass changes comments, tests' descriptions and
RESEARCH_NOTES.md only, committed as b76b8f7e. No candidate was re-measured and the repository-wide
gate was not re-run for it.

What was wrong, and is now fixed in the body above, in RESEARCH_NOTES.md and in the production comments:

1. split4 is not one warp per row. It is one CTA per output row whose four warps split the K dimension
   and reduce their partial sums through shared memory (chunk = threadIdx.x >> 5, s_part[4][kTt]). The
   one-warp-per-row wording belongs to the c4 SIMT tile and is now kept only there.
2. The c4 SIMT tile does not share an activation slab. It stages the quantized weight planes
   (s_nib/s_hi/s_sc) in shared memory and reads activations from the input tensor; the activation slab
   was the removed row-block experiment's mechanism. The "wins because of activation reuse" causal claim is
   dropped — with no counter or ablation evidence, what remains is that complete-Op measurements retained it
   in 11..12.
3. The W=4 fusion loss is no longer explained by a K-split-to-SIMT fallback. The routed form at T=4
   is the launch_t4_pdl pair, whose Q4 side is already SIMT, so that story cannot explain the result;
   what is recorded is the measurement. In the same place, the workspace saving is attributed to the
   Snapshot prototype only: the Record form writes the caller-owned conv_record and its public
   temporary-workspace query is 0 in both builds.
4. The route table is now the public dispatch, not the private launcher's range: GDN Q5 moves 7..10 and
   attention Q5 moves 7..9, row-block was the production route at 7..8 in both projections, 9's production
   route was the c4 tile (it was only a candidate in split4 form), and Q4 is untouched (T=4 PDL pair, Q4
   K-split MMA band 7..12). The two dispatch layers — the fused per-call resolutions and the materialized
   aggregate W x B selection — are stated separately.
5. Verification is split by what each check compares: the complete FP64 reference (Q/K/V printed counts
   with Z and the state checked separately), the sampled FP64 reference (batched helper), the Record
   run_case bit-level relation against the validated Snapshot, the Graph replay (projection helper only),
   and ctest 120 tests with 7 skipped by design for a missing model artifact — the skip is not counted as
   a run.
6. Boundary numbers now carry their timing boundary (complete public Op / the Q5 parent kernel's
   duration inside the complete Op per nsys / one-side probe), and the probes are explicitly not used to
   move a boundary. Where a number's original log is not in the bundle, the note says so rather than
   splicing a different boundary in.

Also in the bundle: the three rejected candidates as minimal patches with their command set, their
per-configuration table and their trial builds' suite logs. The measured prototype sources were reverted
before the final build and are not retained anywhere, so those patches are marked as reconstructions and
the trial-build binaries that do remain are identified by hash.

Deliverable: evidence_r7_review.zip (2159 files, testzip clean, entry set equal to the bundle
directory's file set). It is handed over as a local file; the exact size and sha256 of the delivered
copy are in evidence_r7_review.txt next to it, written after the last rebuild so the number and the
file cannot drift apart. Nothing is merged, and the merge decision stays with the maintainers.

[Neroued]

The code review and local correctness/performance checks look good. Before merging, please squash the four commits into a single commit, rewrite the PR branch history, and force-push the updated branch. Please keep the final code unchanged.

[MOVIBALE]

History squashed. The branch is now a single commit, 32c3d515, on top of 9e163eee. The four commits it replaces are preserved as the branch backup/pre-squash-281 on this fork, so any id cited in earlier review notes still resolves:

before	what it was
718a9ada	route the q4/q5 small-column Q5 parent to split4 at 7..9
61c2c41a	extend the Q5 split4 band to 10 for the GDN parent
cbd24c14	replay the conv record graph at 10 columns too
b76b8f7e	correct the q5 parent shape explanations and route table

The tree is byte-identical to b76b8f7e (tree 70d1ad70), so the diff against master and every measurement in the review notes are unaffected. No code, route, band or tolerance changed in the squash.