Linux prebuilt from CI: CUDA 12.8, sm_80/86/89 + PTX, attached to each release by architectds · Pull Request #136 · Niko1221/Strata

architectds · 2026-09-29T12:45:44Z

The first piece from #103, as you suggested: a Linux prebuilt, built by CI.

The workflow (.github/workflows/linux-prebuilt.yml) compiles strata-linux-x64.zip on Ubuntu 22.04. It needs no GPU.

It builds the engine and the image encoder (strata-vision) with CUDA 12.8.
The architectures are 80-real;86-real;89: sm_80, sm_86 and sm_89 as machine code, plus compute_89 PTX for newer cards.
The CPU code is at the AVX2 baseline (STRATA_PORTABLE=ON). The AVX-512 kernels are still picked at run time where the CPU has them.
It writes the BUILD.json that get_prebuilt reads.

When you publish a release, a second job attaches the zip and its SHA-256 to it. releases/latest/download/strata-linux-x64.zip then sits next to the Windows zip, about 40 minutes after publishing. Until then, a Linux setup compiles, as it does today. Running the workflow by hand, from Actions > linux-prebuilt with other architectures if you want, builds an artifact only. So does a PR that changes the workflow.

setup.py uses CUDA 12 on Linux to match that build:

CUDA 12.8's pip packages, pinned like the CUDA 13 ones. CUDA 12's wheels keep cuBLAS and the runtime in separate folders, which cuda_lib_dirs now finds.
Driver 525 as the minimum.
The 12.8 toolkit when a compile has to install one.

Windows is unchanged (CUDA 13). STRATA_CUDA=12 or 13 picks the other one. A prebuilt built with another CUDA major is not used, because it would look for libraries setup did not install; setup compiles instead.

Tested:

The workflow on my fork: run 36565974211, on upstream main (0.1.23), 37 minutes. The engine for three architectures took 5 min, the image encoder 29 min. The zip is 99.5 MB:
- BUILD.json is {"version": "0.1.23", "archs": [80, 86, 89], "ptx": true, "cuda": "12.8", "vision": "gpu", ...};
- both binaries link libcudart.so.12, libcublas.so.12 and libcublasLt.so.12, which are the pinned pip wheels.
The same build setup (sm_80 only, engine 0.1.20) ran Qwen3.8-Flash-Next IQ3_S on a Colab A100-40G through setup.py --prebuilt all day: 262K context, images, and decode at 61 tok/s.

The RAM-based context cap is the other piece: #134.

…0K-144K The recipe now pins the fork's linux-prebuilt at e416ba5: upstream 0.1.24 plus the two open PRs, Niko1221/Strata#136 (the Linux prebuilt from CI, CUDA 12 on Linux) and #134 (the RAM-based context cap). Its engine is the release that workflow built: CUDA 12.8, sm_80/86/89 + PTX, the image encoder. Measured on the A100-40G on 2026-09-29, on real text with the n-gram table in RAM: - Prefill: 1,448 t/s at 16K, 2,082 at 36K, 2,492 at 70K, 2,497 at 144K. 0.1.20 read 1,301 at 28K and 1,376 at 59K. 0.1.22 put the prompt's attention on tensor cores and 0.1.24 its block selection, the part that grew with length. - Decode: unchanged, 58.8 t/s on a short answer (58.0-61.0 before), 66.1 after a 41K prompt, 72.3 on a next turn. - Images: right in 3.1 s; 0.1.24 thinks before it answers. - READY 18.5 min after the box was assigned. - /v1/status and the reply timings come from upstream now (#107, in since 0.1.23), so the rail needs nothing else. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Niko1221 · 2026-09-29T20:28:20Z

Thanks @architectds, a CI-built Linux engine is very welcome. Could you rebase onto 0.1.26, add native sm_90/sm_120 (the compute_89 PTX needs driver 570+, above your 525 minimum), and have setup avoid a PTX-only prebuilt or a CUDA 13 toolkit when the driver is too old?

…h release - .github/workflows/linux-prebuilt.yml compiles strata-linux-x64.zip on Ubuntu 22.04 with no GPU. It builds the engine and the image encoder with CUDA 12.8, sm_80, sm_86 and sm_89 plus PTX, and the CPU code at the AVX2 baseline (STRATA_PORTABLE). It writes the BUILD.json get_prebuilt reads. A published release gets the zip and its SHA-256 as assets. Running it by hand, or a PR that changes the workflow, builds an artifact only. - setup.py uses CUDA 12 on Linux: - CUDA 12.8's pip packages, pinned like the CUDA 13 ones (cuBLAS and the runtime sit in folders of their own); - driver 525 as the minimum; - the 12.8 toolkit when a compile has to install one. Windows keeps CUDA 13. STRATA_CUDA=12 or 13 picks the other one. A prebuilt built with another CUDA major is not used, since it would look for libraries setup did not install.

…and CUDA 13 - The workflow builds machine code for sm_80, sm_86, sm_89, sm_90 and sm_120, plus PTX for cards newer than those, and writes vision_archs into BUILD.json. CUDA 12.8's machine code runs from driver 525, but its PTX needs 570: a driver compiles the PTX of CUDA releases up to its own. So H100 and RTX 50 now get native code instead of the compute_89 PTX. - setup.py counts a card as covered by an engine when the engine has its architecture. A newer card counts too, but only when the driver can compile that CUDA's PTX (PTX_DRIVER: 12.8 needs 570, 13.0 needs 580). get_prebuilt otherwise compiles instead, and says a driver update would also do. engine_runs_on (and so ensure_engine_for) and the image encoder's vision_archs check use the same rule. - A compile with a driver below 580 passes over a CUDA 13 toolkit on the PC (find_nvcc(below=(13, 0))), and installs 12.8 when there is none, because a CUDA 13 build would not run on that driver.

architectds · 2026-09-29T23:45:15Z

Thanks, done. main moved to 0.1.27 while I was on it, so the branch is now on 0.1.27, in two commits:

The original change, rebased. The only conflict was the MIN_ENGINE line.
Your points:
- Native sm_90 and sm_120. The workflow builds machine code for sm_80, 86, 89, 90 and 120, with compute_120 PTX for newer cards. It also writes vision_archs into BUILD.json.
- PTX only when the driver can compile it. An engine covers a card when it has the card's architecture, or when the card is newer and the driver can compile that CUDA's PTX. PTX_DRIVER gives 570 for 12.8 and 580 for 13.0. Otherwise get_prebuilt compiles instead and says a driver update would also do. engine_runs_on, and so ensure_engine_for, and the new vision_archs check use the same rule.
- No CUDA 13 toolkit below driver 580. A compile passes over a CUDA 13 toolkit on the PC (find_nvcc(below=(13, 0))) and installs 12.8 when there is none.

Tested:

CI on the rebased branch: run 36639718189, 75 min. The engine took 11 min and the image encoder 56. The zip is 149 MB (it was 99.5 with three architectures), with BUILD.json {"version": "0.1.27", "archs": [80, 86, 89, 90, 120], "ptx": true, "cuda": "12.8", "vision_archs": [...]}.

Coverage rules:

card	driver	result
A100	525	prebuilt
H100	535	prebuilt
RTX 5090	550	prebuilt
a card newer than sm_120	560	compiles
a card newer than sm_120	575	prebuilt
B200 (sm_100), RTX 20	any	compiles

The earlier sm_80/86/89 build on an H100 with driver 535, your case, is now refused.

Toolkit choice: with CUDA 13.0 and 12.8 toolkits both installed, find_nvcc picks 12.8 when the driver is older than 580.
Server tests: serve.test_server passes.

The Windows engine gained RTX 20 in 0.1.27, and the Linux build doesn't have sm_75. I can add it if you'd like the two to match; that's roughly 15 more minutes of CI and a slightly bigger zip.

Niko1221 · 2026-09-30T09:18:34Z

Addressed, thanks. Please add sm_75 to match the Windows engine (0.1.27 supports RTX 20). We'll decide on shipping it with a release.

Niko1221 · 2026-09-30T11:30:34Z

Thanks @architectds, the CI build and the driver/PTX rules are well done. It clashes with 0.1.28, though: #220 found that a CUDA 12.8 engine crashes in the prompt path on RTX 50 under Linux, so sm_120 has to come from CUDA 13.0. A new PR on main would need: sm_120 only from a CUDA 13 build, sm_75 added (RTX 20 is supported since 0.1.27), and existing CUDA 13 Linux installs left unchanged. We'll gate the 12.8 build's output against 13.0 before shipping it. We're closing PRs that wait on changes so the list stays manageable; this is not a no. Please open a new PR when it's ready and link this one.

… grouped on the GPU (opt-in STRATA_PF_FUSED=1) Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32 values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync), then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue) and one down launch into the per-slot rows the existing combine reads. Experts are read where they are (cache slot or streamed ring slot, a pointer per expert): no gather. Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU). Everything else - the default, native/IQ packs, small chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch). tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS 1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5. One layer at 2048 tokens x top 10 over 512 experts: MMQ path 11.12 ms, fused 3.13 ms. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…s experts) by default The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch waited for slots its own first batch frees. With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides). RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 -> 2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%. 1024 slots: 4K +24-27%, but the 32K chunk falls from 8192 to 6144 (+6%). The default path's output is identical to 0.1.35's on both prompts. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…ly what the fused kernels need In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the GDN set does). A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for stream_all_min() - 1 tokens. The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default stays 512). Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events. The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert of the layer, routed or not - at a streamed chunk's size nearly all are. RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots): 4K: MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046) 32K: MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153) 128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346 (8192), + 1024-slot ring 2,392 (6144) A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per 8K tokens. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…o layers' experts) With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots, under the 90% cap). RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to 6144 and is still faster. Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656 (+21.5%), 128K 2,128 -> 2,454 (+15.3%). The ring does not change the numbers (the same output as at 512 slots). Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1) src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert - cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of the native packs. The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which multiply at m16n8k16). The activations are int8 per 32 values in natural order (quantize_act_native); fragments by ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add. A work item is 256 weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128). Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers. sm_75, HIP and the default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy tokens, 4K and 32K). tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 / 2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5), IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 / 2.72e-2 (4.4e-5); both tile shapes. One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192 tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6, 12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5. MMQ's products alone (no gathers) at 3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them. (The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0. prefill_fused_moe_test's timing part has the same pattern.) End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K 1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%). The IQ3 packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s of 3.4 s), so faster kernels change little. Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at 162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89, IQ3_S 31 / 256. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…mpt path, the split stages off the RAM budget with Niko1221#467's what-fits fallback Conflicts: - prefill.cpp: upstream's fused int8 experts (Niko1221#136) wrap the host grouping in an else branch; the local CPU tier (STRATA_PREFILL_CPU_ROWS) disables stream_all, so it lives in that branch only. - generate.cpp: Niko1221#467's soft resident fallback takes the stage pairs and every stage's profile too. - gemm.cu, sampler.{hpp,cu}, server.py: both sides kept. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…s own slot count The ring was 384 slots - a number measured on Q2_0, where a slot is one expert blob of 1,382,400 B, so 506 MiB. A slot is a whole blob, so on a pack with bigger blobs the same 384 slots are a different amount of memory: 1,912 MiB on Q8_0, more than half of an 8 GB card's expert cache, and what kept that card at a 1024-token chunk when its own buffers would have fitted 6144. Measured on the 4-way rig, one env var and nothing else: chunk 1024 -> 6144, prefill 87 -> 402 tok/s. So the budget is the constant and the slot count is derived from the pack's max blob: 384 on Q2_0 (pinned) or 96 (not), exactly as before, and the pack it was tuned on does not move. The fused rule (1024 slots, #136) is preserved the same way, and `ring_cap()` still bounds the result. This is also what lets the auto chunk scan treat "the ring is full" as a plain byte comparison. Also exports the three accessors the scan needs: bytes_needed_no_ring (the chunk's own buffers, without the ring), ring_max_slots (the cap in slots) and ring_slots_for (what the ring resolves to for a chunk, after the override, the env var and the pinned-share rule - what `init` lays out and what the INFO line reports). And one fix on the way: `bytes_needed` counted `carve`'s buffers from memory and got two wrong - it took `xn` unconditionally, though `carve` takes it only under STRATA_GR_UNFUSED, and it never took `grs`. Net T*(D-HC)*4 bytes too many: 42 MB at a 1024-token chunk. The direction is safe - the prompt path was told it had less room than it did - but it under-sizes every loan. Co-Authored-By: Claude Code <noreply@anthropic.com>

… grouped on the GPU (opt-in STRATA_PF_FUSED=1) Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32 values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync), then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue) and one down launch into the per-slot rows the existing combine reads. Experts are read where they are (cache slot or streamed ring slot, a pointer per expert): no gather. Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU). Everything else - the default, native/IQ packs, small chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch). tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS 1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5. One layer at 2048 tokens x top 10 over 512 experts: MMQ path 11.12 ms, fused 3.13 ms.

…s experts) by default The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch waited for slots its own first batch frees. With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides). RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 -> 2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%. 1024 slots: 4K +24-27%, but the 32K chunk falls from 8192 to 6144 (+6%). The default path's output is identical to 0.1.35's on both prompts.

…ly what the fused kernels need In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the GDN set does). A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for stream_all_min() - 1 tokens. The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default stays 512). Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events. The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert of the layer, routed or not - at a streamed chunk's size nearly all are. RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots): 4K: MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046) 32K: MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153) 128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346 (8192), + 1024-slot ring 2,392 (6144) A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per 8K tokens.

…o layers' experts) With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots, under the 90% cap). RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to 6144 and is still faster. Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656 (+21.5%), 128K 2,128 -> 2,454 (+15.3%). The ring does not change the numbers (the same output as at 512 slots). Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512).

…IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1) src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert - cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of the native packs. The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which multiply at m16n8k16). The activations are int8 per 32 values in natural order (quantize_act_native); fragments by ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add. A work item is 256 weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128). Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers. sm_75, HIP and the default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy tokens, 4K and 32K). tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 / 2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5), IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 / 2.72e-2 (4.4e-5); both tile shapes. One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192 tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6, 12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5. MMQ's products alone (no gathers) at 3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them. (The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0. prefill_fused_moe_test's timing part has the same pattern.) End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K 1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%). The IQ3 packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s of 3.4 s), so faster kernels change little. Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at 162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89, IQ3_S 31 / 256.

…s own slot count The ring was 384 slots - a number measured on Q2_0, where a slot is one expert blob of 1,382,400 B, so 506 MiB. A slot is a whole blob, so on a pack with bigger blobs the same 384 slots are a different amount of memory: 1,912 MiB on Q8_0, more than half of an 8 GB card's expert cache, and what kept that card at a 1024-token chunk when its own buffers would have fitted 6144. Measured on the 4-way rig, one env var and nothing else: chunk 1024 -> 6144, prefill 87 -> 402 tok/s. So the budget is the constant and the slot count is derived from the pack's max blob: 384 on Q2_0 (pinned) or 96 (not), exactly as before, and the pack it was tuned on does not move. The fused rule (1024 slots, #136) is preserved the same way, and `ring_cap()` still bounds the result. This is also what lets the auto chunk scan treat "the ring is full" as a plain byte comparison. Also exports the three accessors the scan needs: bytes_needed_no_ring (the chunk's own buffers, without the ring), ring_max_slots (the cap in slots) and ring_slots_for (what the ring resolves to for a chunk, after the override, the env var and the pinned-share rule - what `init` lays out and what the INFO line reports). And one fix on the way: `bytes_needed` counted `carve`'s buffers from memory and got two wrong - it took `xn` unconditionally, though `carve` takes it only under STRATA_GR_UNFUSED, and it never took `grs`. Net T*(D-HC)*4 bytes too many: 42 MB at a 1024-token chunk. The direction is safe - the prompt path was told it had less room than it did - but it under-sizes every loan.

…pt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1) Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32 values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync), then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue) and one down launch into the per-slot rows the existing combine reads. Experts are read where they are (cache slot or streamed ring slot, a pointer per expert): no gather. Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU). Everything else - the default, native/IQ packs, small chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch). tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS 1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5. One layer at 2048 tokens x top 10 over 512 experts: MMQ path 11.12 ms, fused 3.13 ms.

…e layer's experts) by default The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch waited for slots its own first batch frees. With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides). RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 -> 2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%. 1024 slots: 4K +24-27%, but the 32K chunk falls from 8192 to 6144 (+6%). The default path's output is identical to 0.1.35's on both prompts.

… hold only what the fused kernels need In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the GDN set does). A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for stream_all_min() - 1 tokens. The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default stays 512). Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events. The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert of the layer, routed or not - at a streamed chunk's size nearly all are. RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots): 4K: MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046) 32K: MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153) 128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346 (8192), + 1024-slot ring 2,392 (6144) A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per 8K tokens.

…ault (two layers' experts) With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots, under the 90% cap). RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to 6144 and is still faster. Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656 (+21.5%), 128K 2,128 -> 2,454 (+15.3%). The ring does not change the numbers (the same output as at 512 slots). Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512).

…gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1) src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert - cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of the native packs. The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which multiply at m16n8k16). The activations are int8 per 32 values in natural order (quantize_act_native); fragments by ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add. A work item is 256 weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128). Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers. sm_75, HIP and the default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy tokens, 4K and 32K). tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 / 2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5), IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 / 2.72e-2 (4.4e-5); both tile shapes. One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192 tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6, 12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5. MMQ's products alone (no gathers) at 3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them. (The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0. prefill_fused_moe_test's timing part has the same pattern.) End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K 1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%). The IQ3 packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s of 3.4 s), so faster kernels change little. Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at 162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89, IQ3_S 31 / 256.

…s own slot count The ring was 384 slots - a number measured on Q2_0, where a slot is one expert blob of 1,382,400 B, so 506 MiB. A slot is a whole blob, so on a pack with bigger blobs the same 384 slots are a different amount of memory: 1,912 MiB on Q8_0, more than half of an 8 GB card's expert cache, and what kept that card at a 1024-token chunk when its own buffers would have fitted 6144. Measured on the 4-way rig, one env var and nothing else: chunk 1024 -> 6144, prefill 87 -> 402 tok/s. So the budget is the constant and the slot count is derived from the pack's max blob: 384 on Q2_0 (pinned) or 96 (not), exactly as before, and the pack it was tuned on does not move. The fused rule (1024 slots, Niko1221#136) is preserved the same way, and `ring_cap()` still bounds the result. This is also what lets the auto chunk scan treat "the ring is full" as a plain byte comparison. Also exports the three accessors the scan needs: bytes_needed_no_ring (the chunk's own buffers, without the ring), ring_max_slots (the cap in slots) and ring_slots_for (what the ring resolves to for a chunk, after the override, the env var and the pinned-share rule - what `init` lays out and what the INFO line reports). And one fix on the way: `bytes_needed` counted `carve`'s buffers from memory and got two wrong - it took `xn` unconditionally, though `carve` takes it only under STRATA_GR_UNFUSED, and it never took `grs`. Net T*(D-HC)*4 bytes too many: 42 MB at a 1024-token chunk. The direction is safe - the prompt path was told it had less room than it did - but it under-sizes every loan.

…pt path, grouped on the GPU (opt-in STRATA_PF_FUSED=1) Strata's own kernels (src/prefill/moe_fused.cu): the layer's input rounded to int8 once per token (a scale per 32 values), the (token, k) pairs grouped by expert on the GPU (counts, offsets, a 64-row tile list - no host sync), then per batch of up to 128 experts one gate/up launch (mma.sync m16n8k32 s8, the 2-bit codes unpacked in registers, Q2_0's offset folded into the activation blocks' sums, SwiGLU and H's int8 rounding in the epilogue) and one down launch into the per-slot rows the existing combine reads. Experts are read where they are (cache slot or streamed ring slot, a pointer per expert): no gather. Used only on the streamed walk (chunks from stream_all_min()) of the canonical Q2_0 layout on sm_80+ CUDA, with STRATA_PF_FUSED=1; it reuses the MMQ buffers (Xq, H, GU). Everything else - the default, native/IQ packs, small chunks, sm_75, HIP - runs the MMQ path unchanged (the old block is only wrapped in the else branch). tests/cuda/prefill_fused_moe_test (exit 77 below sm_80), RTX 5070: vs a double-precision reference, MMQ rel RMS 1.144e-2 (worst row 4.15e-2), fused 1.144e-2 (4.15e-2); fused vs MMQ 9.0e-5. One layer at 2048 tokens x top 10 over 512 experts: MMQ path 11.12 ms, fused 3.13 ms. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…e layer's experts) by default The fused path launches on batches of a layer's streamed experts, so with the 384-slot ring a layer's last batch waited for slots its own first batch frees. With STRATA_PF_FUSED=1 (and the canonical Q2_0 layout) the pinned-RAM ring is 512 slots; without it nothing changes (384 / 96 as before; STRATA_PREFILL_RING still overrides). RTX 5070, Q2_0 (cfg-q2rel), greedy, 3 alternating pairs, prompt tok/s: 4K 1,275 -> 1,516 (+18.9%), 32K 2,176 -> 2,422 (+11.3%); at 384 slots the fused path gave +4-7% / +3%. 1024 slots: 4K +24-27%, but the 32K chunk falls from 8192 to 6144 (+6%). The default path's output is identical to 0.1.35's on both prompts. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

… hold only what the fused kernels need In a layout whose chunks run the fused experts (STRATA_PF_FUSED=1, the Q2_0 pack, a streamed chunk) GU holds the grouping tables, H the int8 H rows and Xq the per-token int8 activations, and Hq nothing - instead of MMQ's FP32 GU / H and per-slot q8_1 rows: ~100 KB a token less in the MoE set, which then no longer sets the shared region's size (the GDN set does). A last chunk below stream_all_min() still runs MMQ in the same buffers, so each keeps MMQ's size for stream_all_min() - 1 tokens. The ring may grow to 1024 slots with the fused experts (STRATA_PREFILL_RING; the default stays 512). Without STRATA_PF_FUSED nothing changes: the same buffers, ring cap (512) and events. The fused walk now counts the experts it reads from the cache in experts_resident (it stayed 0): every resident expert of the layer, routed or not - at a streamed chunk's size nearly all are. RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, 2 alternating pairs, prompt tok/s (borrowed slots): 4K: MMQ 1,305 (1621), fused 1,517 (1749), fused + P3 1,512 (1534), + 1024-slot ring 1,605 (2046) 32K: MMQ 2,177 (3071), fused 2,453 (3199), fused + P3 2,485 (2641), + 1024-slot ring 2,660 (3153) 128K (--max-context 131072 --kv-resident 32768): MMQ 2,108 (chunk 8192), fused 2,140 (6144), fused + P3 2,346 (8192), + 1024-slot ring 2,392 (6144) A 16K chunk still does not fit the 12 GB card: at 8192 tokens the loan is 2,641 of 3,690 slots, ~2.1K of them per 8K tokens. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…ault (two layers' experts) With P3's smaller buffers the 1024-slot ring fits beside an 8192-token chunk at 32K (borrowed 3,153 of 3,690 slots, under the 90% cap). RTX 5070, Q2_0 (cfg-q2rel), --prefill auto, prompt tok/s at 512 / 1024 slots (2 alternating pairs): 4K 1,512 / 1,605, 32K 2,485 / 2,660, 128K (KV streaming) 2,346 / 2,392 - there the chunk falls from 8192 to 6144 and is still faster. Against MMQ (one more pair on this build): 4K 1,299 -> 1,609 (+23.9%), 32K 2,187 -> 2,656 (+21.5%), 128K 2,128 -> 2,454 (+15.3%). The ring does not change the numbers (the same output as at 512 slots). Without STRATA_PF_FUSED nothing changes (384 / 96 slots, cap 512). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…gate/up IQ2_XXS IQ2_XS IQ2_S IQ3_XXS IQ3_S IQ4_XS, down Q2_0 IQ4_NL; opt-in STRATA_PF_FUSED=1) src/prefill/moe_fused_iq.cu: the Q2_0 path's design (moe_fused.cu: the grouping on the GPU, a pointer per expert - cache slot or ring slot, no gather -, SwiGLU and H's int8 rounding in the gate/up epilogue) for the raw GGUF blobs of the native packs. The load stage decodes the i-quant blocks to int8 in shared memory (llama.cpp's codebooks and sign tables, held in shared memory; a 32-value sub-block per thread per 64 values of K, read one stage ahead with 16-bit loads and two 256-value blocks ahead into L2), with each sub-block's scale (per 16 values for IQ2_XS / IQ2_S, which multiply at m16n8k16). The activations are int8 per 32 values in natural order (quantize_act_native); fragments by ldmatrix; the magic bias 1.5 * 2^23 enters as the mma's C, so a dot becomes a float with one add. A work item is 256 weight rows x 64 routed rows, or 128 x 128 when the experts have 56-112 rows each (STRATA_PF_FUSED_TILE=64|128). Wiring (prefill.cpp): a native layer takes these kernels when both its formats are covered, else MMQ (IQ1_M keeps the FP16 path); the 512-slot ring of the fused path now also for a native pack with such layers. sm_75, HIP and the default (no variable) are unchanged: on IQ2_XS, IQ3_XXS and IQ3_S the default output is 0.1.35's (256/256 greedy tokens, 4K and 32K). tests/cuda/prefill_fused_iq_test (exit 77 below sm_80), RTX 5070, 256 tokens x 64 experts against a double-precision reference from ggml's to_float, rel RMS / worst row - MMQ and fused equal to four digits: IQ2_S/Q2_0 1.489e-2 / 2.93e-2 (fused vs MMQ 2.6e-4), IQ2_XXS/Q2_0 1.464e-2 / 2.96e-2 (1.9e-4), IQ2_XS/IQ4_NL 1.465e-2 / 2.41e-2 (3.5e-5), IQ3_XXS/IQ4_NL 1.458e-2 / 2.73e-2 (1.6e-4), IQ3_S/IQ4_NL 1.470e-2 / 2.65e-2 (6.5e-5), IQ4_XS/Q2_0 1.446e-2 / 2.72e-2 (4.4e-5); both tile shapes. One layer (512 experts, top 10), ms, MMQ path -> fused at 2048 / 3584 / 8192 tokens: IQ2_S 10.5 -> 5.2, 12.5 -> 9.3, 21.4 -> 17.5; IQ2_XXS 8.8 -> 4.2, 10.7 -> 7.6, 17.9 -> 13.6; IQ3_S 10.6 -> 5.6, 12.3 -> 8.8, 19.8 -> 18.4; IQ3_XXS 10.2 -> 5.1, 11.9 -> 8.3, 19.1 -> 16.5. MMQ's products alone (no gathers) at 3584 / 8192: IQ2_S 8.9 / 17.9, IQ3_S 8.1 / 15.6: the kernels are about even with MMQ's, the gain is the path around them. (The test syncs the device after its uploads: a cudaMemcpy from pageable memory on the legacy stream can still be queued on Windows while a non-blocking stream runs - the grouping then read zeroed ids, every pair on expert 0. prefill_fused_moe_test's timing part has the same pattern.) End to end, RTX 5070, speed_135.py's settings, greedy, A (MMQ) / B (STRATA_PF_FUSED=1) alternating, 2 pairs, prompt tok/s: IQ2_XS 4K 1,228 -> 1,378 (+12%; a cold first A run of 590 left out), 32K 2,107 -> 2,170 (+3.0%); IQ3_XXS 4K 1,001 -> 1,020 (+1.8%), 32K 1,745 -> 1,739 (-0.4%); IQ3_S 4K 927 -> 925 (-0.2%), 32K 1,382 -> 1,362 (-1.4%). The IQ3 packs' prompts here wait on the host copies of their unpinned experts (IQ3_XXS 4K: 5,062 of 22,573 streamed, ~3.0 s of 3.4 s), so faster kernels change little. Each arm repeats itself (256/256; IQ3_XXS 32K B and IQ3_S 32K A split at 162 / 223, the decode's timing); A vs B greedy agreement, identical prefix of 256: IQ2_XS 204 / 88, IQ3_XXS 182 / 89, IQ3_S 31 / 256. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN

…s own slot count The ring was 384 slots - a number measured on Q2_0, where a slot is one expert blob of 1,382,400 B, so 506 MiB. A slot is a whole blob, so on a pack with bigger blobs the same 384 slots are a different amount of memory: 1,912 MiB on Q8_0, more than half of an 8 GB card's expert cache, and what kept that card at a 1024-token chunk when its own buffers would have fitted 6144. Measured on the 4-way rig, one env var and nothing else: chunk 1024 -> 6144, prefill 87 -> 402 tok/s. So the budget is the constant and the slot count is derived from the pack's max blob: 384 on Q2_0 (pinned) or 96 (not), exactly as before, and the pack it was tuned on does not move. The fused rule (1024 slots, Niko1221#136) is preserved the same way, and `ring_cap()` still bounds the result. This is also what lets the auto chunk scan treat "the ring is full" as a plain byte comparison. Also exports the three accessors the scan needs: bytes_needed_no_ring (the chunk's own buffers, without the ring), ring_max_slots (the cap in slots) and ring_slots_for (what the ring resolves to for a chunk, after the override, the env var and the pinned-share rule - what `init` lays out and what the INFO line reports). And one fix on the way: `bytes_needed` counted `carve`'s buffers from memory and got two wrong - it took `xn` unconditionally, though `carve` takes it only under STRATA_GR_UNFUSED, and it never took `grs`. Net T*(D-HC)*4 bytes too many: 42 MB at a 1024-token chunk. The direction is safe - the prompt path was told it had less room than it did - but it under-sizes every loan. Co-Authored-By: Claude Code <noreply@anthropic.com>

architectds force-pushed the pr/linux-prebuilt branch from 7421b04 to 1663cce Compare September 29, 2026 13:26

Niko1221 mentioned this pull request Sep 29, 2026

Ampere (A100) and CPUs without AVX-512: a Linux prebuilt with CUDA 12, PRs welcome? #103

Closed

architectds force-pushed the pr/linux-prebuilt branch from 1663cce to b1795dd Compare September 29, 2026 22:25

architectds added 2 commits September 29, 2026 17:26

architectds force-pushed the pr/linux-prebuilt branch from b1795dd to 896f75e Compare September 29, 2026 22:27

architectds mentioned this pull request Sep 29, 2026

CI check: linux-prebuilt on 0.1.27 (sm_80/86/89/90/120 + PTX) architectds/Strata#2

Closed

Niko1221 mentioned this pull request Sep 30, 2026

setup: Volta (V100, sm_70) and Turing as an experimental build path #130

Closed

Niko1221 closed this Sep 30, 2026

architectds deleted the pr/linux-prebuilt branch October 2, 2026 15:46

q8atnight mentioned this pull request Oct 2, 2026

multi-GPU: second GPU as an opt-in expert-cache tier (--peer-device), resubmit of #229 on v0.1.36 #531

Closed

sergqwer mentioned this pull request Oct 2, 2026

NVFP4 routed experts on 0.1.35 (experimental): 120a only for the W4A4 unit (rebase of #292) #353

Closed

q8atnight mentioned this pull request Oct 2, 2026

prefill: a group's expert gathers in one launch (bit-identical, long prompts +4%) #439

Closed

gopinath87607 mentioned this pull request Oct 3, 2026

Prefill chunk ring #583

Closed

rafatxf mentioned this pull request Oct 12, 2026

prefill: the fused int8 prompt experts on Turing (sm_75), opt-in (+9 % prefill on RTX 2080 Ti) #2145

Open

Navigation Menu

Repository navigation

Linux prebuilt from CI: CUDA 12.8, sm_80/86/89 + PTX, attached to each release - #136

Linux prebuilt from CI: CUDA 12.8, sm_80/86/89 + PTX, attached to each release#136
architectds wants to merge 2 commits into
Niko1221:mainfrom
architectds:pr/linux-prebuilt

architectds commented Sep 29, 2026

Uh oh!

Niko1221 commented Sep 29, 2026

Uh oh!

architectds commented Sep 29, 2026

Uh oh!

Niko1221 commented Sep 30, 2026

Uh oh!

Niko1221 commented Sep 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

architectds commented Sep 29, 2026

Uh oh!

Niko1221 commented Sep 29, 2026

Uh oh!

architectds commented Sep 29, 2026

Uh oh!

Niko1221 commented Sep 30, 2026

Uh oh!

Niko1221 commented Sep 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants