Strata round 9: performance candidate set on v0.2.9 (source audit, implementation-ready)

Assembled 2026-11-02 from the round-9 source audit of upstream v0.2.9 (fork base branch
r9-base). Research on source and prior evidence only; every candidate below has its
mechanism verified in the v0.2.9 source. This issue is written so a smaller model can
implement each candidate without further research, one branch per candidate off r9-base,
each benched A/B before anything merges.

Machine: one RTX 3090 (SM86), Qwen3.8-Flash-Next Swift IQ3_XXS pack, 256K context.
Baselines: D_bench 171.4 / 164.6 / 160.6 tok/s at 4K / 32K / 128K (geomean 165.5);
D_real geomean 105.1; prompt reads 1,900 tok/s (18.8K fresh), 4K 1,550, 32K 2,440,
128K 2,350 tok/s; TTFT floor ~111 ms.

Implemented already: PR #7 (r9/pr1-render-cache): the C-PS-3 history-render cache,
serve.test_server 360/360 plus 17/17 adjacent serve suites on the mock engine, no GPU
needed to verify the tests.

Common rules for every candidate (from the audit brief): opt-in gate (env or config, off
by default, byte-identical stock path when unset); a host-only test where one can exist;
A/B on identical args at 4K/32K/128K with greedy counting decode plus the Hermes-shaped
fixture; report D_bench, D_real, prompt reads and TTFT deltas with the run table in the
PR; no benchmark number merges into the commit message without a fixture; CUDA changes
must run parity tests on the pack before the speed claim.

Summary

Read-only source audit of upstream v0.2.9, c62d5e3a15b40ea51c96aae9e8d36928965f53ab, engine 0.1.55. Between v0.2.6 and v0.2.9, upstream added CUDA persistent hyper-connection execution (2395c8a53), the wide HC up read (d18d8ba04), sm_80 MMVQ-IL tables and hoisted-load kernels (91aebae63, 8b95bffb0, Niko1221#3779), long-context selection (b434a8aa6, Niko1221#3820), and GP100 staged expert down execution (6a6d58ad0, Niko1221#3867). It fixed QFUSE division and one-token commit defects (7d1016f3f, 789d83903, Niko1221#3948). v0.2.9 further changed the rounding-level GR_V3 path: exact 5 to 8 row specialisations and two rows per warp (21f950b4b), then transposed butterfly reductions for 8 rows and 5 to 7 rows (42aee7a34, 1f05fdc28). These are not improvements to the default sm_86 staged-down path. The sm_86 defaults already include side-mixer branches, daily attention, GR_UP75, fused routed gate/up, QSA_ROWS and B6 rows: src/core/nv_defaults.cpp:26-39. Do not count any of those as new work. The source already has an sm_86 MMVQ-IL table, despite the older wording in docs/MMVQ_IL_TABLE.md: src/kernels/cuda/native_mmvq.cu:3783-3793 and src/kernels/cuda/native_mmvq.cu:3853-3859 are authoritative for dispatch.

The leading implementation candidates are an exact four-warp HC down tile, fusion of router top-10 with its doorbell publication, and a 256-thread variant of the existing fused routed gate/up kernel. All need premise measurements before implementation or promotion. This report specifies 12 bounded candidates, including smaller exact-stack members and one deliberately low-priority state-lifecycle change. No new GPU timing, CUDA build, numerical parity, model execution or speedup was obtained. All gains below are estimates using the brief's older profile, not measured v0.2.9 results. No GPU, SSH, source edit, commit or upstream write was performed.

Scope, evidence and non-negotiable gates

What is and is not being claimed

• Scope: Qwen3.8-Flash-Next ENGINE decode, including verify kernels and the GPU commit path. No model-router service work, prefill redesign, CPU expert optimisation or host serve-loop redesign.
• The supplied configuration remains fixed: RTX 3090 sm_86, 250 W, Broadwell AVX2 host, PCIe 3.0 x16, Swift mixed native quantisation, 256K context, INT8 KV, 32K resident KV, the same sampling parameters and the same MTP settings.
• The source worktree was clean and at the exact tag. The corpus was parsed as 4,277 records. It ends on 2026-10-27. Findings about absence of upstream implementations are provisional after that date. Refresh open, closed and merged upstream work before authorising implementation. Eligible only if no open, closed or merged upstream PR implements the proposed bounded mechanism for this engine. Extend the related thread rather than opening a competing proposal.
• The estimated effects are conditional engineering targets, not confidence intervals. A zero benefit or regression is possible for every candidate. Confidence describes the source mechanism and the likelihood of a useful experiment, not validation.
• With unchanged tokens per window, estimated throughput gain is W / (W - saved_ms) - 1. Arithmetic was calculated with Python using W=24.4 ms for counting and W=22.3 ms for a short real-traffic window. The sampled 22.0 ms and agent 25.3 ms windows bound the real-traffic conversion. This is not a fresh geometric-mean estimate across all tiers. In particular, an overlapped stream's kernel-time reduction need not reduce W at all.
• All candidates target exact arithmetic. That classification is a contract to prove, not an executed result. Failure of byte parity rejects the exact candidate; do not quietly relabel it rounding-level. No lossy routing, dropped experts, changed model, different KV dtype, sampling change or speculative-width change is admissible.

Common implementation contract, applies to every candidate

1. One topic per change, a new STRATA_... switch OFF by default, no default-on architecture table change, no enabled-path allocation during graph capture or replay. Initial support is serial single-CUDA sm_86 only. Unsupported backends and configurations retain stock dispatch. Explicit enablement should report whether the eligible path is supported and selected.
2. Preserve stock kernels and stock allocations when OFF. Build comparison must include old stock versus patched OFF, not merely patched OFF versus ON. Add an OFF-by-default row in docs/DETAILS.md next to the existing decode settings at docs/DETAILS.md:1219 or the relevant CUDA opt-in section. No speed claim belongs in the eventual PR without real paired evidence. Observe CONTRIBUTING.md:63-85 and AGENTS.md:89-100.
3. For enabled paths, compare all written output bytes, including q8_1 scales and sums, routed IDs and weights, state, scratch sentinels and padding where contractual. Run direct calls and captured replay with changing row counts and inputs. Run Compute Sanitizer before full-engine A/B. GPU parity executables are GPU tests even where their driver file ends in .cpp.
4. CPU-only checks cover switch resolution, shape/format eligibility, grid coverage, arithmetic index maps and state-lifecycle planning. Existing src/core/arch_defaults_test.cpp and src/core/window_limits_test.cpp are useful regression routes, not proof of a CUDA result. Add a small pure-host test for each new dispatch/index helper. No existing CPU test proves the new kernels correct.
5. Exact deterministic lane: replay fixed token/window shapes, frozen routing/residency and fixed CPU/GPU ownership with identical initial state; compare logits and full raw token IDs as well as final state. This is a parity-only control, not a proposed production configuration. Also run the unchanged adaptive production lane and explain any timing-induced arithmetic ownership changes. Neither identical greedy text alone nor an uncontrolled sampled seed establishes exactness.
6. Final benchmark protocol A/B-G: A is patched OFF with the current production arguments and all new switches unset. B adds only the named switch. Use four adjacent alternating pairs, fresh engines, same warm-up, the same literal prompts and fixed output budgets. Run counting at 4K/32K/128K and real tiny/natural/sampled/agent tiers. Retain individual results and use the median paired geometric-mean ratio with a bootstrap 95% interval. Standalone adoption bar: D_bench >= 1.015 or D_real >= 1.020, lower interval > 1.000, no gated tier median below 0.98, needles 9 of 9. A candidate forecast below that bar is an exact-stack member, not a standalone success. Validate a combined stack against the same bar and use drop-one comparisons to retain only contributors. Do not add or multiply the estimated percentages below.
7. A/B-G also checks fresh reads of 1K/2K/4K and 128K/256K, follow-up reads, and the 17-token TTFT fixture. These are regression controls, not forecast prefill gains. Preserve all production model/KV/sampling settings. Use the fixed greedy-answer set and the existing teacher-forced tools/quant_quality.py route as a supplementary check if any enabled-path discrepancy appears; a discrepancy still fails an exactness claim.
8. Prove switch delivery from /proc/<pid>/environ on the authorised Linux GPU host, and record executable/source identity, launch arguments, defaults log and an eligible-path counter in a separate diagnostic run. /proc is not available on this audit Mac and was not read. Timing runs have diagnostics disabled. Do not use STRATA_VERIFY_PROFILE=1 as the only critical-path evidence: it disables the shared-expert fork at src/core/verify.cpp:2048.
9. Collect a new stock v0.2.9 profile first. The saved round-8 profile is a ranking prior. The GPU owner must run the premise gate while no other build or timed workload competes for the device. An end-to-end failure stays a failure even if a kernel microbenchmark wins.

Ranked candidates

Ranking is expected exposed saving times confidence divided by implementation effort. Ranges are conditional estimates, not measured gains. 0 for reads means no prompt-read speedup is being claimed. TTFT changes cannot be priced from a four-row window and must be measured separately.

R9-gpu-1: four-warp exact staged HC down tiles

Status and ancestry. A concrete extension of r7-dk-DK2, not a rerun of STRATA_GR_DOWN_TG or GR_V3. Niko1221#315 implements activation staging and finer norms; Niko1221#783 implements exact row-count specialisations; Niko1221#2070 implements token-group splitting. Niko1221#186 and the new v0.2.9 changes implement a different, rounding-level stream-split read. None of the inspected records implements the proposed four-warp row tile for the default staged CUDA read. Post-corpus freshness remains unverified.

Source and change site. src/kernels/cuda/fused_gr.cu:1088-1094 stages activations with i += THREADS; src/kernels/cuda/fused_gr.cu:1098-1107 sets warp = t >> 5, identifies the inject block with blockIdx.x == DOWN_BLOCKS, and maps rows with blockIdx.x * WARPS + warp. Stock dispatch is gr_down_staged_kernel<4, true><<<DOWN_BLOCKS + 1, THREADS, smem, st>>> at src/kernels/cuda/fused_gr.cu:1441. The constants describe 40 main row blocks plus inject at src/kernels/cuda/fused_gr.cu:44.

Implementation handoff. Add STRATA_GR_DOWN_WARPS4=1. Template a separate staged-down variant on four warps/128 threads, leaving the stock instantiations untouched. Parameterise the stage helper's copy stride as well as row mapping and launch bounds. Use LR/4 + 1 blocks: 80 main blocks and one inject block, not a claim of all 82 SMs being busy. Preserve dot8u, each lane's ascending tile/chunk loop, and the same xor reduction at src/kernels/cuda/fused_gr.cu:1135-1167. Do not change global WARPS or the up/norm geometry. Register the enabled variant's dynamic-shared-memory attributes alongside the stock staged attributes. Keep token slicing, q8 weight fallback and inject-null handling. Limit the first version to ordinary BF16 HC, GR_V3 off and default staged dispatch.

Why here, estimated. The brief measures 2.63 ms/window in staged down and 4.70 ms for the HC family. Forty-one blocks leave substantial SM parallelism unavailable in a phase that reads the large HC matrices. Target a 10 to 25% reduction of the 2.63 ms phase, conservatively 0.25 to 0.65 ms/window. More CTAs duplicate activation staging; that is an explicit cost, not free bandwidth. Do not transfer the old GR_V3 down-only gain to this path.

Tests and rejection. Existing GPU routes: src/kernels/gr_parity.cpp, src/kernels/gr_multi_parity.cpp, src/kernels/fused_gr_bench.cpp. Run gr_parity --selftest with STRATA_GR_V3=0, both flag arms, all T=1..8, pending write/inject combinations, direct and graph replay. Add host-only row-ownership and staging-coverage tests; arch_defaults_test confirms OFF remains off. Premise gate: fresh stock down >= 2.3 ms/window and variant saves >= 0.25 ms including staging. Then A/B-G, standalone bar. Reject if barriers or shared-memory pressure negate the gain. Exact proof must include inject output, lo and final mixed output, not only generated text.

R9-gpu-2: fuse top-10 with the residency-aware publication tail

Status and ancestry. Reduced-risk first slice of r8 C8. Keep router GEMV separate. Niko1221#783's route_multi batches top-k but does not publish; Niko1221#646's zero-doorbell graph requires every expert resident, unlike this machine. Niko1221#904 changes graph dependencies, not this tail. The corpus router/doorbell fusion search did not find the proposed combined tail. Do not implement the entire GEMV/top-k/publish chain as a single normal grid: independent GEMV CTAs cannot use a block barrier as a grid barrier.

Source and change site. src/core/verify.cpp:2069-2072 launches MMVF then native_router_top10_multi. src/core/verify.cpp:2127-2130 separately calls doorbell_publish_res. src/kernels/cuda/native_router.cu:100-149 gives each token a warp, preserves winning-expert-lane accumulation and normalises selected weights. The exact tie rule is other > best || (other == best && other_id < expert) at src/kernels/cuda/native_router.cu:138. Publication copies IDs, conditionally copies x, fences and updates the sequence at src/kernels/cuda/elementwise.cu:329-346.

Implementation handoff. Add STRATA_ROUTE_PUBLISH_FUSED=1 and a new entry point for 1..8 tokens. Use one CTA for all active token warps, then all CTA threads participate in the residency reduction, x copy, system fence and sequence update. Refactor top-k into a warp-local device helper without changing its expressions. Replace the current early return for inactive token warps with an active predicate around the warp helper: every thread must reach the later CTA barriers. Publish only after all token IDs and weights exist. Preserve the rule that all-hit calls still publish IDs and ring but need not copy x. In Verifier::forward_range's recording path, suppress only the matching later doorbell call. Keep stock execution for route-prior/resident rewrites, remote helpers, always-publish, device/resident-plan paths, profiling and unsupported group geometry. Do not change the host ring protocol.

Why here, estimated. The brief assigns 0.25 ms to top-k and 0.31 ms to publish over 48 layers. Removing one launch per layer is not itself a 0.31 ms saving: the x stores and fences remain. Target 0.08 to 0.20 ms/window, plus an unpriced benefit if CPU expert work starts sooner. This is normally an exact-stack member.

Tests and rejection. Extend src/kernels/router_top10_parity.cpp, src/kernels/native_multi_parity.cpp and src/kernels/verify_parity.cpp; existing elementwise_parity does not by itself establish doorbell ordering. Add a GPU-mapped host consumer test for all-hit, one miss, all-miss, null residency, ties, T=1..8 and repeated graph replay. Host-only tests enumerate fallback eligibility and exactly one publication per layer-group. Kernel parity must preserve weights as well as IDs. Premise: stock tail >= 0.45 ms/window, net tail saving >= 0.08 ms and no early/stale publication in a long stress run. A/B-G, exact-stack bar, with fresh unprofiled CPU-start timestamps. Reject if reduced copy parallelism delays the doorbell or changes CPU/GPU ownership in the frozen lane.

R9-gpu-3: a 256-thread fused gate/up tile

Status and ancestry. New tile for r8 C7's gate/up half. It is not STRATA_EXP_MODE=8, which missed the bar, and not a repeat of the closed exact Q2_0 down reshape. Niko1221#646 and Niko1221#783 implement existing sub-warps; Niko1221#363 reduces work for the PCIe call. The current fused GU kernel still has 512 threads. No matching 256-thread fused-GU tile was found in the corpus. The 256-thread prompt expert switch from Niko1221#3779 is prefill and is not this kernel.

Source and change site. src/kernels/cuda/iq_kernels.cu:2852-2865 documents and launches the existing fused GU body; src/kernels/cuda/iq_kernels.cu:2873-2876 uses 32 sub-warps and r = blockIdx.x * 32 + sw. The output requires a complete 32-value quantisation block at src/kernels/cuda/iq_kernels.cu:2921-2923. The wrapper always launches 512 threads at src/kernels/cuda/iq_kernels.cu:4263-4270.

Implementation handoff. Add STRATA_GU_SWIGLU_256=1. Keep the grid's 32-hidden-column tile. Sixteen sub-warps each compute hidden columns sw and sw + 16 using the unchanged row_dot2_80_sub16 helper. Fill all 32 shared outputs before any q8_1 store. Do not quantise a half block independently. Keep group slicing, format-specific codebook staging, NVFP4 scales, exact n=1/2/3/4 dispatch and rounded SwiGLU multiplication. The second hidden row can be processed serially to limit live registers. Gate on stock launch_gu_swiglu eligibility, not merely a supported tensor type; preserve the 80-dot contract at src/kernels/cuda/iq_kernels.cu:4258-4262. Initial support excludes trellis, folded and nibble layouts.

Why here, estimated. The existing GU phase is 3.16 ms/window, against a 2.09 ms approximate bandwidth floor for the hit gate/up bytes. A 5 to 14% improvement yields approximately 0.15 to 0.45 ms/window. The hypothesis is less per-CTA thread pressure and better useful residency, not fewer weight bytes. Registers can move the other way because one sub-warp owns two columns. Confidence is limited.

Tests and rejection. Extend src/kernels/native_grouped_parity.cpp, src/kernels/native_expert_parity.cpp, src/kernels/iq_multi_parity.cpp and the real-blob native_expert_bench route. Add q8_1 intermediate byte comparison, all actual mixed gate/up types, n_groups 0/1/full, 1..6 entries per group and strided PCIe grids. Add host-only hidden-column/quantisation coverage tests. Premise: stock GU >= 2.8 ms/window, >= 5% weighted actual-format gain, zero local-memory spills and no material one-entry regression. A/B-G, standalone only if it reaches the common bar, otherwise exact-stack member. Do not extrapolate a single IQ3_S layer to the entire mixed pack.

R9-gpu-4: shape-gated sm_86 hoisted-load IL extension

Status and ancestry. Architecture extension of upstream Niko1221#3779/Niko1221#3712, not a new hoisted-load algorithm and not the old r8 C4 row table alone. Absorbed mechanism: 8b95bffb0. The unimplemented scope is a measured sm_86 selector for the hoisted-load variants. Niko1221#3453/Niko1221#3999 are Ada row-table work. The existing sm_86 ordinary table is already selected. The corpus does not establish a subsequent sm_86 PF implementation; refresh that before proceeding.

Source and change site. src/kernels/cuda/native_mmvq.cu:2207-2222 resolves forced PF variants or a measured table. src/kernels/cuda/native_mmvq.cu:3884-3894 explicitly returns zero when il_cc() != 80. src/kernels/cuda/native_mmvq.cu:2225-2230 begins variant dispatch. The real caller eligibility check is src/core/verify.cpp:1508-1516, and the table-supported types are explicit at src/kernels/cuda/native_mmvq.cu:3930-3933.

Implementation handoff. Before writing a table, run the existing mmvq_il_parity --bench --emit-table --pf 0,3,4 measurement route described in docs/MMVQ_IL_TABLE.md, then repeat with production-equivalent B6 defaults and cache-gap arms. Add STRATA_MMVQ_IL_PF86=1, a separate sm_86 table, and a pure-host resolver that selects it only on sm_86 when this switch is set. An unset switch selects no new cells. Explicit conflicting forced PF settings should keep their established meaning or be rejected clearly; do not silently override them. Store only cells that win in every measured shape of their class; prefer exact-shape eligibility where class-wide evidence is incomplete. Do not paste sm_80 values or invent sm_86 table entries. This is M because the deliverable includes resolver isolation, production-equivalent harness controls, the table and tests, not just a table paste.

Why here, estimated. GDN projections cost 3.35 ms/window and QSA projections 1.17 ms in the old profile; only the subset that actually consults IL is available. New PF selection targets approximately 0.08 to 0.25 ms/window after pack/interleave costs. The prior 3090 row-table result was +0.85%, below its bar, and remains only a stack member. It is not additive measured evidence for PF. The Q5_K MTP head is already near its bandwidth floor and does not justify this candidate.

Tests and rejection. Existing GPU route: src/kernels/mmvq_il_parity.cpp, with direct/graph byte parity and hot/gap controls. Existing CPU route: architecture-resolution tests in arch_defaults_test and the table helper coverage reachable from mmvq_il_parity; add a separate pure-host resolver fixture if needed. Verify OFF, unsupported architectures, unknown shapes, explicit PF overrides and production B6 fallback. Premise: >= 5% weighted gain on eligible real calls and >= 0.12 ms/window estimated pack-inclusive saving. A/B-G as an exact-stack member. Reject if the only win is against the parity harness's non-production plain fallback. No sm_86 PF cells were measured here.

R9-gpu-5: let shared-expert projections use IL

Status and ancestry. A narrower, lower-risk alternative to r8 C16's proposed routed/shared expert merger. Niko1221#3264 reuses already-quantised activations and optimises HIP; Niko1221#783 includes older shared fusions. Neither makes the native shared-expert projections use the IL dispatcher. Niko1221#1437/Niko1221#3779 tune verify dense IL, not this direct shared call chain. Actual applicability to this pack's shared tensor formats must be established before coding.

Source and change site. src/kernels/cuda/shared_expert.cu:223-234 selects an existing q8_1 image, attempts a pair kernel, then calls native_mmvq for gate and up. Down also calls native_mmvq at src/kernels/cuda/shared_expert.cu:242. The pair API handles only GGML type 8 at src/kernels/cuda/native_mmvq.cu:3448-3456. IL supports only the listed types at src/kernels/cuda/native_mmvq.cu:3930-3933. src/core/verify.cpp:2151 deliberately supplies private shared-stream scratch.

Implementation handoff. Add STRATA_SH_MMVQ_IL=1. Add private shared-stream IL scratch at initialisation only when enabled. Interleave once for gate/up, use native_mmvq_il_supported per tensor, and preserve stock pair dispatch when it is available and faster. Build a fresh interleaved image after SwiGLU for down only if that down tensor is eligible. Never reuse the verifier main-stream xil_ without a lifetime/fence proof: the main stream overwrites it while the shared stream is active. Do not fuse gate scaling into routed weights, change quantisation, or allocate a copy of shared weights. Initial support is multi-row native verify, not MTP or prefill.

Why here, estimated. Shared expert is 1.71 ms/window but mostly overlapped; 77 GB/s is not proof of 1.5 ms recoverable wall time. Budget zero until an unprofiled trace finds an exposed join tail. Then target at most 0.15 ms/window. A projection microbenchmark can look much better while the window remains unchanged.

Tests and rejection. Extend src/kernels/shared_expert_parity.cpp and src/kernels/mmvq_il_parity.cpp; compare all outputs and intermediate q8_1 data. Add concurrent-stream graph tests with main-stream scratch being overwritten immediately after the fork. CPU-only eligibility tests enumerate every actual shared type/shape; arch_defaults_test covers OFF behaviour. Premise: eligible types and an exposed shared join tail >= 0.15 ms/window; pack-inclusive IL must shorten that tail. A/B-G, exact-stack bar. Stop immediately if all shared matrices use unsupported formats or no exposed tail exists.

R9-gpu-6: Q2_0 down async row slabs without reshaping weights

Status and ancestry. A different premise from r8 C7's lane reshape and the closed exact Q2_0 down reshape, whose best measured rate was 411 GB/s. It also differs from mode 8 activation staging. Niko1221#3867's mechanism is already present as STRATA_DOWN_WS, compiled and gated for GP100: absorbed by 6a6d58ad0. Do not resubmit that plain staged kernel. The proposed new scope is an sm_86 cp.async double-buffered weight-row slab pipeline using original GGUF bytes and the stock per-row tree. It is a research-risk candidate, not evidence that staging will beat the closed reshape.

Source and change site. Q2_0 down is type 42 here, not type 20. Its exact helper is src/kernels/cuda/iq_kernels.cu:1616-1648: s[c] = __fadd_rn(s[c], s16[c]) then xor 8/4/2/1. Stock dispatch is src/kernels/cuda/iq_kernels.cu:4277-4290. The existing weight-staging rationale and build restriction are src/kernels/cuda/iq_kernels.cu:4292-4309. Down grid sizing and early staged dispatch are src/kernels/cuda/iq_kernels.cu:6444-6453.

Implementation handoff. Add STRATA_DOWN_ASYNC86=1. Implement a separate type-42-only native down variant for n_ff=640, retaining the source GGUF layout. Process two consecutive stock 16-row tiles per CTA. Stage the first contiguous weight slab, issue the next slab's aligned 16-byte async copies while calculating the first, then switch buffers. Copy aligned slab spans, not unaligned 16-byte chunks from individual 180-byte rows. Handle allocation/padding boundaries with explicit valid-byte guards; do not read outside the expert blob. Reuse staged hq and call the same row-dot expressions on shared bytes. Keep entry chunking and zero-group/strided-group behaviour. Do not apply the GP100 emulated integer instruction substitutions on Ampere. Keep type 20 IQ4_NL stock initially since its measured 675 GB/s leaves much less headroom.

Why here, estimated. All routed down costs 2.20 ms/window, with Q2_0 at only 314 GB/s versus the IQ4_NL portion at 675 GB/s. The eligible type-42 component must first be measured separately. Target a net 0.10 to 0.25 ms/window; do not price the entire 1.01 ms gap between all-down time and its bandwidth floor. The premise is latency hiding and coalesced issue on unchanged rows, not a claim that direct down is link-bound.

Tests and rejection. Extend src/kernels/native_grouped_parity.cpp, src/kernels/native_expert_parity.cpp and real-blob native_expert_bench. Compare all output bytes for real type-42 layers, poisoned padding, blob-end slabs, n=1..6 per group, empty groups and PCIe striding. Add a CPU-only slab coverage/alignment test. GPU racecheck/memcheck and captured replay are mandatory. Premise: type-42 >= 1.0 ms/window and async variant saves >= 0.10 ms over stock, also compare against the archived reshape control and stock mode 8. Reject if extra shared memory, copy issue or fewer CTA row tiles causes a regression. A/B-G, exact-stack unless standalone bar is reached.

R9-gpu-7: right-size the residency-aware doorbell block

Status and ancestry. Small exact alternative to R9-gpu-2, not an additive companion. Niko1221#649's store-valued ring is a different HIP protocol; this candidate keeps the existing increment. No 256-thread residency-aware publisher was found in the corpus or in the inspected history.

Source and change site. src/kernels/cuda/elementwise.cu:329-346 already loops over activation stores with i += blockDim.x, performs __syncthreads_or, fences, and increments the sequence. The wrapper hard-codes <<<1, 1024>>> at src/kernels/cuda/elementwise.cu:385-390. The captured decode call has at most eight rows and 80 IDs for the supported lane: src/core/verify.cpp:2127-2130.

Concrete S patch sketch. In doorbell_publish_res, preserve the existing size check, arguments and check_launch. Replace only its launch block size as below. The code uses the file's existing C library environment facilities; add <cstdlib> if not already available directly.

int nt = 1024;
#if !defined(__HIPCC__)
static const bool small = [] {
    const char* v = std::getenv("STRATA_DOORBELL_RES_256");
    return v != nullptr && std::atoi(v) == 1;
}();
if (small && n <= 8 * 2560 && k <= 80) nt = 256;
#endif
doorbell_publish_res_kernel<<<1, nt, 0, (cudaStream_t) stream>>>(
    x, ids, d_res, n_expert, (int) n, (int) k, x_out, ids_out, d_seq);

Why here, estimated. The 0.31 ms/window publisher phase includes mapped writes and a system fence, which this does not remove. Test whether less CTA setup/barrier participation yields 0.02 to 0.08 ms/window. This is explicitly an exact-stack member. More x-store iterations per thread can make it slower.

Tests and rejection. Add publisher ordering cases to src/kernels/verify_parity.cpp or a dedicated GPU doorbell fixture; keep src/kernels/elementwise_parity.cpp as a regression route, not claimed existing coverage of the ring. CPU-only tests cover the k=80/81 and n=20480/20481 boundary and OFF dispatch. Test all-hit, null residency, mixed misses, full misses, T=1..8 and graph replay with a polling host consumer. Premise: >= 0.02 ms/window saving without later CPU readiness. A/B-G, stack-only. Do not combine its forecast with R9-gpu-2, which replaces this call.

R9-gpu-8: one exact launch for QSA K and V

Status and ancestry. Narrow re-costing of r8 C10. Niko1221#1889 already overlaps Q and indexer work; it does not merge K/V. The existing pair API covers Q8_0 only and is not a general IL pair. The corpus paired-K/V projection search returned no matching implementation. Long-context top-k itself is already partly absorbed by Niko1221#3820 and already accelerated by the deployed TOPK_SPLIT setting.

Source and change site. src/core/verify.cpp:1838-1844 forks the Q branch and records its join, then src/core/verify.cpp:1848-1849 invokes mm(wk, ...) and mm(wv, ...) consecutively on the main stream. The mm lambda handles IL and fallback at src/core/verify.cpp:1508-1517. src/kernels/cuda/native_mmvq.cu:3448-3456 limits current pairing to type 8.

Implementation handoff. Add STRATA_QSA_KV_PAIR=1. First reuse native_mmvq_pair only where its exact conditions actually match both stock calls. For other eligible types, add a descriptor/grid mapping that gives disjoint CTA ranges to K and V, dispatching their existing per-row IL or B6 body. Do not concatenate or requantise weight tensors. Derive each half's row-tile count from the same ordinary dispatch; preserve n_in, row count, type and rows-per-warp choices independently. Share the already prepared read-only activation image. K normalisation, rotation, KV append and all existing branch waits stay after the merged launch. Initially fall back for packed/nibble/unsupported tensor formats. If the necessary device-body extraction exceeds 300 lines, retain the L classification rather than pretending this is a small pair-call edit.

Why here, estimated. Twelve QSA layers contribute 1.17 ms of projection time, but K/V are only a subset. Removing 12 dependent launches at roughly the brief's 1.8 us small-kernel scale suggests only about 0.02 ms from launches; the conditional 0.03 to 0.10 ms target requires a measured scheduling/occupancy benefit as well. It cannot justify a 1 to 2% decode claim by itself.

Tests and rejection. Extend src/kernels/mmvq_il_parity.cpp, src/kernels/mmvq_multi_parity.cpp and src/kernels/qsa_parity.cpp. Add different K/V format combinations, unequal row counts, awkward final tiles, no-alias checks and Q-branch concurrent graph replay. Pure-host dispatch tests prove each original row is owned exactly once and unsupported pairs keep the stock path. Premise: save >= 0.03 ms/window with stock SIDE_MIXER=2, not SIDE_MIXER=0. A/B-G, exact-stack bar. Reject if the merged launch delays K availability or yields no gain over the existing Q8_0 pair route.

R9-gpu-9: QSA_ROWS at five and six queries

Status and ancestry. Extend Niko1221#2059's implemented scorer, not a new selection algorithm. Niko1221#2072 uses a different halving scorer and two calls beyond four queries; Niko1221#3820 supplies a separate long-context path. This candidate leaves those paths and TOPK_SPLIT unchanged. It does not increase MTP width or lookup width: only windows that already contain five or six queries are eligible. The corpus did not show a CUDA QSA_ROWS 5/6 specialisation.

Source and change site. src/kernels/cuda/qsa_select.cu:1074-1135 already templates query storage and the exact reduce-scatter score tree. The public wrapper rejects nq > 4 at src/kernels/cuda/qsa_select.cu:2174-2183. The ordinary caller attempts it at src/kernels/cuda/qsa_select.cu:2096-2099 and otherwise launches the multi scorer. The long-context verifier path is separate at src/core/verify.cpp:1775-1786.

Concrete S patch sketch. In qsa_block_scores_rows, replace its four-query limit and final launch branch with the following bounded changes. Keep the geometry checks and existing error handling verbatim.

static const bool wide = [] {
    const char* v = std::getenv("STRATA_QSA_ROWS56");
    return v != nullptr && std::atoi(v) == 1;
}();
if (nq > (wide ? 6 : 4) || s.idx_dim != IDX_DIM ||
    s.idx_n_head != IDX_HEADS || s.idx_block != R) return false;
// Keep the existing n == 1, 2 and 3 branches unchanged, then:
else if (n == 4)
    block_scores_rows_kernel<4><<<256, SCORE_WARPS * 32, 0, st>>>(
        pooled, dead, q_idx, steps, n, max_blocks, scores);
else if (n == 5)
    block_scores_rows_kernel<5><<<256, SCORE_WARPS * 32, 0, st>>>(
        pooled, dead, q_idx, steps, n, max_blocks, scores);
else
    block_scores_rows_kernel<6><<<256, SCORE_WARPS * 32, 0, st>>>(
        pooled, dead, q_idx, steps, n, max_blocks, scores);

Why here, estimated. At 128K the profile gives 0.68 ms/window for the four-query scorer. Do not apply that duration to five/six-query frequency without a histogram. Query registers grow linearly, so these versions may spill. Target 0 to 0.12 ms/window, with nearly zero effect on ordinary four-row counting windows. Possible relevance is existing prompt-lookup verification and other already-wide groups, not a new draft policy.

Tests and rejection. Extend src/kernels/qsa_scores_parity.cpp or the existing QSA rows benchmark fixture with explicit row-kernel calls; the existing scores parity test does not automatically prove this new dispatch. Compare full score bits and subsequent selected IDs against ordinary stock dispatch for nq=1..8, equal keys, dead-tail transitions, mixed n_bid, partial blocks and 4K/32K/128K/256K live extents under a fixed 256K capacity. Add a CPU-only eligibility truth table. Premise: existing 5/6-query calls are >= 5% of selection time and their full calls are >= 10% faster without spills. A/B-G as a stack member; reject if wider-row frequency is too small or if only a synthetic all-six-row workload wins.

R9-gpu-10: heterogeneous GDN qkv plus z grid

Status and ancestry. Re-costed r8 C11, with a substantially narrower claim. Niko1221#1889 already runs the z projection beside qkv/conv and the a/b gates on another branch. Niko1221#783 already specialises the existing kernels. Niko1221#2222 already merges quantisation and interleave, and STRATA_QIL was closed by local measurement. None of those is a heterogeneous qkv/z projection grid. This is not permission to disable SIDE_MIXER and claim the lost overlap as a new baseline.

Source and change site. GDN tensor lookup is src/core/verify.cpp:1522-1529. Stock SIDE_MIXER=2 prepares shared activations, launches qkv on cs and z on side2 at src/core/verify.cpp:1547-1561. It waits for z before recurrence at src/core/verify.cpp:1606-1614. The relevant switch-derived state is src/core/verify.cpp:1234-1237.

Implementation handoff. Add STRATA_GDN_PROJ_GRID=1. Create an initialisation-time descriptor containing the two tensor formats, unchanged weight pointers, independent row counts and output strides. A single launch contains CTA ranges for qkv first, then z. Within each range use the same ordinary IL/B6 per-row body and its original row grouping. The a/b branch remains intact. The experimental group replaces the existing qkv/z fork and join only for eligible native serial windows; batch, split windows, packed variants and unsupported format combinations retain stock. Avoid a host table of mutable device values baked into graph capture. No concatenation or type conversion is allowed.

Why here, estimated. The 3.35 ms GDN projection figure includes more than the two input projections. Fusion saves at most one projection node and its branch bookkeeping per layer, but can delay conv until all z CTAs finish. A conditional 0.05 to 0.15 ms/window is a credible test target, not the old C11 projection-wide estimate. It is lower-ranked because extracting shared exact bodies across mixed quant formats is L work for a likely sub-1% gain.

Tests and rejection. Extend src/kernels/mmvq_il_parity.cpp, src/kernels/mmvq_multi_parity.cpp, and src/kernels/gdn_parity.cpp with separate-stock versus merged output-buffer comparisons. Add a CPU-only descriptor/CTA coverage test with unequal row counts and formats. Run SIDE_MIXER=2 graph replay with the ordinary a/b branch still live, T=1..8, and compare recurrence outputs and untouched verify state. Premise: merged projection-to-conv-to-recurrence dependency span saves >= 0.05 ms/window against stock branches, not merely fewer nodes. A/B-G, exact-stack bar. Reject if qkv's early conv advantage is lost.

R9-gpu-11: four score cells per step in daily INT8 attention

Status and ancestry. A small architecture-specific dispatch experiment using the template from Niko1221#2075/Niko1221#849, not a new attention algorithm. The source explicitly says four was slower on V100. That negative result prevents a blanket default and lowers confidence; it does not measure an 82-SM, 250 W sm_86 card. Niko1221#1888's tensor-core split attention is already present and changes rounding, so it is not the proposed exact mechanism.

Source and change site. The daily attention kernel is src/kernels/cuda/qsa_decode_attn.cu:544-550; src/kernels/cuda/qsa_decode_attn.cu:589-606 reuses shared query slices across NC score cells. The comment at src/kernels/cuda/qsa_decode_attn.cu:593 states CUDA defaults to two and four was slower on V100 because registers cost occupancy. The CUDA macro selects the default two-cell instance at src/kernels/cuda/qsa_decode_attn.cu:902-908. Value loads are already grouped four at a time at src/kernels/cuda/qsa_decode_attn.cu:678-688; do not propose that as new work.

Concrete S patch sketch. In the CUDA, non-HIP arm of the macro definition, add this helper and replace only that CUDA macro. First scope the host-side entry to sm_86 using a cached capability established before capture; the helper below is the switch read, not a replacement for that architecture eligibility check.

bool attn_score4_on() {
    static const bool on = [] {
        const char* v = std::getenv("STRATA_ATTN_SCORE4");
        return v != nullptr && std::atoi(v) == 1;
    }();
    return on;
}
#define STRATA_ATTN_CHUNK_SEL(M, PRE) \
    ((PRE) ? (((M) == 1 && attn_score4_on()) \
        ? attn_chunk_kernel_pre75<M, 4, false> \
        : attn_chunk_kernel_pre75<M>) : attn_chunk_kernel<M>)

The caller must pass the experimental choice only for the ordinary serial decode entry, not use this macro change to broaden MTP, slot, row-pool or prefill eligibility inadvertently. Keep unsupported entry points on the stock macro/helper. Existing STRATA_ATTN_DAILY=1 remains a prerequisite, not another changed arm setting.

Why here, estimated. The chunk phase is 0.55 to 0.57 ms/window at 4K and 128K. A 5 to 10% local reduction gives only about 0.02 to 0.06 ms/window. No changing chunk size, softmax merge order, selected cells or KV dtype. This is stack-only and should be dropped quickly if register occupancy is worse.

Tests and rejection. Extend src/kernels/qsa_merge_daily_parity.cpp and src/kernels/qsa_parity.cpp to compare partial accumulators, maxima, denominators and final output for the two-cell and four-cell instances. Cover 0/1/2/3/63/64/65 selected cells, page -1 masking and graph replay. CPU-only tests prove unsupported modes/entry points never select it. Premise: >= 5% whole chunk-call gain, no spills, no one-row regression. A/B-G as an exact-stack member. If the macro/caller isolation needs more than about 60 lines, classify the implementation M rather than broadening its scope to remain S.

R9-gpu-12: accepted GDN replay inside the next recurrence

Status and ancestry. Re-costed r8 C18, deliberately last. STRATA_COMMIT_BATCH is already default-on: do not propose batching 36 commits as new work. Niko1221#2755/Niko1221#3226 cover other-backend/batched replay. Niko1221#3819 removes a layer-split host wait; Niko1221#3944 rotates pipeline snapshots. Neither implements the proposed serial per-layer replay-plus-next-recurrence body. This candidate may overlap a host-turnaround workstream, so assign one owner before implementation.

Source and change site. src/core/verify.cpp:2761-2785 already commits all GDN conv histories and states in two launches. Verify recurrence loads the original state at src/kernels/cuda/verify_kernels.cu:324-328, advances it with the original per-row FMAs at src/kernels/cuda/verify_kernels.cu:340-357, and writes state only when the commit contract allows it at src/kernels/cuda/verify_kernels.cu:380-383. The conv/hbuf/gate/beta scratch reused by the next window is visible at src/core/verify.cpp:1531-1537 and the current-window producers at src/core/verify.cpp:1588-1606.

Implementation handoff. Add STRATA_GDN_REPLAY_NEXT=1. Keep two bounded generations of the previous window's hbuf/gate/beta data and the accepted count. At the next layer's recurrence, load its old committed state once, replay exactly the accepted previous rows into registers with the stock state-only arithmetic, store that committed state, then calculate the new speculative rows from the same registers without storing rejected speculation. The store of the previous committed state is necessary even if a later cancellation occurs. Conv histories must still be committed before the next conv reads them; this is not permission to defer conv history. Do not allocate a full state per speculative token, which would consume the small remaining VRAM reserve. Compute the extra bytes from actual n_gdn, C, HV and max_t, and reject enablement if the reserve cannot hold them without reducing expert-cache slots or KV residency.

Pending replay must flush before checkpoint/save/park, request completion, prompt reads, session restore/replacement, any external state hash, and destruction/cancellation paths that expose the session. Initial enabled scope excludes batch, pipeline, split windows and any unproven MTP consumer of the target state. Clear generation ownership only after stream-ordered flush completion. Never allow the new window's producers to overwrite pending old hbuf/gate/beta. These are required implementation steps, not details left to a later optimisation pass.

Why here, estimated. The older profile includes about 0.28 ms for GDN commit in the draft range. It is not a proven exposed main-state cost on v0.2.9. The only plausible saving is state store/reload and some replay scheduling, not eliminating replay arithmetic. Conditional net target 0.08 to 0.22 ms/window after extra data/lifecycle costs. Short requests may gain nothing because terminal flushes remain. Confidence is low and the L effort usually loses to earlier exact-stack work.

Tests and rejection. Existing src/kernels/gdn_parity.cpp:558-615 compares batched versus per-layer state and conv commit bytes. Extend it to ordinary commit then next verify versus deferred replay then next verify, accepted counts 0..T, changing T=1..8, and compare full state, conv histories and outputs. Add pure-host generation/flush tests next to src/core/conversation_snapshot_test.cpp and src/core/conversation_cache_test.cpp, including cancel, rewind, restore, prompt transition and destruction. GPU graph replay and sanitizer are mandatory. Premise: first prove exposed target-state commit >= 0.15 ms/window, buffer budget without any slot/KV loss, and no MTP dependency that forces an early flush; require >= 0.08 ms net before A/B-G. Reject rather than narrowing correctness to happy-path consecutive decode. Use the exact-stack bar and full TTFT/read regression suite.

Absorbed, considered and dropped

• r8 C2, old STRATA_GR_V3: its old 4K cancellation result of about 1.003 remains relevant. The new 5..8-row exact specialisations improve that path relative to itself, not exact equivalence to stock HC and not a proven four-row gain. Absorbed by 21f950b4b, 42aee7a34, 1f05fdc28. A separate settings recheck is possible, but not a new implementation candidate and not classed exact against stock.
• CUDA persistent HC: absorbed by 2395c8a53; the old STRATA_HC_PERSIST=1 (opt-in, HIP): the verifier's hyper-connection read as one persistent kernel (gfx1030: read 1.35-2.10x, per window -4.2 % / -2.8 %) Niko1221/Strata#2226 body saying CUDA gets stubs is stale. It changes the reduction structure and has resident-grid barrier risks. Do not file another port or assume CUDA is missing it.
• Wide HC up: absorbed by d18d8ba04; sm_120 default changes do not establish a 3090 gain. Existing STRATA_GR_UP75 is already enabled on sm_86.
• QSA multi-CTA long select: absorbed by b434a8aa6/[Engine] Decode past 64K cells: a multi-CTA top-k and a shuffle-free block scorer in their own window graphs (STRATA_SELECT_LONG=1, opt-in, GP100 builds) Niko1221/Strata#3820. Its special build gate and sm_6x dispatch are not a free improvement on the deployed sm_86 TOPK_SPLIT path. No new top-k gain is credited here.
• Router full GEMV/top-k/publish in one conventional grid: rejected as an M-sized promise. It needs a correct cross-CTA completion mechanism or a single-CTA design that would serialise a 512-row GEMV. R9-gpu-2 captures a bounded safe slice.
• Router load-ahead: not new. src/kernels/cuda/native_bf16.cu:95-107 already preloads up to eight pairs per thread. include/strata/kernels/pdl.hpp:50-57 keeps that prefetch enabled on sm_86 even though PDL execution is sm_90+. A source read that assumes Ampere has PRE=0 is wrong.
• Global restore of __restrict__: 0.1.40.2: decode on a Tesla P40 (sm_61) is about half of 0.1.40's; bisected to 2e4ddf6, and restoring __restrict__ in pdl.hpp (STRATA_PDL_RESTRICT) gives back most of it Niko1221/Strata#1469/Pascal (sm_61): decode halves since 0.1.40.2. 2e4ddf6 drops __restrict__ and peels the MMVQ loop on every CUDA arch Niko1221/Strata#1470 already studied it and the Pascal (sm_61): decode halves since 0.1.40.2. 2e4ddf6 drops __restrict__ and peels the MMVQ loop on every CUDA arch Niko1221/Strata#1470 sm_86 RTX 3080 control reported no measurable difference. Changing the PDL-wide alias contract risks sm_90+ ordering. Not a credible new 3090 candidate without new SASS/counter evidence.
• Exact Q2_0 8-lane reshape: closed by local measurement, best 411 GB/s. No rerun is proposed. R9-gpu-6 must establish a different async issue premise and beat that control.
• Mode 8, QIL, TSUM, GDN_SPLIT, EXPDEC, fused waits, HC_REQ8, IQ256 gather: closed or below-bar in the brief. Their presence in source does not make them new recommendations.
• QFUSE repair: division fix absorbed by 7d1016f3f, grouped parity by 11b049022, one-token commit by 789d83903. The old 15 to 20% real-traffic regression is not proof of remaining gain. Revalidate existing switches separately; no duplicate fix candidate.
• r8 C16 complete shared/routed merger: not costed as if the shared expert's entire 1.71 ms were exposed. Formats, scalar gate semantics and stream overlap differ. R9-gpu-5 is the bounded alternative.
• lm_head/MTP Q5_K vectorisation: the supplied draft head rate of about 867 GB/s is already at the chosen floor. No credible large exact gain is demonstrated by this source pass.
• Norm/quantise blanket fusion: 165 launches at about 1.8 us do not establish a 0.30 ms recoverable gap, especially with existing producer fusion and private stream buffers. QFUSE's recently fixed byte contract must be the oracle for any later work.
• Attention V loads four-at-a-time: already implemented at src/kernels/cuda/qsa_decode_attn.cu:678-688. R9-gpu-11 changes score-cell reuse, not those V loads.
• Expert pruning, defer-CPU and resident-biased routing: [Engine] verify: STRATA_VERIFY_ECAP - cap the verify window's expert set (relaxed MTP, opt-in) Niko1221/Strata#3545, [Engine] verify: STRATA_VERIFY_DEFER_CPU=1, a layer's light CPU expert rows join at the next layer's FFN half (heavy rows on time, 0.1 by default; opt-in) Niko1221/Strata#2690 and route-prior/resident variants change the model's output. Excluded regardless of reported speed.
• Larger speculative width, split tensor-core attention, alternative draft quantisation/vocabulary and adaptive window policies: closed locally or assigned to other workstreams. None is used to price the exact candidates here.

Prefill / serve (prefill-serve audit)

Seeded candidates (mechanism verified in source, value unmeasured)

C-PS-1. The elastic K/V growth point runs a full device sync at prompt start

• Evidence: src/program/generate.cpp:17789-17793 - kvg_ensure(n_prompt + max_new + 64, [] { cudaDeviceSynchronize(); })
  on first use per run.
• Mechanism: precompute the growth target from the request as soon as the token count is known (it is
  known before prefill.run at 17795) and grow before the first kernel launch, overlapping the free-
  list work with model load/prefill start; keep the sync fallback for the late path.
• Gain shape: one full device drain at the worst moment (right before the first prefill kernel); TTFT
  fixture at 17 tokens exists in the brief - measure directly.
• Gate: 17-token TTFT fixture + fresh-read tiers 1K/2K/4K.

C-PS-2. PP telemetry as a first-class gate is present; chunk-boundary checkpoint waits are the target

• Evidence: src/program/generate.cpp:14938-14951 - the batched-prompt path prints per-batch PP rate
  and ver.wait_commit after each batch; commit/wait_commit pair a kernel commit with a CPU-side
  wait per batch.
• Mechanism: audit wait_commit placement in src/core/verify.cpp/conversation path: if the CPU blocks
  on the commit event before issuing the next batch's kernels, the batch boundary pays a launch-latency
  gap; double-buffer the commit (issue batch k+1 kernels before waiting for batch k's commit event)
  where the checkpoint contract allows it.
• Gate: PP lines at 14947 on 32K/128K prompts, paired.

C-PS-3. Effort-tail rendering re-renders the template three times per non-default effort request

• Evidence: serve/server.py:4905-4920 (render_prompt): on effort_end with a non-default effort it
  renders history, full and mine (three full template renders) to splice the effort turn.
• Mechanism: cache the history/prefix rendering per conversation generation (the router's prefix
  keeper already tracks session prefixes), or splice with a jinja block-level render of only the tail;
  the Changing reasoning_effort re-reads the whole prompt: the effort text is the first thing the template renders Niko1221/Strata#458 cache contract (same prompt start) must be preserved byte-exact.
• Gain shape: template render cost is linear in history length (jinja python): at 32K rows this is
  CPU-side per request; measure time.perf_counter around render_prompt in a host-only probe
  (tools/ has template benches).
• Gate: TTFT at 4K/32K/128K with effort=low vs xhigh; byte-identical prompt check (the splice test
  full.startswith(mine) already asserts structure).

C-PS-4. Serving hot path: shared-key validation and body handling per request

• Evidence: serve/server.py:8273 (SHARED_KEYS), 4656-4672 (effort default injection), 7841
  (responses-API template kwargs). Python-side per-request work includes json parse, key validation,
  chat template kwargs build with kwargs dict copying (4911 rebuilds dicts twice).
• Mechanism: profile (host CPU, no GPU) the serve request path for a fixed tiny prompt: measure parse +
  validation + render + first-token latency split; move per-request dict rebuilds into a single build;
  keep validation semantics (the ValueError texts at 8288-8295 are the contract).
• Gate: TTFT fixture + a pure-Python microbenchmark committed as a test, not a production claim.

Queue finished (2026-11-02), verdicts

1. serve/server.py: the render/effort region produced C-PS-3, implemented as fork PR feat(serve): exact cache for the #458 history render (opt-in STRATA_EFFORT_RENDER_CACHE) #7
   (r9/pr1-render-cache, base r9-base = v0.2.9; serve.test_server 360/360,
   serve.test_prompt_encoder + serve.test_request_accounting 17/17 on the mock engine).
   The responses API and shared-key region is C-PS-4 as seeded, no additional candidate found.
   SSE framing is stock generator yields over HTTP/1.0 (serve/server.py:6756, 6014); no win found.
2. conversation_snapshot.cpp park path (690-722): already two pinned pieces pipelined - the next
   piece's DMA runs while the previous piece is memcpy'ed to its segment, blocking-sync events,
   per-piece kParkPiece. No candidate; DROPPED. conversation_*.cpp has no memcpy/sync hot spots
   beyond that park path and startup; the prefix keeper interface is engine-side (strata_prefix,
   serve/server.py:1487-1491) and already token-counted.
3. generate.cpp on_chunk (17773-17781): per-chunk std::vector<int32_t> nxt(T) allocation, a few
   KB per chunk against multi-second chunks; DROPPED as micro. kvg_ensure's full-device-sync
   callback at prompt start (17786-17788) is C-PS-1 as seeded, mechanism confirmed at the call site.
   The lent-slot refill timing print and prefill_batched_ms (17820-17830) confirm C-PS-2's
   telemetry hooks exist; the chunk-boundary syncs are the measurement target, unchanged.
4. Tokenizer: the serve-side fast path is tools/strata_tokenizer.py with prefix reuse already
   counted in serve telemetry (reused tokens, serve/server.py:1036, 1375, 1999-2002): clients
   and the prefix keeper take cached tokens out of the work. No new candidate; DROPPED.
5. Dedupe: r7 prefill.md and serving-host.md cover chunk-size sweeps, the host-side request
   accounting and the tap feed. C-PS-1..4 are disjoint from those (kvg growth sync, chunk-boundary
   gate, render cache, request hot path); no overlap to merge.

Host / speculative decode (host-spec audit)

Seeded candidates (mechanism verified in source, value unmeasured)

C-HS-1. Start the post-prefill lent-slot transfer before the prefill tail, on a dedicated stream

• Evidence: src/program/generate.cpp:17799-17824 - after prefill.run returns, every lent slot is
  refilled serially with cudaMemcpyAsync on stream 0 (src/core/expert_cache.cpp:821-837) and one
  sync_queued wait. Pageable source (blob_stable mmap), so each 4 MB copy runs through driver
  staging chunks that serialise on stream 0; the decode window cannot start until sync_queued clears.
• Mechanism: issue the queued fills from inside the last prefill chunk (rows already committed, the
  blob bytes stable per the blob_stable contract at 2446-2447) on a non-default, non-blocking
  stream, and gate only the first window that touches each slot with a per-slot event.
• Gain shape: hides up to slots x 4 MB / ~10 GB/s of PCIe time behind the prefill tail; on a 24-slot
  loan that is roughly 96 MB ~ 9.6 ms of transfer moved off the critical path (bandwidth arithmetic,
  unmeasured).
• A/B gate: existing log line at 17822 prints refill ms; pair PP/decode windows with STRATA_REFILL_BLOCKING unset.
• Risks: stream-ordered visibility of slot contents for kernels launched on stream 0 (event gates required);
  HIP parity (STRATA_USE_HIP staging path 801-812) stays stock.

C-HS-2. Coalesce the three-piece queued refill into one contiguous destination copy

• Evidence: src/core/expert_cache.cpp:839-860 - fill_slot_queued_parts issues one cudaMemcpyAsync
  per blob part (up to 3, BlobPart parts[3], src/program/generate.cpp:14967-14972) from mmap pieces.
• Mechanism: keep per-part copies only when the destination offsets are non-contiguous; when the slot
  layout is contiguous (the default slot * blob layout, device_slot 754-764), copy the source range
  once. Add a pure-host test that the part map covers the slot with no gaps.
• Gain shape: removes two per-slot API/staging handoffs per refill; with pageable sources each extra
  call costs driver chunking - small per slot, linear in loan size.
• Gate: CPU index-map test + refill-ms log comparison.

C-HS-3. Replace the blocking refill fallback with the pinned ring the engine already owns

• Evidence: src/core/expert_cache.cpp:789-819 (cudaMemcpy sync per slot; HIP stages through one
  4 MB buffer) and src/program/generate.cpp:454 (refill_blocking() env selector) with call sites
  7331, 7474, 14976, 17805.
• Mechanism: the blocking selector exists for correctness cases (mid-window refills at 7331/7474).
  Route those through a pinned double-buffer ring (the pinned.cu budget machinery) with a slot
  watermark: copy into pinned staging on the producer thread, single async H2D per slot, event-gated
  use. Keep STRATA_REFILL_BLOCKING=1 semantics (completion before reuse) by event-blocking on use,
  not on issue.
• Gain shape: a synchronous 4 MB pageable copy blocks ~1 ms+ of CPU per slot inside the decode window;
  ring pipelining overlaps it with compute (bounded, unmeasured; premise gate = time the existing
  blocking refills with tr timestamps at 17801/17822 with STRATA_REFILL_BLOCKING=1).
• Risk: pinned budget eviction interplay with pin_budget tests; residency/ownership rules (14955-14958
  "refilled into the wrong cache is silent") demand a per-cache ownership assertion test.

C-HS-4. Draft upload placement in the MTP round

• Evidence: src/core/mtp.cpp:1294 - comment: the first launch's implicit upload blocked behind a
  device-side spin in verify.cpp; code added an explicit upload. Audit the per-round draft token/
  state uploads for the same implicit-copy pattern at steady state (not just first launch), and the
  progress_at("mtp draft: waiting for the GPU") stall (mtp.cpp:1724) for window accounting.
• Action: map every D2H/H2D in the draft/verify round; look for a round-end sync the CPU pays while
  the GPU idles. Verify against Engine hangs at layer 48 during decode when VRAM headroom drops below ~500 MiB Niko1221/Strata#2065 note before proposing.

Queue finished (2026-11-02), verdicts

1. expert_source.cpp: 149-prefetch is a complete machinery (io_pf_ admission with confidence bar and
   byte budget, jobs/demand/skips/dropped/late counters at 1069-1145, STRATA_IO_PF_BUDGET_MB). The
   host-side serial pread fallback noted at session.cpp:642-652 is the mechanism C-HS-3 names; no
   new candidate beyond C-HS-1..3. No overlap with r7's CPU quantize/doorbell trio.
2. peer_experts.cpp and remote_experts.cpp: the call shape is upload x/meta, run, download output,
   stream sync per call (peer_experts.cpp:288-313, 381-385; remote_experts.cpp:327-381). The per-call
   sync is structural: the output lands host-side for merge. Overlapping a peer call with local
   compute is a real candidate only on multi-device boxes; the 3090 host runs neither lane. Verdict:
   keep as a documented multi-device candidate, not a v0.2.9 3090 gate.
3. mtp.cpp: the draft round is fully graph-captured; the seven host inputs land through one
   multi-mapped copy and the draft ids upload (1294) is inside the graph, so C-HS-4's placement
   question is answered: DROPPED (no host round trip in the round). record_top2 (564) is the
   STRATA_MTP_TOP2 diagnostic, off by default; 600-633 is setup. verify.cpp: the window token
   upload is async (3375), the readbacks at 2521/2905 are profiling-gated, 1191 is a KV-hit slot
   offset copy. Accept path: engineered; no new candidate beyond C-HS-1..2.
4. pinned.cu + pin_budget_test.cpp: the pinned shared arena, section and budget API exists
   (pinned.cu:172-342 surface, pin_budget_test.cpp), so C-HS-3 can reuse it: no new plumbing
   needed, confirmed.
5. Dedupe: r7 expert-tier.md covers the CPU GEMV quantize lanes, the doorbell and expert-tier
   residency; r7 CANDIDATES-8.md covers round-8 items now in the tree. C-HS-1..3 (refill
   transfer scheduling) and C-PS-1..2 (prefill growth sync, chunk-boundary gate) are disjoint
   from both. C-HS-4 DROPPED per verdict 3.