Repository navigation
Conversation
Measured on a CMP 170HX (GA100, 70 SMs): the prompt path dequantizes the dense Q4_K/Q5_K/Q6_K/IQ4_XS weights every chunk. With the wide kernels (STRATA_DEQUANT_WIDE=1, the sm_61 / sm_86 default): GPU time of a 3,259-token prompt 857 -> 792 ms; prompts 3.3K +6.3..+7.2%, 13K +4.7%, 52K +3.9%; greedy output identical (11 of 11 prompts). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Read off mmvq_il_parity --bench --emit-table (the same 3% rule), bitwise equal to native_mmvq. Decode +2.9% on the CMP 170HX (Qwen3.8-Flash-Next IQ3_S). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…6), whole-layer launches (STRATA_PF_BATCH=N), two-pass weight loads in the 256-thread kernels (bitwise) Measured on a CMP 170HX (GA100, 70 SMs), one MoE layer of 3,256 tokens (prefill_fused_iq_test, IQ3_S/IQ4_NL): - STRATA_PF_FUSED_NT=256: 256 threads, 2 blocks an SM, bitwise: 8.27 ms - two-pass weight loads: the weight loads of the next k step are issued before the current step's decode, bitwise. Measured on the 512-thread kernel: 10.09 -> 9.33 ms. Here they run only in the NT = 256 kernels (TWO_PASS), so the default 512-thread kernels keep their load sequence on every card. - STRATA_PF_BATCH=N: experts per launch, one persistent grid a layer. Above kMaxBatch (128) only on the NT = 256 kernels, which take a WideBatch (kBatchRoom = 512 pointers, 4 KB of kernel parameters: CUDA 12.1+, Volta and newer); kBatchRoom is kMaxBatch on HIP, gfx906 and STRATA_EXPERIMENTAL_SM60. All opt-in. Batch, the default kernels (native_kernel WW 2 / 4 at 512 threads, moe_fused.cu, moe_f16_kq.cu) and their launches are unchanged; the NT = 256 instances are new device code in every CUDA build. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…uantizes them on the device 1,440 bytes a token over PCIe instead of 10,240 floats' bytes; the same floats as the host dequant (bitwise). On a PCIe Gen2 x4 card (CMP 170HX) the cmp_bench3 prompts (4K and 16K) were +1.0..+1.7 % faster. Opt-in. - the raw-row buffer's cudaFree and the error clear after a failed cudaMalloc are (void): HIP's hipError_t is [[nodiscard]] - metal: ple_iq4nl_dequant is a stub (stubs_0145.cpp); the Metal prompt runs through its window and never reaches it Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…ir own stream The STRATA_PLE_COMPACT upload sat on the compute stream before layer 1; it now goes on a non-blocking stream with events for the dequant and the host buffer's reuse. Same bytes. The upload stream's synchronize / destroy and its event's destroy are (void) (HIP's [[nodiscard]] hipError_t). Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The table lived in mapped pinned host memory, so every prompt's row gather crossed the link. With the env, when
512 MiB stay free beside it, the table goes to VRAM (322 MiB for an IQ4_XS table) through the calls the existing
cannot-pin path uses; same bytes, same gather kernel. The start-up line says where the table is ("in VRAM" also for
the existing VRAM fallbacks, which it called mapped host memory). The failed copy's cudaFree and the error clear are
(void) (HIP's [[nodiscard]] hipError_t).
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…e stream into a double-buffered scratch Gemm::native ran the dequant serially in front of every dense GEMM on the compute stream. With the env the dequant of the next projection runs on a side stream into the other scratch (64 MiB more of the prompt path's region), with ready/free events; the sliced path stays serial. Same kernels, same bits. Off on HIP. - the side stream's event waits / records and ~Gemm's destroys are (void) (HIP's [[nodiscard]] hipError_t) - metal: gemm_dequant_side() answers false and Gemm::set_dequant_side(nullptr) the serial form (stubs_0145.cpp) Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…16=1) and a warp-specialized kernel (STRATA_PF_FUSED_WS=1) Measured on a CMP 170HX (GA100, 70 SMs), one MoE layer of 3,256 tokens (prefill_fused_iq_test, IQ3_S/IQ4_NL), against the NT = 256 kernels' 8.27 ms: - STRATA_PF_FUSED_F16=1: f16 mma with the block scales folded into the weights and activations (error vs the FP32 reference = MMQ's): 7.90 ms. It runs in the NT = 256 kernels (it implies STRATA_PF_FUSED_NT=256). - STRATA_PF_FUSED_WS=1: producer warps decode into an f16 ring, consumer warps run the mma (named barriers), bitwise = the f16 path: 7.65 ms. Compile-time knobs STRATA_WS_PROD / _SLOTS / _NB / _KS / _HTAB (defaults 256 / 2 / 2 / 2 / 0; _HTAB=1 or 2 turns on the IQ3 decode table). Both opt-in and rounding-level; unset = the kernels as before. Both take a WideBatch, as the NT = 256 kernels do: the F16 kernels are NT = 256 instances, and native_wide() / experts_native(const WideBatch&) also cover the WS kernel (one block a multiprocessor), so STRATA_PF_BATCH above kMaxBatch reaches it; Batch and the default kernels' launches are unchanged. native_ws_kernel and the F16 path are CUDA only: the WideBatch launch()'s WS and F16 arms (WS_THREADS, ws_smem_bytes, f16c_requested: all defined under !__HIPCC__) sit under !__HIPCC__; d.ws is never set on HIP. The two-pass weight loads stay in the NT = 256 kernels, the F16 ones included. The F16 path is sm_80+: the sm_75 pass (STRATA_PF_FUSED_SM75) compiles its native_kernel instances against __trap() stubs of mma16h and i8x4_h2 (dev_info stops before the NT = 256 setup on Turing, so they never run). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…d scan on the tensor cores (sm_80+, opt-in) The chunked WY-form recurrence with every product on the tensor cores, for sm_80+ cards with 48+ SMs and 122 KB of shared memory a block. STRATA_GDN_TC=1 dispatches it before the STRATA_GDN_CHUNKED branch for prompt chunks of 128+ tokens (kGdnTcMin); a card below that, or a failed scratch allocation, takes the kernels as before. Unset = the kernels as before. Rounding-level: the output is not the FP32 kernels' bits. CUDA only (gdn_tc.cuh is under !__HIPCC__; the dispatch sits in kernels.cu's CUDA branch). src/prefill/gdn_tc.cuh (new): the mma.sync / ldmatrix wrappers, the f16 hi/lo split with power-of-two scaling, the dynamic scale exponent, the per-chunk products a warp runs and the two kernels: - gdn_tc_prep_kernel: gdn_chunk_prep_kernel's arithmetic, plus the k and q f16 hi / lo images x 2^8 (built in shared memory from the float4 registers the rows arrive in), the P images x 2^8, T as f16 hi / lo images x 2^et (et from the warp's largest |T|) and the token scalars (beta, exp(gamma), exp(gL - gamma), gamma, the scale 2^-et) as a float4 a token, into gdn_tc_scratch (60.3 MB a device). Its Gram products k_t.k_i and q_t.k_i run as f16x3 mma (48 mma and 32 ldmatrix a warp in place of ~1,150 LDS a thread); after them, warps VPK + j take P and the scalars of value head j while warps j build A and solve with lane i's column of T in registers. - gdn_tc_scan_kernel: a block per value head (48 blocks of 256 threads, one wave); warp w holds value columns 16w..16w+15 of the transposed state (16 m16n8 f32 accumulator tiles) across every chunk of a super-block. X^T / QS^T = S^T K^T / Q^T on f16x3 (the state x 2^es), D^T = Y^T T^T on f16x3 from the prep's T images (6 of 8 tiles, Y^T x 2^ey), o^T = e QS^T + D^T P^T on f16x3 with P's triangle skipped, S^T = eL S^T + Dd^T K on f16x3 through ldmatrix.trans: 324 mma a warp a chunk, one __syncthreads a chunk. Each chunk is staged with 15 cp.async a thread right after the chunk barrier from the prep's images. The warp maxima run through redux.sync over short chains (eight chains of eight, four of four); the warp maximum and the shuffles use the full mask (a lane-divergent L2 prefetch before the first maximum of a chunk made __activemask() unsafe; the kernel has no early return, so the full mask adds no reconvergence code). 255 registers, 0 spills (prep 52). The header is included inside the anonymous namespace of kernels.cu and of the harness, so the harness runs the engine's text verbatim. src/prefill/gdn_rec_parity.cu: check_variant (any launcher of gdn_rec_chunked's shape), the tc variant at T = 1..4099, 8192, 32768 and six stress inputs (state x 0.01 / x 50 / 0, v x 50, slow gates, coherent keys) at 4099 and 32768 against gdn_rec_kh_kernel (gate 1e-4); --selftest-frag (each helper on one warp against a double reference); --bench-tc (the tensor-core series alone, for A / B runs of two builds); --long (T = 51,820). Measured on a CMP 170HX (GA100, 70 SMs): - gdn_rec_parity: output <= 2.3e-6, state <= 2.3e-6 against gdn_rec_kh_kernel over every T and stress case (the coherent-keys case 1.8e-6 / 8.5e-7); --selftest-frag (b) 9.1e-7, (c) 4.2e-7, (d) 3.3e-7, (e) 0 words differ. (--long T = 51,820 at the first cooperative-staging kernel: output 1.17e-06, state 2.58e-07.) - --bench-tc (two builds back to back, twice, medians of alternating runs), us a token at 2048 / 8192 / 32768: v0 0.288 / 0.287 / 0.287 (2.79x gdn_rec_kh_kernel, 0.80 us); redux maxima 0.282 / 0.281 / 0.281 (2.85x); full mask 0.283 / 0.282 / 0.282 (gdn_rec_kh_kernel 1.65 / 6.56 / 26.23 ms); cp.async staging from the prep's images 0.262 / 0.260 / 0.260 (3.07x); short max chains 0.261 / 0.260 / 0.259; the prep's Gram products on the tensor cores 0.238 / 0.238 / 0.238 (3.37x; prep 233 -> 153 us a 2048-token super-block under ncu); the prep's serial phase on six warps 0.219 / 0.218 / 0.218 (3.67x; prep 153 -> 95 us); D^T on f16x3 0.214 / 0.214 / 0.214 (3.75x). - Engine (Qwen3.8-Flash-Next IQ3_S, STRATA_PREFILL_TIMING), with the first engine kernel: the "gdn recurrence" phase of a 3,259-token prompt 105 -> 43 ms, of a 51,820-token prompt 1,705 -> 686 ms; prompts 3,849 -> 4,104 tok/s (3K), 4,163 -> 4,439 (13K), 4,060 -> 4,332 (52K); the 11-prompt greedy set identical to the FP32 kernels. Measured and not kept: - stager warps with named barriers (two more warps staging a chunk ahead, 320 threads): 0.299 / 0.297 / 0.297 us a token against 0.282 / 0.281 / 0.281 (6 % slower): a block of more than 8 warps is capped at 168 registers a thread on GA100 and the consumers spill. - the output norm fused into the scan: correct (y16 within 4.4e-4 of its largest value against gdn_out_norm_kernel), but 0.303 against 0.214 + 0.057 us a token at 32768 (the chunk loop ~2,010 -> ~2,900 instructions at the 255-register cap). - the prep overlapped with the previous super-block's scan on low / high priority streams joined by events (a two-slot scratch): it wins in the harness (0.214 -> 0.215 / 0.191 / 0.184 us a token at 2048 / 8192 / 32768) but loses in the engine: a 3,259-token prompt's GPU timeline 722 -> 760 ms, a 51,820-token one 11.79 -> 12.11 s, the "gdn recurrence" phase 34 -> 37 ms, from the cross-stream scheduling gaps. Both launches stay on the caller's stream. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…t copy) and STRATA_CKPT_VRAM_SLOTS (recent GDN states kept in VRAM) - STRATA_CKPT_ASYNC=1: a turn checkpoint's GDN state is copied device-to-device into a VRAM stage, then a background thread copies it in 4 MiB pieces into a pinned buffer and into the checkpoint; every reader joins first (selective joins: only the copy it needs). Turn checkpoint 72 -> 1 ms on a PCIe Gen2 x4 card (118 MB of GDN state at 1.64 GB/s). - STRATA_CKPT_VRAM_SLOTS=N: the last N saved GDN states stay in VRAM (tag-checked; only allocated while 1.5 GB stays free). Restoring a parked conversation 116 -> 28-47 ms. - STRATA_TRACE lines carry a steady-clock timestamp. Both opt-in; outputs bitwise identical (conversation A/B/A check). HIP: ckpt_async_on() is false under STRATA_HIP_GFX906 as well, and the runtime calls whose result the cache does not check are cast to (void) (the HIP shims declare hipError_t [[nodiscard]]), so this commit carries its own HIP guard and builds alone for gfx1100 and gfx906. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…nd and moves the dropped checkpoints Parking the outgoing conversation copied its K/V device to host and deep-copied its checkpoints before the next prompt could start: 213 ms at a 3K context, 334 ms at 13K on a PCIe Gen2 x4 card (CMP 170HX). - STRATA_PARK_ASYNC=1 (CUDA, opt-in): conversation_kv_save copies the device K/V device to device into a VRAM stage (only while 1.5 GB stays free beside it) and a background thread moves the stage into the host buffers (4 MiB pieces through two pinned buffers). conversation_park_join waits for it: every reader of a parked image's K/V, the cache before it frees an entry, and the image itself (SavedConversation::pending) before its buffers go. ConversationBuffer::resize_for_overwrite leaves the added bytes unfilled for that copy. - With it, the checkpoints this request drops right after the park (all of them when another state comes in, else the ones off this prompt's path) move into the parked image instead of being copied, and a background allocation refills STRATA_CKPT_POOL for the next save. - conversation_snapshot_save takes dropped checkpoints' storage (STRATA_CKPT_POOL) for its copies. Park 213 -> 64-75 ms (3K), 334 -> 68 ms (13K); outputs identical (greedy 11 of 11, the A/B/A conversation check). HIP: park_async_on() is false under STRATA_HIP_GFX906 as well, on_device() returns false in the shims (none maps cudaPointerGetAttributes; the slot then takes no part, as when the query fails), and the runtime calls whose result the cache does not check are cast to (void), so this commit carries its own HIP guard and builds alone for gfx1100 and gfx906. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…copies gathered, helpers off the host core, park traces Measured on a CMP 170HX (PCIe Gen2 x4): a conversation switch's park still paid 38-52 ms for the live state - the STRATA_CKPT_POOL pool was empty at every park (a switch takes two checkpoints and the refill added one), so the 118 MB GDN vector was allocated fresh and page-faulted. Now: - STRATA_CKPT_POOL=N: the pool keeps N (default 2) ready checkpoints, pre-sized to the session's state sizes, topped up on a background thread after any take. - STRATA_CKPT_GATHER=1: the 37 small synchronous device-to-host copies of a checkpoint (ple, tails, dead, block_pos) gathered device-to-device into one buffer and fetched with one copy through a pinned buffer. - The STRATA_PARK_ASYNC stage is allocated with 64 MiB of headroom (no cudaFree/cudaMalloc as the conversation grows) and the pointer-attribute query is skipped for a VRAM-resident K/V. - The three background helpers (pool top-up, the checkpoint DMA, the park DMA) leave the host's pinned core. - A parked image whose live GDN state is still being copied joins before its buffers are freed. - STRATA_PARK_TRACE prints the joins, the live-state vectors' readiness and the K/V sub-timers; the d2d time is measured right. Park at a 3K context 64-75 -> 21-40 ms; outputs identical (greedy 11 of 11, the A/B/A conversation check). HIP: ckpt_gather_on() is false under STRATA_HIP_GFX906 as well, and the runtime calls whose result the cache does not check are cast to (void), so this commit carries its own HIP guard and builds alone for gfx1100 and gfx906. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…' K/V in VRAM in front of the host copy With STRATA_PARK_ASYNC, a park copies the K/V device-to-device into a VRAM slot (64 MiB granules, up to N MiB in all, only while 1.5 GB stays free) that the background DMA reads from; the host image is still filled. A switch back restores the K/V device-to-device from the slot (~1 ms instead of 27-107 ms over a PCIe Gen2 x4 link) and needs no join on the DMA. Each K/V layer carries its own join-and-release token, so any order of destruction is safe; the elastic K/V's growth failure path releases every slot and retries once. Restore of a 3K conversation 28-47 -> 7-12 ms on a CMP 170HX; outputs identical. HIP: the runtime calls whose result the cache does not check are cast to (void) (the HIP shims declare hipError_t [[nodiscard]]), so this commit carries its own HIP guard and builds alone for gfx1100 and gfx906. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…C, STRATA_CKPT_VRAM_SLOTS, STRATA_PARK_ASYNC, STRATA_CKPT_POOL=N, STRATA_CKPT_GATHER, STRATA_PARK_VRAM_MIB; CUDA, opt-in) One section in DETAILS.md for the switches the four conversation-cache commits add or extend: what each copies where, its VRAM / host memory cost, that the first three stay off on HIP, gfx906 and Metal, the new cost of an existing STRATA_CKPT_POOL=1 / STRATA_SLOT_CKPT_MOVE=1 setup (N x 118 MB held from start-up), and the times measured on the CMP 170HX the commits name. STRATA_PARK_TRACE and the timestamp on the serve loop's STRATA_TRACE lines are listed as diagnostics. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…des their stream's priority Opt-in (unset = as before) and bitwise by construction (no kernel changes, the same launches on the same inputs; only the dispatch order moves): STRATA_GRAPH_PRIO=1: the side mixer's branch stream (GDN gates, QSA indexer chain) is created at the device's highest priority, cs_ and side2_ in the middle, the shared expert's stream at the lowest, and the window graph is instantiated with cudaGraphInstantiateFlagUseNodePriority so each kernel node runs at its capturing stream's priority instead of the launch stream's. The 1-block indexer kernels are then dispatched ahead of the q GEMV's queued blocks and the shared expert behind the experts. A log line per captured window prints the node priority histogram; if a driver did not copy the stream priority into the nodes, the capture-time notes (the branch forks and the priority streams' tails) set the attribute per node (not needed on driver 610 / CUDA 13.3: the 3-token window shows 120 x -5, 1063 x -2, 336 x 0, exactly the expected 36 + 84 high and 7 x 48 low). Measured on the CMP 170HX (IQ3_S, --spec 3, 3K context, cmp_bench3 dec tok/s, interleaved in one lend, each arm with STRATA_QIL=1): ref 125.3/126.6 | GRAPH_PRIO 127.8/129.0 | ref 125.0/126.8 = +1.8 % (both GRAPH_PRIO runs above every reference run). Greedy 11-prompt ids 11/11; the window node count is unchanged (1519 at T = 3 with QIL). HIP: the prio_forks_ / prio_tails_ members (cudaGraphNode_t, not in the wave32 shim), capture()'s node-priority block and the note_prio_fork / note_prio_tail / apply_node_priorities bodies sit under !STRATA_USE_HIP && !STRATA_HIP_GFX906 && !STRATA_GPU_METAL; init()'s cudaDeviceGetStreamPriorityRange query under !STRATA_USE_HIP && !STRATA_HIP_GFX906 (gprio = false otherwise); graph_prio_on() is false on HIP, gfx906 and Metal and [[maybe_unused]]. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…ab), the clock warm-up and MMVQ_IL_TMIN A go/no-go harness for an L2 weight prefetch, with its measurement on the CMP 170HX (GA100, 70 SMs, 32 MiB L2), T = 3. Harness only; no engine code changes. mmvq_il_parity --bench gains --hot (one weight copy: the L2-hot ceiling of a prefetch), --gap MB (MB of OTHER weights read between the timed calls, as the window does between a prefetch and its consumer), --persist (a 20 MiB persisting L2 set-aside + an access-policy window on the stream) and --ab (every arm per case inside ONE process: cold / hot / hot+gap / hot+gap+persist / cold - the clocks ramp over the first ~100 ms of a run, so arms in separate runs are not comparable; a memset gap was tried first and its write-backs made every arm slower than cold), plus a clock warm-up and MMVQ_IL_TMIN. Engine row choices (kIlRows_80), cold -> hot, us: z 2560x6144: IQ4_XS 16.0 -> 14.7 (-8 %), Q4_K 20.8 -> 17.7 (-15 %), Q5_K 22.3 -> 18.9 (-15 %), Q6_K 27.9 -> 20.7 (-26 %) out 6144x2560: IQ4_XS 16.1 -> 13.9 (-14 %), Q4_K 21.9 -> 19.2 (-12 %, multi), Q5_K 23.6 -> 20.5 (-13 %, multi), Q6_K 25.1 -> 20.4 (-19 %) ncu (caches kept vs flushed): L2 sector hit rate 100 / 93 / 96 % hot vs 24 / 21 / 16 % cold (Q6_K out, Q6_K z, IQ4_XS z), but at ncu's locked clock the duration moves only -4.5 / -10 / -3 %: the IL kernel is issue-bound once the lines are in L2 (sm__throughput 47-51 %). With 12, 16, 24 or 40 MB of other weights read between the prefetch and the consumer, plain LRU keeps NOTHING (hot+gap = cold or 1-2 us worse for every type at every gap), and the 20 MiB set-aside keeps a fraction: Q6_K z 27.9 -> 26.0, Q6_K out 25.1 -> 24.6, Q4_K / Q5_K unchanged, IQ4_XS z 16.0 -> 17.6 (slower). So an L2 prefetch (z forked 150 us ahead behind the experts' 58 MB, out 25-60 us ahead behind z's 9-13 MB) cannot meet its -15 % floor on this card: no-go. gfx906 (the harness builds there with tests): persist_window's L2 access policy window and the --persist / --ab set-aside query sit under !STRATA_HIP_GFX906; the harness prints "not available in this build". The new timing code ignores the runtime calls' results explicitly ((void)), so the harness adds no new nodiscard warnings to the gfx906 build. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… (STRATA_MMVQ_IL_PF=1) native_mmvq_il_pf_kernel is native_mmvq_il_kernel with every weight load of a group (4 chunks x R rows) issued before any of the group's dots, in a branch-free main loop over the full groups (the old kernel's four `if (kbx < bpr)` sections each sat behind their own branch, so a warp had one 2-row x 1-block group in flight; ptxas also sinks a load whose only use is behind a branch into it), one pointer set per row with the other three blocks at immediates, and the tail group clamped and predicated. The dots use P2's exact integer forms: __vsubss4(x, 0x20) for Q6_K's [0, 63] codes as (x + 0x60606060) ^ 0x80808080 (2 instructions for ptxas's 8), and (float)(dp * sc) for |dp * sc| < 2^22 as the 1.5 * 2^23 magic add on the FMA pipe instead of the quarter-rate I2F (same for Q4_K / Q5_K's dot1 * sc, dot2 * m, and IQ4_XS's sumi then times (float)(ls - 32), exact under 2^24). The per-lane accumulation order, the v-slot sums and the xor tree are untouched, so every output is bitwise native_mmvq's: mmvq_il_parity checks each variant on every case, T 1-8, rows 1/2/4. The main loop is 441 SASS instructions per 8 row-blocks (55 a row-block, from 89), and ncu's long-scoreboard stall drops from 3.0 to 1.1-1.8 cycles an instruction; the win is then bounded by registers (128 at 4 blocks a SM), so the choice of variant and rows a warp is a measured sm_80 table (kIlPf_80 / kIlPf58_80, variant << 4 | rows, from --emit-table --pf 0,3,4): variant 3 (4 blocks a SM) or 4 (5 blocks a SM: the 2560-row shapes' 320 blocks fit one wave). Only Q6_K is in the table: Q4_K / Q5_K / IQ4_XS are bitwise and measured too but tie today's kernel within noise. Tried and rejected (measured slower): hoisting without P2 (latency hidden but 16% more instructions, issue-bound), a whole-group double buffer (168 registers, 3 blocks a SM) and a per-section ring (spills at 2 rows a warp). mmvq_il_parity --pf a,b,.. (interleaved medians, 5 rounds of 100), cold, engine clocks, T x rows a warp: Q6_K 2560x10240 T2 r2 37.6 -> 30.0 T3 r2 39.6 -> 37.0 T4 r2 45.8 -> 39.6 T5 r2 55.5 -> 48.5 us Q6_K 6144x2560 T2 r2 27.8 -> 19.1 T3 r1 25.1 -> 22.7 T4 r1 27.9 -> 23.8 T5 r2 45.5 -> 35.1 Q6_K 2560x6144 T2 r2 25.3 -> 20.1 T3 r2 28.0 -> 23.9 T4 r2 32.5 -> 25.1 T5 r2 35.1 -> 30.9 Engine (cmp_bench3, 3K context, interleaved ref / on / ref / on): dec 125.5 126.7 | 127.0 127.9 | 124.4 126.0 | 127.1 127.9 tok/s (+1.4%); the 11-prompt greedy set 11 / 11. Unset or 0: today's kernels; 2-6 force one variant for every type (A/B; 6 = hoisted loads without P2). HIP: the inline-PTX load helpers (ldg_nc_*, madw) and the Q6_K dot forms only the CUDA-only il_pf kernels use sit under !__HIPCC__, and the traits' TYPE constants are [[maybe_unused]], so this commit carries its own HIP guard and builds alone for gfx1100 and gfx906 with no new warnings in native_mmvq.cu. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…witches) DETAILS.md gets a section with the two switches, what each does, that both are bit-exact, and their decode gain on the CMP 170HX (+1.8 % with STRATA_QIL=1, +1.4 %). MMVQ_IL_TABLE.md gets a paragraph on the kIlPf_80 / kIlPf58_80 tables (variant << 4 | rows, Q6_K only) and on measuring a card with mmvq_il_parity --bench --emit-table --pf 0,3,4. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
GPU runs on this revision. Everything below ran on one CMP 170HX:
How the runs were made:
Overall
Per featureEach pair is the PR's switches off/on on that branch's own build. PR 1's defaults are compared against stock main. Percentages are on/off means. The pair ratios are in brackets.
Notes:
Not run on this revision: teacher-forced KL for PR 3's two rounding-level switches, any AMD or SYCL execution, and other sm_80 cards. 🤖 Generated with Claude Code |
linxuhao
changed the title
|
Teacher-forced KL of PR 3's two rounding-level switches (CMP 170HX, PR 3's build on 011c9e6, 250 W). Method:
🤖 Generated with Claude Code |
|
Independent data for
One question from reading the code (unconfirmed, not measured): does the VRAM stage come out of the expert cache when I am not opening a competing pull request. If it helps, I can run this PR's |
|
Thanks, this is useful. Your 72.4 ms matches what we measured with reused vectors (71.4 ms), so the copy dominates on PCIe 3.0 x16 as well as on our Gen2 x4. On your question (from the code, not measured on your card): the stage does not come out of the expert cache. It is allocated lazily by the first turn checkpoint save. That happens after the engine is ready, when The flip side is that the stage has no headroom check of its own. Yes, please run it. The branch is based on 011c9e6 (0.2.2). The checkpoint change is the single commit 0d51292, which applies without conflicts onto v0.2.4.1 (checked textually, not built there). Your pinned-only variant is also a good data point: on a fast link, the VRAM stage may not be worth its 113 MiB. 🤖 Generated with Claude Code |
|
Thanks. I ran 0d51292 on its own, cherry-picked onto main 4b9e05a (clean, same patch-id) and built with setup's CUDA line for this card (sm_86), on one RTX 3090 24 GB Risk case: completed. Fresh engine, four turns that extend each other (24K tokens), then a fresh 253,915-token prompt, a fresh 200,034-token prompt
Identity. With Speed, release defaults, 4 pairs (OFF ON ON OFF, then ON OFF OFF ON). Tiny-prompt TTFT (engine prompt ms, median of 6 chats) OFF / ON: 353.8 / 339.8, Why it is small here ( Build break. With Not tested: the sampled thinking-on tier, 128K and 256K speed tiers, |
…, not in conversation_file.cpp 0d51292 (STRATA_CKPT_ASYNC) and 3e0c141 (STRATA_PARK_ASYNC) made conversation_file.cpp include conversation_snapshot.hpp and call the joins, so conversation_file_test and session_save_reclaim_test (built from conversation_file.cpp without CUDA) stopped compiling with STRATA_BUILD_CONVERSATION_TESTS=ON. The joins move to the start of the SAVE/RESTORE handler in generate.cpp. That is also earlier: the save copies the checkpoints (session_save_checkpoints) and may reclaim some before put_checkpoint ran, and the join there only matched the copied vector, not the one the copy in flight writes. conversation_file.cpp is upstream's again. Reported by bitaccesscomau in Niko1221#3779. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
Thank you, this answers the open question and found a real bug. Headroom. With a minimum of 147 MiB free and no fallback in 42 of 42 saves, I will not add a headroom check to the stage. Thanks for correcting the free-at-ready figure: I had not counted the 452 MiB the driver keeps reserved. Build break: fixed in 33a8146 (on top of the branch, so
Speed on a fast link. Your reading of the code matches mine:
The K/V copy getting slower during a park is expected: both copies share the one copy engine. I have no fix for that here. 🤖 Generated with Claude Code |
… chunked scan on the tensor cores, sm_80 (#3779, #3712) Problem: the prompt path's DeltaNet recurrence is a large part of the prompt time of a GA100 card. Change, opt-in and rounding-level (the output is not the FP32 kernels' bits): src/prefill/gdn_tc.cuh holds the chunked WY-form recurrence with every product on the tensor cores (f16x3 splits with power-of-two scaling, one block per value head) and its prep kernel; kernels.cu dispatches it before the STRATA_GDN_CHUNKED branch for prompt chunks of 128+ tokens when STRATA_GDN_TC is set. Where it runs and where it does not: - gdn_rec_tc returns "not supported" and the existing kernels run unchanged on a card below sm_80, with fewer than 48 SMs, or with less than 122 KB of shared memory a block (the kernel needs 124,944 B). That excludes sm_86, sm_89 and sm_120 (99 KB opt-in): the code is compiled into those builds and never launched there. It runs on sm_80 and sm_90. A failed scratch allocation (60.3 MB a device) also falls back. - Device code for cards below sm_80 compiles to trapping stubs behind STRATA_GDN_TC_SM80; the whole header sits under !__HIPCC__. - Unset, the dispatch costs one cached flag test per call. Tests: gdn_rec_parity gains the tc variant (T = 1..4099, 8192, 32768, six stress inputs against gdn_rec_kh_kernel, gate 1e-4; "not on this card" where it cannot run), --selftest-frag (each tensor-core helper on one warp against a double reference) and --bench-tc. The new ctest entry gdn_rec_parity_selftest_frag runs the helper self-test, which any sm_80+ card can run (77 below sm_80). The full tc variant can only run on an sm_80 or sm_90 card. Reporter's numbers (linxuhao, CMP 170HX): recurrence 105 -> 34 ms at 3.3K tokens and 1,705 -> 547 ms at 52K; prompts +3.9 % (4K), +6.0 % (16K), +6.7 % (52K) over the same base; within 2.3e-6 of the FP32 kernel in gdn_rec_parity; teacher-forced KL against the bitwise base mean 0.0157, median 0.0011, max 0.098, top-1 100 %; 1 of 5 conversation checks identical and decode -2.6 % in their pair (they think draft acceptance moves with the rounding; not checked). Theirs, not measured here. Original change by linxuhao (PR #3779, commit f1cabc9).
…-to-host copy) and STRATA_CKPT_VRAM_SLOTS (#3779, #3712) Problem: on a slow host link a turn checkpoint copies the whole GDN state to the host on the request path (118 MB at 1.64 GB/s on a PCIe Gen2 x4 card, 72 ms; an independent RTX 3090 on PCIe 3.0 x16 showed 73 ms to 62 ms with the switch). Restoring a parked conversation copies it back. Change, both opt-in and CUDA only (off on HIP, gfx906 and Metal): - STRATA_CKPT_ASYNC=1: the GDN state is copied device to device into a VRAM stage (one buffer the size of the state, allocated at the first checkpoint save, plus a pinned host buffer of the same size; if either allocation fails the save falls back to the synchronous copy), and a background thread moves it in 4 MiB pieces into the checkpoint's host vector. Every reader joins first and only for the copy it needs (restore, snapshot save and restore, the parked copies, the serve loop at each request start, the SAVE/RESTORE session handler). - STRATA_CKPT_VRAM_SLOTS=N (with ASYNC): the last N saved GDN states stay in VRAM, tag-checked, only allocated while 1.5 GB stays free; a restore of a checkpoint whose tag a slot holds copies device to device. - STRATA_TRACE lines carry a steady-clock timestamp; STRATA_PARK_TRACE prints the timing of the save steps. - Unset: the code adds a mutex-guarded no-op join per checkpoint save, restore and snapshot, a tag field that stays 0 and a few clock reads; the copies are the synchronous ones as before. Differs from the PR as submitted: the PR put a join into conversation_file.cpp (put_checkpoint), which broke the conversation_file_test and session_save_reclaim_test builds without CUDA and matched only the copied vector, not the one the copy in flight writes, so a SAVE within one checkpoint copy after a turn could write an incomplete GDN state (found by bitaccesscomau; fixed in the PR's last commit 33a8146). This commit carries that fix: conversation_file.cpp stays as it is and the join runs first in the SAVE/RESTORE handler of generate.cpp. The commit 33a8146 is therefore not a separate commit here. One more join, found by running the new ctest entries: conversation_snapshot_sources (the SAVE path of a session file) saves the running state with the background copy and the caller writes it to the file at once, after the handler's join has already run. It now joins the copy of that checkpoint before it returns; without it a session file written with STRATA_CKPT_ASYNC=1 could hold an incomplete GDN state (conversation_snapshot_test failed at "both files read" before the join). Tests: conversation_snapshot_test now joins after each save and, with STRATA_TEST_BIG_GDN=1 (a GDN state over 1 MiB, so the async path is taken), runs as the new ctest entries conversation_snapshot_test_ckpt_async and conversation_snapshot_test_ckpt_vram_slots: every checkpoint, capture and restore must give the same bytes as the plain run. Not covered: no test destroys a checkpoint while its copy is in flight (readers join; the cache's eviction paths do not); the reporter saw 3 of about 35 engine starts stop answering after a park with evictions, on the paths with the switches off, and could not reproduce it in 15 reruns. The VRAM stage has no headroom check of its own (an RTX 3090 with 147 MiB minimum free showed no failure in 42 of 42 saves; the allocation failure falls back silently). Reporter's numbers (linxuhao, CMP 170HX): turn checkpoint 71.4 -> 0.9 ms, restore 114-126 -> 28-47 ms. Independent data (bitaccesscomau, RTX 3090 24 GB, PCIe 3.0 x16, release defaults, 4 pairs): checkpoint save 73.2 -> 62.4 ms, tiny-prompt time to first token median ratio 1.038, outputs byte-identical. Theirs, not measured here. Original change by linxuhao (PR #3779, commits 0d51292 and 33a8146).
…ackground and moves the dropped checkpoints (#3779, #3712) Problem: parking the outgoing conversation copies its K/V device to host and deep-copies its checkpoints before the next prompt can start: 213 ms at a 3K context, 334 ms at 13K on a PCIe Gen2 x4 card (the reporter's CMP 170HX). Change, opt-in and CUDA only (off on HIP, gfx906 and Metal): - STRATA_PARK_ASYNC=1: conversation_kv_save copies the device K/V device to device into a VRAM stage (only while 1.5 GB stay free beside it; otherwise the save copies as before), and a background thread moves the stage into the host buffers in 4 MiB pieces through two pinned buffers. conversation_park_join waits for it: every reader of a parked image's K/V, the cache before it frees an entry, and the image itself (SavedConversation::pending) before its buffers go. ConversationBuffer::resize_for_overwrite leaves the added bytes unfilled for that copy; the buffer's segments use an allocator that value-initializes only when a fill value is given, and resize(n, value) is unchanged for every other caller. - With it, the checkpoints this request drops right after the park (all of them when another state comes in, else the ones off this prompt's path) move into the parked image instead of being copied, and a background allocation refills STRATA_CKPT_POOL for the next save. Only the main session's checkpoints move, not an idle batch slot's (the park_source of #2526 serves both). - conversation_snapshot_save takes dropped checkpoints' storage (the existing STRATA_CKPT_POOL) for its copies; with the pool off, nothing is taken. - Unset: park_drops is null, no storage is taken, the K/V copies are the synchronous ones; the added work is the K/V size sum for the stage request and a few join calls. Differs from the PR as submitted: the PR's join in conversation_file.cpp (put_payload) is not ported; the SAVE/RESTORE handler of generate.cpp joins the park first (see the previous commit, which carries the PR's 33a8146 fix). Hand-merged into the park_source / ParkSource refactor of #2526. Not covered: if the background copy fails after resize_for_overwrite, the parked image keeps unfilled bytes; the failure is logged. A restore of such an image would read them. No test of the park path with a failing copy. Tests: the PARK_ASYNC path has no unit test in this commit (it needs a card with 1.5 GB free for the stage); the reporter's A/B/A conversation check covers it, see the checklist. Reporter's numbers (linxuhao, CMP 170HX): park 213 -> 64-75 ms (3K), 334 -> 68 ms (13K); greedy answers identical on their 11 prompts and the A/B/A conversation check. Theirs, not measured here. Original change by linxuhao (PR #3779, commit 3e0c141).
…DY=N, gathered small checkpoint copies STRATA_CKPT_GATHER=1, park traces (#3779, #3712) Problem (reporter, CMP 170HX, PCIe Gen2 x4): a conversation switch's park still paid 38-52 ms for the live state, because the STRATA_CKPT_POOL pool was empty at every park (a switch takes two checkpoints), so the 118 MB GDN vector was allocated fresh and page-faulted; and a checkpoint's 37 small synchronous device-to-host copies queue behind the 4 MiB DMA pieces of STRATA_CKPT_ASYNC. Change: - STRATA_CKPT_POOL_READY=N (opt-in, 1..16, with STRATA_CKPT_POOL on): the pool keeps N ready checkpoints, pre-sized to the session's state sizes with their pages in, topped up on a background thread after any take (not with a layer split). - STRATA_CKPT_GATHER=1 (CUDA, opt-in): the small running-state vectors (ple, and each layer's indexer tail, dead row and block position) are gathered device to device into one buffer and fetched with one copy into a pinned buffer. - The park stage is allocated rounded up to 64 MiB (a growing conversation does not cudaFree/cudaMalloc it every time) and a VRAM-resident K/V skips the pointer attribute query. - The background helpers (pool top-up, the checkpoint DMA, the park DMA) leave the host's pinned core (aux_cpus.hpp); no output depends on where a thread runs. - A parked image whose live GDN state is still being copied joins before its buffers are freed (SavedConversation::live_pending). - STRATA_PARK_TRACE prints the joins, the live-state vectors' readiness and the K/V sub-timers. Differs from the PR as submitted: the PR made STRATA_CKPT_POOL=1 allocate two ready checkpoints at startup (+236 MB of host RSS) and read STRATA_CKPT_POOL=N as a count. That changes what an existing switch does, so here STRATA_CKPT_POOL=1 keeps its meaning (the pool holds dropped checkpoints only, nothing is allocated ahead of time) and the ready pool has its own variable, STRATA_CKPT_POOL_READY. A user of the PR's STRATA_CKPT_POOL=2 sets STRATA_CKPT_POOL=1 STRATA_CKPT_POOL_READY=2. The PR's one-checkpoint refill after a park with STRATA_PARK_ASYNC is replaced by the top-up, as in the PR, and so is only active with the new variable. Merged by hand into #2526's park_source. With none of the new variables set: the stage size is rounded up only when the park stage is on, the extra code is trace output behind STRATA_PARK_TRACE and the join calls. Reporter's numbers (linxuhao, CMP 170HX): park at a 3K context 64-75 -> 21-40 ms (all park switches together, A/B/A conversation check identical, greedy 11 of 11 identical). Theirs, not measured here. Original change by linxuhao (PR #3779, commit 56088e7).
…sations' K/V in VRAM in front of the host copy (#3779, #3712) Problem: restoring a parked conversation copies its K/V from the host over the link: 27-107 ms over a PCIe Gen2 x4 link. Change, opt-in and only with STRATA_PARK_ASYNC=1 (CUDA only): a park also keeps the whole K/V of its image in a VRAM slot of its own (64 MiB granules, up to N MiB in all, only while 1.5 GB of VRAM stays free beside it); the background copy to the host reads from the slot, and a restore of that image copies device to device from the slot (about 1 ms) without waiting for the background copy. Each K/V layer carries its own join-and-release token, so any order of destruction is safe (the layers outlive their image when they are retained for the next park). Under VRAM pressure the elastic K/V's growth releases every slot and the stage, then retries once (both growth sites in generate.cpp). With the variable unset the slot code is not reached: the budget is 0, a restore finds dev_tag 0 and takes the host path, and the growth sites see conversation_park_vram_release_all() return 0 and behave as before. Tests: none of its own (it needs a card with the VRAM for the slot and a parked conversation); see the checklist for the reporters' A/B/A conversation check with STRATA_PARK_VRAM_MIB set. Reporter's numbers (linxuhao, CMP 170HX): restore of a 3K conversation 28-47 -> 7-12 ms; outputs identical. The first restore of a run is slower (41.6 / 56.9 ms) than later ones (6.5-12.6 ms) and they have not found why. Theirs, not measured here. Original change by linxuhao (PR #3779, commit bafca41).
…C and the park switches (#3779, #3712) One section in DETAILS.md for the switches the four conversation-cache commits add: what each copies where, its VRAM and host memory cost, that they stay off on HIP, gfx906 and Metal, and the measured times. Rewritten from the PR's section: the times are labelled as the reporter's (linxuhao, one CMP 170HX on PCIe Gen2 x4; we have no sm_80 card and measured none of them), the independent RTX 3090 data of bitaccesscomau is added, and the pool row describes STRATA_CKPT_POOL_READY=N instead of the PR's change to the meaning of STRATA_CKPT_POOL=N (STRATA_CKPT_POOL=1 is unchanged). Original section by linxuhao (PR #3779, commit a04f439).
…rnel nodes their stream's priority (#3779, #3712) Problem: the decode window graph runs every kernel node at the launch stream's priority, so the one-block indexer kernels of the side branch wait behind the queued blocks of the big q GEMV, and the shared expert competes with the routed experts. Change, opt-in (bitwise by construction: no kernel changes, the same launches on the same inputs; only the dispatch order moves): with STRATA_GRAPH_PRIO=1 the side mixer's branch stream (GDN gates, QSA indexer chain) is created at the device's highest priority, cs_ and side2_ in the middle, the shared expert's stream at the lowest, and the window graph is instantiated with cudaGraphInstantiateFlagUseNodePriority. If a driver did not copy the stream priority into the nodes, the capture-time notes (the branch forks and the priority streams' tails) set the attribute per node. One log line per captured window prints the node priority histogram. A card whose priority range is 0..0 changes nothing. CUDA only: the members and the three methods sit under !STRATA_USE_HIP && !STRATA_HIP_GFX906 && !STRATA_GPU_METAL; graph_prio_on() is false there. Merged by hand with #3702 (the window-graph eviction): instantiate_evicting now takes both keep_t and the flags, and the retry after freeing other window sizes keeps the flags (the merged code would otherwise have dropped the node priorities on that retry). With the variable unset: every stream is created as before and the capture-time notes return at once. Tests: none of its own (a priority is a scheduling hint with no value to compare). The window graphs are exercised by the engine; see the checklist for the reporter's greedy-answer identity check. Reporter's numbers (linxuhao, CMP 170HX, IQ3_S, --spec 3, 3K context, interleaved in one session, each arm with STRATA_QIL=1): reference 125.3 / 126.6, GRAPH_PRIO 127.8 / 129.0, reference 125.0 / 126.8 tok/s, +1.8 %. Greedy 11 of 11 identical. In the later per-feature pair together with STRATA_MMVQ_IL_PF: decode +1.3 %. Theirs, not measured here. Original change by linxuhao (PR #3779, commit 68c487c).
…ab), the clock warm-up and MMVQ_IL_TMIN (#3779, #3712) Harness only; no engine code. The bench part of mmvq_il_parity times the interleaved matvec cold (24 copies of the weights rotated past the L2), hot (one copy), hot with other weights' reads in between (--gap MB, what a window's GEMVs do between a prefetch and its consumer), and with the timed weights in a persisting L2 set-aside (--persist); --ab runs the five arms per case in one process, because the clocks ramp over the first ~100 ms of a run. A 150 ms warm-up of launches runs before the first timed case, and MMVQ_IL_TMIN bounds T from below. The parity part is unchanged: without --bench the same checks run in the same order (every rows choice against native_mmvq, bitwise). The bench code is CUDA only (the test is built with STRATA_ENABLE_CUDA); the persisting-L2 calls are left out under STRATA_HIP_GFX906. Original change by linxuhao (PR #3779, commit 1f4b813).
…eaved GEMV kernels for Q6_K (#3779, #3712) Problem: the interleaved small-batch matvec (native_mmvq_il) has one weight block in flight per warp: its four `if (kbx < bpr)` sections each sit behind their own branch, so on a GA100 (ncu: 432 GB/s on the qkv Q6_K matrix) the loads of a group are not issued together. Change, opt-in (every variant is bitwise native_mmvq): native_mmvq_il_pf_kernel issues every weight load of a group (4 virtual warps x R rows) before any of the group's dots, with the addresses clamped to the row's last block for a tail, and with the ring variants the next group's loads before this group's dots. The per-lane accumulation order is the plain kernel's. P2's forms replace conversions by exact ones where the value is a small integer: (float)p becomes p + 1.5 * 2^23 minus 1.5 * 2^23 for |p| < 2^22, and the saturating byte subtract of Q6_K's codes (0..63, so nothing saturates) becomes two instructions. Q4_K, Q5_K, IQ4_XS and Q6_K have the variants; Q3_K (the Q2_0 pack's type) keeps the plain kernel. - STRATA_MMVQ_IL_PF=1: the measured table (kIlPf_80 / kIlPf58_80: Q6_K only, read only when the card is sm_80; on every other card 1 does nothing); =2..6: one variant for every type and shape (A/B only, any CUDA card). - Unset or 0: the plain kernels through the same launch code (the variant argument is 0, and the added code is one cached environment read and a switch). - CUDA only; HIP has no interleaved path. Tests: mmvq_il_parity --pf a,b,.. bit-checks each listed variant against native_mmvq for every type and shape, T = 1..8 and rows 0/1/2/4, and with --bench times them interleaved (PFBENCH lines) and prints the joint (variant, rows) table in the source form of kIlPf_<cc>. New ctest entry mmvq_il_parity_pf runs it with --pf 0,2,3,4,5,6 on whatever card is present. Reporter's numbers (linxuhao, CMP 170HX): decode +1.4 % (two reference / on / reference / on sets) over the plain kernels; with STRATA_GRAPH_PRIO decode +1.3 % in their per-feature pair (4K prompts +3.9 %, 16K +1.2 %, 52K +0.2 % in the same pair, which the reporter reads as noise). Q4_K, Q5_K and IQ4_XS variants tie the plain kernels there. Theirs, not measured here. Original change by linxuhao (PR #3779, commit 116992a).
…TA_GRAPH_PRIO and STRATA_MMVQ_IL_PF (#3779, #3712) DETAILS.md gets two sections: the decode switches (GRAPH_PRIO, MMVQ_IL_PF) with the two sm_80-only defaults named, and the prompt switches (PF_FUSED_NT / BATCH / F16 / WS, GDN_TC, PLE_COMPACT / PLE_COPY_STREAM, EMBED_VRAM, DEQUANT_SIDE) with what is bitwise and what is rounding-level. MMVQ_IL_TABLE.md gets the sm_80 row of the table list and a paragraph on the hoisted-load tables (kIlPf_80 / kIlPf58_80, Q6_K only) and on measuring a card with mmvq_il_parity --bench --emit-table --pf 0,3,4. Rewritten from the PR's text: every figure is labelled as the reporter's (one CMP 170HX; we have no sm_80 card and measured none of them), and the prompt-switch section is new (the PR documented only the decode switches and the conversation cache). Original docs by linxuhao (PR #3779, commit 6d6f0f3).
… tables and the opt-ins) # Conflicts: # CMakeLists.txt # docs/MMVQ_IL_TABLE.md # src/kernels/cuda/native_mmvq.cu # src/kernels/mmvq_il_parity.cpp # src/program/generate.cpp
…llow STRATA_DEFAULTS=026: back to the 0.2.6 kernels and the shared table; explicit =1 and =0 still win; arch_defaults_test checks the rule; release-notes line
P0 #3911 #1503; P1 fixes; file-02 items; 22 PR ports (authors kept) incl. #3779 (sm_80 defaults under STRATA_DEFAULTS=026) and #4121 (STRATA_PIPED_MTP opt-in); audit fixes #3602 #3948 #3638 #3797 #2409; owner decisions #3935 (DeepSeek sm_89 packages), #3844, #3453 (sm_89 tables on), #3976 (layer-major prefill on by default for a 2-card CUDA layer split; revert 1e41fa269 to undo). Final gate: sm_60, SYCL, z440 sm_86 CPU 38/38 GPU 23/23, boston CPU 24/24 GPU 23/23, owner-PC MSVC CPU 36/36 GPU 20/20, snapshot 10/10 on z440 and boston.
|
Thank you. Fixed in 0.2.7: The sm_80 defaults and the opt-in sm_80 switches of #3779, STRATA_GRAPH_PRIO and STRATA_MMVQ_IL_PF (#3779, #3712). Your name is on the commit. We could check it with unit tests and builds, but not on your hardware or with your workload. Please try 0.2.7 and tell us how it goes. Release: https://github.com/Niko1221/Strata/releases/tag/v0.2.7 |
… forwarded to the workers - The coordinator's CONFIG carries "env": the pool-wide opt-ins set on it (kPoolForwardEnv: Niko1221#4005's STRATA_PREFILL_COPY_EVENT_EVERY / _COPY_WAIT_DEDUPE, Niko1221#3779/Niko1221#3712's STRATA_PARK_ASYNC / STRATA_CKPT_ASYNC / STRATA_CKPT_GATHER, and STRATA_ROUTE_PRIOR). A worker applies each one it has not set itself, before its first prompt, park, checkpoint or verifier reads it, and logs "strata pool: from the coordinator: ...". STRATA_DEFAULTS=026 is not forwarded (a worker reads it before CONFIG) and the VRAM/RAM-sized switches stay per PC. - <config>.pool.json takes "env": {STRATA_*: value} (checked: up to 32, STRATA_* names, short values without ; or newlines), applied to the engine of either split role through pool_env; the Pool tab's save keeps it. - Niko1221#4010's STRATA_PIPE_ADAPT_EVERY start-up line also on a pool coordinator (the pipelined groups drain the workers' stages before a round; the workers already swap between the windows they serve). - Checked for a worker: the async park's device copy is synchronised before the park returns (only the copy to RAM runs on), so a following prompt cannot overwrite what is being copied. - docs/POOL.md: "Opt-in switches for a split" (what each does, pool-wide or per PC, how to set them). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0157fjixQD2nRxvCvtXq66DD
|
Thank you for porting this into 0.2.7. Here are the results from the CMP 170HX (sm_80, PCIe Gen2 x4, 250 W, Qwen IQ3_S, every expert in VRAM), using the release tag built with the production CMake options. Each arm was run twice, interleaved. The sm_80 defaults are bitwise identical. 0.2.7 against 0.2.7 with
The opt-in switches work as ported. 0.2.7 defaults against 0.2.7 with the switches we run in production (the #3779 set plus upstream's
The greedy answers were identical (11 of 11). As expected, the multi-turn conversation check differs in wording, because 🤖 Generated with Claude Code |
|
Following your request in #3712 (2026-10-27) to try 0.2.7 on sm_80 and post the results here. All figures below are one A100 80GB (PCIe); multi-card numbers are not included. A100 (sm_80) results: 0.1.40.3 → 0.1.52 on Qwen3.8-Flash-Next, plus a self-built sm_80 DeepSeek engineTest box: 7× A100 80GB PCIe, Linux, driver 580.126.09, CUDA 12.8 and 13.1 toolchains, 128 cores, 1 TB RAM. One card is used per service (a dedicated card for Qwen, a 40 GB slice of another card for DeepSeek when sharing with production). 1. Confirming the CUDA-12 requirement (for #4461)
So on A100 the CUDA 12 package is the only choice — that matches what you said in #4461. It runs on driver 580.126.09. One nuance worth adding to the docs: the source tree does build sm_80 with CUDA 13.1. I compiled 2. Qwen3.8-Flash-Next IQ3_S: 0.1.40.3 vs 0.1.52 (same bench, same pack, same card)The 0.1.40-era pack loaded in 0.1.52 unchanged ( Method: Config:
A measurement hazard others may hit: the first request after loading ran at 116 tok/s (old engine) and 90 tok/s (new engine) because of cache effects. Averaging it with the next two runs inflates the mean to 84 tok/s and makes the new engine look 22% slower. Compare steady-state runs only. Two traps when moving an existing 0.1.40.3 setup to 0.1.52:
For completeness: 3. DeepSeek-V4.1-Flash-1017 on A100 (self-built sm_80 engine)
DSpark: Context: 131,072 verified. A 130,543-token prompt was accepted, first token after 212 s, and a needle placed at 2/3 of the document was answered exactly. The model's own limit is 1,048,576 (yarn factor 16 on 65,536) but we did not test longer. One config gotcha that cost me a day (may deserve a line in DEEPSEEK41.md): 4. Not tested here
Happy to run any specific measurement you want on this box. |
|
Thank you for running this on a real A100. One thing stands out: your absolute numbers are below what we get on the CMP 170HX. That card is a cut-down GA100 on a PCIe Gen2 x4 link (1.5 TB/s HBM, 250 W). Our results on 0.2.7 (engine 0.1.52), Qwen IQ3_S with every expert in VRAM:
An A100 should be at least as fast, so we think the gap comes from the configuration rather than the card. These are the differences from our setup that we have measured on our side:
If you have time, could you run two more arms on the same card and the same bench? Our "0.2.7 defaults" column above is arm A, and the switches column is arm B. A: same arguments as yours, except:
B: A plus these switches: With all of them on, our 11 greedy answers were identical to the defaults. Our multi-turn check differed in wording, because If your host has two sockets, One question about the side checks: 🤖 Generated with Claude Code |
|
Thanks — you are right that it was configuration, and worse: two of my numbers were measurement artifacts. Correcting them first, because they invalidate part of what I posted. Two corrections to my earlier comment
So the honest picture is the one you predicted: decode is level with your CMP; the prompt path is ~35% behind. And to answer your question directly: Arms, same card, same bench, corrected methodMedian of 3 × 512-token decodes with unique prompts; prompt = median of 2, new tokens only, unique documents. One A100 80GB PCIe (2.0 TB/s HBM, x16 Gen4), driver 580.126.09, engine 0.1.52 (cuda12 package), Qwen IQ3_S, 24,576 expert slots (= every expert), vision on.
For reference, your CMP: decode 109–111 / 124–128; prompt 3,260 / 4,290 at 52K.
Two notes on the flags you listed
What is left of the prompt gap, and a questionWith your switches we are at 2,370–2,479 against your 3,260–4,290, so ~30% is still unexplained, and I do not think it is expert residency. What differs from your setup on our side:
I tried the obvious version of it: Happy to run more arms; the box is a single card for Qwen so restarts are ~50 s. |
|
Thanks for the careful rerun. The build is not the gap. We ran your package (
"Stock" means only How we measure the prompt rate. We take it from the engine's own log line, for example
Our KV settings: Agreed on 🤖 Generated with Claude Code |
|
Both corrections accepted — thank you.
Your exact argument line, on our A100Same package (
So: with identical flags and an identical engine, our decode lands on your numbers, and our prompt is still ~35% below. I agree this is not build and not configuration any more. A caution on our numbers before you trust the shape of them: decode on this box has outliers we cannot reproduce. Single arms gave 209.7, 236.8 and 277.8 tok/s, and repeating the same arm gave 120.0 / 123.1 on one and 277.8 / 139.1 within one arm. Prompt is stable within ~5%; decode needs repeats. Everything above is a repeat, not a single run. Where I think the remaining 35% is, and a questionThe difference from your box that we have not eliminated: two sockets and 1 TB of RAM that is ~50% busy with other jobs. Two things that would settle it:
Adopted your KV settings meanwhile: |
|
The engine-side line is in the child engine's stdout/stderr log, not the server's progress stream. A clean run on my RTX 3090 emitted:
That line appeared without |
|
@bitaccesscomau's pointer was the missing piece: the line is in the config's Engine-side numbers with your exact argument lineSame package, one A100 80GB PCIe,
Decode on the same arms: 165.6 / 168.2 / 181.7 tok/s, drafts accepted 97/110, 145/157, 159/168. So on filler prompts PLE I/O mode makes no difference here, exactly as you predicted; and the rate is ~5,000 tok/s, i.e. above the CMP's 4,333 at 52K, with decode 166–182 against 124–128. The 35% "gap" was our environment, not the cardHere is what produced the low numbers I posted before, and I should flag it because other A100 reporters may hit it too: this box also ran a second model service holding 52 GB of VRAM and ~106 GB of host RAM. With that service running, our engine-side prompt line read So: under host-memory pressure the prompt path loses about half its speed on this machine, and it is invisible in the client-side numbers we were reporting. Your 3,257 / 4,333 and our 5,002 are consistent once the hosts are compared at similar occupancy. I also tried the pinning arm you suggested ( Two small things back to you: (1) |
|
Thanks, that settles it. Host memory pressure halving the prompt rate is worth knowing, and the engine-side line makes it visible. One note for anyone comparing against our numbers: our CMP is not at stock clocks. It runs with a +250 MHz V/F offset (an undervolt), a 1,650 MHz ceiling and a 250 W limit for Strata. Going from a 1,350 to a 1,650 MHz ceiling gave Strata +14% on prompts and +21% on decode on this card. Your A100 at stock clocks has 108 SMs against the CMP's 70, so it is no surprise that it comes out ahead once the host is quiet. 🤖 Generated with Claude Code |
|
That clock note explains the CMP's numbers well (and it makes the 108 vs 70 SM comparison honest). Two things from our side, both with PLE on A100: 0.3–0.5%, not 18.8%
Host side: So on this card PLE is not a factor at all — the prompt is ~36–38% expert GEMMs ( For anyone reporting numbers from a shared box: state the host's free RAMSame flags, same engine, same 78K filler prompt, four runs today:
That is a 2.2× spread from host occupancy alone, with |
Problem and result
sm_80 has no entry in the engine's per-card tables, so on a GA100 card most of the prompt and decode path runs on choices made for other architectures. #3712 has the measurements from a CMP 170HX (GA100, sm_80, 70 SMs, PCIe Gen2 x4) running Qwen3.8-Flash-Next IQ3_S with every expert resident. On our working branch, with these features and the switches turned on, prompts run up to 2.0x the speed of stock 0.2.1.4, decode +20 %, and conversation park/restore 5-10x faster.
This PR puts all of it in front of you in one place, rebased on current main. It is not meant to be merged as it is. Each feature is its own commit (or a few consecutive ones), so you can cherry-pick, or ask us to split it into the five thematic PRs listed in #3712, or into finer ones.
Related discussion
#3712 (scope proposal; no agreement yet). Also #3252 (A100 slowdowns, where sm_80 also falls back to the shared MMVQ-IL table) and #3453 (the sm_89 MMVQ-IL table, which touches the same table list).
Approach
Every switch is opt-in. With all of them unset, the default path is byte-identical to main on every card, with two exceptions on sm_80 only. Both exceptions are bitwise: the output stays identical.
STRATA_DEQUANT_WIDE=0turns it off)STRATA_MMVQ_IL_SM80=0restores the shared table)STRATA_PF_FUSED_NT=256256-thread fused IQ prompt experts,STRATA_PF_BATCH=Nwhole-layer launches (WideBatch)STRATA_PLE_COMPACT=1raw IQ4_NL PLE rows dequantized on the device;STRATA_PLE_COPY_STREAM=1their upload on its own streamSTRATA_EMBED_VRAM=1token-embedding table in VRAM (only while 512 MiB stay free)STRATA_DEQUANT_SIDE=1dense dequant on a side streamSTRATA_PF_FUSED_F16=1/STRATA_PF_FUSED_WS=1f16 mma fused experts (rounding-level)STRATA_GDN_TC=1DeltaNet prompt recurrence as a chunked tensor-core scan (rounding-level; sm_80+, 48+ SMs)STRATA_CKPT_ASYNC=1,STRATA_CKPT_VRAM_SLOTS=Nconversation checkpointsSTRATA_PARK_ASYNC=1,STRATA_CKPT_POOL=N,STRATA_CKPT_GATHER=1background park, ready checkpoint poolSTRATA_PARK_VRAM_MIB=Nparked K/V in VRAMSTRATA_GRAPH_PRIO=1verify-window graph nodes keep their stream prioritySTRATA_MMVQ_IL_PF=1hoisted-load, exact-integer MMVQ-IL kernels for Q6_KCompatibility effects:
STRATA_CKPT_POOL=1now also pre-allocates 2 ready checkpoints (+236 MB of host RSS).GRAPH_PRIOare CUDA only. On HIP, gfx906 and Metal they are off.Validation
Passed:
tools/quant_quality.py prompt/pcompare, set v1) of the rounding-level switches against the bitwise base. GDN_TC: mean 0.0157, median 0.0011. F16 / WS: mean 0.0062, median 0.00006. Both on: mean 0.0165. Top-1 agreement is 100 % in every arm. Details: [Engine] sm_80 (GA100) tuning from #3712, all five features in one PR: one commit each, to split or cherry-pick as you prefer #3779 (comment).Failed: none observed in builds. During the GPU runs, in 3 of about 35 engine starts the engine stopped answering after a park with evictions, each time on upstream's conversation-cache path (the PR 4 switches off). It did not reproduce in 15 reruns of the same sequences (5 greedy runs, 10 full benchmark runs on the stock build); details are in the results comment.
Not run:
Limits
🤖 Generated with Claude Code