Repository navigation
Q36: release the GGUF mapping and stop zeroing the overwritten arena - #2823
Closed
ggeorgovassilis wants to merge 21 commits into
Closed
ggeorgovassilis wants to merge 21 commits into
ggeorgovassilis wants to merge 21 commits into
Conversation
Quantize the resident K/V cache of the Qwen3.6 path (src/qwen35/q36_engine.cu) from fp32 to q8_0 (32 int8 codes + one fp16 scale per block) when --kv int8 is passed; the default --kv fp32 path is byte-identical. Adds the quantize/ dequantize primitive (include/strata/kernels/q36_kv_q8.hpp, src/kernels/cuda/q36_kv_q8.cu) and a GPU parity test (src/kernels/q36_kv_q8_parity.cpp). Implements docs/Q36_INT8_KV.md. Measured on an RTX 4060 Max-Q Laptop 8GB / Ryzen 7 7735HS / 62GB with Qwen3.6-35B-A3B-UD-Q4_K_XL.gguf (see docs/Q36_INT8_KV.md): - greedy decode is bit-for-bit identical to fp32 on short and long prompts; the 248,320-logit dump on the long prompt differs by at most 0.79. - the context window doubles: fp32 16,384 vs int8 32,768 (the 131,072 target is not reachable; the ceiling is dense weights + expert-streaming staging + the attention-score buffer, not K/V). - decode is ~4% slower at 4K (16.0 -> 15.4 tok/s); prefill unchanged. CUDA verified; the HIP build of the shared source is still to be run. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
- STRATA_Q36_SINGLE_STAGE=1 halves the streamed-prompt staging VRAM (one buffer instead of two double-buffered) - STRATA_Q36_NO_MMQ=1 is now numeric (MMQ stays on by default) - STRATA_Q36_PLACE_DEBUG logs the staging allocation - shrink the attention-scores buffer (256 -> 64 MB) when the staging buffers only just miss, so placement still succeeds Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
- docker/build.sh builds strata:local for this machine's GPU architecture - docker/Dockerfile bakes the engine, server and tokenizer tool from pinned sources - docker/entrypoint.py prepares the tokenizer and run config from the mounted model - docker/run.sh resolves and mounts the model and forwards the STRATA_* switches - tools/test_docker_run.py asserts run.sh's mode selection and env forwarding Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
…red, dropped) Qwen3.6-35B-A3B-UD-Q4_K_XL is Q4_K gate/up with Q5_K down, so its down projection was the last one in the CPU pool with no multi-token AVX2 kernel. A q5k_dot was written on the q5k-down-kernel branch - bit-exact against ggml at 1..8 tokens, wired into native_down_rows behind STRATA_KQ256 - and measured in the engine: 18.14 tok/s with the switch off against 17.72 with it on (18 and 12 greedy generations, 32K context, int8 KV, +-3% spread within one setting), with the replies byte-identical either way. No gain, because the verify windows hold about two drafts, where the kernel is only ~1.08x on the rows. Not enabled, and the branch is parked: bench/results/2026-10-16-q5k-down-kernel/ has the numbers and the harness. This also settles the plan document's premise - the "+3 t/s" it credits to the sibling Q4_K kernel is in no document, and UNSLOTH_Q4.md already says that kernel measured no faster.
Records the diagnosis and the phased implementation plan for the host's kcompactd churn while the engine runs: - the GGUF mapping (MAP_PRIVATE PROT_READ, ~22.85 GB) is read once and then kept mapped for the process lifetime; nothing needs it after place_experts - the expert arena is one cudaHostAlloc of ~18.75 GiB and cannot be reclaimed - /proc/vmstat showed 71,744,161 failed page migrations against 2,373,960 successful ones with compact_stall at 146: no shortage, just futile work The plan is five opt-in phases (memory report, mapping release, pageable arena, large-page arena, padding-only memset) plus a defaults phase, each with its gates, its image tag and its rollback. The default path stays byte-identical until a phase has measured better, per AGENTS.md. Also records the host sysctls and THP settings applied on 2026-10-17, the values they replaced, how to revert them, and the fact that they are not persisted yet.
…y memset
Implements docs/Q36_MEMORY_COMPACTION.md.
- Phase 0: the memory/compaction counters in the engine's own log, behind
STRATA_MEM_REPORT (off by default): include/strata/platform/mem_report.{hpp,cpp}
plus a fixture test.
- Phase 1/5: release the GGUF mapping after place_experts() has read every expert
(madvise(MADV_DONTNEED) then posix_fadvise(POSIX_FADV_DONTNEED) - fadvise cannot
drop pages that are still mapped). Measured ~18.5 GiB of cache back to the OS
with decode unchanged, so it is on by default; STRATA_Q36_RELEASE_MAPPING=0 turns
it off.
- Phase 4: zero only the arena padding (and an NVFP4 scale tail), not the
18.75 GiB expert_blob() overwrites. Byte-identical by construction; 14/14
native-expert parity tests pass.
- Phases 2 and 3 are moot: the arena is memfd/shm-backed, not mlocked (measured on
a RTX 4060 Max-Q, driver 580.178.04), so "un-pin the arena" fixes nothing.
Built: CPU suite 44/44, CUDA build, parity suite on the card. HIP/SYCL not built
here (no ROCm/SYCL on this box).
What would make Strata faster on the RTX 4060 Max-Q / Ryzen 7 7735HS / 62 GB box, on the qwen36 path and on the main engine: the PRs that are new, applicable and measured, plus the ones already merged after the history rewrite and the ones for other hardware.
…n this PC (prompt phases, decode window split, hit rate per turn) It starts the server with the user's settings plus STRATA_PREFILL_TIMING=1 and STRATA_DECODE_TIMING=1, sends three short questions, fresh 2K / 16K / 32K prompts (then 128 tokens at the long context) and a five-turn agent session through the HTTP API, reads the engine's own lines for each request, and writes bench/results/<date>-profile-<gpu>- <model>/README.md: the prompt read's GPU time grouped (experts over PCIe, dequant, expert GEMM, attention, GDN, other), the decode window split (the card, the experts computed on the CPU, the draft, the adaptive tier), the hit rate and PCIe share per request, the agent session turn by turn, and findings in words that name the setting each wait depends on. --gpu-stages adds STRATA_VERIFY_PROFILE=1; --url with --engine-log reads a running server. The config is read, never written; --expert-profile-save is dropped for the run as the bench does. tools/test_strata_profile.py: the parsers on real engine lines, the grouping, the findings, the run against a fake server that writes an engine log (25 tests, no GPU). An example report from an RTX 5070 Ti + Ryzen 7 5800X3D (Swift 1.5 IQ3_S) is in bench/results/. Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
A build folder configured while the option's default was OFF keeps STRATA_Q6K_EXPERTS=OFF in its CMakeCache.txt even after the default flipped to ON (465cd61): CMake does not rewrite an option already in the cache, and source_hash() fingerprints the source files but not the CMake option values, so setup recompiles the engine with the Q6_K gate/up expert kernels still disabled. A model whose routed experts are Q4_K/Q6_K (Unsloth UD-Q4_K_XL) then fails to load with: strata-q36: blk.34.experts are Q4_K/Q6_K, which the native expert kernels do not take (build with -DSTRATA_Q6K_EXPERTS=ON for Q6_K) Passing -DSTRATA_Q6K_EXPERTS=ON explicitly overrides the stale cache entry, so the engine always has the kernels its default build promises. Same pin in the Dockerfile build for consistency. Not hardware-specific: it is a CMake cache-staleness issue, independent of GPU/platform.
…65cd61 flipped both options the same way) 465cd61 (2026-10-12) moved STRATA_MMQ_KQUANTS from OFF to ON exactly like STRATA_Q6K_EXPERTS, and CMake caches it the same way: an explicit -D on one option does not rescue the other (qiuzan2001's cmake 4.4.3 repro: configure A/B at OFF, flip both defaults to ON, reconfigure with only -DA=ON -> the cache ends A=ON, B=OFF). So pinning only Q6K_EXPERTS left an old build/ still compiling without the strata_kq K-quant MMQ library (the UD-Q4_K_XL prompt kernels), and build-hip had the same stale-cache case: the HIP ON default never reaches an existing cache. Pin -DSTRATA_MMQ_KQUANTS=ON in the CUDA engine configure, the Dockerfile and the HIP configure, and -DSTRATA_Q6K_EXPERTS=ON in the HIP configure too. Both flags match the CMakeLists.txt defaults on main, so a clean build is byte-identical; the pin only overrides a stale cache entry. Tests pin both flags in the CUDA and HIP engine def lists.
… a Zen 3) bf16_rows_dot_multi held its per-token sums in __m256 acc[MAXT] indexed by a variable. With eight 256-bit accumulators live GCC also unrolled the column loop and left the whole array in memory, so every token's FMA read and wrote a stack slot (vmovaps 0xc0(%rsp) ... vfmadd213ps 0xe0(%rsp)). On a Ryzen 7 5700X3D one layer's router [512 x 2560] at 4 tokens took 0.41 ms against 0.25 ms for the AVX1 copy in kq_avx1.cpp that kept its accumulators in registers, so every AVX2 CPU ran the slower of its two kernels (router_dot_parity --cpu 2, median of 3). Pragma GCC unroll 1 on the column loop changes nothing; the variable-indexed array is what forces the array into memory. Four tokens per pass, in four named accumulators. A row above four tokens is read a second time, from L1: nt=8 measures 0.37 ms, the same as the spilled form. On the same CPU the AVX2 kernel goes 0.41 -> 0.19 ms at 4 tokens and nt=2 0.36 -> 0.16 ms, and it now runs 1.3x ahead of the AVX1 copy instead of 1.6x behind it. The loop body carries no stack slot any more (was 21 memory ops per iteration). Untouched: the arithmetic and the reduction. Each token's FMA chain still runs over the same columns in the same order and ends in the same hsum_float_8, so every sum is bit-identical - the output for nt 1/2/4/8 at cols 2560/5120 hashes the same before and after (FNV-1a over the whole out[]), and router_dot_parity reports 0 failures with the AVX2 rel error unchanged at 3.16e-07. Left alone: the expert row kernels (iq_avx2.cpp, q2_avx2.cpp, kq256_gu_rows, kq256_rows) and the AVX1 path. The iq_avx2_parity bench over four format pairs x 1..8 tokens moved within 1.6% (noise), and the CPU parity tests (8 targets) pass. The router stays what it was, an estimate for which experts to prefetch. CUDA and HIP do not compile this file; the SYCL port does, and it builds there under icpx 2026.1 with no new warnings and no spills either. The AVX1 copy in kq_avx1.cpp kept a comment saying it reads its rows "exactly as the AVX2 kernel does"; with the row loop chunked four tokens at a time that is no longer the AVX2 kernel's shape, so the clause is dropped there. Comment only - no code change.
q5_1_dot and q8_0_dot in kq_avx2.cpp held their per-token sums in the same runtime-indexed __m256 acc[MAXT] that made bf16_rows_dot_multi spill: with the array live the compiler leaves it in memory, and every token's FMA reads and writes a stack slot. Both are now template <int NT>, dispatched 1..4 and, past that, run four tokens per pass with the row's bytes re-read from L1 - the same idiom the router change above and kq256_gu_rows already use. Ryzen 7 9800X3D (Zen 5), GCC 15.2.0, build/native_expert_parity --synthetic (ggml-quantized weights and activations), one thread pinned to core 0, at the verify window's four tokens: Q5_1 down rows [2560 x 640]: 190-197 us -> 155-156 us (1.24x) Q8_0 down rows [2560 x 640]: 135 us -> 104 us (1.30x) Same sums: 0 of 36 token-sets differ in any bit against ggml's own vec_dot over 1..8 tokens, both before and after; native_expert_parity reports 0 failures. Left alone: q4k_dot keeps its runtime-nt array. Its token accumulators are not what spills - the unpatched function carries 8 ymm stack ops for the sums against 24 (q5_1) and 59 (q8_0); what q4k_dot spills is its weight-side unpack arrays (q4l/q4h/scale_l/scale_h), which token chunking does not touch and re-does. The same change measured +5% at four tokens and -4..-7% at six to eight, so it is a wash at the window and a small regression above it, and q4k is the gate/up format that runs most. The AVX1 path, the iq_avx2/q2_avx2 rows and the engine's dispatch are untouched. Backends: CUDA and HIP do not compile this file. The SYCL port does; not built on this machine (no icpx here) - the router change in this branch built there. Tests: native_expert_parity --synthetic q4_K/q5_1 and q4_K/q8_0 (0 failures, 0 bits differ), iq_avx2_parity, iq_avx1_parity, q8k_quant_parity, router_dot_parity - all exit 0.
…hout the barrier between them; pool_layer_bench run_split_multi_native ran a layer's CPU experts in two phases (all gate/up rows, the host's quantize, all down rows). Mode 7: each expert's gate/up rows in 8 parts and its down rows in 4; the thread finishing an expert's last gate/up part quantizes its intermediate and flags it; its down parts wait for that flag only, queued one expert behind (STRATA_POOL_FUSED_LAG). The same native_gu_rows / native_down_rows calls, so the same bits. Q2_0 layers keep the phases. Opt-in: measured on one CPU. Why: on Zen 2 the i-quant gate/up rows are compute-bound (7 threads 24.7-27.3 GB/s) and IQ4_NL down rows run at RAM speed (~39 GB/s), so each phase left the other's resource idle. pool_layer_bench (new; the real ExpertPool, one layer's 6-7 CPU experts a call, a quarter 2 tokens, 4 GB of blobs, output hash over 300 fixed layers), Ryzen 7 3800X, 7 workers + host: IQ2_S / IQ3_XXS / IQ3_S with IQ4_NL down 0.420 -> 0.387, 0.500 -> 0.456, 0.526 -> 0.469 ms a layer (-8 / -9 / -11%), the same hashes. generate.cpp (CUDA/HIP and SYCL) notes the fused layer calls in the pool stats. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…down without the barrier between them; pool_layer_bench" This reverts commit f4ddbf4.
… driver #2775, #1181, #1262, #1989, #1838 (and #1581) touch only the main engine (generate/expert_source/weights/layer/verify/prefill/native_dense/fused_gr). This box's only MoE model runs the q36 driver, which shares none of those files, and there is no Flash-Next pack to measure them against, so a 32K serve-mode A/B would be flat by construction. Deferred, not rejected.
… with alloc/copy phase timing The load's hot loop is place_experts() copying 18.75 GiB of routed experts from the GGUF mapping into the pinned host arena, one expert at a time on a single thread. Measured on the RTX 4060 Laptop / Ryzen 7 7735HS: 64 s serial (warm and cold within 2%), ~22.5 s of cudaHostAlloc pinning, so ~1m50s total. - STRATA_Q36_PARALLEL_PLACE: 0/unset = serial (the pre-change default, byte-identical); N = copy on N threads. Disjoint chunks, so the placed bytes are identical. - the 'hybrid experts:' log now reports 'arena alloc X s, expert copy Y s, N copy threads', so the lever's movement is visible without a profiler. Measured (same binary, controlled cache): copy 63.8s -> 16.2s at 8 threads (3.9x), 13.4s at 16 threads (4.8x); total load ~104s -> ~58s. Greedy-decode tokens identical flag on vs off. See docs/Q36_LOAD_TIME.md.
…NA A/B arm - STRATA_Q36_PARALLEL_PLACE: unset now copies the experts on one thread per logical core (was serial). =0 keeps the serial path. Measured 4.8x faster, byte-identical. - STRATA_Q36_PAGEABLE_ARENA=1: malloc the expert arena instead of cudaHostAlloc (pinning) it. Diagnosis (compaction counters frozen during the alloc, 47 GiB free) shows the ~14-22 s pin is the pin itself, not the OS rearranging RAM. - Measured: pageable arena cuts place_experts 27.8s -> 13.8s and decode is byte- identical, but prompt throughput halves (295 -> 138 tok/s) because staging goes through the driver's bounce buffer - so it stays opt-in; a deferred pin is the right follow-up. See docs/Q36_LOAD_TIME.md.
…the DMA The pageable arena cut the ~14 s pin but halved prompt throughput (295 -> 138 tok/s): cudaMemcpyAsync from pageable memory falls back to the driver's bounce buffer. CUDA requires page-locked host memory for the fast H2D DMA the prompt staging and slot refills do, so the pin is a necessary cost, not a lever. The parallel copy default stays (it is byte-identical and has no such tradeoff).
Owner
|
Thanks for the pull request. We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to review properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected. To reopen it takes about two minutes:
The rules are in CONTRIBUTING.md. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Implements the Q36 memory-compaction plan (
docs/Q36_MEMORY_COMPACTION.md): the engine reports its own memory/compaction counters, releases the GGUF mapping after load, and stops zeroing the 18.75 GiB expert arena it immediately overwrites. The original goal - quiet thekcompactdchurn - was already met by host settings; the measured engine win is ~18.5 GiB of RAM returned to the OS at no decode cost.What changed
include/strata/platform/mem_report.{hpp,cpp}(new) +src/platform/mem_report_test.cpp: read/proc/meminfo,/proc/vmstatand the THP mode into the engine's own log, behindSTRATA_MEM_REPORT(off by default). A row the kernel does not export prints-, not 0.include/strata/artifact/gguf_reader.hpp:GgufFile::release_pages()/released()andGgufModel::release_pages()/released()—madvise(MADV_DONTNEED)thenposix_fadvise(POSIX_FADV_DONTNEED)(fadvise cannot drop pages that are still mapped).src/qwen35/q36_engine.cu:release_model_mapping()afterplace_experts(), on by default;STRATA_Q36_RELEASE_MAPPING=0reverts. A guard makesexpert_blob()fail loudly if it is ever called after the release. The arena memset now zeroes only the stride padding and thekMmqTailtail (plus an NVFP4 scale tail) instead of the whole 18.75 GiB.tests/core/expert_blob_padding_test.cpp(new): the padding-only zeroing is byte-identical to the whole memset.Extra Notes
Measured on an RTX 4060 Max-Q 8 GiB, Ryzen 7 7735HS, 62 GB, 32K,
--kv int8 --mtp 2:Mlockedstayed 104 KiB throughout -cudaHostAllocon this driver (580.178.04) is memfd/shm-backed, not mlocked, so the plan's "un-pin the arena" phase and its dependent were dropped.Built: CPU suite 44/44, CUDA build, and the native-expert parity suite on the card (14/14). HIP and SYCL builds were not run here (no ROCm/SYCL on this box).