Repository navigation
Conversation
medking82
marked this pull request as draft
October 4, 2026 14:07
medking82
changed the title
medking82
marked this pull request as ready for review
October 4, 2026 23:05
…ost OOM skips parking Co-authored-by: zengyishou <zengyishou@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…er cards --vram-reserve-mib is every card's reserve, but the card that drives the monitors needs more headroom than one that drives none: one value has to fit the display card, and the other cards then keep that headroom for nothing. --vram-reserve-later-mib N is the reserve of every card after the first (default: --vram-reserve-mib's value), in the split search and in each later card's cache sizing. RTX 4060 Ti (no display, layers 0-19) + RTX 5080 (display, 20-47), Swift 1.5 IQ3_XXS at 160K: --vram-reserve-mib 1000 gives the 4060 Ti's cache 4551 slots; --vram-reserve-mib 300 --vram-reserve-later-mib 1800 gives it 4867, and the 5080 keeps 1.3 GB free with everything loaded.
--resident-experts (the low-RAM mode's resident variant) ran as plain --mmap-experts as soon as a layer split was set. It now keeps its RAM copy on a split: - the experts every stage's cache holds are left out of the RAM copy, not only CUDA0's; - when the rest does not fit whole (#467), the copy keeps the hottest by the whole expert profile. Each stage only keeps its own layers' share of the profile, and ranked by CUDA0's share alone the copy held layers 0..K-1's experts and read every later layer's from the files (4.0 GiB in RAM instead of 21.2 on the PC below); - an adaptive swap copies the evicted expert back into RAM from the cache of the card that owns its layer, on that card's stream. host_res holds slot numbers of whichever cache owns the layer: reading a later stage's slot out of CUDA0's cache put another expert's bytes in RAM, and after the first swaps the answers turned to garbage. (Latent until now: the guard kept the mode off a split.) - the decode profile prints every stage's GPU profile, not only CUDA0's. Remote expert caches (--expert-cache-remote) are still refused with it. Swift 1.5 IQ3_XXS at 160K, q4_0 KV, RTX 4060 Ti (layers 0-19) + RTX 5080 (20-47), i9-14900KF, 32 GB of RAM, Windows 11, four greedy prompts, decode tok/s (first round / fourth): - 5080 alone, --resident-experts: 25 / 29 - split, --mmap-experts (what --resident-experts became): 32 / 64 - split, --resident-experts (22 GiB of experts locked in RAM): 71 / 69 With --pcie-frac 0 --adapt-every 0 the split's greedy output is the same as --mmap-experts' (4/4 prompts). With adaptive swaps on, 33 varied requests at temperature 0.6 (thinking on and off, five languages, code, each checked for degenerate output) exchanged 81,783 experts with the two caches and gave no degenerate answer.
With an engine that keeps --resident-experts on a layer split, setup no longer steers a low-RAM PC with two cards to one of them: - at install, the low-RAM mode on several GPUs takes the resident variant by the same rule as on one card (what no card holds fits the RAM with LOW_RAM_HEADROOM_GB beside it), without the one-card question; - at a start, a resident config on several GPUs keeps --resident-experts (split_mmap leaves it alone); - a resident config started on a PC with two cards is offered both, "y" by default, as any other config. It is gated on MIN_ENGINE >= RESIDENT_SPLIT_ENGINE (0.1.40 here, a guess at the release that would carry the engine change), so until then setup behaves as before and the golden configs are unchanged. Two tests run the new paths with MIN_ENGINE patched; tools/test_setup_*.py pass (260 tests).
…r stages say which stage Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…non-blocking streams read the table Every upload of the expert residency table (host_res -> d_res, and each stage's) was a plain cudaMemcpy: a pageable copy on the legacy stream, which may return before its DMA has landed, while the kernels that read the table run on streams created cudaStreamNonBlocking, which do not wait for the legacy stream. A kernel enqueued right after could read a table half old, half new: an expert computed by both the GPU and the CPU, or by neither. The ~100 KB copy lands in microseconds, so this rarely fires, but nothing ordered it. The uploads now go through res_put(), which waits for the copy on the current device's legacy stream only. Same values, same output. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…M; a -1 raises a plan error flag (#871) A prompt loan, a VRAM shrink or an adaptive swap marks experts -1 in the residency table. The zero-doorbell graph planned from the device table and returned silently on a -1, so the window ran a stale plan (!!!! at 100% residency). refresh_ar() checks the host table before each window and runs the doorbell graph (its own captured variant) while any expert is out; resident_plan also writes a mapped error flag the host checks after the window. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ill::init (#828 test switch) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…abort latency) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…s and an incomplete checkpoint (#752) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…bundle) onto 0.1.39 Merges f-halo-int (36fa455 + 56 commits: the gfx1151 target, WMMA prompt attention / GEMM / scorer, GDN recurrence, fused gfx11 prompt experts, decode switches, drafting) with f-halo-int-S (STRATA_HCD_EXACT, 16bc95a) and f-halo-int-W (table16, float4 combine, LDS alias + occupancy, d23d4e2). Conflicts resolved against 0.1.39's own all-resident verify graph, forked shared-expert stream, fused swiglu q8_1, PLE batch, GDN commit split and peer prompt path; the gfx1151-only code stays gated (build fixes and the arch defaults follow in separate commits). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…NT is 0.1.39's all-resident graph 0.1.39 already skips the host round trip when every expert is in VRAM (all_resident_, on by default, #871's refresh_ar), so the port keeps that graph and adds only what halo built on top of it: the LFUSE gated combine. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…LE_HIP (it includes hip_runtime.h; the CUDA build broke) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… zeros on gfx1103 / gfx1152), cached per device WmmaTraits::mma's #else returned its accumulator (zeros) for a gfx11 part outside the __gfx1100__..__gfx1151__ list; it traps now, and the runtime gates of the WMMA GEMM, the prompt attention, the block scorer and the HC up-projection take exactly gfx1100 / 1101 / 1102 / 1150 / 1151 (strata/kernels/gfx_arch.hpp) so a gfx1103 or gfx1152 falls back to hipBLASLt instead of launching code that is not there. The gate is cached per device (a process-wide static answered for the first GPU only; hipGetDeviceProperties on every call was the review's other complaint). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…1151 only (arch_defaults), the rounding-level ones stay opt-in 16 switches whose output was byte-identical to the switch off on Aurora (GDN head / PP / CONVL2 / NOY / SPLIT, QFUSE, CVEC_FUSE, HCD_EXACT, PF_PAD, Q8 / Q6 packed, MMVF_ROWS, ATTN_LANECELL, EXPERT_V2, TSUM, LFUSE, PLE_BATCH) are set at start-up when device 0 is gfx1151 and the user has not set them; STRATA_GFX1151_DEFAULTS=0 turns the table off. The GDN head kernels are AMD-only now (HIP builds; CUDA keeps its key-head kernel), and verify.cpp's namespace-scope env reads became lazy so the table reaches them. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…, the DONE line's newline escapes Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…o tuning table, or a solution other than 1176 / 1177), and hip_prefill_hcd_exact_parity checks the kernel against hipBLASLt The kernel reproduces one hipBLASLt solution's summation order; the test compares both on the same inputs for T from 64 to 16384 (bitwise) and SKIPs without STRATA_HIPBLASLT_TUNING. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
PleTable had one bool per format (q5_0, fp8) and an if-chain in decode(), open() and format(). Now one list, ple_formats(), holds one entry per table type: its GGUF type name, the bytes of a 160-value row, whether it carries a scale and its dequantizer. open(), format(), the row size, the error string and the row buffers all read it, so the next formats (Q8_0, Q5_1, Q4_0, BF16) are one entry each. No behaviour change: IQ4_NL, Q5_0 and FP8 decode with the same functions, and the "not IQ4_NL, Q5_0 or FP8 (I8)" error text is built from the list and reads the same. ple_reader_test --selftest now writes a PLE-only GGUF for every format in the list and checks both readers (direct and mapped) against the format's own dequantizer, for single rows and a 16-row token gather. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Ordinary GGUFs ship per_layer_token_embd.weight in other formats than IQ4_NL: Unsloth's UD-Q6_K_XL and Swift-1.5 Q4_K_L as Q8_0 (170-byte rows, ~54 GB), a Q5_K_M finetune as Q5_1 (120-byte rows). The engine refused both at start. - dequantize_q5_1 in strata/artifact/dequant.hpp (ggml's dequantize_row_q5_1: unsigned 5-bit codes times d plus the block minimum m); Q8_0 rows use the existing dequantize_q8_0 - two entries in the PLE format list (the table-driven commit before this one): Q5_1 and Q8_0, on both readers; PLE_ROW_BYTES_MAX grows to 170 - tests: dequant_q5_1_test (synthetic blocks against ggml, no model, no GPU), dequant_bf16_test knows Q5_1, ple_q5_parity takes Q5_0, Q5_1 or Q8_0 tables and checks both the mapped and the direct reader Ported onto the format table by the maintainer from anon761's PR #651 (the --ple-io ram fault-in is the next commit). Quality of the Q8_0 table: ksanislo measured 0.53% error against the BF16 table, FP8 2.66%, IQ4_NL 7.59%. Co-authored-by: ksanislo <ksanislo@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
PleTable::open compared rows x row_bytes against the file after the 64-bit product could already have wrapped, and mapped the tensor before checking it fits. Now: refuse a row count that overflows rows x row_bytes and a data section that starts past the end of the file, check offset and size without an addition that can wrap (offset > have, need > have - offset), require a single-tensor shard to start at offset 0 and fill the file, and map the table only after all of that. close() also resets the FP8 scale and the byte counter. ple_reader_test --selftest writes tables that claim one row too many, one too few and a wrapping row count, for every format, and checks both readers refuse each with a message and stay closed. Checks taken from CC-David-CC's #865 (its opt-in Q8_0 table is #651's; the direct reader reads Q8_0, so the experimental switch and the mapped-only limit are not needed). Co-authored-by: CC-David-CC <CC-David-CC@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
#651) mlock alone brings a cold file in one page at a time from one thread, ~0.47 GiB/s (about two minutes for a 54 GB Q8_0 table); 16 threads reach ~2 GiB/s on anon761's machine. The pages are then resident and mlock only pins them. When mlock fails, the pages stay faulted in (the message says they may be reclaimed) instead of being touched again one by one. ple_reader_test gets --make-table PATH FORMAT (a big synthetic table, with --rows) and --ram-time PATH (open mapped and locked, print the seconds) to time it cold. Ported by the maintainer from anon761's PR #651. Co-authored-by: ksanislo <ksanislo@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
A plain llama-quantize Q4_0 file of this model (bartowski's Qwen3.8-Flash-Next-Q4_0 and Swift-1.5 Q4_0: gate/up Q4_0, down Q4_0 on 42 layers and Q4_1 on 6, a Q4_0 PLE table and token embedding) was refused at start: no GPU kernels for Q4_0/Q4_1 experts, and the PLE reader took only IQ4_NL, Q5_0 and FP8. Follows #222 (Q4 to Q6 quantizations). - Q4_0: Fmt<2> with the codes centred before dp4a (q - 8 as signed bytes), the exact integer form ggml-cpu's q4_0 x q8_0 dot uses, rather than llama.cpp's CUDA chain (uncentred codes minus 8 x the q8_1 float sum): on a real Q4_0 expert the CUDA chain measured 1.4e-2 relative error against the float product, the centred form 5.3e-3. - Q4_1: Fmt<3>, the Q5_1 chain without the fifth bit and the same min-term choice. - dq_q4_0 / dq_q4_1 for the prompt path's dequant fallback; row bytes, is_iq and the GU/D/MMVQ format lists. - PLE: Q4_0 rows (90 B, as IQ4_NL) through the same reader: one entry in the PLE format list (ple_format_test and ple_reader_test cover it). - docs/UNSLOTH_Q4.md: a Q4_0 section with the measured run. native_expert_parity --synthetic (now in the CMake matrix): q4_0/q4_0 CPU 1.29e-2, GPU 1.29e-2, CPU-GPU 1.1e-7; q4_0/q4_1 CPU-GPU 3.2e-4; GPU dequant bit-exact against to_float. The existing pairs are unchanged. On bartowski's file (revision 928589f), layers 0 and 5 (Q4_1 down) and 20 and 47: CPU and GPU 1.1-1.3e-2, CPU-GPU 1.1e-7 (Q4_0 down) and 5e-4 to 1.1e-3 (Q4_1 down). Served on one RTX 3060 (all experts in RAM, 131K context): 24,461-token prompt 787 tok/s read, 24.0 tok/s out; 67,934 tokens 737 / 26.4, the hidden fact found. Ported onto the format table by the maintainer (the PLE part of the PR is the Q4_0 table entry). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The full checkpoint stores the table as BF16 (320-byte rows, 102.4 GB), the source of record; the engine gets it as IQ4_NL (about 8% off per row) or FP8. A bfloat16 is the top half of a float32, so widening is a shift: exact for every value, no rounding and no scale. One entry in the PLE format list (type BF16, self-describing: no scale, no metadata) and bf16_dequant_row; the direct reader takes the 320-byte row like any other width, and PLE_ROW_BYTES_MAX is 320. ple_reader_test checks the widening is exact over the patterns a lossy path breaks on (subnormals, +-inf, both NaNs); the generic every-format test covers a synthetic BF16 GGUF through both readers. Not in this commit: the PR's rename of tools/ple_fp8_pack.py (nothing needs it) and any setup option. Whether a table format other than IQ4_NL is worth offering is a quality question (see the measurement in the 0.1.40 notes). Ported by the maintainer from constantindjonkam's PR #586. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…range) so v0.1.39 compiles The sycl port went in as 5047172 ("sycl: port 0.1.39"). Two engine-side changes were merged into the release branch after it (cb97e6c, 1189218) and nothing in sycl/ was updated for them, so the v0.1.39 tag does not compile: - #626 (ec35cf0) made the thread affinity a value: pin_current_thread() returns strata::kernels::cpu::ThreadAffinity and restore_thread_affinity() takes it. sycl/include/strata/core/session.hpp still stored a `long long` (no conversion exists, and it did not include the header that names the type); sycl/src/core/session.cpp still assigned the returned struct to it and passed the `long long` back. - #559 (ba707d8) added the two layer-range arguments to NativeDense::load. The sycl out-of-line definition was still the 4-argument one and no longer matched the declaration in include/strata/core/native_dense.hpp. Both are fixed by mirroring the CUDA sources: the same member type, the same `outside()` predicate, and the same place it is applied in the tensor loop. `pinned` now follows CUDA as well: it is pinned_core.valid, so a pin that failed is not "restored". Nothing changes when the pin succeeds. No caller passes layer_lo/layer_hi yet, so `outside()` is never true today: this restores the build and does not change what the engine computes. Measured here: Arch Linux, Intel oneAPI DPC++ 2026.0.0 + oneMKL 2026.0, SPIR-V (JIT), cmake -S sycl -B build-sycl -G Ninja -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx, then `cmake --build build-sycl --target strata -- -k 0` -> 0 errors, links. The two errors are a type and a declaration mismatch, so the compiler version does not matter. Not built here: the other targets (kernel parity tests, benches) and the AOT path. Not run on an Arc in this change.
The migration turned `q = cudaStreamQuery(cs)` into `DPCT_CHECK_ERROR(cs->ext_oneapi_empty())`. That macro yields 0 whenever
the call does not throw, whatever the queue holds. The guard `q != 1 && *seq < want` (verify.cpp, session.cpp) therefore
fired on any ring that took more than 2 ms and reported "never rang (graph finished)" while the graph was still running.
With every expert resident the host does not wait per layer (STRATA_VERIFY_NO_HOST), so the path measured on the B70 never
reached it. With host-mirrored experts the first layer whose GPU time passes 2 ms does: layer 15 of the Coder on one Arc Pro
B60, every time.
q is now 0 when the queue is empty and 1 while it is running, as the guard and the trace comment ("1 = still running") expect.
fixups.py gets the same rewrite (entry 8), so a fresh dpct run of session.cpp produces it too.
Measured on one Arc Pro B60 (4,042 of 12,288 Coder IQ1_M experts mirrored in RAM, 128 greedy tokens, AOT bmg-g21): before,
exit 1 at layer 15 in 5 of 5 tries; after, exit 0 in 7 of 7 runs, at 11.2-11.9 tok/s with coherent text.
…tomic in kv_stream resolve (#783 PR-a) A CTA of the fixed 256-CTA grid whose first block is past the batch's largest n_bid returns before staging the queries. Same scores, same bits. STRATA_QSA_EARLY_EXIT=0 restores the old staging. kv_stream resolve_kernel: only threads that looked something up add to the shared counter. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…with an all-outputs path, and the gdn_parity T=1..8 memcmp check (#783 PR-b) gdn_ab_multi is instantiated per T=1..6 (generic above) with the weight unpack hoisted out of the token loop; gdn_step_norm_multi prefetches the next token's q/k while the current one runs, and when no state is written and every token is an output (the verify half) it skips the per-token bookkeeping. gdn_step_norm_kernel splits its reduction buffer so one barrier goes. gdn_parity gains section 5: for every T = 1..8, gdn_ab_multi against fused_gdn_ab, and gdn_step_norm_multi (verify half, commit half, t_out_begin, state-only commit) against the single-token kernel chained, compared with memcmp. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…sigmoid scale (#783 PR-c) native_router_top10_multi runs one warp per token (8 per CTA) instead of one CTA per token; native_moe_combine(_multi) takes a float4 kernel for k=10 with a shared row on aligned buffers; the shared expert's sigmoid-scale is vectorised. Same arithmetic order, so same bits. New native_multi_parity (memcmp): the multi router against the single call for n=1..19 on random, tied and NaN logits, and the combine against a host replay of its contract for the float4 and scalar paths. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ommit graph and the MTP catch-up (#783 PR-d) kv_append_q8_steps / kv_append_q4_steps append n cells in one 3D-grid launch per pool, and native_qsa_indexer_append_steps walks the n indexer cells in order inside one CTA (one barrier between cells). Verify record_window, the commit graph and MtpDrafter::record_forward use them; a multi-slot batch keeps the per-token loop (each row has its own K/V). STRATA_NO_BATCH_KV_STEP=1 restores the per-token calls. The kv single-cell calls are now the n=1 case of the same kernel. kv_steps_parity (new, memcmp) covers INT8, Q4_0 and the hybrid's folded single-plane calls for n=1..8; qsa_parity gains append_steps against the sequential indexer append at strides 1 and 5 under none/linear/YaRN. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
PR-e) The plan kernel grouped duplicate experts and numbered the groups with thread 0 walking all entries; every entry now finds its group in parallel and a packed (entries << 16 | groups) warp prefix scan numbers them. The residency check, the empty plan and *plan_err of #871 are kept as they are. static_asserts pin the 4-warp / packing limits. resident_plan_parity (new, memcmp): n=1..80 entries over pools of 3..512 experts with and without a slot table, with and without a skip word, plus the non-resident case, against a host replay of the plan's contract. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…he MTP drafter, off on HIP (#783 PR-f) native_qsa_rms_norm_rope does norm<256> and the rope in one launch with the same FMA order as the pair (register barriers keep nvcc from contracting across the norm/rope boundary); rope_parity check 6 compares it with native_qsa_rms_norm_weighted + native_rope_apply by memcmp for head_dim 128/256, strided q|gate and contiguous inputs, unscaled and YaRN. The in-place rope also launches one 32-thread block per row. Verify's norm_rope and q split, and the MTP drafter's q path (all T tokens at once; K for a one-token window) use it. native_norm_rope_usable() gates it: on the CUDA build it is on unless STRATA_NO_NORM_ROPE=1; on the HIP build it is off until rope_parity check 6 passes on AMD hardware, and STRATA_NORM_ROPE=1 turns it on there. The MTP q path's split, FWHT and gate-apply now run over all T tokens in one call each (row-independent). Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The 4-token accumulator variant of the fused GR down read keeps the same tile, block size, accumulation order and variant as the kFusedGrMaxT one, so its bits are the same; #832 measured +2.8-3.5% decode with it on an RTX 5070. Unset, the engine now turns it on for compute capability 12.0 only (decided once per device), and off on every other card and on the HIP build, where it was not measured. STRATA_GR_DOWN_MAX4=0 or =1 still decides explicitly. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…6 projections (#783 PR-g) gr_read_multi / gr_write_multi (and native_gr_rms_norm_weighted_multi, native_gr_pre_gated_multi, native_gr_post_multi under them) run T tokens in one launch per step; the verify window's PLE writes, last-layer write and final mixer, and the MTP drafter's write and final mixer use them. The BF16 MMVF multi kernel and the GR up / staged down kernels are instantiated per exact T (1..6, generic above), the BF16 weight tile is unpacked once for all tokens, with the same FMA order as before. STRATA_NO_MULTI_GR=1 restores the per-token calls and the generic kernels. STRATA_GR_DOWN_MAX4 keeps its meaning from the previous commit (sm_120 default on, explicit setting decides); this commit does not change that default. gr_multi_parity (new, memcmp): gr_read_multi (with and without injection) and gr_write_multi (in place and out of place) against the single-token calls at the real geometry, T=1..8. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…MVQ columns, decoded-once small-format tables, two rows per block (#783 PR-h) iq_kernels.cu: row_dot_80_sub16 (two rows per warp, all lanes active) for the n_embd == 2560 gate/up of every split format, exact-N instantiations of the multi-column MMVQ (launch_mmvq_multi_nc), the 3-entry tail in the grouped gate/up and down kernels. native_mmvq.cu: SmallTraits for Q4_0, Q5_0, Q8_0 and IQ4_NL decode the weight chunks and block scale once per (row, block) and apply them to every activation column with the arithmetic of small_q8_dot, and the multi-column launch runs two rows per block for any K. Kill switches: STRATA_NO_SUB16_GU=1 (sub-warp gate/up/down), STRATA_OLD_IQ_MMVQ=1 (per-column kernels, existing), STRATA_NO_MMVQ_ROWS2=1 (one row per block at large K). mmvq_multi_parity, iq_multi_parity, native_grouped_parity and native_expert_parity pin the columns to the single-column results. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…PR-i) MtpDrafter::draft now runs the round for T = a + 1: the catch-up writes the K/V of the accepted cells [0, a) and the full layer covers row a, instead of writing all T cells and rewriting cell a. No later attention reaches a cell past the one being drafted and the next round writes it first, so the draft tokens and probabilities are unchanged. draft_first (a = 0) no longer writes its repeated rows at all. STRATA_MTP_CATCHUP_ALL=1 restores the whole-window catch-up (and the old staging sizes). Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…erts are in the window (#783 PR-j) STRATA_SH_FORK_LATE=0 forks at the top as before. The other two commit-4 changes of the PR are not taken: the shared-expert Q8_1 reuse is already in f-140 as the opt-in STRATA_VERIFY_QDEDUP (and the fused q8_1 store), and the PLE native-key batch is already in f-140 as the opt-in STRATA_PLE_BATCH=1; both would have changed output on the native packs by default. Dropped on purpose: STRATA_VERIFY_PRELAUNCH and the dead copy_rows_from_mapped<ZERO_HITS>. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…GLU + Q8_0 in the grouped S2 experts (#783 PR-k) quantize_q8_0_scaled_kernel reduces the block's max with a warp shuffle (one warp per 32-value block) instead of one thread walking the block; fmax ignores NaN the same way in both. moe_grouped_s2 fuses SwiGLU with that quantize and writes the activation back to gate_up, so the grouped parity checks see the same buffer. STRATA_S2_SWIGLU_Q8=0 keeps the two launches; a static_assert pins FF % 32 == 0 for the fused kernel. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…MTP wiring) For a window of more than one token the drafter routes with one multi-column router GEMV and one multi-row top-10 and combines all rows in one launch (verify already did; each row is bitwise the single-token call, native_multi_parity). Needs the native router (512 experts, k = 10) and the native combine; otherwise the per-token calls stay. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Found by native_multi_parity on gfx1151 (Aurora): the float4 combine differed from the scalar kernel there, because the first pair's other product was contracted into the FMA. Written as fmaf per expert it is the pinned contract (first product rounded, FMA per expert in order) on every compiler. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…her the state does Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…== 0 branch (#783 PR-b) On gfx1151 the shuffle reduction of the sum of squares inside the divergent-looking branch rounded differently for ~3% of the outputs than the single-token kernel (gdn_parity section 5, Aurora); the base kernel had the shuffles outside it, and with them back out the multi kernel is bitwise the single one for T=1..8 on gfx1151. No change for CUDA. Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…kernel is not enabled (#783 PR-f) On the HIP build the fused kernel is off until check 6 passes there, so check 6 reports SKIPPED unless STRATA_NORM_ROPE=1 asks for it. (gfx1151 on Aurora, 0.1.40: with STRATA_NORM_ROPE=1 it does not pass - about 7% of the rotated elements differ - which is why the kernel stays off on HIP.) Ported from #783 (stuchapin909, head 4327807), split per piece. Co-authored-by: stuchapin909 <178119976+stuchapin909@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…t of 6d60dfc1) Bitwise on the 5070 (identity gate: Q2_0 and IQ3_XXS identical 10/10 with it on against off, same exe), but the 10-pair interleaved decode A/B (Q2_0, default config) shows no win: B/A medians 0.97 / 0.97 / 0.99 on the three greedy requests. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… keep f-140's file (#783 PR-k) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…_MULTI_GR flag (#783 PR-g) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The default STRATA_HIP_ARCHS of tools/hip/build_windows.bat gains gfx1151. The ROCm 10.2.0a20260930 Windows SDK compiles the strata target for gfx1151 (stream 140-j built it on this PC) and its rocm-sdk-device-gfx1151 wheel carries the rocBLAS kpack and the hipBLASLt kernels package_windows.py copies, as for gfx1201. Unvalidated on a Windows Strix Halo. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…emory, recommend UD-IQ4_XS gfx1151 is now a supported AMD arch (TheRock gfx1151 index; HIP build on Linux, docs/STRIX_HALO.md). The check is exact (gfx_arch_is, as the engine's gfx_arch.hpp): gfx1150 / gfx1152 / gfx1103 are named as Strix Point / Krackan / Phoenix-Hawk Point and stay unsupported. PCI 1002:1586 (one id for the 8060S / 8050S / 8040S) and the Windows adapter names map to gfx1151; the KFD simd_count names the 8060S / 8050S / 8040S. Unified memory: the BIOS carve-out + the shared pool (GTT; Windows half the RAM) size the context; only the carve-out counts as extra room in the low-RAM mode. A discrete card is picked before an APU. A Windows zip without gfx1151 code is refused with the reason. UD-IQ4_XS is the recommended model from ~80 GB of memory; setup changes no host setting. tools/test_setup_strix_halo.py uses Aurora's real sysfs. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… wheel's library/gfx1151/) instead of crashing Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
…e-VRAM readout after the release vmm_test failed on the RTX 5070 / driver 610.47: the pageable host-to-device copy was still in flight when a.unmap() pulled the chunks out from under it, an illegal memory access that stuck to every later call (the last two checks, the 8-chunk free-VRAM gain and cudaGetLastError, failed as a result). It was not a VRAM leak: released memory comes back at once (18 MiB gained for 9 chunks of 2 MiB, polled 10 s). generate.cpp already synchronizes before it unmaps. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
…n Linux rejects its std::array locals) dma_batch_parity.cu failed to compile on Linux/WSL (nvcc 13.0 with g++): 'template argument 2 is invalid' at the std::array<T, cap> locals of copy_expert_blobs, whatever the form of the size (int, size_t, a literal). The engine objects and the Windows build were not affected; the Linux tests build was. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
Owner
|
Heads-up: today we cleaned up this repository's commit history, and GitHub closed this PR automatically when GitHub cannot reopen it, because its branch is built on the old history. If you want to keep it going: start a new branch from the new Sorry for the noise, and thank you for the contribution. |
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
19 participants
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.
An opt-in memory policy resizes expert RAM/VRAM caches while the model remains loaded, preserving the native process, generation session and KV state. This contribution also adds Automatic, Full, Daily, Busy and Off resource controls to the existing full Monitor. The policy and resource presets remain disabled by default; request-boundary reload mode and the original percentage policy remain available.
Desktop resource presets
Full targets 2 GiB free RAM and a 256 MiB VRAM reserve; Daily targets 4 GiB / 700 MiB; Busy targets 8 GiB / 1,536 MiB. These are reserves, not actual cache allocations or throughput guarantees.
Automatic starts in Daily, selects Daily while Codex/ChatGPT is present, selects Busy for sustained external application CPU/RSS load, and selects Full after confirmed absence and low external load. Conservative transitions require 8 seconds; relaxed transitions require 60 seconds of valid advancing readings. Manual selection persists until Automatic is restored; Off restores the configured percentage policy.
The optional sampler uses process identity, CPU counters and RSS in the existing telemetry loop, excludes the server/native process family, and exposes only sanitized preset status. Absent or disabled presets do not sample processes. Missing, stale, incomplete or replayed readings cannot earn a relaxed transition.
Stable fixed-preset growth requires 30 seconds of fresh headroom and proposes at most +2 GiB RAM / 128 MiB VRAM reserve reduction per step. The configured sustained-pressure window still wins. Retargeting preserves actual allocation, pending acknowledgements and the native-error growth retry gate; no selection unloads or restarts the model.
GET/POST /v1/resourcesretains API-key admission, own-page JSON origin checks, bounded request bodies and no wildcard CORS for resource controls. Only a complete validatedenabled/selectionblock is saved, atomically, before live retargeting; the existing backup is preserved. A failed save cannot change the live selection. The UI distinguishes saved targets from actual, pending, limited and failed native allocations, and guards stale metrics responses after a selection change.Existing live-policy behavior retained
Matching applied acknowledgements own actual cache sizes. Sustained pressure never grows the other cache; stale capabilities, failed writes and partial/error acknowledgements cannot fabricate successful recovery. Live mode rejects parallel batch capacity owners and enabled
--vram-elastic; those upstream features remain available with the policy disabled. Safe reader boundaries, guarded prefill borrowing, STOP/EOS and complete continuation reuse remain documented in LIVE_MEMORY.md.Optional
recovery_secondsdefaults to 0; nonzero values must be 30–3,600 seconds. Under the original percentage policy, only matching applied RAM-pressure relief arms recovery toward the actual pre-pressure ceiling, at most +2 GiB per fresh qualified window. Limited acknowledgements stop fast recovery, and native errors clear the episode without bypassing the growth retry deadline. Ordinary percentage-policy expansion retains its 120-second debounce and 600-second cooldown. Fixed presets use their separate bounded 30-second path.The omitted
min_ram_headroom_gibdefault remains 5.5 GiB; explicit values from 2 GiB are accepted. The pressure default remains 60 seconds; explicit values from 2 seconds are accepted. Percentage RAM targets retain both their percentage and absolute floors. An explicit VRAM floor of 0 retains percentage headroom and native admission/workspace guards, including live startup's 256 MiB minimum. The policy and presets do not change the model, quantization, context/KV allocation or sampling settings.Upstream compatibility
Based on upstream main
1735d6471df29b42c26170efaac1f1446a58640f. The integration preserves upstream exchange rotation, resident layer-split support, lookup-chain drafting, fatal-engine recovery and slot-save configuration. Live RAM blocks keep their own ownership path, including the public resident accessor. Live mode explicitly rejects--adapt-async 1, whose split exchange path assumes a fixed resident arena; that upstream mode remains available without live resizing. Chained drafts remain bounded by the request's remaining output allowance, preserving the emitted continuation prefix.Validation
Contribution head:
15a59d4785f492b4df3fb78862373a5383696450.generate.cpp,expert_source.cppandlive_memory_test.cpptranslation units compiled.Earlier daily runtime observation
The following measurements are historical evidence from the separately deployed daily source, not performance validation of the refreshed contribution head.
Functional GPU/browser acceptance used separately deployed daily source
f1183b586ba1a25bb88485b50471b6ee0f2f907b, with its separate sidebar/context features and unchanged native artifact. Eight model runs retained the same native PID/create-time and 64K context without native error; HTTPS Monitor selections Daily/Full/Busy and Automatic saved and read back correctly, and Automatic was restored. With Thinking Off, approximately 6,765 input tokens and two 512-output runs per preset, output-weighted decode rates were Full 63.4, Daily 62.0 and Busy 53.7 tokens/s. These capped Off samples have no matched pre-change baseline and do not establish a controller speedup or whole-request High throughput.The uncapped High-thinking Auto/Busy request completed at natural EOS: 33,159 input / 5,741 generated tokens, averaging 32.2 decode tokens/s over the whole request (208.502 seconds wall time). The requested 50+ whole High target was not achieved. Native resizing cost, cache geometry/warming, file misses, output length, speculative acceptance and desktop load remain relevant; successful functional selection is not a universal performance claim. Earlier real-request limitations and measurements remain documented with their original inputs. Only aggregate observations are published; no private conversation, credentials or runtime configuration is included. The daily branch must not be replaced directly by this contribution branch.