perf(cuda): fused decode/verify/MTP kernels and graph launch reductions on 0.1.39 by stuchapin909 · Pull Request #783 · Niko1221/Strata

stuchapin909 · 2026-10-04T14:51:01Z

Summary

This is the follow-up to #646 rebased cleanly onto 0.1.39 (6f32ec0). Since #646 merged only the initial commit on that branch (884c14e), this PR brings in the remaining verify-window, MTP, and batch-decode kernel and CUDA-graph fusions split into 4 self-contained commits by subsystem.

Measured apples-to-apples on 2x RTX 3090 24 GB (IQ2_XS, --layer-split 25/23, 100% of all 24,576 experts resident in VRAM on both baseline and PR):

Metric (2x RTX 3090, `IQ2_XS`, 100% VRAM Resident)	Pre-#646 (100% Resident)	`0.1.39` (#646 `884c14e`)	`0.1.39` + This PR (#783)	Pure Kernel/Graph Delta
Verify Window Latency (`lru_code`, `T = 5..8`)	`15.92 ms/win`	`~15.10 ms/win`	`13.07 ms/win` (`13.07–15.02` across suite)	-13.4% vs `0.1.39` (-17.9% vs pre-#646)
3-Prompt Benchmark (`lru_code` / `skiplist` / `prose`)	`164.1 / 156.4 / 151.0 tok/s`	`~170.2 / 162.0 / 156.5 tok/s`	`186.6 / 177.1 / 170.4 tok/s`	+9% to +10% vs `0.1.39` (+13.7% vs pre-#646)
Live Coding Decode (53k–62k context, 4,790 tokens)	—	—	`210.7 – 217.5 tok/s`	Sustained high-acceptance coding session

(Note: Compared to 0.1.38's default 26/22 split where GPU 0 held 491/512 experts at 18.31 ms/win / 149.2 tok/s, combining 0.1.39's 25/23 all-resident split with #646 + #783 reduces verify window time from 18.31 ms/win to 13.07 ms/win.)

All changes are scoped strictly to include/strata/{core,kernels} and src/{core,kernels/cuda} (32 files), preserve 0-ULP bitwise parity across all 25 GPU parity tests, and maintain full HIP (gfx906 wave64 + RDNA wave32) compatibility.

Commits & What's Included

1. Multi-token GR / RMSNorm / post-op kernels and exact-`T` specializations (`fused_gr.cu`, `gr.cu`, `native_gr_norm.cu`, `native_gr_postops.cu`, `native_bf16.cu`, `quantize_act.cu`)

fused_gr.cu:
- Adds pre-unpacked BF16 8-wide dot helpers (Bf16x8, unpack8, dot8u, dot8u_ptr) so weight tiles are unpacked once across tokens.
- Templates gr_up_multi_kernel<MAX_T, EXACT_T> and gr_down_staged_kernel<MAX_T, EXACT_T> to eliminate dynamic loop bounds for T = 1..8 (STRATA_GR_DOWN_MAX4=0 disables the MAX_T=4 specialization for ct <= 4).
- Preserves f16_from_f32, STRATA_DP4A, and STRATA_LDG from f166564 for HIP/ROCm portability.
gr.cu (gr_read_multi, gr_write_multi): Multi-token hyper-connection read and write across n_tok tokens in one launch.
native_gr_norm.cu (native_gr_rms_norm_weighted_multi): Multi-token weighted RMSNorm across n_tok rows.
native_gr_postops.cu (native_gr_pre_gated_multi, native_gr_post_multi): Multi-token pre-gated mixing and residual post-ops.
native_bf16.cu (bf16_f32_mmvf_multi_kernel<BLOCK_SIZE, NT, EXACT_T>): Compile-time token-count dispatch in bf16_gemv_fp32_mmvf_multi.
quantize_act.cu: Warp-shuffle reduction in quantize_q8_0_scaled_kernel.

2. IQ/MMVQ sub-warp & 2-row kernels and MoE / shared-expert fusions (`iq_kernels.cu`, `native_mmvq.cu`, `native_moe.cu`, `native_router.cu`, `s2_expert_grouped.cu`, `shared_expert.cu`, `elementwise.cu`)

iq_kernels.cu:
- Adds 16-lane sub-warp row_dot_80_sub16<TY, NC, EXACT_N, STAGE_GRID> and templates mmvq_multi_kernel<TY, NC, STAGE_GRID, EXACT_N> (launch_mmvq_multi_nc).
- Adds fused gate+up handling (gu_split) in native_gu_multi_kernel alongside 0.1.39's down_iq4nl_multi_kernel.
native_mmvq.cu: Adds SmallTraits<Q40Block, 4>, SmallTraits<Q50Block, 4>, SmallTraits<Q80Block, 8>, and SmallTraits<IQ4NLBlock, 4> 2-row-per-warp (ROWS = 2) multi-column MMVQ specialization in launch_multi_n when !multi_exact.
native_moe.cu (combine_k10_vec4): float4-vectorized $k=10$ expert output combination in native_moe_combine and native_moe_combine_multi.
native_router.cu (route_multi, native_router_top10_multi): Multi-token top-10 router kernel across n_tok tokens in a single launch.
s2_expert_grouped.cu (swiglu_quantize_q8_0_scaled_kernel): Fuses SwiGLU and scaled Q8_0 quantization in moe_grouped_s2 while writing back gate_up[idx] so s2_expert_grouped_parity remains 100% bitwise identical (STRATA_OLD_GROUPED=1 selects the previous path).
shared_expert.cu (sigmoid_scale_rows_vec4_kernel, launch_sigmoid_scale_rows): float4-vectorized sigmoid gating in shared_expert and shared_expert_multi.
elementwise.cu (copy_rows_from_mapped_kernel<ZERO_HITS>): Allows copy_rows_from_mapped (zero_hits = false) to skip redundant zero-fill when all experts are resident.

3. Batched KV append, QSA indexer/RoPE fusion, and parallel `resident_plan_kernel` (`kv_q4.cu`, `kv_q8.cu`, `native_qsa_indexer.cu`, `native_rope.cu`, `verify_kernels.cu`, `fused_gdn.cu`, `qsa_select.cu`, `kv_stream.cu`)

kv_q4.cu (kv_append_q4_steps) & kv_q8.cu (kv_append_q8_steps): Append K and V across n_tok tokens in a single 3D grid launch.
native_qsa_indexer.cu (native_qsa_indexer_append_steps): Multi-token QSA indexer append in a single launch.
native_rope.cu (norm_rope_kernel<TAB>, native_qsa_rms_norm_rope): Fuses per-head QSA RMSNorm and RoPE application.
verify_kernels.cu:
- resident_plan_kernel<ALL_RESIDENT>: Keeps kResidentPlanMax = 128 (f945515) and replaces thread-0's serial $O(n^2)$ bubble sort with parallel shared-memory grouping and a single packed (my_cnt << 16) | (is_first ? 1 : 0) prefix scan using __shfl_up_sync(0xffffffffu, pref, d). Using width-32 __shfl_up_sync without __ballot_sync keeps the scan identical on both 32-lane warps (CUDA / RDNA) and 64-lane wavefronts (gfx906).
- Templates gdn_ab_multi_kernel<MAX_T, EXACT_T> and gdn_step_norm_multi_kernel<ALL_OUT>.

4. Wire multi-token kernels into `Verifier` and `MtpDrafter` CUDA graphs (`verify.cpp`, `mtp.cpp`, `verify.hpp`)

verify.cpp:
- Replaces per-token loops in Verifier::record_window with native_gr_rms_norm_weighted_multi, native_gr_pre_gated_multi, native_gr_post_multi, native_qsa_rms_norm_rope, kv_append_q8_steps, kv_append_q4_steps, native_qsa_indexer_append_steps, and native_router_top10_multi.
- Skips redundant copy_rows_from_mapped zero-fill on the all-resident direct_parts path.
- Note on environment flags: STRATA_FUSE_HEAD_GR=1 (fused LM-head GR read) and STRATA_VERIFY_PRELAUNCH=1 (cross-stage prelaunch) are both opt-in (default off) so default --layer-split commit and stage ordering remain identical to 0.1.39 (deee447).
mtp.cpp:
- Dispatches gr_read_multi, gr_write_multi, native_gr_rms_norm_weighted_multi, native_qsa_rms_norm_rope, kv_append_q8_steps, kv_append_q4_steps, native_qsa_indexer_append_steps, and native_router_top10_multi inside MtpDrafter::record_forward.

Verification

GPU Parity Suite (CTest — 25/25 Passed)

Built and ran the GPU parity test suite on Windows 11 / CUDA 13.2 (2x NVIDIA GeForce RTX 3090 24 GB):

mmvq_multi_parity — Passed
iq_multi_parity — Passed
iq_parity_fixtures / iq_parity — Passed
native_grouped_parity — Passed
shared_expert_parity — Passed
s2_expert_grouped_parity — Passed (bitwise identical across all 22 checks)
s2_gemv_parity / s2_gemv_q8_parity — Passed
gr_parity — Passed
gdn_parity — Passed
qsa_parity / qsa_topk_parity — Passed
rope_parity — Passed
quantize_act_parity — Passed
route_window_parity / router_top10_parity — Passed
elementwise_parity — Passed
bf16_gemv_parity — Passed
sampler_parity — Passed
cvec_parity — Passed
kv_q8_parity / kv_q4_parity / kv_hybrid_parity / kv_stream_parity — Passed

GonzMartinez · 2026-10-04T17:09:19Z

2x RTX 3090, Swift 1.5 IQ3_XXS, partially resident, layer split: +6.5-7.8% decode, but greedy output not identical to 0.1.39

Another data point, for the configuration nobody has posted yet: the IQ3_XXS pack on a 2-GPU layer split with only part of the experts in VRAM.

Setup. 2x RTX 3090 24 GB (sm_86, PCIe x8, 13.4 GB/s measured, no NVLink), Ryzen 9 9950X3D, 156 GB RAM, Ubuntu 24.04 (container), driver 595.71.05, CUDA 12.0. Swift 1.5 IQ3_XXS (GSQ-RCO). Both arms are source builds with identical CMake flags (Release, CMAKE_CUDA_ARCHITECTURES=86): main at 6f32ec0 vs this PR's head e1f292a in a separate worktree; same serve/server.py.

Engine args (setup's own): --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 --mtp <rt> --max-context 262144 --kv int8 --kv-resident 32768 --vision --vram-reserve-mib 700 --remote-expert-opt, layer split auto (K=26), plus "parallel": 2 and --conversation-cache-mib 32768 --conversation-cache-slots 8 --prompt-cache 16 for the speed runs. Expert cache 33.0 GiB / 20,091 slots with 2 batch slots (82% resident), 35.2 GiB / 21,394 slots without; decode hit rate ~99%. So all_resident_ is false here.

Speed

8 real prompts (code, explanations, prose), 1,024 output tokens each, one request at a time, thinking on. Engine-side reciprocal mean TPOT. Arms alternated stock, PR, stock, PR with a fresh engine start each.

	stock run 1	stock run 2	PR run 1	PR run 2	change
default sampling (tok/s)	114.7	117.4	123.8	123.5	+6.5%
temperature 0 (tok/s)	113.4	116.0	122.9	124.7	+7.8%
draft acceptance, default / T=0	0.715 / 0.710	0.739 / 0.728	0.709 / 0.722	0.709 / 0.717	same

Both PR runs beat both stock runs in each mode, with the same acceptance, so the windows really are cheaper. Prompt reading was not measured.

Exactness: not identical on this configuration

Gate: 5 greedy prompts (chat, code, math, prose, and a 21,348-token prompt), 256 tokens each, through the HTTP server, thinking text included in the comparison. Exactness settings on both arms: STRATA_IQ_MT_MIN=1 --pcie-frac 0 --adapt-every 1000000, no batch slots, a fresh engine per arm.

Stock vs stock (two fresh starts): 5 of 5 IDENTICAL, drafts accepted identical.
PR vs stock: only math identical; the others differ:

prompt	first difference	drafts accepted stock / PR
chat	char 220	56 / 54
code	char 272	144 / 154
math	identical	173 / 173
prose	char 468	134 / 123
long (21K)	char 160	147 / 146

For example (chat): stock ...focus on code not author. Need concise. Ensure brief. / PR ...focus on code not author / constructive comments. Need concise.; (code): stock P[n]Y[n]M[n]DT[n]H[n]M[n]S. Need parse into seconds / PR P[n]Y[n]M[n]D[T[n]H[n]M[n]S], also weeks?.

PR with STRATA_GR_DOWN_MAX4=0 (a second fresh start): the same output hashes as the PR's default on all 5, so the PR build is deterministic and the divergence is not in that path.

The texts stay coherent, and the first difference is never at token 0, so this looks like a changed accumulation order in a path that only runs on a layer split or with partial residency (your parity runs were IQ2_XS, fully resident). If you add kill switches for the four commits, or point me at the most likely kernels, I can bisect on this box; the gate takes about 4 minutes per arm.

…lizations - fused_gr.cu: add pre-unpacked BF16 8-wide dot helpers (Bf16x8, unpack8, dot8u, dot8u_ptr) and template gr_up_multi_kernel / gr_down_staged_kernel on <MAX_T, EXACT_T> to eliminate dynamic loop bounds for T=1..8 (STRATA_GR_DOWN_MAX4=0 opts out of the MAX_T=4 cap for ct <= 4; STRATA_NO_MULTI_GR=1 falls back to 0.1.39 dispatch). Preserves f16_from_f32 / STRATA_DP4A / STRATA_LDG portability from f166564. - gr.cu / gr.hpp: add gr_read_multi and gr_write_multi for multi-token hyper-connection read/write in a single launch (STRATA_NO_MULTI_GR=1 selects per-token gr_read / gr_write). - native_gr_norm.cu / .hpp: add native_gr_rms_norm_weighted_multi. - native_gr_postops.cu / .hpp: add native_gr_pre_gated_multi and native_gr_post_multi. - native_bf16.cu: template bf16_f32_mmvf_multi_kernel on <BLOCK_SIZE, NT, EXACT_T> for exact token counts in bf16_gemv_fp32_mmvf_multi. - quantize_act.cu: use warp shuffle reduction in quantize_q8_0_scaled_kernel.

…sions - iq_kernels.cu: add 16-lane sub-warp row_dot_80_sub16<TY, NC, EXACT_N, STAGE_GRID>, template mmvq_multi_kernel on <TY, NC, STAGE_GRID, EXACT_N> with launch_mmvq_multi_nc dispatch, and add 16-lane / 8-lane sub-warp grouped expert gate+up and down kernels in native_gu_multi_kernel / native_down_multi_kernel (STRATA_NO_SUB16_GU=1 falls back to 32-lane row_dot_multi). - native_mmvq.cu: add SmallTraits<Q40Block, 4>, SmallTraits<Q50Block, 4>, SmallTraits<Q80Block, 8>, and SmallTraits<IQ4NLBlock, 4> 2-row-per-warp (ROWS = 2) multi-column MMVQ dispatch in launch_multi_n when !multi_exact. - native_moe.cu: add float4-vectorized combine_k10_vec4 for k=10 expert combination in native_moe_combine and native_moe_combine_multi. - native_router.cu: add route_multi kernel for multi-token top-10 routing in native_router_top10_multi. - s2_expert_grouped.cu: add swiglu_quantize_q8_0_scaled_kernel to fuse SwiGLU and scaled Q8_0 quantization in moe_grouped_s2 while writing back gate_up[idx] so s2_expert_grouped_parity remains bitwise identical (STRATA_OLD_GROUPED=1 selects the previous path). - shared_expert.cu: add float4-vectorized sigmoid_scale_rows_vec4_kernel and launch_sigmoid_scale_rows. - elementwise.cu / .hpp: template copy_rows_from_mapped_kernel<ZERO_HITS>.

…resident_plan - kv_q4.cu / .hpp & kv_q8.cu / .hpp: add kv_append_q4_steps and kv_append_q8_steps to append K and V across multiple tokens in a single 3D grid launch. - native_qsa_indexer.cu / .hpp: add native_qsa_indexer_append_steps for multi-token indexer append. - native_rope.cu / .hpp & rope_parity.cpp: add norm_rope_kernel<TAB> and native_qsa_rms_norm_rope to fuse per-head RMSNorm and RoPE application with exact FMA evaluation order matching norm<256> + apply (0 ULP verified by Check 6 in rope_parity). - verify_kernels.cu: - Template resident_plan_kernel<ALL_RESIDENT> (kResidentPlanMax = 128, preserving f945515) and replace thread-0's serial O(n^2) bubble sort with parallel shared-memory grouping and a single packed (my_cnt << 16) | (is_first ? 1 : 0) prefix scan over __shfl_up_sync (width 32, no __ballot_sync, wave32 and gfx906 wave64 safe). - Template gdn_ab_multi_kernel<MAX_T, EXACT_T> and gdn_step_norm_multi_kernel<ALL_OUT>. - fused_gdn.cu, qsa_select.cu, kv_stream.cu: register/loop staging refinements.

…A graphs - verify.cpp / verify.hpp: - Dispatch gr_read_multi, gr_write_multi, native_qsa_rms_norm_rope, kv_append_q8_steps, kv_append_q4_steps, and native_qsa_indexer_append_steps across verify/batch tokens in Verifier::record_window and Verifier::capture_commit. - Add runtime env kill switches STRATA_NO_MULTI_GR=1, STRATA_NO_NORM_ROPE=1, and STRATA_NO_BATCH_KV_STEP=1 for isolated A/B bisection without recompiling. - Preserve 0.1.39's hit_out + moe_hit_add on the !all_resident_ path and synchronous --layer-split commit/stage ordering by default (deee447). - mtp.cpp: - Dispatch gr_read_multi, gr_write_multi, native_qsa_rms_norm_rope, kv_append_q8_steps, and kv_append_q4_steps inside MtpDrafter::record_forward (honoring STRATA_NO_MULTI_GR, STRATA_NO_NORM_ROPE, and STRATA_NO_BATCH_KV_STEP).

stuchapin909 · 2026-10-04T17:58:33Z

Thanks for running the 4-minute gate and sharing the detailed output, @GonzMartinez — that caught a subtle --use_fast_math SASS difference that wasn't covered by the existing rope_parity CTest!

Root Cause of the Greedy Exactness Divergence (`chat` / `code` / `prose` / `long`)

norm_rope_kernel in src/kernels/cuda/native_rope.cu (Commit 3):
Both native_qsa.cu (norm<256>) and native_rope.cu (apply) are compiled with --use_fast_math.
- In 0.1.39, norm<256> computes output[col] = (scale * input[col]) * gamma[col] using FMUL.FTZ (mul.ftz.f32), stores to global memory, and then apply loads a, b and lowers a * c - b * s / a * s + b * c to bs = b * s; bc = b * c; fmaf(a, c, -bs); fmaf(a, s, bc).
- In our fused norm_rope_kernel, we had written __fmul_rn(__fmul_rn(scale, x_tid), gamma[pair]) — which emitted mul.rn.f32 (disabling FTZ on the intermediate product) and allowed ptxas to contract across the norm/RoPE boundary. That caused a 1-ULP difference on a fraction of rotated Q/K elements, which accumulated across layers until flipping a near-tied greedy argmax around char 160–468.
- Fix: norm_rope_kernel now uses (scale * x_tid) * gamma[pair] with PTX register barriers (mov.f32) on a, b and bs, bc before fmaf(a, c, -bs) / fmaf(a, s, bc), producing the exact same SASS instruction sequence (FMUL.FTZ -> FFMA.FTZ) as norm<256> + apply.
- New CTest coverage: Added Check 6 to rope_parity comparing native_qsa_rms_norm_rope (both strided 2 * head_dim Q/gate split and contiguous in-place head_dim, across hd=128 and hd=256, unscaled and YaRN) against native_qsa_rms_norm_weighted + native_rope_apply:
```
native_qsa_rms_norm_rope hd=128 nr=64      bit-identical (0 of 49152 elements differ)
native_qsa_rms_norm_rope hd=256 nr=64      bit-identical (0 of 98304 elements differ)
```
!all_resident_ expert buffer path in src/core/verify.cpp (Commit 4):
Restored 0.1.39's exact hit_out + copy_rows_from_mapped (zero_hits = true) + moe_hit_add path when !all_resident_ (such as 82% resident on 2x RTX 3090 24 GB).

Runtime Kill Switches Added (No Recompile Needed)

All four env-var switches you requested are now wired in:

STRATA_NO_MULTI_GR=1 — disables gr_read_multi / gr_write_multi and exact-T fused_gr template specializations (Commit 1)
STRATA_NO_SUB16_GU=1 — disables the 16-lane / 8-lane sub-warp grouped expert gate/up/down kernels and falls back to 32-lane row_dot_multi (Commit 2)
STRATA_NO_NORM_ROPE=1 — disables fused native_qsa_rms_norm_rope in verify.cpp and mtp.cpp and falls back to cudaMemcpy2DAsync + native_qsa_rms_norm_weighted + native_rope_apply (Commit 3)
STRATA_NO_BATCH_KV_STEP=1 — disables kv_append_q8_steps / kv_append_q4_steps / native_qsa_indexer_append_steps and falls back to per-token loops (Commit 3/4)

Force-pushed the updated 4-commit stack (d18d8b7..4327807) — ready for a re-run whenever you have a moment!

…P kernels and graph launch reductions

…#783, Niko1221#742, Niko1221#743, Niko1221#731, Niko1221#764, k1, k2b, the sm_75 BF16 GEMM path), with the measured numbers

q8atnight · 2026-10-04T19:44:19Z

Reply to PR #783 (stuchapin909) — posted via gh as q8atnight

Another 2-card row, for the config class GonzMartinez flagged (layer split + partial
residency): 2× RTX 5060 Ti 16 GB (sm_120, both at x8 / 16 GT/s), 32 GB RAM, Linux, native build
(CUDA 13.4, CMAKE_CUDA_ARCHITECTURES=120). IQ2_XS, layer split auto (0–23 / 24–47), 262,144
ctx, --kv int8 --kv-resident 32768, --mmap-experts, --prefill 6144, MTP spec 4, one stream.
Expert cache 7,176 VRAM slots — all_resident_ false, decode hit 96.8–99.4 %. Baseline = 0.1.39
(+ our #699 carry); PR = same base + your 4327807 (the only conflict was our 3-hunk #699 delta
in mtp.cpp, re-applied cleanly on your version).

Decode: 600-token story, temp 0, 6 runs per arm, page cache dropped before each arm's
start, arms alternated, first (cold) run excluded: baseline 78.0 and 79.9 tok/s, your head
83.0 and 80.8 → mean +3.7 %, which is not significant at this sample size
(10 samples per side, pooled std ≈ 5.8). The 3-run rounds looked like +5–8 %; more samples ate
the ordering effect. So: real but small, and only visible above the noise with many runs.
Prefill flat: 32K 1,919/1,887 (base) vs 1,923/1,876 (PR); 128K ~2,070–2,150 both; expert
slots (7,176) and the prompt-path borrow (2,846 MiB) identical; ready time ~39 s both.
Divergence spot (not a full gate): two temp-0 prompts × 2 repeats per binary — a step-by-step
arithmetic answer byte-identical within AND across the two binaries; a 120-word prose prompt was
cut by thinking before content in all four runs (no signal either way). No divergence found at
IQ2_XS with 7,176 resident slots.

Conclusion from this rig: neutral-to-mildly-positive speed (decode +3–4 % at best, prefill 0),
parity spot-clean in the partial-residency class. On a 32 GB RAM box where decode is already
70–85 tok/s and the page cache is the tight resource, that's not enough for us to carry an
unmerged kernel PR — but the row is yours for the table.

T-Crypt · 2026-10-04T20:39:04Z

Single RTX 4090 (sm_89, Linux, 32 GB RAM), Flash-Next GSQ-RCO IQ2_XS, --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 --mtp --max-context 262144 --kv q4_0 --resident-experts. Base: v0.1.39 (6f32ec0) + #700; candidate: the same + this PR. Fresh process per build, alternated base/PR twice, three prompt shapes, 3 runs each (mean decode tok/s):

shape	base r1	base r2	PR r1	PR r2	mean delta
code	138.5	138.5	141.9	140.3	+1.9%
reason	155.3	151.0	158.6	159.8	+3.9%
agent	134.1	134.3	141.8	139.1	+4.7%

PR faster in both rounds on all three shapes; draft acceptance and expert-tier hits look unchanged. Parity: all *_parity targets pass (qsa_topk_parity --selftest, iq_parity with generated fixtures); native_expert_parity --synthetic q4_K q5_1 reports 2 failures on stock v0.1.39 as well, so not from this PR. Short contexts only, I didn't measure decode at depth. We're now running this build day to day.

T-Crypt · 2026-10-04T21:42:50Z

Follow-up from the single-4090 report above: decode at depth, same rig and config (262K, q4_0, resident experts). Haystack prompt with a recall question, then the same prompt plus a 300-token summary for decode. Fresh process per build, two rounds:

	119K decode	238K decode	119K / 238K prompt read
v0.1.39	106.2 / 106.6 t/s	126.7 / 127.3 t/s	4,046 / 3,898 t/s
+ this PR	121.8 / 121.5 t/s	138.9 / 139.3 t/s	4,036 / 3,893 t/s

That's +14% decode at 119K and +9.5% at 238K, bigger than the +2 to +5% I saw on short prompts. Prompt reading is unchanged, and recall was correct in every run.

GonzMartinez · 2026-10-05T00:28:11Z

Re-run on `4327807`: still not identical on IQ3_XXS + 2-GPU layer split, and the four switches don't cover it

Thanks for the quick turnaround and the switches. Same box, same gate and stock reference as my report above; only the PR head changed (rebuilt from 4327807, same CMake flags; the running engine was confirmed to be that binary and to carry the switches in its environment).

4327807 vs stock 0.1.39 (6f32ec0): the same result as e1f292a, byte for byte. Every arm below produced the same five output hashes as e1f292a did:

arm	chat	code	math	prose	long (21K)
`4327807` default	char 220	char 272	identical	char 468	char 160
all four off (`STRATA_NO_MULTI_GR=1 STRATA_NO_SUB16_GU=1 STRATA_NO_NORM_ROPE=1 STRATA_NO_BATCH_KV_STEP=1`)	char 220	char 272	identical	char 468	char 160
`STRATA_NO_MULTI_GR=1` only	same	same	identical	same	same
`STRATA_NO_SUB16_GU=1` only	same	same	identical	same	same
`STRATA_NO_NORM_ROPE=1` only	same	same	identical	same	same
`STRATA_NO_BATCH_KV_STEP=1` only	same	same	identical	same	same

Stock reproduced itself in my earlier runs (5/5 identical across fresh starts), and all six PR runs here gave the same hashes. So on this configuration the divergence is in a part of the stack that none of the four switches turns off, and the norm/RoPE fix doesn't reach it. Candidates by elimination, from the commit messages: the 2-row MMVQ kernels and MoE/shared-expert fusions in commit 2 beyond the gate/up/down sub-warp path, the parallel resident_plan in commit 3, or the graph wiring in commit 4. A switch for any of those would let me narrow it down; each gate run takes about a minute here.

One housekeeping note: the force-push converted 23 of the 26 touched files to CRLF line endings, so the PR diff against e1f292a shows ~14,000 changed lines; with --ignore-cr-at-eol it's 180 lines in 7 files.

q8atnight · 2026-10-05T06:19:55Z

Our IQ3_S gate row for the config nobody has covered — IQ3_S + partial residency on a 2× 3090 rig
(NVLink pair + PCIe, peer expert tier), fresh processes, same base, outputs compared word-for-word against
the stock v0.1.39 arm on the identical config and env.

Same-box arms: stock 6f32ec0 vs 6f32ec0 + this PR, gate env
(STRATA_IQ_MT_MIN=1 STRATA_MMQ_NO_STREAMK=1 STRATA_PEER_HOT=0), temp 0, fixed prompts.

Single-GPU (gS, expert cache 6,000): chat DIFFERS, code DIFFERS, math DIFFERS (3/3), long gate FAIL.
Peer tier (gD, 3,000 + 3,950 slots, partial residency): same result — 3/3 + long gate FAIL, at the same
character positions as the single-GPU run, so the divergence is the PR itself, not the peer path.
The four switches did not cover it here either.
Speed (gate context, decode at 3,000 tokens, ×3): stock 49.0 / 50.3 / 50.4 → PR 52.1 / 52.4 / 52.3 t/s = +6 %
on the peer-tier path. Prefill unchanged. So the gain is real on this hardware class too — it just comes with
the answer change, which is what we cannot carry.

Conclusion on our side: not adopted. If the fused kernels get an exact path for partially-resident IQ3_S (or a
gate-passing default with the switches), we will re-run this same reference in a maintenance window — the
window is cheap now that the harness is set up (worktrees strata-w039 / strata-ab783 + saved refs
bench/pr783ref-gD{,-long}.json).

@GonzMartinez

…tributes The force-push that produced 4327807 rewrote the line endings of 23 source files (.cu/.cpp/.hpp) from LF to CRLF. Nothing in their content changed, but the PR diff against v0.1.39 read as ~12,900 phantom changed lines -- a reviewer had to run `--ignore-cr-at-eol` to find the real 180 lines in 7 files. This commit converts those files back to LF with no content change (`git diff --ignore-cr-at-eol` against the parent is empty except for .gitattributes), and adds `* text=auto eol=lf` so a Windows checkout cannot reintroduce it. The existing exceptions still win: *.bat stays CRLF, bench/prompts/** and tests/data/native_experts/** stay untouched. Reported by @GonzMartinez on Niko1221#783.

stuchapin909 · 2026-10-06T00:46:55Z

Closing this. The four-commit stack is too much to review or bisect: @GonzMartinez and @q8atnight both still see greedy mismatch vs 0.1.39 on partial-resident / layer-split IQ3, and the kill switches do not cover it. Speed on fully resident boxes is real; that is not enough to keep an open PR that changes answers.

Thanks @GonzMartinez, @q8atnight, and @T-Crypt for the gates. We will not ask for another re-run on this branch. If we come back it will be one testable change at a time, default-exact against 0.1.39.

Branch stays on the fork for the record.

…tomic in kv_stream resolve (#783 PR-a) A CTA of the fixed 256-CTA grid whose first block is past the batch's largest n_bid returns before staging the queries. Same scores, same bits. STRATA_QSA_EARLY_EXIT=0 restores the old staging. kv_stream resolve_kernel: only threads that looked something up add to the shared counter. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…with an all-outputs path, and the gdn_parity T=1..8 memcmp check (#783 PR-b) gdn_ab_multi is instantiated per T=1..6 (generic above) with the weight unpack hoisted out of the token loop; gdn_step_norm_multi prefetches the next token's q/k while the current one runs, and when no state is written and every token is an output (the verify half) it skips the per-token bookkeeping. gdn_step_norm_kernel splits its reduction buffer so one barrier goes. gdn_parity gains section 5: for every T = 1..8, gdn_ab_multi against fused_gdn_ab, and gdn_step_norm_multi (verify half, commit half, t_out_begin, state-only commit) against the single-token kernel chained, compared with memcmp. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…sigmoid scale (#783 PR-c) native_router_top10_multi runs one warp per token (8 per CTA) instead of one CTA per token; native_moe_combine(_multi) takes a float4 kernel for k=10 with a shared row on aligned buffers; the shared expert's sigmoid-scale is vectorised. Same arithmetic order, so same bits. New native_multi_parity (memcmp): the multi router against the single call for n=1..19 on random, tied and NaN logits, and the combine against a host replay of its contract for the float4 and scalar paths. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…ommit graph and the MTP catch-up (#783 PR-d) kv_append_q8_steps / kv_append_q4_steps append n cells in one 3D-grid launch per pool, and native_qsa_indexer_append_steps walks the n indexer cells in order inside one CTA (one barrier between cells). Verify record_window, the commit graph and MtpDrafter::record_forward use them; a multi-slot batch keeps the per-token loop (each row has its own K/V). STRATA_NO_BATCH_KV_STEP=1 restores the per-token calls. The kv single-cell calls are now the n=1 case of the same kernel. kv_steps_parity (new, memcmp) covers INT8, Q4_0 and the hybrid's folded single-plane calls for n=1..8; qsa_parity gains append_steps against the sequential indexer append at strides 1 and 5 under none/linear/YaRN. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

PR-e) The plan kernel grouped duplicate experts and numbered the groups with thread 0 walking all entries; every entry now finds its group in parallel and a packed (entries << 16 | groups) warp prefix scan numbers them. The residency check, the empty plan and *plan_err of #871 are kept as they are. static_asserts pin the 4-warp / packing limits. resident_plan_parity (new, memcmp): n=1..80 entries over pools of 3..512 experts with and without a slot table, with and without a skip word, plus the non-resident case, against a host replay of the plan's contract. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…he MTP drafter, off on HIP (#783 PR-f) native_qsa_rms_norm_rope does norm<256> and the rope in one launch with the same FMA order as the pair (register barriers keep nvcc from contracting across the norm/rope boundary); rope_parity check 6 compares it with native_qsa_rms_norm_weighted + native_rope_apply by memcmp for head_dim 128/256, strided q|gate and contiguous inputs, unscaled and YaRN. The in-place rope also launches one 32-thread block per row. Verify's norm_rope and q split, and the MTP drafter's q path (all T tokens at once; K for a one-token window) use it. native_norm_rope_usable() gates it: on the CUDA build it is on unless STRATA_NO_NORM_ROPE=1; on the HIP build it is off until rope_parity check 6 passes on AMD hardware, and STRATA_NORM_ROPE=1 turns it on there. The MTP q path's split, FWHT and gate-apply now run over all T tokens in one call each (row-independent). Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

…iko1221#783 PR-e) The plan kernel grouped duplicate experts and numbered the groups with thread 0 walking all entries; every entry now finds its group in parallel and a packed (entries << 16 | groups) warp prefix scan numbers them. The residency check, the empty plan and *plan_err of Niko1221#871 are kept as they are. static_asserts pin the 4-warp / packing limits. resident_plan_parity (new, memcmp): n=1..80 entries over pools of 3..512 experts with and without a slot table, with and without a skip word, plus the non-resident case, against a host replay of the plan's contract. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…he MTP drafter, off on HIP (Niko1221#783 PR-f) native_qsa_rms_norm_rope does norm<256> and the rope in one launch with the same FMA order as the pair (register barriers keep nvcc from contracting across the norm/rope boundary); rope_parity check 6 compares it with native_qsa_rms_norm_weighted + native_rope_apply by memcmp for head_dim 128/256, strided q|gate and contiguous inputs, unscaled and YaRN. The in-place rope also launches one 32-thread block per row. Verify's norm_rope and q split, and the MTP drafter's q path (all T tokens at once; K for a one-token window) use it. native_norm_rope_usable() gates it: on the CUDA build it is on unless STRATA_NO_NORM_ROPE=1; on the HIP build it is off until rope_parity check 6 passes on AMD hardware, and STRATA_NORM_ROPE=1 turns it on there. The MTP q path's split, FWHT and gate-apply now run over all T tokens in one call each (row-independent). Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…6 projections (Niko1221#783 PR-g) gr_read_multi / gr_write_multi (and native_gr_rms_norm_weighted_multi, native_gr_pre_gated_multi, native_gr_post_multi under them) run T tokens in one launch per step; the verify window's PLE writes, last-layer write and final mixer, and the MTP drafter's write and final mixer use them. The BF16 MMVF multi kernel and the GR up / staged down kernels are instantiated per exact T (1..6, generic above), the BF16 weight tile is unpacked once for all tokens, with the same FMA order as before. STRATA_NO_MULTI_GR=1 restores the per-token calls and the generic kernels. STRATA_GR_DOWN_MAX4 keeps its meaning from the previous commit (sm_120 default on, explicit setting decides); this commit does not change that default. gr_multi_parity (new, memcmp): gr_read_multi (with and without injection) and gr_write_multi (in place and out of place) against the single-token calls at the real geometry, T=1..8. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…MVQ columns, decoded-once small-format tables, two rows per block (Niko1221#783 PR-h) iq_kernels.cu: row_dot_80_sub16 (two rows per warp, all lanes active) for the n_embd == 2560 gate/up of every split format, exact-N instantiations of the multi-column MMVQ (launch_mmvq_multi_nc), the 3-entry tail in the grouped gate/up and down kernels. native_mmvq.cu: SmallTraits for Q4_0, Q5_0, Q8_0 and IQ4_NL decode the weight chunks and block scale once per (row, block) and apply them to every activation column with the arithmetic of small_q8_dot, and the multi-column launch runs two rows per block for any K. Kill switches: STRATA_NO_SUB16_GU=1 (sub-warp gate/up/down), STRATA_OLD_IQ_MMVQ=1 (per-column kernels, existing), STRATA_NO_MMVQ_ROWS2=1 (one row per block at large K). mmvq_multi_parity, iq_multi_parity, native_grouped_parity and native_expert_parity pin the columns to the single-column results. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…o1221#783 PR-i) MtpDrafter::draft now runs the round for T = a + 1: the catch-up writes the K/V of the accepted cells [0, a) and the full layer covers row a, instead of writing all T cells and rewriting cell a. No later attention reaches a cell past the one being drafted and the next round writes it first, so the draft tokens and probabilities are unchanged. draft_first (a = 0) no longer writes its repeated rows at all. STRATA_MTP_CATCHUP_ALL=1 restores the whole-window catch-up (and the old staging sizes). Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…erts are in the window (Niko1221#783 PR-j) STRATA_SH_FORK_LATE=0 forks at the top as before. The other two commit-4 changes of the PR are not taken: the shared-expert Q8_1 reuse is already in f-140 as the opt-in STRATA_VERIFY_QDEDUP (and the fused q8_1 store), and the PLE native-key batch is already in f-140 as the opt-in STRATA_PLE_BATCH=1; both would have changed output on the native packs by default. Dropped on purpose: STRATA_VERIFY_PRELAUNCH and the dead copy_rows_from_mapped<ZERO_HITS>. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…GLU + Q8_0 in the grouped S2 experts (Niko1221#783 PR-k) quantize_q8_0_scaled_kernel reduces the block's max with a warp shuffle (one warp per 32-value block) instead of one thread walking the block; fmax ignores NaN the same way in both. moe_grouped_s2 fuses SwiGLU with that quantize and writes the activation back to gate_up, so the grouped parity checks see the same buffer. STRATA_S2_SWIGLU_Q8=0 keeps the two launches; a static_assert pins FF % 32 == 0 for the fused kernel. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…PR-c, MTP wiring) For a window of more than one token the drafter routes with one multi-column router GEMV and one multi-row top-10 and combines all rows in one launch (verify already did; each row is bitwise the single-token call, native_multi_parity). Needs the native router (512 experts, k = 10) and the native combine; otherwise the per-token calls stay. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…1#783 PR-c) Found by native_multi_parity on gfx1151 (Aurora): the float4 combine differed from the scalar kernel there, because the first pair's other product was contracted into the FMA. Written as fmaf per expert it is the pinned contract (first product rounded, FMA per expert in order) on every compiler. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…== 0 branch (Niko1221#783 PR-b) On gfx1151 the shuffle reduction of the sum of squares inside the divergent-looking branch rounded differently for ~3% of the outputs than the single-token kernel (gdn_parity section 5, Aurora); the base kernel had the shuffles outside it, and with them back out the multi kernel is bitwise the single one for T=1..8 on gfx1151. No change for CUDA. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…kernel is not enabled (Niko1221#783 PR-f) On the HIP build the fused kernel is off until check 6 passes there, so check 6 reports SKIPPED unless STRATA_NORM_ROPE=1 asks for it. (gfx1151 on Aurora, 0.1.40: with STRATA_NORM_ROPE=1 it does not pass - about 7% of the rotated elements differ - which is why the kernel stays off on HIP.) Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

… keep f-140's file (Niko1221#783 PR-k)

…_MULTI_GR flag (Niko1221#783 PR-g)

A decode window is a long chain of small kernels, and on a fast card many of them are over before the next one has fetched its weights. With programmatic dependent launch the window's quantizations and dense projections (multi-column MMVQ, multi-row BF16 GEMV) start while the kernel before them is still running: they load their first weights, then wait (griddepcontrol.wait) for its output. Nothing is written before the wait, so every value is the one an ordinary launch computes. - pdl.hpp: pdl_wait / pdl_trigger and launch_pdl. The attribute is used only inside the verify window's capture (pdl_scope), with a CUDA 12.3+ runtime, on sm_90+ (STRATA_EMULATE_CC counts), for a kernel whose loaded code is sm_90+ (ptxVersion and binaryVersion >= 90: older PTX, JIT-compiled, has no wait and would race), and only when every node the capturing stream's next node depends on is a kernel. So there is never a programmatic edge after a memcpy, memset, host or event node. After a join of branches each joined kernel gets a programmatic edge (the wait covers them all); STRATA_DF_PDL=2 allows a single predecessor only. - On CUDA the PDL kernels' activations and outputs are no longer __restrict__ (the compiler could otherwise load them before the wait); the weights keep it. For every architecture: the host pass must see the device pass's signature, or MSVC rejects the template stubs. - HIP: ordinary launches, no-op device functions, the kernels' plain loops. - copy_rows_strided replaces the q/gate split's cudaMemcpy2DAsync on the batched QSA path: a memcpy node would end the PDL chain there. A copy: the same bytes, PDL on or off. - The profiler's stamps pass the early launch on inside the window. - STRATA_VERIFY_NODES also counts the programmatic edges and those that are not kernel to kernel (0 expected). - STRATA_DF_PDL=0: ordinary launches. Byte-identical: PDL changes when a kernel starts, not what it computes. Off on pre-sm_90 cards (the 4060 Ti) and on HIP. Checked: pdl_parity (new test) captures a verify-shaped chain (quantize, small-K and large-K MMVQ, BF16 GEMV, the shared Q8_1 buffer rewritten, a memcpy node, a fork and a join) with and without PDL: bitwise equal outputs over 200 replays with every buffer NaN-poisoned before each; programmatic edges only kernel to kernel, and never into a node that also follows a non-kernel node; present on sm_90+, absent elsewhere. In our fork the same kernels gave identical windows over full requests (per-layer hashes of the MoE inputs and the final residual). Measured together with the next commit (PDL is on the 5080 only: the 4060 Ti, sm_89, has none); the numbers are there. PR Niko1221#904's PDL commit ported onto 0.1.40 (8 conflicts resolved against 0.1.40's EXACT_T bf16 kernels, Niko1221#783's rows rule, TSUM and the w2/y2 pair outputs; launch_pdl needs the defaulted arguments spelled out). pdl_parity passes on sm_120; no window-time change here.

…py (#1188, #1169) #783's batched append of a verify window's cells passed no host copy for the hybrid K8V4 format (int8 and q4_0 pass st.host, and the per-token path passes K8V4's two halves). With --kv-resident, a block that was not resident when its cells were appended came back from the host copy without them, so the tokens a window wrote were lost: free generation degenerated into repetition (#1188, #1135) and parked conversations lost their decode-written turns (#1169). The batched path now passes the same host halves as the per-token path.

…y path, PLE post-ops in one pass (all opt-in) Ported from #904 onto the 0.1.40 verify window. Every switch is OFF by default, so default output and the default launch sequence are unchanged: - STRATA_DF_PDL=1: programmatic dependent launch (pdl.hpp) for the window's quantizations and dense projections (multi-column MMVQ, multi-row BF16 GEMV); sm_90+ only, CUDA 12.3+; HIP and older cards keep ordinary launches. - STRATA_DF_PLE=1: the window's hyper-connection write and PLE post-ops in one pass, which also writes the commit's history snapshots (native_gr_post_multi with strides, native_ple_postops_batch_snap). - STRATA_DF_QB_Q4=1: a Q4_0 KV cache takes the batched QSA query path too. - STRATA_DF_BRANCH=1: the mixer's side branches inside the window graph. Left opt-in because the author saw a stall on Linux with NVML queries. The q/gate split now uses copy_rows_strided (a kernel) instead of a cudaMemcpy2DAsync so a PDL chain stays kernel to kernel; the bytes are the same. Left out: STRATA_DF_KVAPP (kv_append_q4_step_multi). 0.1.40 already appends a window's K/V cells in one launch (kv_append_q4_steps, #783 PR-d). Merged by hand with the 0.1.40 forms of the same code: the exact-T BF16 multi-token kernel, the PAIR/TS MMVQ variants, QFUSE (q8_attn), the fused norm+rope and the batched residual write. Co-authored-by: Francesco Albano <119073350+Hardin22@users.noreply.github.com> (PR #904)

…tomic in kv_stream resolve (Niko1221#783 PR-a) A CTA of the fixed 256-CTA grid whose first block is past the batch's largest n_bid returns before staging the queries. Same scores, same bits. STRATA_QSA_EARLY_EXIT=0 restores the old staging. kv_stream resolve_kernel: only threads that looked something up add to the shared counter. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…with an all-outputs path, and the gdn_parity T=1..8 memcmp check (Niko1221#783 PR-b) gdn_ab_multi is instantiated per T=1..6 (generic above) with the weight unpack hoisted out of the token loop; gdn_step_norm_multi prefetches the next token's q/k while the current one runs, and when no state is written and every token is an output (the verify half) it skips the per-token bookkeeping. gdn_step_norm_kernel splits its reduction buffer so one barrier goes. gdn_parity gains section 5: for every T = 1..8, gdn_ab_multi against fused_gdn_ab, and gdn_step_norm_multi (verify half, commit half, t_out_begin, state-only commit) against the single-token kernel chained, compared with memcmp. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…sigmoid scale (Niko1221#783 PR-c) native_router_top10_multi runs one warp per token (8 per CTA) instead of one CTA per token; native_moe_combine(_multi) takes a float4 kernel for k=10 with a shared row on aligned buffers; the shared expert's sigmoid-scale is vectorised. Same arithmetic order, so same bits. New native_multi_parity (memcmp): the multi router against the single call for n=1..19 on random, tied and NaN logits, and the combine against a host replay of its contract for the float4 and scalar paths. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…ommit graph and the MTP catch-up (Niko1221#783 PR-d) kv_append_q8_steps / kv_append_q4_steps append n cells in one 3D-grid launch per pool, and native_qsa_indexer_append_steps walks the n indexer cells in order inside one CTA (one barrier between cells). Verify record_window, the commit graph and MtpDrafter::record_forward use them; a multi-slot batch keeps the per-token loop (each row has its own K/V). STRATA_NO_BATCH_KV_STEP=1 restores the per-token calls. The kv single-cell calls are now the n=1 case of the same kernel. kv_steps_parity (new, memcmp) covers INT8, Q4_0 and the hybrid's folded single-plane calls for n=1..8; qsa_parity gains append_steps against the sequential indexer append at strides 1 and 5 under none/linear/YaRN. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…iko1221#783 PR-e) The plan kernel grouped duplicate experts and numbered the groups with thread 0 walking all entries; every entry now finds its group in parallel and a packed (entries << 16 | groups) warp prefix scan numbers them. The residency check, the empty plan and *plan_err of Niko1221#871 are kept as they are. static_asserts pin the 4-warp / packing limits. resident_plan_parity (new, memcmp): n=1..80 entries over pools of 3..512 experts with and without a slot table, with and without a skip word, plus the non-resident case, against a host replay of the plan's contract. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…he MTP drafter, off on HIP (Niko1221#783 PR-f) native_qsa_rms_norm_rope does norm<256> and the rope in one launch with the same FMA order as the pair (register barriers keep nvcc from contracting across the norm/rope boundary); rope_parity check 6 compares it with native_qsa_rms_norm_weighted + native_rope_apply by memcmp for head_dim 128/256, strided q|gate and contiguous inputs, unscaled and YaRN. The in-place rope also launches one 32-thread block per row. Verify's norm_rope and q split, and the MTP drafter's q path (all T tokens at once; K for a one-token window) use it. native_norm_rope_usable() gates it: on the CUDA build it is on unless STRATA_NO_NORM_ROPE=1; on the HIP build it is off until rope_parity check 6 passes on AMD hardware, and STRATA_NORM_ROPE=1 turns it on there. The MTP q path's split, FWHT and gate-apply now run over all T tokens in one call each (row-independent). Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

…6 projections (Niko1221#783 PR-g) gr_read_multi / gr_write_multi (and native_gr_rms_norm_weighted_multi, native_gr_pre_gated_multi, native_gr_post_multi under them) run T tokens in one launch per step; the verify window's PLE writes, last-layer write and final mixer, and the MTP drafter's write and final mixer use them. The BF16 MMVF multi kernel and the GR up / staged down kernels are instantiated per exact T (1..6, generic above), the BF16 weight tile is unpacked once for all tokens, with the same FMA order as before. STRATA_NO_MULTI_GR=1 restores the per-token calls and the generic kernels. STRATA_GR_DOWN_MAX4 keeps its meaning from the previous commit (sm_120 default on, explicit setting decides); this commit does not change that default. gr_multi_parity (new, memcmp): gr_read_multi (with and without injection) and gr_write_multi (in place and out of place) against the single-token calls at the real geometry, T=1..8. Ported from Niko1221#783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com>

beastrobin · 2026-10-13T14:16:07Z

One more data point from a fully-resident box, closing out the speed question here rather than reopening the review: RTX PRO 6000 Blackwell 96 GB, IQ3_S, engine 0.1.43.1 (which carries your STRATA_DUMMY_NODES=k probe hook). Injecting k=32 empty kernel nodes per layer - 1,536 extra nodes per verify window, ~1.5-3x the real node count of this chain - changed measured decode by less than the machine's noise band (five suite groups, signs flip, pooled high-n medians +/-3 %). So on this tier one graph/kernel node costs <= ~0.5 us and removing launches alone can never reach the +9-10 % seen on 2x3090 rigs; it caps near 2-3 % even if every real node were fused away. The parts of this stack that pull experts or touch host traffic are a separate story, but as pure launch-overhead work at full residency: dead end confirmed on Blackwell too. Full method and tables: comment dated 2026-10-13 on #2268; baseline suite for the box in #2280.

stuchapin909 force-pushed the perf/cuda-decode-wave2-6 branch 2 times, most recently from d37990f to e1f292a Compare October 4, 2026 15:41

stuchapin909 added 4 commits October 4, 2026 10:56

stuchapin909 force-pushed the perf/cuda-decode-wave2-6 branch from e1f292a to 4327807 Compare October 4, 2026 17:58

DarthCY-K pushed a commit to DarthCY-K/Strata that referenced this pull request Oct 4, 2026

Merge upstream PR Niko1221#783 (v0.1.39 base): fused decode/verify/MT…

3e9534b

…P kernels and graph launch reductions

T-Crypt mentioned this pull request Oct 4, 2026

bench: community report, RTX 4090 + 32 GB RAM at 512K context (IQ2_XS) #834

Closed

jomcgi mentioned this pull request Oct 4, 2026

inference-oom: profile multi-token verification launch fusion jomcgi-org/homelab#6851

Open

T-Crypt mentioned this pull request Oct 4, 2026

generate: one startup VRAM plan before the expert cache is committed (#765) #796

Closed

q8atnight mentioned this pull request Oct 5, 2026

perf(cuda,decode): zero-doorbell resident verify graph, sub-warp expert packing & shared-mem staging (+39-72% tok/s) #646

Closed

poakes28 mentioned this pull request Oct 5, 2026

Keep MTP drafting in concurrent batch slots #846

Closed

stuchapin909 closed this Oct 6, 2026

anon761 pushed a commit to anon761/Strata-upstream that referenced this pull request Oct 6, 2026

perf(quantize): quantize_act.cu is already one warp a block in f-140;…

da6782a

… keep f-140's file (Niko1221#783 PR-k)

anon761 pushed a commit to anon761/Strata-upstream that referenced this pull request Oct 6, 2026

fix(mtp): a NUL byte had replaced the backslash-zero of the STRATA_NO…

18fac29

…_MULTI_GR flag (Niko1221#783 PR-g)

merbanan mentioned this pull request Oct 6, 2026

verify: the batched K/V append writes a streamed K8V4 state's host copy (#1188, #1169) #1264

Closed

blange48 mentioned this pull request Oct 7, 2026

Isolate concurrent image positions and batch-MTP draft paths #1242

Closed

stuchapin909 mentioned this pull request Oct 8, 2026

verify: remove the unused resident_plan_par_kernel (CUDA and SYCL); INTEL.md no longer names STRATA_PLAN_PARALLEL #1554

Closed

johnmiu mentioned this pull request Oct 9, 2026

Windows HIP (RX 7900 XT, gfx1100): every generation fails with "mtp prefill: unspecified launch failure" on a GSQ-RCO IQ3_XXS native pack (0.1.41 and 0.1.40.4) #1644

Closed

k0valik mentioned this pull request Oct 15, 2026

replay (review only): hub.lumenfield.work-eelgaev-Strata-AC922-ac922 — filtered squash of 920 commit(s) from https://github.com/eelgaev/Strata-AC922 k0valik/Strata#4

Open

Navigation Menu

Repository navigation

perf(cuda): fused decode/verify/MTP kernels and graph launch reductions on 0.1.39 - #783

perf(cuda): fused decode/verify/MTP kernels and graph launch reductions on 0.1.39#783
stuchapin909 wants to merge 5 commits into
Niko1221:mainfrom
stuchapin909:perf/cuda-decode-wave2-6

stuchapin909 commented Oct 4, 2026 •

edited

Loading

Uh oh!

GonzMartinez commented Oct 4, 2026

Uh oh!

stuchapin909 commented Oct 4, 2026

Uh oh!

q8atnight commented Oct 4, 2026

Uh oh!

T-Crypt commented Oct 4, 2026

Uh oh!

T-Crypt commented Oct 4, 2026

Uh oh!

GonzMartinez commented Oct 5, 2026

Uh oh!

q8atnight commented Oct 5, 2026 •

edited

Loading

Uh oh!

stuchapin909 commented Oct 6, 2026

Uh oh!

beastrobin commented Oct 13, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

Conversation

stuchapin909 commented Oct 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Commits & What's Included

1. Multi-token GR / RMSNorm / post-op kernels and exact-T specializations (fused_gr.cu, gr.cu, native_gr_norm.cu, native_gr_postops.cu, native_bf16.cu, quantize_act.cu)

2. IQ/MMVQ sub-warp & 2-row kernels and MoE / shared-expert fusions (iq_kernels.cu, native_mmvq.cu, native_moe.cu, native_router.cu, s2_expert_grouped.cu, shared_expert.cu, elementwise.cu)

3. Batched KV append, QSA indexer/RoPE fusion, and parallel resident_plan_kernel (kv_q4.cu, kv_q8.cu, native_qsa_indexer.cu, native_rope.cu, verify_kernels.cu, fused_gdn.cu, qsa_select.cu, kv_stream.cu)

4. Wire multi-token kernels into Verifier and MtpDrafter CUDA graphs (verify.cpp, mtp.cpp, verify.hpp)

Verification

GPU Parity Suite (CTest — 25/25 Passed)

Uh oh!

GonzMartinez commented Oct 4, 2026

2x RTX 3090, Swift 1.5 IQ3_XXS, partially resident, layer split: +6.5-7.8% decode, but greedy output not identical to 0.1.39

Speed

Exactness: not identical on this configuration

Uh oh!

stuchapin909 commented Oct 4, 2026

Root Cause of the Greedy Exactness Divergence (chat / code / prose / long)

Runtime Kill Switches Added (No Recompile Needed)

Uh oh!

q8atnight commented Oct 4, 2026

Reply to PR #783 (stuchapin909) — posted via gh as q8atnight

Uh oh!

T-Crypt commented Oct 4, 2026

Uh oh!

T-Crypt commented Oct 4, 2026

Uh oh!

GonzMartinez commented Oct 5, 2026

Re-run on 4327807: still not identical on IQ3_XXS + 2-GPU layer split, and the four switches don't cover it

Uh oh!

q8atnight commented Oct 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

stuchapin909 commented Oct 6, 2026

Uh oh!

beastrobin commented Oct 13, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

stuchapin909 commented Oct 4, 2026 •

edited

Loading

1. Multi-token GR / RMSNorm / post-op kernels and exact-`T` specializations (`fused_gr.cu`, `gr.cu`, `native_gr_norm.cu`, `native_gr_postops.cu`, `native_bf16.cu`, `quantize_act.cu`)

2. IQ/MMVQ sub-warp & 2-row kernels and MoE / shared-expert fusions (`iq_kernels.cu`, `native_mmvq.cu`, `native_moe.cu`, `native_router.cu`, `s2_expert_grouped.cu`, `shared_expert.cu`, `elementwise.cu`)

3. Batched KV append, QSA indexer/RoPE fusion, and parallel `resident_plan_kernel` (`kv_q4.cu`, `kv_q8.cu`, `native_qsa_indexer.cu`, `native_rope.cu`, `verify_kernels.cu`, `fused_gdn.cu`, `qsa_select.cu`, `kv_stream.cu`)

4. Wire multi-token kernels into `Verifier` and `MtpDrafter` CUDA graphs (`verify.cpp`, `mtp.cpp`, `verify.hpp`)

Root Cause of the Greedy Exactness Divergence (`chat` / `code` / `prose` / `long`)

Re-run on `4327807`: still not identical on IQ3_XXS + 2-GPU layer split, and the four switches don't cover it

q8atnight commented Oct 5, 2026 •

edited

Loading