hc: the multi-token hyper-connection read split finer by BlueKingMuch · Pull Request #315 · Niko1221/Strata

BlueKingMuch · 2026-09-30T23:49:18Z

The multi-token hyper-connection read (fused_gr_read_multi: every verify window and the MTP drafter) keeps v1's arithmetic exactly - every sum over the same elements in the same order - and only splits the work differently. On an RTX 4080 SUPER it takes the two HC reads of a verify window from 5.01 to 4.21 ms (-16 %), about 0.9 ms of GPU time per round, with the same tokens.

What changes

Norm (v2, v3): one block per token and stream instead of one per token. Each thread visits the elements of its stream in the order it did in v1, so every sum of squares and rs is v1's; R' * w_norm is recomputed and scaled once (the same two roundings as v1's store and in-place scale).
Down projection (v3): v1's rows per block and lane order, but the activations arrive by cp.async in half-stream tiles (1280 floats: 160 = 5 x 32 chunks, so a lane still accumulates its chunks lane + 32 q in ascending order), two in flight, in a bank-conflict-free two-plane layout, and the next tile's weights load while the current one is used. Before sm_80 and on HIP the staging is a plain copy (the same bits).
The down kernels carry as many tokens per launch as the card's opt-in shared memory allows (v3: 10 KB per token), sliced like v1 since 0.1.27; pre-Volta cards use the per-block limit.

Exactness

Verifier::init runs fused_gr_check() once per card: v1, v2, v3 and the single-token read on random bf16 weights and inputs (1..8 tokens, with and without the pending write), every output compared bit for bit. The newest variant that agrees runs on that card, and the log says which:
strata hc: CUDA0: the hyper-connection read runs as v3 (...); checked bit for bit against v1 on this card.
A card that has not been checked runs v1. STRATA_HC_V2=0 keeps v1, =2 stops at v2.
End to end: requests without anything timing-dependent (no prompt-lookup drafts, no expert swaps, pcie_frac=0) give the same 256 tokens with v3 and v1, greedy and sampled (T 1.0, top_p 0.95, top_k 20, seed 1), in the same number of rounds.

It is on by default where the check passes, since it changes no bit. If you prefer it opt-in, that is one line in env_variant().

Measured

RTX 4080 SUPER 32 GB (sm_89), Ryzen 7 5800X3D, 64 GB DDR4, PCIe 3.0 x16, Windows 11, CUDA 13.3; IQ3_S, --max-context 262144, int8 KV, 11055 experts in VRAM. A local test build of 0.1.30 (30ec18e) with this commit, #242, #284, and #186 and #258 behind switches; every comparison is within that one build (STRATA_HC_V2 unset = v3 against STRATA_HC_V2=0 = v1), so everything else is identical.

The deterministic requests above (8K context, 256 tokens, the same tokens and 100 / 99 rounds in both):

	v1	v3
greedy: ms per round	26.43	25.71	-2.7 %
sampled: ms per round	25.00	24.28	-2.9 %
GPU wait per round (greedy / sampled)	12.08 / 12.46	11.20 / 11.57	-0.88 ms
tokens/s (greedy / sampled)	96.9 / 103.4	99.6 / 106.5	+2.8 % / +3.0 %

With prompt-lookup drafts and expert swaps (3 x 256 tokens per cell, 8K and 64K context, greedy and sampled) tokens/s scatter by several percent from request to request, but the GPU wait per round drops in every cell: 14.67 -> 13.63, 14.90 -> 13.77, 15.94 -> 14.91, 16.00 -> 15.03 ms.

GPU stages (STRATA_VERIFY_PROFILE=1, 8K, ms per verify window, DeltaNet + attention layers summed; the stamps split only the first read of each layer):

	v1	v2	v3	#186 (`STRATA_GR_V3=1`)	v1 + #258's 1280 tile, forced on sm_89
read 1: norm	0.55	0.42	0.42	0.27	0.55
read 1: down	1.02	1.02	0.76	1.16	0.92
read 1: up	0.64	0.64	0.63	0.68	0.64
read 2 + router	2.80	2.65	2.40	2.70	2.70
sum	5.01	4.73	4.21	4.81	4.81

On this card v3 is 0.6 ms per window faster than #186, and it keeps v1's bits.

To reproduce on the same build: STRATA_HC_V2=0 (v1) against unset (v3); STRATA_VERIFY_PROFILE=1 for the stages.

Tested on sm_89 only. The pre-sm_80 and HIP paths compile the same kernel with plain copies instead of cp.async; there the startup check decides as well.

…r bit on each card (STRATA_HC_V2) The multi-token read (verify windows, the MTP drafter) keeps v1's arithmetic exactly - every sum over the same elements in the same order - and only splits the work differently: - the norm (v2, v3): one block per token and stream instead of one per token. Each thread visits the elements of its stream in the order it did in v1, so every sum of squares and rs is v1's; R' * w_norm is recomputed and scaled once (the same two roundings as v1's store and in-place scale). On an RTX 4080 SUPER: 0.58 -> 0.23 ms per 4-token round. - the down projection (v3): v1's rows per block and lane order, but the activations arrive by cp.async in half-stream tiles, two in flight, in a bank-conflict-free two-plane layout, and the next tile's weights load while the current one is used. A tile is 160 = 5 x 32 chunks, so a lane still accumulates its chunks lane + 32 q in ascending order. Before sm_80 and on HIP the staging is a plain copy (the same bits). Verifier::init runs fused_gr_check once per card: v1, v2, v3 and the single-token read on random bf16 weights and inputs (1..8 tokens, with and without the pending write), every output compared bit for bit; the newest variant that agrees runs on that card and the log says which. A card that has not been checked runs v1. STRATA_HC_V2=0 keeps v1, =2 stops at v2. The down kernels carry as many tokens per launch as the card's shared memory allows (v3: 10 KB per token), as v1 did. Claude-Session: https://claude.ai/code/session_01FKZxC8sUotEzdWXws2FTUC Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Niko1221 · 2026-10-01T06:21:00Z

Thanks! This is in the 0.1.32 batch. We take it with its variant/env names changed so they don't clash with STRATA_GR_V3 from #186, and the self-check comparing against the single-token read; every change goes through our identity gate (byte-identical output on four models) and a speed A/B before it ships.

…ed, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 main already has an opt-in "v3" hyper-connection read (#186, STRATA_GR_V3=1, another summation order); #315's "v2"/"v3" and STRATA_HC_V2 read as the same thing. Now: the plain read (0.1.31's default), split (the norm per token and stream) and staged (split + the down projection's activations staged by cp.async), kHcPlain/kHcSplit/ kHcStaged, gr_norm_split_kernel / gr_down_staged_kernel, and STRATA_HC_SPLIT: unset or 2 = the newest the card's check accepts, 1 = at most split, 0 = the plain read. The check compares against the plain read (launch_multi with kHcPlain) and the single-token read, never against what STRATA_GR_V3 selects; with STRATA_GR_V3=1 that opt-in read still runs first, unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

Niko1221 · 2026-10-01T17:29:31Z

Released in 0.1.32 with its variants renamed (STRATA_HC_SPLIT) and the self-check against the plain read. Merged under your name (GitHub doesn't show it as merged because it went in by hand). Thank you!

A decode window on this model 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. These changes cut launches and let the next kernel start early. Every one keeps the arithmetic of the path it replaces (checked with per-layer hashes of the MoE inputs and the final residual over full requests: identical windows), and each has an off switch for A/B runs. - Programmatic dependent launch on sm_90+ (STRATA_DF_PDL, default on): the verify kernels trigger their dependents early, and the dependents prefetch their weight rows into L2 before waiting for the activations. Off on older cards. - Two side streams inside the verify graph (STRATA_DF_BRANCH): the query side (q, the indexer query) runs beside the K/V side and joins before attention. - A Q4_0 KV cache batches too (STRATA_DF_QB_Q4, STRATA_DF_KVAPP): the query rotation is row-wise, so the window's rows go in one launch, and the window's K/V cells are appended in one launch. - PLE: the per-token post-ops of a window in one kernel (STRATA_DF_PLE). - A layer split's hand-off copies (3 x T) in one launch (STRATA_DF_HANDOFF). - native mmvq: a warp per row for large outputs (STRATA_DF_MMVQW): 10-30% less time on 6144..248320-row IQ4_XS on both cards (slower at 640 rows, so it starts at 2048). - Attention loops with their loads batched (STRATA_DF_ATTN), SwiGLU folded into the shared expert's down quantization (STRATA_DF_SHARED). - copy_rows_strided replaces a per-token cudaMemcpy2DAsync for the q/gate split. - An L2 prefetch of the next layer's weights on a side stream, opt-in (STRATA_DF_L2PF_MB=n, STRATA_DF_L2PF_MIN_SM). It did not help on the 5080 + 4060 Ti. Upstream's staged hyper-connection read (Niko1221#315) and decode-once multi-column IQ kernels (Niko1221#242) do what two earlier parts of this change did, so those parts were dropped in favour of upstream's.

…ed, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 main already has an opt-in "v3" hyper-connection read (#186, STRATA_GR_V3=1, another summation order); #315's "v2"/"v3" and STRATA_HC_V2 read as the same thing. Now: the plain read (0.1.31's default), split (the norm per token and stream) and staged (split + the down projection's activations staged by cp.async), kHcPlain/kHcSplit/ kHcStaged, gr_norm_split_kernel / gr_down_staged_kernel, and STRATA_HC_SPLIT: unset or 2 = the newest the card's check accepts, 1 = at most split, 0 = the plain read. The check compares against the plain read (launch_multi with kHcPlain) and the single-token read, never against what STRATA_GR_V3 selects; with STRATA_GR_V3=1 that opt-in read still runs first, unchanged.

…t / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 main already has an opt-in "v3" hyper-connection read (Niko1221#186, STRATA_GR_V3=1, another summation order); Niko1221#315's "v2"/"v3" and STRATA_HC_V2 read as the same thing. Now: the plain read (0.1.31's default), split (the norm per token and stream) and staged (split + the down projection's activations staged by cp.async), kHcPlain/kHcSplit/ kHcStaged, gr_norm_split_kernel / gr_down_staged_kernel, and STRATA_HC_SPLIT: unset or 2 = the newest the card's check accepts, 1 = at most split, 0 = the plain read. The check compares against the plain read (launch_multi with kHcPlain) and the single-token read, never against what STRATA_GR_V3 selects; with STRATA_GR_V3=1 that opt-in read still runs first, unchanged.

…t / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 main already has an opt-in "v3" hyper-connection read (Niko1221#186, STRATA_GR_V3=1, another summation order); Niko1221#315's "v2"/"v3" and STRATA_HC_V2 read as the same thing. Now: the plain read (0.1.31's default), split (the norm per token and stream) and staged (split + the down projection's activations staged by cp.async), kHcPlain/kHcSplit/ kHcStaged, gr_norm_split_kernel / gr_down_staged_kernel, and STRATA_HC_SPLIT: unset or 2 = the newest the card's check accepts, 1 = at most split, 0 = the plain read. The check compares against the plain read (launch_multi with kHcPlain) and the single-token read, never against what STRATA_GR_V3 selects; with STRATA_GR_V3=1 that opt-in read still runs first, unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the paths this fork still mirrors: 26 file(s) under include/strata -- 25 clean, 1 with conflict markers. Conflict-marked files -- a pull request with conflicts: resolve them, do not discard the work: include/strata/kernels/verify_kernels.hpp Original authors: 4EverBuilder,Andy,Anon,BlueKingMuch Original commits touching these paths (152 subject(s)): - Niko1221#253: the 8 GiB arena pin cap only under WDDM; STRATA_ARENA_PIN_GIB; the split's prompt chunk steps down - Niko1221#284 follow-up: wait for the async commit on an event, at every point that touches the session afterwards - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 - Niko1221#340 prototype: a split's streamed ring is 96 slots when most experts are resident (STRATA_SPLIT_RING); STRATA_SPLIT_SMALL_OWN=S keeps an S-token chunk's buffers per lending stage for the session; --no-prefill-borrow on a split books the ring (it no longer starves the draft head) - Niko1221#465 engine: batch slots for one GPU - warn instead of refuse, windows over the active slots, slot conversation cache, prompts read beside decoding slots, BYIELD - Niko1221#583 default on (owner): STRATA_RING_BYTES=0 opts out; DETAILS.md with the measured quality - Niko1221#583: the byte-budget ring only where it gains; short prompts keep 0.1.39's ring - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base) Review artifact produced by fork-replay.sh (a spike adapted from sync-upstream.sh): a filtered, squashed replay for reading -- not built, not tested, and not evidence that any of it works. Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922 Fork-Replay-Branch: ac922 Fork-Replay-Base: fb58e0d Fork-Replay-Tip: 1f029aa Fork-Replay-Group: include/strata Fork-Replay-Automated: true

Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the paths this fork still mirrors: 21 file(s) under src/kernels -- 14 clean, 7 with conflict markers. Conflict-marked files -- a pull request with conflicts: resolve them, do not discard the work: src/kernels/cpu/native_expert.cpp src/kernels/cpu/pool.cpp src/kernels/cuda/dequant_bf16.cu src/kernels/cuda/fused_gr.cu src/kernels/cuda/iq_kernels.cu src/kernels/cuda/verify_kernels.cu src/kernels/ngram.cpp Original authors: 4EverBuilder,Andy,BlueKingMuch,Chanyeong Lim Original commits touching these paths (151 subject(s)): - Niko1221#152 as an opt-in (the owner's choice): STRATA_IQ_MT_MIN=1 makes CPU expert rounding independent of the drafting; the default stays 0.1.29's rule (IQ3_S decoded 1-3% slower with it on); docs; the width test runs the opt-in - Niko1221#289 follow-up: a placeholder GPU name when the properties query fails; HIP builds compile (no device_code_error probe) - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 - Niko1221#646 on 0.1.39b: the layer split and the CPU path as before, the resident plan for 8-token windows - Niko1221#646: the single-matrix mmvq does not stage the IQ grids (10-20% slower on the RTX 5070 at 3-8 columns) - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base) - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested. - C-1, C-2: the prompt path's QSA select grid reaches only the active blocks; the indexer appends a chunk at once Review artifact produced by fork-replay.sh (a spike adapted from sync-upstream.sh): a filtered, squashed replay for reading -- not built, not tested, and not evidence that any of it works. Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922 Fork-Replay-Branch: ac922 Fork-Replay-Base: fb58e0d Fork-Replay-Tip: 1f029aa Fork-Replay-Group: src/kernels Fork-Replay-Automated: true

Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the paths this fork still mirrors: 26 file(s) under include/strata -- 25 clean, 1 with conflict markers. Conflict-marked files -- a pull request with conflicts: resolve them, do not discard the work: include/strata/kernels/verify_kernels.hpp Original authors: 4EverBuilder,Andy,Anon,BlueKingMuch Original commits touching these paths (152 subject(s)): - Niko1221#253: the 8 GiB arena pin cap only under WDDM; STRATA_ARENA_PIN_GIB; the split's prompt chunk steps down - Niko1221#284 follow-up: wait for the async commit on an event, at every point that touches the session afterwards - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 - Niko1221#340 prototype: a split's streamed ring is 96 slots when most experts are resident (STRATA_SPLIT_RING); STRATA_SPLIT_SMALL_OWN=S keeps an S-token chunk's buffers per lending stage for the session; --no-prefill-borrow on a split books the ring (it no longer starves the draft head) - Niko1221#465 engine: batch slots for one GPU - warn instead of refuse, windows over the active slots, slot conversation cache, prompts read beside decoding slots, BYIELD - Niko1221#583 default on (owner): STRATA_RING_BYTES=0 opts out; DETAILS.md with the measured quality - Niko1221#583: the byte-budget ring only where it gains; short prompts keep 0.1.39's ring - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base) Review artifact produced by fork-replay.sh (a spike adapted from sync-upstream.sh): a filtered, squashed replay for reading -- not built, not tested, and not evidence that any of it works. Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922 Fork-Replay-Branch: ac922 Fork-Replay-Base: fb58e0d Fork-Replay-Tip: 1f029aa Fork-Replay-Group: include/strata Fork-Replay-Automated: true

Squashed from 920 commit(s) fb58e0d..1f029aa on https://github.com/eelgaev/Strata-AC922, restricted to the paths this fork still mirrors: 21 file(s) under src/kernels -- 14 clean, 7 with conflict markers. Conflict-marked files -- a pull request with conflicts: resolve them, do not discard the work: src/kernels/cpu/native_expert.cpp src/kernels/cpu/pool.cpp src/kernels/cuda/dequant_bf16.cu src/kernels/cuda/fused_gr.cu src/kernels/cuda/iq_kernels.cu src/kernels/cuda/verify_kernels.cu src/kernels/ngram.cpp Original authors: 4EverBuilder,Andy,BlueKingMuch,Chanyeong Lim Original commits touching these paths (151 subject(s)): - Niko1221#152 as an opt-in (the owner's choice): STRATA_IQ_MT_MIN=1 makes CPU expert rounding independent of the drafting; the default stays 0.1.29's rule (IQ3_S decoded 1-3% slower with it on); docs; the width test runs the opt-in - Niko1221#289 follow-up: a placeholder GPU name when the properties query fails; HIP builds compile (no device_code_error probe) - Niko1221#315 follow-up: its variants and switch renamed (plain / split / staged, STRATA_HC_SPLIT) - no clash with STRATA_GR_V3 - Niko1221#646 on 0.1.39b: the layer split and the CPU path as before, the resident plan for 8-token windows - Niko1221#646: the single-matrix mmvq does not stage the IQ grids (10-20% slower on the RTX 5070 at 3-8 columns) - AMD HIP backend for gfx1100 (PR Niko1221#121, squashed onto its 0.1.21 base) - AMD HIP: RDNA2 gfx1030 (tested), gfx(1031, 1032) untested. - C-1, C-2: the prompt path's QSA select grid reaches only the active blocks; the indexer appends a chunk at once Review artifact produced by fork-replay.sh (a spike adapted from sync-upstream.sh): a filtered, squashed replay for reading -- not built, not tested, and not evidence that any of it works. Fork-Replay-Source: https://github.com/eelgaev/Strata-AC922 Fork-Replay-Branch: ac922 Fork-Replay-Base: fb58e0d Fork-Replay-Tip: 1f029aa Fork-Replay-Group: src/kernels Fork-Replay-Automated: true

…nfig - On upstream 0.1.32: what upstream took, what the fork keeps on (the opt-ins), the rebased pull requests, the conflicts, and the checks (the fork without its own defaults = upstream + Niko1221#353 byte for byte; the same first token and greedy tokens as 0.1.31-nvfp4.3 at a fixed cache; the same speed in interleaved runs). - The speed table re-measured to each answer's end (--stop-eos). The old "161-173 tokens/s after a 32K prompt" was the loop after a 22-token answer; the adaptive tier's numbers carry the same caveat (half of each run). - GR_V3 stays on: Niko1221#315's staged read measured 18.49 against 18.23 ms a round. - The bundle's config sets "anthropic_thinking": "on_request" (0.1.31-nvfp4.x's behaviour; upstream's default thinks). - Switches: STRATA_ROPE_TABLE=0 (STRATA_ROPE_LEGACY is gone), --prefill auto:N, anthropic_thinking. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio

BlueKingMuch mentioned this pull request Oct 1, 2026

decode: hyper-connection read in 2 kernels, stream-split (-1.5 ms per verify window) #186

Closed

Niko1221 closed this Oct 1, 2026

sergqwer mentioned this pull request Oct 1, 2026

gr_parity: compare the split read (STRATA_GR_V3=1) within float rounding (the test fails with it on) #379

Closed

bitaccesscomau mentioned this pull request Nov 1, 2026

Strata round 9: v0.2.9 performance candidate set (source audit, implementation-ready) bitaccesscomau/Strata#8

Open

Navigation Menu

Repository navigation

hc: the multi-token hyper-connection read split finer - #315

hc: the multi-token hyper-connection read split finer#315
BlueKingMuch wants to merge 1 commit into
Niko1221:mainfrom
BlueKingMuch:hc-read-bitwise

BlueKingMuch commented Sep 30, 2026 •

edited

Loading

Uh oh!

Niko1221 commented Oct 1, 2026

Uh oh!

Niko1221 commented Oct 1, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

BlueKingMuch commented Sep 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What changes

Exactness

Measured

Uh oh!

Niko1221 commented Oct 1, 2026

Uh oh!

Niko1221 commented Oct 1, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

BlueKingMuch commented Sep 30, 2026 •

edited

Loading