Repository navigation
Conversation
…n one kernel (SM60 build, no caller yet) swiglu_kernel (n_ff 640), then quantize_mmq_q8_1<D4>'s rows of its output, in one pass: a block per row (128 threads, 4 values each, blockIdx.y the row's 512-value half), H never stored. Each value is swiglu_kernel's expression and each 32-value block quantize_mmq_q8_1's (the same unit and flags: the same bytes in hq). swiglu_quant_ok(type) says where it applies (a q8_1 D4 down type, no FP4 activations). Built only in the STRATA_EXPERIMENTAL_SM60 CUDA build; HIP and the other builds get swiglu_quant_ok() == false and a no-op. Nothing calls it in this commit, so every launch is the one before. Test: prefill_swiglu_quant_test (STRATA_BUILD_TESTS, byte for byte against swiglu + quantize; exit 77 where the kernel is not built). Partial port of a larger prompt-path fusion change: only the swiglu + q8_1 kernel, which the weight-stationary expert path uses for all of a layer's rows at once. The per-group wiring, the MoE combine + hyper-connection write fusion and the FP16 dequantization part are left out.
…ch per layer and product; not used yet) On GP100 (sm_60, no dp4a) the prompt path's MoE experts are a gate/up, swiglu and down MMQ launch per 16-expert group, each MMQ tile decoding its weight blocks again, and every expert of a group padded to the group's largest row count. When a layer's experts are all resident in VRAM they can be read where they are and run weight-stationary. moe_ws: the Swift pack's expert products as CTAs of 128 (or 256) weight rows of one expert against up to 64 of its activation rows, a thread per weight row: each weight block is decoded once per pass instead of once per MMQ tile, and an expert runs its own row count. Gate/up IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S [1280 x 2560]; down Q2_0, IQ4_NL [2560 x 640], the K loop stopping at 640. Every product keeps MMQ's int8 values, scales and per-32-value epilogue in its K order (a 32-value block's integer sum is exact, so only the float epilogue's order matters), so GU, H and Dm are the per-group MMQ's bits. Weights are read in place (a per-layer offset table, + the down offset); the CTAs take 64-row chunks of the experts, largest first (ws::chunks). - Built only in the STRATA_EXPERIMENTAL_SM60 build (CMake adds the four files to strata_mmq there; not under HIP). - ws::mode(): 0 unless STRATA_MMQ_WS=1, and then only where every visible GPU is cc 6.0 (the kernels were tuned and checked on GP100 only). - ws::faster carries the measured ws / per-group MMQ ratios per format (one GP100 die, the 0.1.41-based tree): 0.25-0.35 at 2.5 rows per routed expert, 0.38-0.52 at 10, 0.52-0.76 at 35, 0.61-0.82 at 160. Nothing calls these yet; the prompt path takes them in the next change.
… product (STRATA_MMQ_WS=1, GP100) A layer whose routed experts are ALL resident in the expert cache, on this GPU alone - no peer share, no CPU share, no streamed entries of the layer, not the fused layout, not a direct-read (STRATA_MMQ_DIRECT) layer - and whose formats have the moe_ws kernels runs: quantize_scatter -> gate/up as one weight-stationary launch -> ONE swiglu_quant over all T*K rows into Hq -> down as one launch into Dm's rows, instead of a gate/up, swiglu and down MMQ launch per 16-expert group. The experts are read where they lie: the layer's offsets from the cache's base go up with the bounds (after them in the mapped table, or through the pinned upload ring), and the CTAs' chunk table after them, 16-byte aligned. A layer is taken only where ws::faster says both products beat the per-group MMQ at its rows per routed expert, so the weight-stationary gate/up never declines a layer it was picked for. Every other layer, pack and GPU keeps its path; kPfGemmGU / kPfGemmD mark the new launches; the first such layer says so on stderr. Opt-in and GP100 only: STRATA_PREFILL_WS is defined only in the STRATA_EXPERIMENTAL_SM60 CUDA build (elsewhere a stub says mode 0), ws::mode() is 0 unless STRATA_MMQ_WS=1 and every visible GPU is cc 6.0, and the offset and chunk buffers (the mapped table's extra room, carve's and bytes_needed_impl's takes) exist only where it is not 0. Unset, nothing is allocated and every launch is the one before. GU, H (all 8 q8_1 blocks of every row) and Dm keep MMQ's bits: on the 0.1.41-based tree, layer-level parity on a GP100 die for all 10 (gate/up, down) pairs of the Swift pack at T = 128 / 512 / 1792 / 8192. There the prompt read was 24% faster at ~1K tokens and 14% at 30K.
ruibeikaa
marked this pull request as ready for review
October 22, 2026 18:57
… instructions per pass (opt-in, same bits)
The weight-stationary kernels issue as many instructions as GP100 can, so their time follows the instruction count,
and part of it is per-pass overhead rather than the inner MACs. nl4::dn_kernel_fast, picked only with
STRATA_WS_FAST=1 for the tier that runs at u >= 20 rows per routed expert (prompt chunks of ~1024 tokens and up):
DUFF a switch on the number of token pairs falls through the 16 unrolled pairs (a jump table) instead of a guard
per pair and block
DYT dy as [block][token]: a pair's two dy in one 8-byte shared load (122 -> 102 registers, still two CTAs per SM)
Every token keeps its own operations in their order (the same XMAD chains, the same epilogue), so Dm is the default
kernel's bytes. Below u 20, for Q2_0 and for the gate/up kernels nothing changes. Unset, the kernels are picked as
before; one log line when on.
moe_ws_down_bench (EXCLUDE_FROM_ALL, the SM60 build): a layer's 512 real down experts from the GGUF, lognormal and
flat-32 routings, ws::down with and without the switch compared byte for byte and timed.
This was referenced Oct 23, 2026
Niko1221
pushed a commit
that referenced
this pull request
Oct 27, 2026
…n one kernel (#3752) Problem: the prompt path stores H (swiglu of the gate/up rows) as floats and then quantizes it to the q8_1 rows the down product reads: two kernels and an H round trip per MMQ group. Change: swiglu_quant_kernel, a block per row (128 threads, 4 values each, blockIdx.y the row's 512-value half), computes swiglu_kernel's expression for each value and quantize_mmq_q8_1<D4>'s result for each 32-value block, so hq holds the same bytes and H is never stored. swiglu_quant_ok(type) says where it applies (a q8_1 D4 down type, no FP4 activations). Built only in the CUDA build with STRATA_EXPERIMENTAL_SM60; HIP and the other builds get swiglu_quant_ok() == false and a no-op. Nothing calls it in this commit; the weight-stationary path of the next commit does. prefill.cpp gets the matching stubs for builds without the MMQ sources. Tests: prefill_swiglu_quant_test (new, registered under STRATA_BUILD_TESTS in the SM60 build, skips with 77 without a device) compares the fused kernel with swiglu + quantize byte for byte for Q2_0, IQ4_NL and Q8_0 down types, split and interleaved gate/up rows, a zero row and a row with one huge value. Run on z440 (RTX A4000, sm_86, CUDA 13.0, SM60 build, gpu lock): all 18 cases print "0 of N bytes differ" (1, 37 and 4113 rows) and the test ends "prefill_swiglu_quant_test: ok". sm_60 compile of the tree (CUDA 12.9, P1004): exit 0. This is a partial port of another contributor's prompt-path fusion change; only the kernel is taken, as in PR 3752.
Niko1221
pushed a commit
that referenced
this pull request
Oct 27, 2026
…n STRATA_MMQ_WS=1 (#3752) Problem: on a GP100 die (compute capability 6.0: no dp4a, no tensor cores) the prompt pass's routed experts run as MMQ launches per 16-expert group: a gate/up, a swiglu, a quantize and a down launch per group, each group's experts gathered and padded to the group's largest row count, every MMQ tile decoding its weight blocks again. With every expert resident the per-group overhead, not the weight bytes, is a large share of a prompt read on this card. Change: with STRATA_MMQ_WS=1, a layer whose routed experts are all resident on this GPU runs quantize_scatter, then gate/up as ONE weight-stationary launch over all its experts (read where they lie in the expert cache), ONE swiglu_quant over all its rows, then the down as ONE launch. New src/prefill/moe_ws*.cu(h): a CTA holds 128 (or 256) weight rows of one expert and runs them against up to 64 of that expert's activation rows, a thread per weight row, so each weight block is decoded once per pass; gate/up for IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS and IQ3_S, down for Q2_0 and IQ4_NL. Every product keeps MMQ's int8 values, scales and per-32-value epilogue in its K order, so GU, H and the down rows are MMQ's bits. ws::faster carries the author's measured ws / MMQ time ratios per format and picks a layer only where both products win at its rows per routed expert. prefill.cpp decides per layer after quantize_scatter (all experts resident, this GPU alone, no peer or CPU share, not the fused layout, not STRATA_MMQ_DIRECT, no Hadamard, no NVFP4 tail, no FP4 activations, aligned, not STRATA_DBG_NAN); every other layer keeps its path. STRATA_WS_FAST=1 (second opt-in, same bits) selects a down kernel for the IQ4_NL 32-token tier with fewer instructions per pass. Gating: the kernels and STRATA_PREFILL_WS exist only in the CUDA build with STRATA_EXPERIMENTAL_SM60 (not HIP, not Metal; prefill.cpp links stubs that answer mode 0 elsewhere). At run time ws::mode() is 1 only with STRATA_MMQ_WS=1 and every visible GPU at compute capability 6.0. Unset, nothing is allocated (the offset and chunk buffers, under 1 MiB per die, exist only when the mode is not 0), the arena layout and the mapped table are the ones before, and the per-layer test short-circuits on the cached mode. Not ported: moe_ws_down_bench (needs a model shard) and its CMake target. Added: ws::set_mode (tests only) and tests/cuda/prefill_ws_parity_test.cpp, a synthetic layer-level check: random experts of all 5 gate/up x 2 down formats, random routing, the per-group MMQ (mmq::Context per expert) against ws::gate_up + swiglu_quant + ws::down, GU, Hq and Dm compared byte for byte, the down also alone on the reference H and with STRATA_WS_FAST on. It gives its verdict on a compute capability 6.0 card only (the kernels follow MMQ's sm_60 SASS); other cards run it for information with --any. Tests: prefill_ws_parity_test on boston (2x RTX A6000, sm_86, CUDA 13.1, gpu lock, SM60 build, --any because it is not a GP100): 40 cases (all 5 gate/up x 2 down formats, four routings each), 0 declined; the largest GU and down-alone error against that card's MMQ is 5.2e-07 of the largest value (float rounding: MMQ's sm_86 epilogue is not the sm_60 one the kernels follow), the full chain's Dm differs up to 1e-3 where int8 rounding of H flips on those rounding differences, and STRATA_WS_FAST output equals the default kernel's in all cases, 0 differing bytes. The byte-for-byte verdict against MMQ is the P100's. sm_60 compile of the whole tree (CUDA 12.9, P1004): exit 0; sm_86 build without the SM60 option (z440): exit 0. SASS of the 10 kernels moe_mmq.cu had before this change is identical in the sm_60 build. Not covered: no GP100 run here. The contributor's numbers (prompt reads 31-37% shorter on four GP100 dies with every expert resident: 240 tokens 2,156 -> 1,364 ms, a fresh 26K read 41,231 -> 28,555 ms; one run per arm) were not re-measured. The model gate is phaseA-tests/gh027/gates/3752-mmq-ws.sh.
Owner
|
Thank you. Fixed in 0.2.7: A resident layer's experts weight-stationary on GP100, opt-in STRATA_MMQ_WS=1. Your name is on the commit. Release: https://github.com/Niko1221/Strata/releases/tag/v0.2.7 |
sakuumomo
added a commit
to sakuumomo/Strata
that referenced
this pull request
Oct 28, 2026
…1#3752 (GP100 weight-stationary experts) landed with the 0.2.7 rebase, after this file was written. The kernels are pure PTX (prmt/mad.wide/f16x2 asm): on the CUDA side they are strata_mmq sources only under STRATA_EXPERIMENTAL_SM60, gated at run time to machines whose every GPU is cc 6.0. A hip-src copy of them fails on 'invalid output constraint =f' / 'invalid instruction' in the gfx1150 pass, and could never dispatch on an AMD device - so they belong beside the gdn_chunk/mmq_nvfp4/mmq_kq_lib exclusions. moko flavor build was the witness: moe_ws_down/gu_iq2/gu_iq3 were the only FAILED targets.
sakuumomo
added a commit
to sakuumomo/Strata
that referenced
this pull request
Oct 28, 2026
…f the HIP glob - upstream Niko1221#3752 (GP100 weight-stationary experts) landed with the 0.2.7 rebase, after this file was written. The kernels are pure PTX (prmt/mad.wide/f16x2 asm): on the CUDA side they are strata_mmq sources only under STRATA_EXPERIMENTAL_SM60, gated at run time to machines whose every GPU is cc 6.0. A hip-src copy of them fails on 'invalid output constraint =f' / 'invalid instruction' in the gfx1150 pass, and could never dispatch on an AMD device - so they belong beside the gdn_chunk/mmq_nvfp4/mmq_kq_lib exclusions. moko flavor build was the witness: moe_ws_down/gu_iq2/gu_iq3 were the only FAILED targets.
sakuumomo
added a commit
to sakuumomo/Strata
that referenced
this pull request
Oct 30, 2026
…1#3752 (GP100 weight-stationary experts) landed with the 0.2.7 rebase, after this file was written. The kernels are pure PTX (prmt/mad.wide/f16x2 asm): on the CUDA side they are strata_mmq sources only under STRATA_EXPERIMENTAL_SM60, gated at run time to machines whose every GPU is cc 6.0. A hip-src copy of them fails on 'invalid output constraint =f' / 'invalid instruction' in the gfx1150 pass, and could never dispatch on an AMD device - so they belong beside the gdn_chunk/mmq_nvfp4/mmq_kq_lib exclusions. moko flavor build was the witness: moe_ws_down/gu_iq2/gu_iq3 were the only FAILED targets.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Problem and result
On a GP100 die (compute capability 6.0: 732 GB/s HBM2, no
__dp4a, no tensor cores) the prompt pass's routed experts run as MMQ launches per 16-expert group: a gate/up, a swiglu, a quantize and a down launch per group, each group's experts gathered and padded to the group's largest row count, and every MMQ tile decoding its weight blocks again. With every expert resident in VRAM, that per-group overhead rather than the weight bytes is a large share of a prompt read on this card: avoiding it removes about a third of the read time below.With
STRATA_MMQ_WS=1, a layer whose routed experts are all resident on this GPU runsquantize_scatter-> gate/up as one weight-stationary launch over all its experts, read where they lie in the expert cache -> oneswiglu_quantover all its rows -> down as one launch, instead of the per-group launches. On four GP100 dies with every expert resident, prompt reading drops by 31-37% (240 tokens 2,156 -> 1,364 ms, 1.15K 4,210 -> 2,760 ms, a fresh 26K read 41,231 -> 28,555 ms, a 1.1K append to a 26K conversation 4,873 -> 3,257 ms; one run per arm, details below); decode is unchanged. The products keep MMQ's int8 values, scales and epilogue order, so GU, H and the down rows are MMQ's bits: every reply and every engine state hash stays the same. Unset (the default), nothing is allocated and every launch is the one before.Related discussion
This is the prompt-path counterpart of the GP100 decode work in #3455 (scope issue #3454): that PR keeps the prompt path's MMQ gather unchanged; this one replaces the gather and per-group launches for fully resident layers. The two are independent: this branch is main plus three commits and does not need #3455. If #3455 is merged first, its 4-bit gate/up layers must be excluded from this path (one condition, since their slots no longer hold GGUF blocks) or served by a 4-bit variant of the weight-stationary gate/up, which I have measured and will send as a follow-up once both are in. The scope is the same as #3454's (opt-in, GP100 only, default path byte-identical); I am happy to open a separate scope issue if you prefer.
Approach
mmq::swiglu_quant): swiglu and the down product's q8_1 rows in one kernel, byte for byteswiglu+quantize's output, with a test (prefill_swiglu_quant_test, underSTRATA_BUILD_TESTS). It is a partial port of another contributor's prompt-path fusion change (kept under their name): only this kernel, no caller of its own. Built only in theSTRATA_EXPERIMENTAL_SM60CUDA build.moe_ws, kernels only): newsrc/prefill/moe_ws.{hpp,cu,cuh},moe_ws_down.cu,moe_ws_gu_iq2.cu,moe_ws_gu_iq3.cu. A CTA is 128 (or 256) weight rows of one expert against up to 64 of that expert's activation rows, a thread per weight row: each weight block is decoded once per pass instead of once per MMQ tile, and each expert runs its own row count. Gate/up for IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS, IQ3_S; down for Q2_0 and IQ4_NL. Every product keeps MMQ's int8 values, scales and per-32-value epilogue in its K order (a 32-value block's integer sum is exact, so only the float epilogue's order matters).ws::fastercarries the measured ws / per-group MMQ time ratios per format (one GP100 die) and picks a layer only where both products are faster at its rows per routed expert. They need only ggml'scommon.cuh/mmq.cuhat the pinned commit.prefill.cpp): the per-layer decision afterquantize_scatter: all routed experts resident, this GPU alone (no peer or CPU share, no streamed entries of the layer), not the fused layout, not aSTRATA_MMQ_DIRECTlayer, no Hadamard, no NVFP4 tail, no FP4 activations, aligned. The layer's expert offsets from the cache's base and the CTAs' chunk table go up with the group bounds (in the mapped table, or through the pinned upload ring). Every other layer keeps its path.ws::mode()is 1 only withSTRATA_MMQ_WS=1and every visible GPU at compute capability 6.0.STRATA_PREFILL_WSis defined only in theSTRATA_EXPERIMENTAL_SM60CUDA build (not under HIP); elsewhere a stub says mode 0. The offset and chunk buffers (under 1 MiB per die) are allocated only when the mode is not 0, so unset, the arena layout and the mapped table are the ones before.Validation
Environment note: everything below was measured with every expert resident in VRAM on four GP100 dies (two PH402 SKU 200 boards, 32 GB HBM2 per die; the Swift 1.5 IQ3_S pack fits whole: no CPU share, no PCIe share, no adaptive tier). The path declines a layer that is not fully resident, so a single 16 GB P100 that keeps experts in RAM takes the per-group MMQ as before; that setup was not measured. HIP, SYCL and Metal were not built.
Rig: four GP100 dies (application clocks 715 MHz memory / 1050 MHz graphics, 140 W cap), Windows 11, driver 581.80 (TCC), CUDA 12.9, VS 2022. Build:
-DSTRATA_EXPERIMENTAL_SM60=ON -DCMAKE_CUDA_ARCHITECTURES=60, ggml at the pinned 3cf03257f, unpatched. Model: Swift 1.5 IQ3_S (abliterated GSQ-RCO GGUF, dense matrices in Q8_0). Engine arguments, the same on every arm:--serve --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.85 --mtp <drafter> --max-context 1048576 --kv int8 --layer-split 12,24,36 --mtp-window 8192 --pipeline-windows 2 --conversation-cache-mib 32768 --conversation-cache-slots 8 --rope-scaling yarn --rope-scale 4. Environment, the same on every arm:STRATA_Q8_SM60=1 STRATA_STAGE_TRIM=1 STRATA_PREFILL_PIPE=70 STRATA_MMQ_RESIDENT_SORT_NE=1 STRATA_DF_BRANCH=1 STRATA_FUSE_DECODE=1 STRATA_FUSE_WAIT=2 STRATA_SIDE_MIXER=2 STRATA_GR_UP75=1 STRATA_B6_MMVQ_ROWS=1 STRATA_ADAPT_LAG=2 STRATA_ADAPT_NOWAIT=1 STRATA_DECODE_TIMING=1 STRATA_STATE_HASH=1.STRATA_MMQ_WSwas the only difference between the two branch arms.Passed:
STRATA_MMQ_WS=1. Ten requests per arm: five short prompts (400-512 tokens out), a fresh 26K-token read + 400 out, prompt reads of 240 and 1.15K tokens (16 out), a 26K-token conversation and a 1.1K append to it (16 out). Prompt chunks were the same in every arm (512 below 1.2K tokens, 768 above).STRATA_STATE_HASHitem (GDN, PLE, tail, pooled, KV, MTP) equal on all ten requests. X1's log hasa resident layer's experts weight-stationary, one launch per product (STRATA_MMQ_WS=1); X0's does not.--pipeline-windows 3; a 256-token prompt chunk floor; fused prompt-path swiglu + quantize and combine + write; changes of mine not posted yet), turning the weight-stationary path on took prompt reads 240 tokens 2,847 -> 2,084 ms, 1.15K 3,857 -> 2,163, 7.7K 15,283 -> 9,587, 33K 49,687 -> 34,139, a fresh 26K 41,932 -> 27,810, an 88K append 125,258 -> 91,682 ms, with every reply hash and state item equal to the arm before on 15 requests including a 250K-token read. Those numbers depend on that configuration and are not this PR's claim.STRATA_EXPERIMENTAL_SM60=ON(strata) and sm_86 without it (strata), both on this branch, both clean. In the sm_86 build none of the new kernels is compiled.Failed: none observed.
Not run: HIP, SYCL and Metal builds (no such machine here; the new code is under the
STRATA_EXPERIMENTAL_SM60CUDA-only guards). The PR is a draft for that reason: a HIP, SYCL or Metal build by someone who has one would close it.prefill_swiglu_quant_teston this branch (my builds hadSTRATA_BUILD_TESTS=OFF; the kernel is covered by the end-to-end hashes, since every weight-stationary layer goes through it). A layer-level parity tool for the weight-stationary kernels on this revision (it was done on an older 0.1.41-based tree, all 10 gate/up and down format pairs of the Swift pack at T = 128 / 512 / 1792 / 8192; here the evidence is the end-to-end reply and state hashes). Partial residency, peer or CPU shares (declined by condition). Packs other than Swift 1.5 IQ3_S. Cards other than GP100.Limits
STRATA_EXPERIMENTAL_SM60builds only.Addendum (2026-10-23): fewer instructions per pass in the IQ4_NL down's 32-token tier,
STRATA_WS_FAST=1(opt-in, default off)One more commit.
nl4::dn_kernel_fastinmoe_ws_down.cuis picked only withSTRATA_WS_FAST=1, for the IQ4_NL down's 256 x 32 tier (u >= 20 rows per routed expert, i.e. prompt chunks of about 1024 tokens and up). The weight-stationary kernels issue close to all of GP100's instruction slots, so their time follows the instruction count; two changes cut the per-pass overhead without touching any token's arithmetic: aswitchon the number of token pairs falls through the 16 unrolled pairs instead of a guard per pair and block, and dy is stored [block][token] so a pair's two dy are one 8-byte shared load (122 -> 102 registers, still two CTAs per SM). Every token's XMAD chains and epilogue are the default kernel's, so Dm is the same bytes. Below u 20, for the Q2_0 down and for every gate/up nothing changes. The commit also addsmoe_ws_down_bench(EXCLUDE_FROM_ALL, sm_60 build): a layer's 512 real down experts from the GGUF, lognormal and flat-32 routings,ws::downwith and without the switch compared byte for byte and timed.The gate/up variants I measured with the same ideas are written for #3455's 4-bit slots; this branch's GGUF-layout gate/up kernels have no measured counterpart, so none is included.
Environment note: PH402 SKU 200 boards, four GP100 dies of 32 GB HBM2 each (128 GB in all), every expert resident in VRAM, no CPU share, no PCIe share. The single-die numbers below use one of those dies (715 / 1050 MHz).
Bitwise, one GP100 die, this branch:
moe_ws_down_bench <IQ3_S shard> --layers 0,2,1 --t 768,1024,1536,2048,4096 --models ln0.5,ln1.0,flat32 --reps 15withSTRATA_MMQ_WS=1: 45 checks, 0 failures (every Dm row memcmp-equal; layer 1 is a Q2_0 down, the default kernel either way, as a control).Timing, same run (IQ4_NL layer 0; layer 2 within 0.6% of it; median of 15 interleaved; ms per layer, default ->
STRATA_WS_FAST=1):(flat at T >= 2048 is capped at 512 experts x 32 rows.) The Q2_0 control is x0.998-1.001 everywhere.
Default off:
cuobjdumpofmoe_ws_down.cu,moe_ws_gu_iq2.cuandmoe_ws_gu_iq3.cu, sm_60 build: the 20 kernels this branch had before the commit have identical SASS after it;dn_kernel_fast<256, 32, 2, 2>is the only addition, and its SASS is identical to the build used in the end-to-end run below. The sm_86 build compiles none of the WS kernels and builds clean.End to end, four GP100 dies, not this branch's revision: engine 0.1.47 (main 011c9e6) plus 26 opt-in GP100 commits of mine ([Engine] GP100 (sm_60): 4-bit slots for resident IQ2_S / IQ3_XXS / IQ3_S gate/up experts, decode +9.7% (opt-in, STRATA_EXPERT_NIBBLE=1) #3455's 4-bit slots, this PR with a 4-bit variant of its gate/up,
--pipeline-windows 3and others),STRATA_MMQ_WS=1andSTRATA_SM60_TABLE=0in both arms, one fresh engine per arm:STRATA_WS_FAST=1took a fresh 26K read 27,555 -> 27,235 ms (-1.16%), a fresh 33K read 34,426 -> 34,028 ms (-1.16%) and an 88K-context 1.5K append 84,897 -> 84,113 ms (-0.92%); decode unchanged (-0.11%); every reply hash identical on 9 requests. That arm also ran faster 4-bit gate/up kernels for [Engine] GP100 (sm_60): 4-bit slots for resident IQ2_S / IQ3_XXS / IQ3_S gate/up experts, decode +9.7% (opt-in, STRATA_EXPERT_NIBBLE=1) #3455's slots, which are not in this PR, so the down's own share of those numbers is smaller (on one die the two parts were of similar size). One run per arm.Not run: HIP, SYCL and Metal builds (no such machine here; the new code is inside the
STRATA_EXPERIMENTAL_SM60CUDA-only WS sources). An end-to-end run on this branch's own revision.CI note: the red
metal-linkcheck is not caused by this PR. The job fails while compiling upstream'smetal/src/mtl_rt.mm(lines 69 and 92:MTLGPUFamilyMetal4,MTLCompileOptions.mathMode,MTLMathModeFast/MTLMathModeSafe), which themacos-14runner's SDK does not provide. The workflow runs on any pull request that touchessrc/orinclude/; this PR changes no file undermetal/.