[Hardin22]

Stacked on #905 (and through it on #859 and #876). Only the last two commits are new.

These are the last pieces of the two-GPU work. With them, main reaches what our fork does on the same PC.

1. Decode attention: a merge bounded by the selection width (STRATA_ATTN_MERGE_V2=1, opt-in since 9d27f7e: bitwise the same, but not yet measured apart from the other switches).

   • It is bitwise equal to the current merge. attn_merge_parity checks 768 cases (fp16 / int8 / q4_0 / k8v4 KV) on an RTX 5080 and an RTX 4060 Ti.
   • Faster or equal everywhere in the kernel bench, e.g. one query at full width on the 5080: 98 → 21 µs.
   • Off on HIP, where it has never run.

   Same commit: STRATA_MTP_KV=f16 (opt-in), the draft layer's KV in f16. It changes only the drafts, never the verified tokens, and costs ~230 MiB of VRAM.

2. Pipelined decode (opt-in switches, only inside --pipeline-windows):

   • STRATA_PL_PLE_PREFETCH reads the next window's PLE rows ahead, as the chain's outputs land.
   • STRATA_PL_PLE_LATE gives each stage-0 window its own read handle, so service never blocks on another window's reads. Layer 0's experts are served while its flag waits for the rows.
   • STRATA_PL_EARLY_CHAIN launches the chain before the verdict's bookkeeping.

   The serial loop, run() and the batch windows are unchanged.

Measured

Same setup as #905: 4060 Ti + 5080 split, resident RAM mode (#848), --pipeline-windows 2 --adapt-async 1, --trim-stage-weights, Swift 1.5 IQ3_XXS at 160K (q4_0 KV). Greedy, five prompts × two rounds, decode tok/s:

	run 1	run 2
off (STRATA_ATTN_MERGE_V2=0, nothing else set)	117.5	117.5
on (STRATA_PL_PLE_PREFETCH=1 STRATA_PL_PLE_LATE=1 STRATA_PL_EARLY_CHAIN=1 STRATA_MTP_KV=f16, merge v2 by default)	132.1	130.6

That is +12%; I have not split it per switch. For reference, our fork on the same prompts measured 131.7 and 133.7, so the stack is at parity with it.

Checks

• Exactness: the three pipelined switches give the same greedy text on as off (4/4 prompts, 500 tokens, STRATA_IQ_MT_MIN=1 --pcie-frac 0 --adapt-every 0, blocking tier).
• Stress: with every switch on and STRATA_PIPELINE_FORCE_MISS=2, 11 varied requests gave no degenerate answer and no stall.

The whole two-GPU stack, on this PC

RTX 4060 Ti 16 GB + RTX 5080 16 GB, i9-14900KF, 32 GB of RAM, Swift 1.5 IQ3_XXS at 160K:

	decode tok/s
main today, the 5080 alone (what setup picks for 32 GB)	~29
main today, both cards (--mmap-experts once the file cache is warm)	~64
#848 (resident RAM mode on the split)	~70
+ #859 pipelined windows	~84
+ #876 / #905 async tier beside the pipeline, + #904, --trim-stage-weights	~118
+ this	~131

The last two rows use a fine-tuned draft layer, --spec-min-p 0.65 and Windows power throttling off; the rows above them use the stock draft layer. So read this table as an outline. Each PR has its own A/B with the same settings on both sides.

[CC-David-CC]

absolute king
what a banger

[CC-David-CC]

got any other ideas, king? (cant sleep. up pushing for the roofline https://www.youtube.com/watch?v=xG_Z90EAZZs )

[justxiami]

sm_70 data on the merge commit alone (b1f03d333): bitwise confirmed, kernel faster, end to end flat to slightly worse on 4x V100.

We cherry-picked only that commit onto v0.1.39 (5 files, clean, no rejects) and built it for CMAKE_CUDA_ARCHITECTURES=70, CUDA 12.8. Rig: 4x Tesla V100-SXM2-16GB (sm_70), --layer-split 12,24,36 --split-device 1,2,3 = 4 stages, IQ3_XXS native pack, --kv int8, --spec 3 --spec-min-p 0.35.

The bitwise claim holds on Volta. attn_merge_parity on a V100:

attn_merge_parity on device 0 (Tesla V100-SXM2-16GB, 7.0)
[fp16] 192 cases, bit-identical (0 differ)   [int8] 192 (0)   [q4_0] 192 (0)   [k8v4] 192 (0)
attn_merge_parity: 768 cases, 0 differ, 15225012 nonzero output values in all -> PASS

The kernel is faster here too (--bench, 200 calls, CUDA events, cap 32768):

fmt	n_q	width	merge, old	merge, v2	saved
fp16	1	32768	217.8 us	51.7 us	166.0 us
fp16	3	32768	222.1 us	65.6 us	156.5 us
fp16	7	32768	276.4 us	121.3 us	155.1 us
fp16	3	2051	33.1 us	5.3 us	27.9 us
q4_0	1	64	26.9 us	2.8 us	24.1 us

End to end we cannot see it. Same binary, only the environment differs (STRATA_ATTN_MERGE_V2 1 vs 0), each arm with an engine restart, five fixed prompts at temperature 0 that produce identical window counts in both arms, plus a 32K / 120K / 65K / 8K decode ladder with 3 reps each:

prompt case	windows	verify ms/window V2=0	V2=1	delta
table	53	23.51	24.42	+0.91
json	62	23.14	24.09	+0.95
zh prose	75	22.09	22.93	+0.84
code	74	22.59	22.84	+0.25
tool JSON	37	22.39	22.40	+0.01
mean				+0.59 ms = +2.6%

Decode ladder, per-rep pairing (same rung, same rep): +3.20 / +6.50 / -0.10 (8K), +0.80 / +0.10 / +3.40 (32K), -1.20 / -2.30 / -0.90 (65K), -1.40 / -0.10 / -3.60 (120K) tok/s - 5 faster, 7 slower, pooled +0.42%, which is inside our run-to-run spread. Shape: short prompts slightly positive, long prompts slightly negative - the opposite of the bench table, where wider selections save more.

Two questions:

1. Does the merge win depend on something sm_70 does not have (async copy, L2 residency behaviour, occupancy at HD)? On our rig a window is 22-27 ms with 3-4 rows, and 156 us per call is under 1% of it if the merge runs only a few times per window. How many merge calls does one verify window make per stage on a 48-layer model? If it is once per stage, the arithmetic explains our result.
2. Is the steps read (__ldg(steps + blockIdx.y * kStepCount + kStepWidth)) meant to keep the width on the same card, and does it hold when the chunk pass ran on a different stage? Our split hands off between cards each window.

We did not promote it; our production build is unchanged. We can re-run with more reps, at 2 stages, or with --spec 4 (more rows per window, where the bench says the saving grows) if that would help you decide.

[Hardin22]

Thanks for the careful numbers. Here's what the code says.

Where v2 runs. It runs in every decode attention call:

• the main QSA layers in verify (verify.cpp:883);
• the drafter (mtp.cpp:594);
• plain decode and the prompt path (layer.cpp:980, prefill.cpp:2226).

Calls per window: three per stage.

• The QSA layers are 3, 7, …, 47 (layout.hpp:63), so 12,24,36 puts three on each card.
• Each does one merge launch over all T rows (--spec-split is opt-in, verify.cpp:566): 12 per window.
• They run at cap 2051, i.e. 33 chunks (verify.cpp:355). Once the context passes 2051 tokens the width equals the cap, and only the load batching changes.
• The 512-chunk shape your bench timed (attn_merge_parity is fixed at cap 32768, :254) is the drafter's. That is 1-2 calls per window at --spec 3, on the last card only, and it is timed as draft, not verify.

So your arithmetic holds: on the main layers the merge can't move a ~23 ms window much either way.

Why slower? I can't find added work.

• Same single launch, same grid and block (qsa_decode_attn.cu:516-523). No pre-pass, no sync, no new graph node.
• The only differences are one dependent 4-byte width load per block, and more registers for the 16-chunk staging.
• Nothing needs sm_80.

+0.59 ms would need ~50 µs more on each of 12 small merges, so I suspect noise, but I haven't measured sm_70.

steps. It is always read on the same card. The merge reads the step record the chunk kernel just read, in the same call and on the same stream (:602-603), so a layer's two passes never span cards. Each stage computes the width from its own pos0 (verify.cpp:1346, :1560) into its own device buffer; only the residual crosses cards.

Default. You're right to question it. Our +12% was the bundle (merge v2 plus STRATA_MTP_KV=f16 and the three pipeline switches), never isolated, and our fork had v2 off. So it's now opt-in (9d27f7e, STRATA_ATTN_MERGE_V2=1). I'll post an isolated A/B here: same binary, interleaved, with the draft column from STRATA_DECODE_TIMING=1, which is where the bound should actually skip work. If it only pays in the drafter, I'll limit it to that call.

[Hardin22]

The isolated A/B I promised: same binary (the whole stack, pipelined plus async), STRATA_ATTN_MERGE_V2 0 vs 1, interleaved, five prompts × two rounds, 4060 Ti + 5080, Swift 1.5 IQ3_XXS at 160K. Decode tok/s:

	pair 1	pair 2
v2 off	116.7	133.6
v2 on	129.5	126.0

The two pairs point opposite ways. The mean is +2% for v2, inside our run-to-run spread, so like your V100 result it's flat end to end. It stays opt-in. If I find a case where the drafter's wide merge shows up in the draft time, I'll come back with that alone.

[Niko1221]

Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when main was force-pushed (that is the "base branch was force-pushed" event above). It is not a rejection.

GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new main, cherry-pick your own commits onto it (git fetch upstream && git switch -c my-fix upstream/main && git cherry-pick <your commits>), and open a new PR. Or simply leave it: when we take your change for a release we port it onto the new main ourselves and keep you as its author.

Sorry for the noise, and thank you for the contribution.