Repository navigation
Conversation
… and sum finite too 0.1.39 clamped two q8_1 activation quantizers (Niko1221#606: an fp16 sum or scale past 65504 rounds to inf, the dot products then read inf * 0 = NaN, and a NaN in the residual ends as one token repeated). Two more store the block unclamped: - native_swiglu_quantize_q8_1_kernel (native_mmvq.cu), the shared expert's fused SwiGLU + q8_1, on by default (STRATA_FUSED_SWIGLU_Q81) in decode and the multi-token pass; - native_gu_fused_kernel (iq_kernels.cu), the routed experts' fused SwiGLU + q8_1 (gfx906 builds, STRATA_EXP_MODE=8). Both now go through q8_1_finite / q8_1_quant / q8_1_ds as their siblings do: every finite block is stored bit for bit as before. iq_multi_parity's Niko1221#606 check covers the shared expert's one (gate = 32, up = x / 32: its activations are x exactly). rwkeyes found the first one (Niko1221#606). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio
… moe_combine, LDS alias + occupancy) Bitwise-identical to f4e96d5 (generated ids identical on UD-IQ4_XS, UD-Q4_K_XL, IQ3_S, IQ3_XXS at 8K/4K/64K; harness fused-output hash unchanged): - table16: IQ4_NL / IQ4_XS decode through the direct-selector v_perm form (X5; exhaustive 2^32 check). -DSTRATA_W_NO_T16 restores the old form. - moe_combine: four columns per thread with float4 loads, same fma order (X5); falls back to the scalar kernel if Dm / shared / bo are not 16-byte aligned. -DSTRATA_W_NO_COMB. - native_w11_kernel: the H tile aliased onto the weight buffers (gate/up LDS 46.5 -> 30.2 KB), amdgpu_waves_per_eu(8), 4 blocks per WGP (X1), only for the types that do not spill under the 192-VGPR cap (gate/up IQ2_*, IQ3_*, IQ4_XS; down Q2_0, Q8_0, IQ4_NL). -DSTRATA_W_NO_OCC. - X3 on1 (skip the second 16-row n tile of the down kernel) NOT included: no gain alone (-0.1%), and with the occupancy cap it makes IQ4_NL down spill 62 VGPRs and the whole bundle 2.3% slower. Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01QJJwUv9g3cD2ztX6KWS8rk
…): a host 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 (Niko1221#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).
…; later 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
…221#550) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…M; a -1 raises a plan error flag (Niko1221#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 (Niko1221#828 test switch) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…1221#879 abort latency) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…s and an incomplete checkpoint (Niko1221#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, Niko1221#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>
…iko1221#409 review), not in a process-wide static 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 Niko1221#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 Niko1221#865 (its opt-in Q8_0 table is Niko1221#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>
Niko1221#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 Niko1221#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 Niko1221#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 Niko1221#586. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…nse::load layer 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: - Niko1221#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. - Niko1221#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.
…6 projections (Niko1221#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 Niko1221#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 (Niko1221#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 Niko1221#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>
…o1221#783 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 Niko1221#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 (Niko1221#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 Niko1221#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 (Niko1221#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 Niko1221#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-c, 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 Niko1221#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>
…1#783 PR-c) 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 Niko1221#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 (Niko1221#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 Niko1221#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 (Niko1221#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 Niko1221#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 (Niko1221#783 PR-k) Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…_MULTI_GR flag (Niko1221#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
|
Large (3,400 lines, a new launcher). Not merged for 0.1.40; the setup half is taken from #945. Experimental launcher stays a separate review. |
A preset is what the page shows instead of a wall of flags: the family and size, the context, the KV precision, images and the rest of the choices setup.py takes. Nothing here invents a setting - every field names the setup.py flag it writes, and a field left empty is not passed at all, so this PC's default decides; that is what makes a preset portable to another machine. The form stops at what a person actually changes: context scaling and the draft token subset are not offered, because setup picks them for the context it is given, and a page that shows a flag should be a page someone can decide about. The built-ins are not stored either: they are derived from setup's tables and its recommendation for this PC, so a repository that adds a size adds it to the page too. Presets are plain JSON in .strata-launcher/presets.json, written whole or not at all (like setup's configs, Niko1221#459), so a launcher closed half-way never leaves an empty file behind. Sizes are read the way setup reads them: a size that names no families in its table belongs to the original model and to its fine-tunes (setup's own default), which is what Swift 1.5's IQ2_XS and IQ3_XXS are. Tests: 22 without a GPU or a download - the choices a user cannot make are refused with the reason, a context past the trained one is setup's to scale, and the flags a preset writes are the flags setup would be given. Co-Authored-By: Strata qwen3.8-flash-next-iq3_s <strata@localhost> AI-developed: written by an agent (pi) running on Strata's own local model, qwen3.8-flash-next-iq3_s; tested by hand on an RTX 5060 Ti 16 GB / Ryzen 7 5700X3D / 64 GB PC with engine 0.1.39
…ncher.sh) One window over a Strata folder: the sizes setup offers and what each needs on this PC, the models this folder has and what their own strata-<model>.json says they run with, presets you save and run again, and setup's calibration as a job you can watch and stop. It is a front-end, not a second engine: every answer comes from tools/strata_mcp.py, the controller the MCP server uses, and installing, applying and starting run setup.py and serve/server.py the way SETUP.bat and run-<model>.bat do. Nothing is downloaded or started that those cannot do, and the only files the page writes of its own are in .strata-launcher/. It listens on 127.0.0.1 only, the same rule the model server follows; closing the window leaves a started model running, exactly like the start scripts. Next to the preset the page puts what the model is on this PC, not a speed number: its own config (which is what starting it runs), the files it misses, what setup's calibration recorded for it, and which model the server is serving now - the last one is what makes the Start button say "Running" instead of offering a second model. A model runs what its own config says, so starting a preset that differs applies it first: setup rewrites the config and the files already there are only checked again. That difference is compared with what setup's rules decide, not with the words - a preset saying "auto", or saying what auto decides on this PC, is the model as it runs and never re-runs setup to change nothing. Tests: 29 with the real Strata folder but a fixed fake PC, no port answering and the launcher's state in a temporary folder - a preset is stored and comes back, a plan carries the preset's flags once each, starting or measuring something that is not there says so plainly, and a path cannot escape the page's folder. Co-Authored-By: Strata qwen3.8-flash-next-iq3_s <strata@localhost> AI-developed: written by an agent (pi) running on Strata's own local model, qwen3.8-flash-next-iq3_s; tested by hand on an RTX 5060 Ti 16 GB / Ryzen 7 5700X3D / 64 GB PC with engine 0.1.39
…ADME translations, and the .bak setup leaves stays out of git docs/LAUNCHER.md: what the page is, what each part shows, what a preset is and what it writes, what installing over an installed model changes, what Start does when the preset and the config differ, the calibration, and where its own files live. README, AGENTS.md, AI_SETUP, INSTALL, MCP_SERVER and MODELS each get the one line that points here, and the README line goes into its six translations as the other README changes do; AGENTS.md also says that launcher/ is a front-end over tools/strata_mcp.py and how to run its tests, so the next change there starts from the same rule. .gitignore: setup's Model settings card (Niko1221#564) keeps the config it replaces as strata-<model>.json.bak, and that is not source. Co-Authored-By: Strata qwen3.8-flash-next-iq3_s <strata@localhost> AI-developed: written by an agent (pi) running on Strata's own local model, qwen3.8-flash-next-iq3_s; tested by hand on an RTX 5060 Ti 16 GB / Ryzen 7 5700X3D / 64 GB PC with engine 0.1.39
pipeob0
force-pushed
the
feature/launcher-ui
branch
from
October 6, 2026 08:28
ab6bc49 to
40490fa
Compare
Contributor
Author
Rebased onto 0.1.40Rebased onto Your k8v4 change caught one of our labels. The page described k8v4 as "hybrid K8/V4 (23% less than 8-bit; it does not stream its KV cache)", which was true of 0.1.39 and is not true of 0.1.40 (#711). The label now says it streams its KV cache too, and a test reads setup's own rule for each KV precision on a PC with room and fails if a label claims a streaming setup does not write: 55 launcher tests and 44 setup tests green on top of 0.1.40 ( |
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.
Continued in #1092
This PR closed by itself when
mainwas force-pushed during the history cleanup (its branch was built on the old history), and that branch is now deleted. The same 3 commits, cherry-picked onto the newmainwith no conflicts, are in #1092. The setup half is inmainfrom #945.Experimental - please read first
The launcher is experimental: it is a new surface over
setup.pyandserve/server.py, it has been run on one PC (Windows 11, RTX 5060 Ti 16 GB, Ryzen 7 5700X3D, 64 GB, engine 0.1.39; the branch is rebased onto 0.1.40'smainand the suites below were run there), and it will have bugs. What has not been exercised: Linux (./launcher.shis written and marked executable but never run on a Linux box), AMD/ROCm, several GPUs on one model, WSL, and a preset saved by an older version of the page (the presets file has no schema version yet). Nothing it does goes outside this PC, and nothing it does is impossible forSETUP.bat/run-<model>.batto do; the only files it writes of its own are in.strata-launcher/. Treat it as a preview, not as a replacement for the scripts.Also AI-developed: written by an agent (pi) running on Strata's own local model (
qwen3.8-flash-next-iq3_s), reviewed and tested by hand on the PC above - see theAI-developed:trailer on each commit.What it is
LAUNCHER.bat(Linux:./launcher.sh) serves a page on127.0.0.1for the caseSTART-HERE.batis not meant for: "the Coder today with 32K, leaving 2 GB of VRAM for the game", without remembering flags. It shows the sizes setup offers and what each needs on this PC, the models this folder has and what their ownstrata-<model>.jsonsays they run with, presets you save and reuse, and setup's--calibrateas a job you can watch and stop.(The same file is in the PR at
docs/media/launcher.png;docs/LAUNCHER.mdshows it too.)A front-end, not a second engine: every answer comes from
tools/strata_mcp.py, the controller the MCP server already uses; installing, applying and starting runsetup.pyandserve/server.pythe waySETUP.batandrun-<model>.batdo. No second config format, no second install path, no dependency beyond the standard library, andsetup.py/START-HERE.batdo not change: if you never open the launcher, nothing about your install changes.A preset is setup's own choices with a name: each field names the
setup.pyflag it writes, and a field left empty is not passed so this PC's default decides - that is what makes a preset portable to another machine. Starting a preset applies it first when the model's config differs, and that difference is compared with what setup's rules decide, not with the words: a preset sayingauto, or saying whatautodecides on this PC, is the model as it runs and never re-runs setup to change nothing. Those rules are the ones #945 put insetup.py, now inmain.Relation to the open work
mainand carries only the launcher: 3 commits, 30 files, +3443/-14. The setup half is yours, with yourkv_streaming_wanted(the k8v4 change included) - the launcher calls it, it does not copy it.--kv k8v4stream its KV. The page's label for k8v4 said it does not, which was true of 0.1.39 and is not true now; the label follows setup, and a test reads setup's rule for each KV precision and fails if a label claims a streaming setup does not write.--vram-reserve-mib): the web app exists once a model is running, the launcher is for choosing and installing before anything runs.python -m unittest launcher.test_presets launcher.test_api(25 + 30).tools.test_setup_riskandtools.test_setup_goldenstay green on top of this: 44 tests.Docs
docs/LAUNCHER.md, plus one line in README (and in its six translations), AGENTS.md, AI_SETUP, INSTALL, MCP_SERVER and MODELS..gitignoregains/strata-*.json.bak, which setup's settings card leaves behind.Known rough edges
launcher/,LAUNCHER.bat,launcher.shand.strata-launcher/.