Repository navigation
Conversation
… in shared memory (STRATA_EXPDEC=1, opt-in, bitwise); RTX 2080 Ti gate/up kernel -11/-23% at T 1-4, decode +1.1% on a 131k layer split native_expert_grouped already decodes each weight part once per block, but the apply re-reads every entry's q8_1 block from global memory with 9 scattered 4-byte loads, so an extra row of the window costs ~60% of a whole expert. expdec_gu_kernel stages the group's entries in shared memory (two conflict-free 16-byte loads per call, the fp16 scale as a float), takes rb rows per block and stages the IQ3 codebooks too; the same lane does the same calls in the same order with the same operations (apply_u = apply on the block's words), so the output is bitwise the default kernels' (expdec_bench: identical on every cell, 2 cards). Gate/up IQ2_XXS/IQ2_XS/IQ3_XXS/IQ3_S/IQ2_S at n_embd 2560; the down (Q2_0, IQ4_NL) behind STRATA_EXPDEC_DOWN=1 (in the engine the Q2_0 down was ~10% slower). Knobs STRATA_EXPDEC_R/_DR/_MINB/_RB/_DRB/_MINBLK/_D20; one log line per (role, format) at the first launch. docs/EXPERTDEC.md. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
@Niko1221 This is the first of two small opt-in decode levers from our RTX 2080 Ti + RTX 3060 layer-split setup on Windows (the second is #2059). STRATA_EXPDEC=1 keeps the verify window's group of q8_1 rows in shared memory for the routed experts' gate/up; same kernels order and same sum trees, so the output is bit-exact (expdec_bench checks every cell with memcmp, and the engine's greedy text is unchanged). Measured with nsys in the engine (gate/up kernel -11/-23% at T 1-4 on the 2080 Ti) and end to end on a forced 131k bench (+1.1%); the down stays behind STRATA_EXPDEC_DOWN because in the engine it was slower. Ported and re-checked on v0.1.43. Off by default, CUDA only; happy to adjust names, defaults or the scope if you prefer something else, or to run more measurements. |
|
Tried this together with #2059 ( Rig: 4x RTX 5080 16 GB (PCIe Gen3), i9-10920X, 123 GiB RAM, Linux, CUDA sm_120. Qwen3.8-Flash-Next GSQ-RCO IQ3_S,
There is no measurable change on this rig, and no errors. Most of our routed experts are not VRAM-resident (16 GB cards, a 4-stage split), and each window's time goes to the stage hand-offs and the CPU or host experts, so the resident gate/up and the QSA block scores are presumably too small a share to show. I did not compare outputs token by token. |
|
Thanks a lot for testing both on that rig, very useful. That matches what we'd expect: the gain comes from VRAM-resident routed experts on sm_75, where gate/up is a larger share of the window (on our 2080 Ti the resident gate/up kernel is −11/−23% at T 1-4 and the window +1.1% / +2.9% with #2059). With most experts off-GPU on 16 GB cards, the hand-offs and host experts dominate, so a neutral result is the expected outcome. I'll note in the description that the levers target Turing/Ampere setups with mostly resident experts. If you happen to have the per-stage |
|
Hardware: 2x RTX 5060 Ti 16 GB layer split (both x8) + RTX 2080 Super for vision, Ryzen 7 3700X (AVX2, no AVX-512), ~32 GB RAM. IQ3_XXS, One row:
+5.0 t/s at 1K decode, prefill flat; the 32K decode delta is within run-to-run spread here. Answers: our gate (3 greedy prompts x 256) was already 1/3 identical on the stack before this flag — chat and code differ from earlier stack items, math identical. With the flag on the count did not change and chat/math are unchanged, but the code prompt's first diff against the stock reference moved (char 786 → 535), i.e. the code answer differs between the two arms on this box — so the bit-exact claim did not hold on our IQ3_XXS row. A 10-question verifiable probe scored 10/10 in both arms. |
|
Thanks, imanu86. We ported this as an opt-in (STRATA_EXPDEC=1). On an RTX 3060 (IQ3_XXS, 30K context) it was +0.5% decode, on an RTX A4000 +1 to +5% at short context, greedy text identical to the default in our identity gate (10 prompts and 4K/20K prompt dumps, with the switch on). Together with your other levers it gave +7% at 30K on both cards. It stays off by default until a second card class shows the same. Credit is kept in the commit. |
|
Thanks for the pull request. Your change is already in If part of your work did not make it in and you want it considered, open a new pull request on top of |
Title
Issue: none (an opt-in speed change).
Summary
STRATA_EXPDEC=1(opt-in, off by default) runs the verify window's VRAM-resident routed experts' gate/up through akernel that keeps the group's q8_1 activation rows in shared memory and computes several rows per lane. Same kernels
otherwise, same launch order, and the output is bit for bit the default kernels' (
expdec_benchchecks it withmemcmpon every cell). CUDA only; HIP and the prompt path are not touched.Why
native_gu_multi_kernelalready decodes each weight part once per block and applies it to the group's 1-4 entries,but the apply re-reads each entry's q8_1 block from global memory with 9 scattered 4-byte loads (36-byte blocks: a
warp's load touches 16-32 sectors). Measured with nsys on a 131k decode (RTX 2080 Ti, stage 1 of a layer split), an
extra row of the window costs about 60% of a whole expert - ~1.6 us against ~2.6 us per expert and layer - although it
reads nothing from DRAM. That is most of the gate/up kernel's slope with the window's rows (+0.86 ms per row).
What changed
src/kernels/cuda/iq_kernels.cu:expdec_gu_kernel<TG, R, MINB>(gate/up IQ2_XXS, IQ2_XS, IQ3_XXS, IQ3_S, IQ2_S atn_embd 2560) and
expdec_down_kernel<TD, R>(down Q2_0 at n_ff 640; IQ4_NL withSTRATA_EXPDEC_D20=1). The group'sentries are staged 4 at a time in a layout where a call's 8 words are two conflict-free 16-byte loads (IQ4_NL: two
8-byte loads) and the fp16 block scale is already a float; R rows per lane share the activation words; a block takes
rbrows (auto from the launch's groups) instead of 16; the IQ3_XXS / IQ3_S codebooks are staged in shared memory likethe IQ2 ones.
Split<TG>::apply_uisapplyon the block's words;Split<18>/<21>::load<true>read the staged grid.native_expert_groupedtakes the new kernels when the lever is on and the layout matches, else the default ones.include/strata/kernels/iq_kernels.hpp:native_expert_set_expdec,native_expert_set_expdec_params/native_expert_expdec_params(NativeExpertExpdecParams),native_expert_expdec_takes.src/kernels/expdec_bench.cpp(EXCLUDE_FROM_ALL): the default kernels against EXPDEC at T 1-4 on 10 format pairs,synthetic blobs read cold from a >= 96 MB pool,
memcmpof the output, median of 5 xrepscalls, per phase.docs/EXPERTDEC.md.Knobs (environment, read at startup):
STRATA_EXPDEC_R(gate/up rows per lane, 1 | 2, default 1),STRATA_EXPDEC_DR(down, 1 | 2 | 4, default 4),
STRATA_EXPDEC_MINB(gate/up__launch_bounds__min blocks per SM, 4 | 2, default 4),STRATA_EXPDEC_RB/_DRB(rows per block, 0 = auto),STRATA_EXPDEC_MINBLK(the auto rule's minimum blocks, 800),STRATA_EXPDEC_DOWN(the down too, default 0),STRATA_EXPDEC_D20(IQ4_NL down too, default 0). The first launch ofeach (role, format) logs one line (
strata expdec: CUDA1 gate/up type 21: expdec_gu_kernel R 1 MINB 4, ... captured into a graph).The down stays separate on purpose. In the micro-bench the Q2_0 down with EXPDEC is 0.65-0.85x of the default, but
in the engine (nsys, two captures with two block geometries) it was about 10% slower than the default kernel, which
already keeps the group's h rows in shared memory; IQ4_NL showed no gain in either. So
STRATA_EXPDEC=1takes thegate/up only and
STRATA_EXPDEC_DOWN=1is a separate opt-in for whoever wants to measure it on another card.Bit for bit
For every (row, entry) the sum is the default kernel's tree (
row_dot_80_sub16: lane t of a 16-lane sub-warp doescalls t, t+32, t+64 into
sand t+16, t+48 intos16,s = __fadd_rn(s, s16), butterfly xor 8-4-2-1;row_dot_40_sub8/
row_dot_20_sub16for the downs). The same lane does the same calls in the same order with the same integer andfloat operations on the same words; the entries of a group are independent, so the staging and the rows-per-lane /
rows-per-block choices only change which lane computes what.
expdec_bench:identicalon all 1,280 cells (10 pairs xT 1-4 x 2 window shapes x 8 choices x 2 cards). In the engine the greedy text is unchanged.
Measurements
Hardware: RTX 2080 Ti 22 GB (sm_75) + RTX 3060 12 GB (sm_86), Windows 11, CUDA 12.6, layer split (3060: layers 0-15,
2080 Ti: 16-47 + head),
--pipeline-windows 2, Qwen3.8-Flash-Next IQ3_XXS pack, 131k context.Micro-bench, gate/up phase, us per call with 20 distinct experts (the shape of a real window),
def-> ratio with EXPDEC:Per-expert / per-entry cost (us per layer call, RTX 2080 Ti, IQ3_S): 5.20 / 1.26 default -> 3.73 / 0.64.
nsys in the engine (96 verify windows, median kernel time per window at T = 1 / 2 / 3 / 4):
RTX 2080 Ti gate/up 1.60 / 2.58 / 3.51 / 4.16 ms -> 1.42 / 2.07 / 2.75 / 3.20 ms (-11 / -20 / -22 / -23%), slope per
row +0.86 -> +0.60 ms; RTX 3060 gate/up 0.77 / 1.43 / 1.97 / 2.23 -> 0.80 / 1.21 / 1.83 / 1.84 ms.
End to end, deterministic 131k bench (forced tokens and window sizes, 1024 decode tokens, three interleaved runs
lever on / off / on): +1.1% decode tokens/s (+0.0 and +2.1 in the two B arms, no B arm below A), text identical.
With the companion QSA scores lever (
STRATA_QSA_ROWS=1, separate PR) on top: +2.9% together. The expert kernels are~20% of the critical card's window and EXPDEC takes ~20% of that, so the end-to-end number is small by construction.
Ported and validated on v0.1.43 (micro-bench bit-exact on both cards; our full stack on 0.1.43 matches our 0.1.42
daily on the forced 131k bench).
Risks / limits
STRATA_OLD_IQ_MMVQ=1turns it off as it turns the multi kernels off. Captured graphs keep the kernels they captured.
MINB4 (116-128 with2). Checked withcuobjdump -res-usageon sm_75 and sm_86.untouched and is the next step; T=1 gains less (the flat cost is the weight side).
Scope. The gain comes from VRAM-resident routed experts / decode QSA scores on Turing/Ampere setups where these kernels are a large share of the verify window. On rigs where most routed experts are off-GPU (e.g. 4x RTX 5080 16 GB, see the test in #2058) the window is dominated by stage hand-offs and host experts, and the levers measure neutral.