Merge remote-tracking branch 'macm5us/deepseek4-pro-correctness' into glm5.2

# Conflicts:
#	ds4_metal.m

.gitattributes

@@ -0,0 +1,3 @@
+# Preserve hosted-model continuation bytes exactly; these fixtures may contain
+# model-generated trailing whitespace that is meaningful to tokenizer scoring.
+gguf-tools/quality-testing/data/*/continuations/*.txt -diff

Makefile

@@ -10,6 +10,7 @@ endif
 DEBUG_FLAGS ?= -g
 CFLAGS ?= -O3 -ffast-math $(DEBUG_FLAGS) $(NATIVE_CPU_FLAG) -Wall -Wextra -std=c99
 OBJCFLAGS ?= -O3 -ffast-math $(DEBUG_FLAGS) $(NATIVE_CPU_FLAG) -Wall -Wextra -fobjc-arc
+QUALITY_CFLAGS ?= -O3 $(DEBUG_FLAGS) $(NATIVE_CPU_FLAG) -Wall -Wextra -std=c11
 
 LDLIBS ?= -lm -pthread
 METAL_SRCS := $(wildcard metal/*.metal)
@@ -67,6 +68,9 @@ ds4-eval: ds4_eval.o ds4_help.o $(CORE_OBJS)
 ds4-agent: ds4_agent.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o $(CORE_OBJS)
 	$(CC) $(CFLAGS) -o $@ ds4_agent.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o $(CORE_OBJS) $(METAL_LDLIBS)
 
+gguf-tools/quality-testing/score_official: gguf-tools/quality-testing/score_official.c ds4.h $(CORE_OBJS) rax.o
+	$(CC) $(QUALITY_CFLAGS) -I. -o $@ gguf-tools/quality-testing/score_official.c $(CORE_OBJS) rax.o $(METAL_LDLIBS)
+
 cpu: ds4_cli_cpu.o ds4_server_cpu.o ds4_bench_cpu.o ds4_eval_cpu.o ds4_agent_cpu.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o rax.o $(CPU_CORE_OBJS)
 	$(CC) $(CFLAGS) -o ds4 ds4_cli_cpu.o ds4_help.o linenoise.o $(CPU_CORE_OBJS) $(LDLIBS)
 	$(CC) $(CFLAGS) -o ds4-server ds4_server_cpu.o ds4_help.o ds4_kvstore.o rax.o $(CPU_CORE_OBJS) $(LDLIBS)
@@ -128,6 +132,9 @@ ds4-eval: ds4_eval.o ds4_help.o $(CORE_OBJS)
 ds4-agent: ds4_agent.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o $(CORE_OBJS)
 	$(DS4_LINK) -o $@ $^ $(DS4_LINK_LIBS)
 
+gguf-tools/quality-testing/score_official: gguf-tools/quality-testing/score_official.c ds4.h $(CORE_OBJS) rax.o
+	$(DS4_LINK) $(filter-out -ffast-math,$(QUALITY_CFLAGS)) -I. -o $@ gguf-tools/quality-testing/score_official.c $(CORE_OBJS) rax.o $(DS4_LINK_LIBS)
+
 cpu: ds4_cli_cpu.o ds4_server_cpu.o ds4_bench_cpu.o ds4_eval_cpu.o ds4_agent_cpu.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o rax.o $(CPU_CORE_OBJS)
 	$(CC) $(CFLAGS) -o ds4 ds4_cli_cpu.o ds4_help.o linenoise.o $(CPU_CORE_OBJS) $(LDLIBS)
 	$(CC) $(CFLAGS) -o ds4-server ds4_server_cpu.o ds4_help.o ds4_kvstore.o rax.o $(CPU_CORE_OBJS) $(LDLIBS)
@@ -241,4 +248,4 @@ q4k-dot-test: tests/test_q4k_dot.c
 	./tests/test_q4k_dot
 
 clean:
-	rm -f ds4 ds4-server ds4-bench ds4-eval ds4-agent ds4_cpu ds4_native ds4_server_test ds4_test ds4_agent_test tests/test_q4k_dot *.o tests/cuda_long_context_smoke tests/cuda_long_context_smoke.o
+	rm -f ds4 ds4-server ds4-bench ds4-eval ds4-agent ds4_cpu ds4_native ds4_server_test ds4_test ds4_agent_test gguf-tools/quality-testing/score_official tests/test_q4k_dot *.o tests/cuda_long_context_smoke tests/cuda_long_context_smoke.o

QA_BEFORE_RELEASES.md

@@ -54,7 +54,43 @@ in this system.
   `./ds4_test --server`.
 - Run `./ds4-eval --self-test-extractors`.
 
-## 3. Metal Flash Path
+## 3. Official Continuation Quality Gates
+
+These tests are release-blocking after tokenizer, template, KV-cache, attention,
+MoE routing, quantization, logit, or model-graph changes.  They are
+teacher-forced continuation checks against hosted-model output and API
+top-logprob slices, so do not replace them with one sampled chat answer.
+
+- Build the scorer:
+  `make -C gguf-tools quality-score`.
+- Run the tracked DeepSeek V4 Flash smoke vectors:
+  `./ds4_test --logprob-vectors`.
+  This covers the official Flash API vectors in `tests/test-vectors/`, including
+  short prompts and long-prompt attention cases.
+- Run the 100-case DeepSeek V4 Flash fixture for every released Flash GGUF:
+  `gguf-tools/quality-testing/score_official /path/to/deepseek-v4-flash.gguf gguf-tools/quality-testing/data/flash/manifest.tsv /tmp/flash.tsv 4096`.
+- Run the 100-case GLM 5.2 OpenRouter fixture for every released GLM GGUF:
+  `gguf-tools/quality-testing/score_official models/GLM-5.2-UD-Q4_K_XL.gguf gguf-tools/quality-testing/data/glm52-openrouter-100/manifest.tsv /tmp/glm52-q4.tsv 4096`.
+  Current Q4 XL reference band: first-token match `95/100`, API top-1 agreement
+  about `0.942`, and API pair-order agreement about `0.880`.
+- Run the same GLM fixture for reduced-precision GLM release files.  The Q2
+  routed reference is lower quality but should stay near first-token match
+  `92/100`, API top-1 agreement about `0.890`, and API pair-order agreement
+  about `0.800` unless the quantization changed deliberately.
+- Run the 100-case DeepSeek V4 PRO fixture for every released PRO GGUF:
+  `gguf-tools/quality-testing/score_official /path/to/deepseek-v4-pro.gguf gguf-tools/quality-testing/data/pro/manifest.tsv /tmp/pro.tsv 4096`.
+- For SSD streaming, run the same official-continuation scorer once with full
+  residency and once with `--ssd-streaming` for the release model.  The summary
+  and API agreement should stay in the same quality band.
+- Compare any candidate against the previous release or last-known-good output:
+  `python3 gguf-tools/quality-testing/compare_scores.py /tmp/old.tsv /tmp/new.tsv`.
+  Treat a large first-token-match drop, a clear NLL regression, or a material
+  API top-1/pair-order regression as a blocker unless the release notes call out
+  an intentional quality tradeoff.
+- Keep the raw `summary` and `api_summary` lines in the release notes or QA log.
+  Do not use stale manifests from `misc/` as release evidence.
+
+## 4. Metal Flash Path
 
 Use the normal Flash GGUF that 128 GB users run.
 
@@ -73,7 +109,7 @@ Use the normal Flash GGUF that 128 GB users run.
   generation speed, and KV bytes with the last known good numbers for the same
   machine.
 
-## 4. Metal PRO Path
+## 5. Metal PRO Path
 
 PRO support is experimental, but release builds must not break it silently.
 
@@ -84,7 +120,7 @@ PRO support is experimental, but release builds must not break it silently.
   touching model shape, tensor lookup, routed expert mapping, template logic,
   and KV payload compatibility.
 
-## 5. SSD Streaming
+## 6. SSD Streaming
 
 SSD streaming is a capacity path, so test both correctness and user experience.
 
@@ -103,7 +139,7 @@ SSD streaming is a capacity path, so test both correctness and user experience.
 - If streaming cache internals changed, test the same prompt twice and compare
   first-token/logprob sanity between runs.
 
-## 6. CUDA / DGX Spark
+## 7. CUDA / DGX Spark
 
 Before a release, ask the user for CUDA access if it is not already configured.
 Use the DGX Spark / GB10 host `toor@192.168.0.180`.  Do not claim CUDA is
@@ -121,7 +157,7 @@ release-ready without this pass.
 - Verify that any CUDA-only warning fixes are also clean on macOS and do not
   change Metal behavior.
 
-## 7. ROCm / Strix Halo
+## 8. ROCm / Strix Halo
 
 Use the Strix Halo Framework Desktop via the VPN hostname `strixhalo`
 (`antirez@strixhalo`).  This host validates the ROCm backend; do not use it as
@@ -142,7 +178,7 @@ a substitute for CUDA or Metal release testing.
 - Record startup memory/cache messages, prefill speed, generation speed, and
   whether the backend reports `ROCm backend initialized`.
 
-## 8. Distributed Inference
+## 9. Distributed Inference
 
 Distributed code has regressed around route setup, KV snapshots, request IDs,
 and split model loading.  Test it whenever distributed, KV, session, or model
@@ -160,7 +196,7 @@ loading code changes.
 - If CUDA distributed is relevant, test across the CUDA hosts and record
   generation speed, not just "it works".
 
-## 9. Disk KV Cache
+## 10. Disk KV Cache
 
 Disk KV cache bugs are high impact for server users.
 
@@ -174,7 +210,7 @@ Disk KV cache bugs are high impact for server users.
 - Test stripped agent sessions: `/strip <id>` then `/switch <id>` should rebuild
   by prefill and render sane history.
 
-## 10. Server APIs
+## 11. Server APIs
 
 The server must keep compatibility across OpenAI, Responses, and Anthropic
 clients.
@@ -187,7 +223,7 @@ clients.
 - Test `--trace` and confirm rendered prompts, cache decisions, generated text,
   and tool-parser events are useful without leaking unrelated state.
 
-## 11. ds4-agent
+## 12. ds4-agent
 
 The agent is the most stateful component.  Test it manually, not only by build.
 
@@ -219,7 +255,7 @@ The agent is the most stateful component.  Test it manually, not only by build.
   status bar fill to terminal width, syntax highlighting in Markdown/code blocks,
   and SSH/remote terminal flicker.
 
-## 12. Download Script And Model Files
+## 13. Download Script And Model Files
 
 - Test `download_model.sh` in a temporary directory so local weights are not
   overwritten.
@@ -228,7 +264,7 @@ The agent is the most stateful component.  Test it manually, not only by build.
 - Verify legacy removed targets fail clearly.
 - Verify README model names match the script and Hugging Face repository.
 
-## 13. Performance And Power
+## 14. Performance And Power
 
 - Run `ds4-bench` on the release machine and compare with tracked CSV baselines.
 - Test `--power 100` is not throttled.
@@ -237,11 +273,12 @@ The agent is the most stateful component.  Test it manually, not only by build.
 - Confirm context buffer size, raw KV rows, compressed KV rows, and mmap behavior
   match expectations for 32k, 100k, and any release-advertised context size.
 
-## 14. Release Sign-off
+## 15. Release Sign-off
 
 Do not sign off until:
 
 - macOS Metal Flash passed.
+- Official continuation quality gates passed for every released model family.
 - CUDA was tested on the CUDA machine or the release notes explicitly say CUDA
   was not validated.
 - ROCm was tested on Strix Halo or the release notes explicitly say ROCm was

README.md

@@ -139,6 +139,20 @@ but must be enabled explicitly with `--mtp`. The current MTP/speculative
 decoding path is still experimental: it is correctness-gated and currently
 provides at most a slight speedup, not a meaningful generation-speed win.
 
+GLM 5.2 support is currently limited to the GGUF files tested by this branch:
+
+```sh
+./download_model.sh glm-unsloth-q4  # Unsloth UD-Q4_K_XL, 11 shards
+./download_model.sh glm-antirez-q2  # antirez routed Q2_K single-file GGUF
+./download_model.sh glm-antirez-q4  # antirez routed Q4_K single-file GGUF
+```
+
+The supported GLM layout keeps dense/model-control tensors in the existing
+Q8/F32 paths and supports routed expert gate/up tensors in `Q2_K`, `Q4_K`, or
+`Q5_K`; routed expert down tensors are supported in `Q2_K`, `Q4_K`, `Q5_K`, or
+`Q6_K`. Other GLM GGUF quant layouts should be treated as unsupported until they
+are added deliberately and scored against the official 100-case fixture.
+
 Then build:
 
 ```sh
@@ -206,14 +220,19 @@ more memory for context, set the routed expert cache explicitly:
 ./ds4 -m ./ds4flash.gguf --ssd-streaming --ssd-streaming-cache-experts 32GB
 ```
 
-The `32GB` value is a memory budget for complete routed experts, not a generic
-byte cache. DwarfStar converts it to the number of full experts that fit for the
-current GGUF. Non-routed weights, KV cache, graph scratch, and activations need
-additional memory. Only the automatic cache budget does the subtraction for you:
-it takes 80% of the Metal recommended working set, subtracts non-routed weights,
-then uses the rest for routed experts. Leave the hot expert preload enabled for
-normal use; use `--ssd-streaming-cold` and `--ssd-streaming-preload-experts N`
-only for measurements.
+The `32GB` value is a routed-expert memory budget, not a generic byte cache.
+DwarfStar first reserves headroom for the two full routed layers used by
+overlapped streaming prefill, then converts the remaining bytes to the number of
+dynamic cached experts that fit for the current GGUF. Explicit `NGB` budgets may
+also be capped after context/KV accounting so the Metal working set stays out of
+the slow pressure zone. A plain number such as
+`--ssd-streaming-cache-experts 4000` is different: it means exactly 4000 dynamic
+expert slots, with no extra accounting. Non-routed weights, KV cache, graph
+scratch, and activations need additional memory. The automatic cache budget takes
+80% of the Metal recommended working set, subtracts non-routed weights, then
+applies the same routed-prefill headroom before sizing the dynamic cache. Leave
+the hot expert preload enabled for normal use; use `--ssd-streaming-cold` and
+`--ssd-streaming-preload-experts N` only for measurements.
 
 ### Practical SSD streaming examples
 
@@ -245,15 +264,11 @@ and occasional work when you accept slow generation. Start with `--nothink`:
 
 On an M5 Max with 128GB of RAM, a short PRO q2 streaming decode benchmark found
 the automatic budget best: it selected about `59GB` of routed expert cache.
-Manual `64GB` to `75GB` caches were close on that machine. Larger explicit
-`NGB` requests are capped before inference so the expert buffers remain
-lockable instead of falling into macOS paging. If the system is under extra
-memory pressure and `mlock` still fails, ds4 refuses to install pageable
-expert-cache entries and releases a locked-cache margin before continuing with
-the measured lockable cache size. Prefer the automatic budget; if setting the
-cache manually on this class of machine, start around `48GB` to `64GB`, then
-increase only while the startup log reports a lockable cache. Once the machine
-is stable, re-enable thinking with a conservative generation limit:
+Manual `64GB` to `75GB` caches were close on that machine. Prefer the automatic
+budget; if setting the cache manually on this class of machine, start around
+`48GB` to `64GB`, then increase only while the machine remains responsive and
+the startup log shows the requested dynamic cache. Once the machine is stable,
+re-enable thinking with a conservative generation limit:
 
 ```sh
 ./ds4 \

download_model.sh

@@ -1,7 +1,10 @@
 #!/bin/sh
 set -e
 
-REPO="antirez/deepseek-v4-gguf"
+DEEPSEEK_REPO="antirez/deepseek-v4-gguf"
+GLM_UNSLOTH_REPO="unsloth/GLM-5.2-GGUF"
+GLM_ANTIREZ_REPO="antirez/GLM-5.2-GGUF"
+REPO=$DEEPSEEK_REPO
 Q2_IMATRIX_FILE="DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix.gguf"
 Q4_IMATRIX_FILE="DeepSeek-V4-Flash-Q4KExperts-F16HC-F16Compressor-F16Indexer-Q8Attn-Q8Shared-Q8Out-chat-v2-imatrix.gguf"
 Q2_Q4_IMATRIX_FILE="DeepSeek-V4-Flash-Layers37-42Q4KExperts-OtherExpertLayersIQ2XXSGateUp-Q2KDown-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix-fixed.gguf"
@@ -9,6 +12,12 @@ PRO_Q2_IMATRIX_FILE="DeepSeek-V4-Pro-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-Instruct-
 PRO_Q4_LAYERS00_30_FILE="DeepSeek-V4-Pro-Q4K-Layers00-30.gguf"
 PRO_Q4_LAYERS31_OUTPUT_FILE="DeepSeek-V4-Pro-Q4K-Layers-31-output.gguf"
 MTP_FILE="DeepSeek-V4-Flash-MTP-Q4K-Q8_0-F32.gguf"
+GLM_UNSLOTH_Q4_REMOTE_BASE="UD-Q4_K_XL/GLM-5.2-UD-Q4_K_XL"
+GLM_UNSLOTH_Q4_LOCAL_BASE="GLM-5.2-UD-Q4_K_XL"
+GLM_UNSLOTH_Q4_FIRST_FILE="$GLM_UNSLOTH_Q4_LOCAL_BASE-00001-of-00011.gguf"
+GLM_ANTIREZ_IQ2XXS_FILE="GLM-5.2-UD-IQ2_XXS_RoutedIQ2XXS_blk78Q2K.gguf"
+GLM_ANTIREZ_Q2_FILE="GLM-5.2-UD-Q2_K_RoutedQ2K.gguf"
+GLM_ANTIREZ_Q4_FILE="GLM-5.2-UD-Q4_K_RoutedQ4K.gguf"
 
 ROOT=$(CDPATH= cd -- "$(dirname -- "$0")" && pwd)
 OUT_DIR=${DS4_GGUF_DIR:-"$ROOT/gguf"}
@@ -20,7 +29,7 @@ TOKEN=${HF_TOKEN:-}
 
 usage() {
     cat <<EOF
-DeepSeek V4 GGUF downloader
+DwarfStar GGUF downloader
 
 Usage:
   ./download_model.sh q2-imatrix [--token TOKEN]
@@ -31,6 +40,10 @@ Usage:
   ./download_model.sh pro-q4-layers31-output [--token TOKEN]
   ./download_model.sh pro-q4-split [--token TOKEN]
   ./download_model.sh mtp [--token TOKEN]
+  ./download_model.sh glm-unsloth-q4 [--token TOKEN]
+  ./download_model.sh glm-antirez-iq2xxs [--token TOKEN]
+  ./download_model.sh glm-antirez-q2 [--token TOKEN]
+  ./download_model.sh glm-antirez-q4 [--token TOKEN]
 
 Targets:
 
@@ -69,6 +82,22 @@ Targets:
        It is useful with q2-imatrix, q2-q4-imatrix, and q4-imatrix, but must be
        enabled explicitly with --mtp when running ds4 or ds4-server.
 
+  glm-unsloth-q4
+       GLM 5.2 Unsloth UD-Q4_K_XL quant from unsloth/GLM-5.2-GGUF.
+       Downloads all 11 shards and links ./ds4flash.gguf to the first shard.
+
+  glm-antirez-iq2xxs
+       GLM 5.2 antirez routed IQ2_XXS GGUF from antirez/GLM-5.2-GGUF.
+       Includes Q2_K block 78 and is intended for reduced-memory testing.
+
+  glm-antirez-q2
+       GLM 5.2 antirez routed Q2_K GGUF from antirez/GLM-5.2-GGUF.
+       About 262 GB on disk.
+
+  glm-antirez-q4
+       GLM 5.2 antirez routed Q4_K GGUF from antirez/GLM-5.2-GGUF.
+       About 434 GB on disk.
+
 Options:
   --token TOKEN  Hugging Face token. Otherwise HF_TOKEN or the local HF token
                  cache is used if present.
@@ -87,8 +116,9 @@ Then the default commands work:
 After downloading mtp, enable it explicitly, for example:
   ./ds4 --mtp <download directory>/$MTP_FILE --mtp-draft 2
 
-PRO files are downloaded with the official Hugging Face downloader because
-they are too large for the curl path used by the smaller GGUF files.
+PRO and GLM files are downloaded with the official Hugging Face downloader
+because they are too large, sharded, or nested for the curl path used by the
+smaller DeepSeek Flash GGUF files.
 EOF
 }
 
@@ -101,6 +131,8 @@ MODEL=$1
 shift
 MODEL_FILES=
 LINK_MODEL=1
+FORCE_HF_DOWNLOAD=0
+FLATTEN_DOWNLOADS=0
 
 case "$MODEL" in
     q2-imatrix) MODEL_FILE=$Q2_IMATRIX_FILE ;;
@@ -114,6 +146,31 @@ case "$MODEL" in
         LINK_MODEL=0
         ;;
     mtp) MODEL_FILE=$MTP_FILE; LINK_MODEL=0 ;;
+    glm-unsloth-q4)
+        REPO=$GLM_UNSLOTH_REPO
+        MODEL_FILE=$GLM_UNSLOTH_Q4_FIRST_FILE
+        MODEL_FILES=
+        for part in 00001 00002 00003 00004 00005 00006 00007 00008 00009 00010 00011; do
+            MODEL_FILES="$MODEL_FILES $GLM_UNSLOTH_Q4_REMOTE_BASE-${part}-of-00011.gguf"
+        done
+        FORCE_HF_DOWNLOAD=1
+        FLATTEN_DOWNLOADS=1
+        ;;
+    glm-antirez-q2)
+        REPO=$GLM_ANTIREZ_REPO
+        MODEL_FILE=$GLM_ANTIREZ_Q2_FILE
+        FORCE_HF_DOWNLOAD=1
+        ;;
+    glm-antirez-iq2xxs)
+        REPO=$GLM_ANTIREZ_REPO
+        MODEL_FILE=$GLM_ANTIREZ_IQ2XXS_FILE
+        FORCE_HF_DOWNLOAD=1
+        ;;
+    glm-antirez-q4)
+        REPO=$GLM_ANTIREZ_REPO
+        MODEL_FILE=$GLM_ANTIREZ_Q4_FILE
+        FORCE_HF_DOWNLOAD=1
+        ;;
     -h|--help|help)
         usage
         exit 0
@@ -149,6 +206,9 @@ if [ -z "$TOKEN" ] && [ -s "$HOME/.cache/huggingface/token" ]; then
 fi
 
 needs_hf_download() {
+    if [ "${FORCE_HF_DOWNLOAD:-0}" -eq 1 ]; then
+        return 0
+    fi
     case "$1" in
         "$PRO_Q2_IMATRIX_FILE"|"$PRO_Q4_LAYERS00_30_FILE"|"$PRO_Q4_LAYERS31_OUTPUT_FILE")
             return 0
@@ -167,12 +227,22 @@ find_hf_command() {
     return 1
 }
 
+local_download_name() {
+    if [ "${FLATTEN_DOWNLOADS:-0}" -eq 1 ]; then
+        basename "$1"
+    else
+        printf '%s\n' "$1"
+    fi
+}
+
 download_one_hf() {
     file=$1
-    out="$OUT_DIR/$file"
+    local_file=$(local_download_name "$file")
+    out="$OUT_DIR/$local_file"
+    hf_out="$OUT_DIR/$file"
     part="$out.part"
 
-    mkdir -p "$OUT_DIR"
+    mkdir -p "$(dirname "$out")"
 
     if [ -s "$out" ]; then
         echo "Already downloaded: $out"
@@ -188,7 +258,7 @@ download_one_hf() {
 
     HF_CMD=$(find_hf_command || true)
     if [ -z "$HF_CMD" ]; then
-        echo "PRO downloads require the official Hugging Face CLI." >&2
+        echo "Large GGUF downloads require the official Hugging Face CLI." >&2
         echo "Install it with:" >&2
         echo "  python3 -m pip install -U huggingface_hub hf_xet" >&2
         exit 1
@@ -205,6 +275,11 @@ download_one_hf() {
         "$HF_CMD" download "$REPO" "$file" --repo-type model --local-dir "$OUT_DIR"
     fi
 
+    if [ "$hf_out" != "$out" ] && [ -s "$hf_out" ]; then
+        mv "$hf_out" "$out"
+        rmdir "$(dirname "$hf_out")" 2>/dev/null || true
+    fi
+
     if [ ! -s "$out" ]; then
         echo "Hugging Face download finished but expected file is missing: $out" >&2
         exit 1
@@ -213,7 +288,8 @@ download_one_hf() {
 
 download_one() {
     file=$1
-    out="$OUT_DIR/$file"
+    local_file=$(local_download_name "$file")
+    out="$OUT_DIR/$local_file"
     part="$out.part"
     aria2_part="$out.aria2"
     url="https://huggingface.co/$REPO/resolve/main/$file"
@@ -223,7 +299,7 @@ download_one() {
         return
     fi
 
-    mkdir -p "$OUT_DIR"
+    mkdir -p "$(dirname "$out")"
 
     if [ -e "$aria2_part" ]; then
         echo "Found incomplete aria2 download sidecar: $aria2_part" >&2

ds4.h

@@ -96,6 +96,7 @@ typedef struct {
     const char *mtp_path;
     ds4_backend backend;
     int n_threads;
+    int context_size;
     uint32_t prefill_chunk;
     int mtp_draft_tokens;
     float mtp_margin;
@@ -106,13 +107,17 @@ typedef struct {
     int power_percent;
     uint32_t ssd_streaming_cache_experts;
     uint64_t ssd_streaming_cache_bytes;
+    uint32_t ssd_streaming_full_layers;
     uint32_t ssd_streaming_preload_experts;
     uint64_t simulate_used_memory_bytes;
     bool warm_weights;
     bool quality;
     bool ssd_streaming;
     bool ssd_streaming_cold;
+    bool ssd_streaming_full_layers_set;
     bool inspect_only;
+    bool first_token_test;
+    bool metal_graph_test;
     bool load_slice;
     uint32_t load_layer_start;
     uint32_t load_layer_end;
@@ -154,14 +159,24 @@ const char *ds4_engine_model_name(ds4_engine *e);
 int ds4_engine_layer_count(ds4_engine *e);
 uint32_t ds4_engine_layer_compress_ratio(ds4_engine *e, uint32_t layer);
 uint64_t ds4_engine_hidden_f32_values(ds4_engine *e);
+bool ds4_engine_glm_layer_payload_bytes(ds4_engine *e,
+                                        uint32_t layer,
+                                        uint32_t full_live,
+                                        uint32_t key_dim,
+                                        uint32_t value_dim,
+                                        uint32_t compact_live,
+                                        uint32_t index_live,
+                                        uint64_t *out);
 /* Stable id for cache compatibility.  0 is the original Flash shape, so old
  * KV files with the previously-zero reserved byte remain Flash-compatible;
  * Pro and later shapes must use nonzero ids. */
 int ds4_engine_model_id(ds4_engine *e);
+bool ds4_engine_is_glm_dsa(ds4_engine *e);
 const char *ds4_backend_name(ds4_backend backend);
 bool ds4_think_mode_enabled(ds4_think_mode mode);
 const char *ds4_think_mode_name(ds4_think_mode mode);
 const char *ds4_think_max_prefix(void);
+const char *ds4_glm_reasoning_effort_text(ds4_think_mode mode);
 uint32_t ds4_think_max_min_context(void);
 ds4_think_mode ds4_think_mode_for_context(ds4_think_mode mode, int ctx_size);
 /* Uses the active model shape selected by ds4_engine_open(); call after opening
@@ -171,6 +186,11 @@ ds4_context_memory ds4_context_memory_estimate_with_prefill(
         ds4_backend backend,
         int ctx_size,
         uint32_t prefill_chunk);
+ds4_context_memory ds4_context_memory_estimate_with_prefill_mode(
+        ds4_backend backend,
+        int ctx_size,
+        uint32_t prefill_chunk,
+        bool ssd_streaming);
 bool ds4_log_is_tty(FILE *fp);
 void ds4_log(FILE *fp, ds4_log_type type, const char *fmt, ...);
 int ds4_engine_generate_argmax(ds4_engine *e, const ds4_tokens *prompt,
@@ -214,6 +234,11 @@ void ds4_chat_append_assistant_prefix(ds4_engine *e, ds4_tokens *tokens, ds4_thi
 
 char *ds4_token_text(ds4_engine *e, int token, size_t *len);
 int ds4_token_eos(ds4_engine *e);
+bool ds4_token_is_stop(ds4_engine *e, int token);
+bool ds4_token_is_thinking_control(ds4_engine *e, int token);
+bool ds4_token_is_stop_for_think_mode(ds4_engine *e,
+                                      int token,
+                                      ds4_think_mode mode);
 int ds4_token_user(ds4_engine *e);
 int ds4_token_assistant(ds4_engine *e);
 

ds4_bench.c

@@ -6,10 +6,10 @@
  *
  * The benchmark walks one fixed token sequence to configurable context
  * frontiers, measuring only the newest prefill interval at each frontier.  It
- * then snapshots the live session in memory, performs a fixed greedy decode
- * run without allowing EOS, restores the snapshot, and continues to the next
- * frontier.  Snapshot save/restore time is intentionally outside both timing
- * windows.
+ * then snapshots the live session in memory when the payload is small enough,
+ * performs a fixed greedy decode run without allowing EOS, restores the
+ * snapshot or replays the prefix, and continues to the next frontier.  Snapshot
+ * save/restore time is intentionally outside both timing windows.
  */
 
 #include <errno.h>
@@ -22,6 +22,8 @@
 #include <string.h>
 #include <time.h>
 
+#define DS4_BENCH_DEFAULT_SNAPSHOT_MAX_BYTES (UINT64_C(1) << 30)
+
 typedef struct {
     const char *model_path;
     const char *prompt_path;
@@ -40,6 +42,7 @@ typedef struct {
     uint32_t prefill_chunk;
     uint32_t ssd_streaming_cache_experts;
     uint64_t ssd_streaming_cache_bytes;
+    uint32_t ssd_streaming_full_layers;
     uint32_t ssd_streaming_preload_experts;
     uint64_t simulate_used_memory_bytes;
     double step_mul;
@@ -49,6 +52,7 @@ typedef struct {
     bool quality;
     bool ssd_streaming;
     bool ssd_streaming_cold;
+    bool ssd_streaming_full_layers_set;
 } bench_config;
 
 static double bench_now_sec(void) {
@@ -57,6 +61,29 @@ static double bench_now_sec(void) {
     return (double)ts.tv_sec + (double)ts.tv_nsec / 1000000000.0;
 }
 
+static uint64_t bench_snapshot_max_bytes(void) {
+    const char *env = getenv("DS4_BENCH_SNAPSHOT_MAX_BYTES");
+    if (!env || env[0] == '\0') return DS4_BENCH_DEFAULT_SNAPSHOT_MAX_BYTES;
+    if (!strcmp(env, "unlimited") || !strcmp(env, "UNLIMITED") ||
+        !strcmp(env, "inf") || !strcmp(env, "INF")) {
+        return UINT64_MAX;
+    }
+    char *end = NULL;
+    unsigned long long v = strtoull(env, &end, 10);
+    if (env[0] == '\0' || !end || *end != '\0') {
+        fprintf(stderr,
+                "ds4-bench: invalid DS4_BENCH_SNAPSHOT_MAX_BYTES=%s; using default %llu\n",
+                env,
+                (unsigned long long)DS4_BENCH_DEFAULT_SNAPSHOT_MAX_BYTES);
+        return DS4_BENCH_DEFAULT_SNAPSHOT_MAX_BYTES;
+    }
+    return (uint64_t)v;
+}
+
+static double bytes_to_gib(uint64_t bytes) {
+    return (double)bytes / (1024.0 * 1024.0 * 1024.0);
+}
+
 static void usage(FILE *fp, const char *topic) {
     ds4_help_print(fp, DS4_HELP_BENCH, topic);
 }
@@ -260,6 +287,10 @@ static bench_config parse_options(int argc, char **argv) {
             }
             c.ssd_streaming_cache_experts = experts;
             c.ssd_streaming_cache_bytes = bytes;
+        } else if (!strcmp(arg, "--ssd-streaming-full-layers")) {
+            int v = parse_nonnegative_int(need_arg(&i, argc, argv, arg), arg);
+            c.ssd_streaming_full_layers = (uint32_t)v;
+            c.ssd_streaming_full_layers_set = true;
         } else if (!strcmp(arg, "--ssd-streaming-preload-experts")) {
             int v = parse_int(need_arg(&i, argc, argv, arg), arg);
             if (v <= 0) {
@@ -441,11 +472,13 @@ static int next_frontier(const bench_config *c, int cur) {
 
 static void log_context_memory(ds4_backend backend,
                                int         ctx_size,
-                               uint32_t    prefill_chunk) {
+                               uint32_t    prefill_chunk,
+                               bool        ssd_streaming) {
     ds4_context_memory m =
-        ds4_context_memory_estimate_with_prefill(backend,
-                                                 ctx_size,
-                                                 prefill_chunk);
+        ds4_context_memory_estimate_with_prefill_mode(backend,
+                                                      ctx_size,
+                                                      prefill_chunk,
+                                                      ssd_streaming);
     fprintf(stderr,
             "ds4-bench: context buffers %.2f MiB (ctx=%d, backend=%s, prefill_chunk=%u, raw_kv_rows=%u, compressed_kv_rows=%u)\n",
             (double)m.total_bytes / (1024.0 * 1024.0),
@@ -512,9 +545,11 @@ int main(int argc, char **argv) {
         .model_path = cfg.model_path,
         .backend = cfg.backend,
         .n_threads = cfg.threads,
+        .context_size = cfg.ctx_alloc,
         .prefill_chunk = cfg.prefill_chunk,
         .ssd_streaming_cache_experts = cfg.ssd_streaming_cache_experts,
         .ssd_streaming_cache_bytes = cfg.ssd_streaming_cache_bytes,
+        .ssd_streaming_full_layers = cfg.ssd_streaming_full_layers,
         .ssd_streaming_preload_experts = cfg.ssd_streaming_preload_experts,
         .simulate_used_memory_bytes = cfg.simulate_used_memory_bytes,
         .power_percent = cfg.power_percent,
@@ -522,6 +557,7 @@ int main(int argc, char **argv) {
         .quality = cfg.quality,
         .ssd_streaming = cfg.ssd_streaming,
         .ssd_streaming_cold = cfg.ssd_streaming_cold,
+        .ssd_streaming_full_layers_set = cfg.ssd_streaming_full_layers_set,
         .expert_profile_path = cfg.expert_profile_path,
         .distributed = cfg.dist,
     };
@@ -532,7 +568,10 @@ int main(int argc, char **argv) {
     }
     ds4_engine *engine = NULL;
     if (ds4_engine_open(&engine, &opt) != 0) return 1;
-    log_context_memory(cfg.backend, cfg.ctx_alloc, cfg.prefill_chunk);
+    log_context_memory(cfg.backend,
+                       cfg.ctx_alloc,
+                       cfg.prefill_chunk,
+                       cfg.ssd_streaming);
 
     char *text = read_file(cfg.prompt_path ? cfg.prompt_path : cfg.chat_prompt_path);
     ds4_tokens prompt = {0};
@@ -581,12 +620,14 @@ int main(int argc, char **argv) {
             return 1;
         }
     }
-    fprintf(out, "ctx_tokens,prefill_tokens,prefill_tps,gen_tokens,gen_tps,kvcache_bytes\n");
+    fprintf(out, "ctx_tokens,prefill_tokens,prefill_tps,gen_tokens,gen_tps,gen_first_ms,gen_steady_tokens,gen_steady_tps,kvcache_bytes\n");
     fflush(out);
 
     const int eos = ds4_token_eos(engine);
     const bool distributed = cfg.dist.role == DS4_DISTRIBUTED_COORDINATOR;
     ds4_session_snapshot snap = {0};
+    const uint64_t snapshot_max_bytes = bench_snapshot_max_bytes();
+    bool warned_large_snapshot = false;
     char err[256];
     int previous = 0;
     int rc = 0;
@@ -613,15 +654,37 @@ int main(int argc, char **argv) {
             break;
         }
 
-        if (cfg.gen_tokens > 0 && !distributed) {
-            if (ds4_session_save_snapshot(session, &snap, err, sizeof(err)) != 0) {
-                fprintf(stderr, "ds4-bench: snapshot at %d failed: %s\n", frontier, err);
-                rc = 1;
-                break;
+        const bool need_restore_after_generation =
+            cfg.gen_tokens > 0 && frontier < cfg.ctx_max;
+        bool have_snapshot = false;
+        if (need_restore_after_generation && !distributed &&
+            getenv("DS4_BENCH_DISABLE_SNAPSHOT") == NULL) {
+            const uint64_t payload_bytes = ds4_session_payload_bytes(session);
+            const bool large_snapshot_forced =
+                getenv("DS4_BENCH_FORCE_SNAPSHOT") != NULL;
+            if (payload_bytes > snapshot_max_bytes && !large_snapshot_forced) {
+                if (!warned_large_snapshot) {
+                    fprintf(stderr,
+                            "ds4-bench: session payload snapshot is %.2f GiB, above the %.2f GiB benchmark limit; "
+                            "replaying prefixes instead (set DS4_BENCH_FORCE_SNAPSHOT=1 to force snapshots)\n",
+                            bytes_to_gib(payload_bytes),
+                            bytes_to_gib(snapshot_max_bytes));
+                    warned_large_snapshot = true;
+                }
+            } else if (payload_bytes > 0) {
+                if (ds4_session_save_snapshot(session, &snap, err, sizeof(err)) != 0) {
+                    fprintf(stderr, "ds4-bench: snapshot at %d failed: %s\n", frontier, err);
+                    rc = 1;
+                    break;
+                }
+                have_snapshot = true;
             }
         }
 
         const double gen_t0 = bench_now_sec();
+        double gen_first_sec = 0.0;
+        double gen_steady_sec = 0.0;
+        int gen_done = 0;
         for (int i = 0; i < cfg.gen_tokens; i++) {
             if (ds4_session_pos(session) + 1 >= ds4_session_ctx(session)) {
                 fprintf(stderr, "ds4-bench: generation would exceed allocated context at frontier %d\n", frontier);
@@ -634,20 +697,25 @@ int main(int argc, char **argv) {
                 rc = 1;
                 break;
             }
+            const double token_t0 = bench_now_sec();
             if (ds4_session_eval(session, token, err, sizeof(err)) != 0) {
                 fprintf(stderr, "ds4-bench: decode at frontier %d failed: %s\n", frontier, err);
                 rc = 1;
                 break;
             }
+            const double token_t1 = bench_now_sec();
+            if (i == 0) gen_first_sec = token_t1 - token_t0;
+            else gen_steady_sec += token_t1 - token_t0;
+            gen_done++;
         }
         const double gen_t1 = bench_now_sec();
         if (rc != 0) break;
 
-        if (cfg.gen_tokens == 0) {
-            /* Pure prefill benchmark: leave the live session at the frontier. */
-        } else if (distributed) {
+        if (!need_restore_after_generation) {
+            /* Nothing later depends on the frontier state. */
+        } else if (distributed || !have_snapshot) {
             if (ds4_session_sync(session, &prefix, err, sizeof(err)) != 0) {
-                fprintf(stderr, "ds4-bench: distributed replay restore at %d failed: %s\n", frontier, err);
+                fprintf(stderr, "ds4-bench: replay restore at %d failed: %s\n", frontier, err);
                 rc = 1;
                 break;
             }
@@ -660,14 +728,18 @@ int main(int argc, char **argv) {
         }
 
         const double gen_sec = gen_t1 - gen_t0;
+        const int gen_steady_tokens = gen_done > 1 ? gen_done - 1 : 0;
         fprintf(out,
-                "%d,%d,%.2f,%d,%.2f,%llu\n",
+                "%d,%d,%.2f,%d,%.2f,%.3f,%d,%.2f,%llu\n",
                 frontier,
                 prefill_tokens,
                 prefill_sec > 0.0 ? (double)prefill_tokens / prefill_sec : 0.0,
-                cfg.gen_tokens,
-                gen_sec > 0.0 ? (double)cfg.gen_tokens / gen_sec : 0.0,
-                (unsigned long long)(distributed ? 0 : snap.len));
+                gen_done,
+                gen_sec > 0.0 ? (double)gen_done / gen_sec : 0.0,
+                gen_first_sec * 1000.0,
+                gen_steady_tokens,
+                gen_steady_sec > 0.0 ? (double)gen_steady_tokens / gen_steady_sec : 0.0,
+                (unsigned long long)(have_snapshot ? snap.len : 0));
         fflush(out);
 
         previous = frontier;

ds4_cli.c

@@ -5,7 +5,7 @@
 
 /* ds4 CLI.
  *
- * One-shot mode builds a single DeepSeek chat prompt and exits.  Interactive
+ * One-shot mode builds a single model-family chat prompt and exits.  Interactive
  * mode keeps a rendered token transcript plus one ds4_session, so follow-up
  * turns reuse the live Metal KV checkpoint just like the server does.  The CLI
  * deliberately keeps policy here and leaves graph/cache mechanics inside the
@@ -28,16 +28,21 @@
 typedef struct {
     const char *prompt;
     const char *system;
+    bool raw_prompt;
     int n_predict;
     int ctx_size;
     float temperature;
     float top_p;
     float min_p;
+    bool temperature_set;
+    bool top_p_set;
+    bool min_p_set;
     uint64_t seed;
     bool dump_tokens;
     const char *dump_logits_path;
     const char *dump_logprobs_path;
     int dump_logprobs_top_k;
+    int decode_consistency_tokens;
     const char *perplexity_file_path;
     const char *imatrix_dataset_path;
     const char *imatrix_output_path;
@@ -142,6 +147,16 @@ static int parse_int(const char *s, const char *opt) {
     return (int)v;
 }
 
+static int parse_nonnegative_int(const char *s, const char *opt) {
+    char *end = NULL;
+    long v = strtol(s, &end, 10);
+    if (s[0] == '\0' || *end != '\0' || v < 0 || v > INT32_MAX) {
+        fprintf(stderr, "ds4: invalid value for %s: %s\n", opt, s);
+        exit(2);
+    }
+    return (int)v;
+}
+
 static uint64_t parse_u64(const char *s, const char *opt) {
     char *end = NULL;
     unsigned long long v = strtoull(s, &end, 10);
@@ -191,19 +206,40 @@ static ds4_backend default_backend(void) {
 
 static void log_context_memory(ds4_backend backend,
                                int         ctx_size,
-                               uint32_t    prefill_chunk) {
+                               uint32_t    prefill_chunk,
+                               bool        ssd_streaming) {
     ds4_context_memory m =
-        ds4_context_memory_estimate_with_prefill(backend,
-                                                 ctx_size,
-                                                 prefill_chunk);
+        ds4_context_memory_estimate_with_prefill_mode(backend,
+                                                      ctx_size,
+                                                      prefill_chunk,
+                                                      ssd_streaming);
+    const uint64_t kv_bytes = m.raw_bytes + m.compressed_bytes;
+    const uint64_t total_bytes = kv_bytes + m.scratch_bytes;
+    const bool color = ds4_log_is_tty(stderr);
+    const char *green = color ? "\x1b[32m" : "";
+    const char *bright_green = color ? "\x1b[1;32m" : "";
+    const char *reset = color ? "\x1b[0m" : "";
     fprintf(stderr,
-            "ds4: context buffers %.2f MiB (ctx=%d, backend=%s, prefill_chunk=%u, raw_kv_rows=%u, compressed_kv_rows=%u)\n",
-            (double)m.total_bytes / (1024.0 * 1024.0),
+            "%sds4: memory: KV %.2f GiB (raw %.2f + compressed %.2f) "
+            "+ buffers %.2f GiB = %s%.2f GiB context%s\n",
+            green,
+            (double)kv_bytes / 1073741824.0,
+            (double)m.raw_bytes / 1073741824.0,
+            (double)m.compressed_bytes / 1073741824.0,
+            (double)m.scratch_bytes / 1073741824.0,
+            bright_green,
+            (double)total_bytes / 1073741824.0,
+            reset);
+    fprintf(stderr,
+            "%sds4: memory detail: ctx=%d prefill_cap=%u raw_kv_rows=%u "
+            "compressed_kv_rows=%u backend=%s%s\n",
+            green,
             ctx_size,
-            ds4_backend_name(backend),
             m.prefill_cap,
             m.raw_cap,
-            m.comp_cap);
+            m.comp_cap,
+            ds4_backend_name(backend),
+            reset);
 }
 
 static ds4_think_mode cli_effective_think_mode(const cli_generation_options *gen) {
@@ -244,7 +280,8 @@ static void cli_prefill_progress_cb(void *ud, const char *event, int current, in
     (void)total;
     cli_prefill_progress *p = ud;
     if (!p || !event || p->input_tokens <= 0) return;
-    const bool is_display = strcmp(event, "prefill_display") == 0;
+    const bool is_display =
+        strcmp(event, "prefill_display") == 0;
     if (strcmp(event, "prefill_chunk") && !is_display) return;
     if (is_display && !p->use_color) return;
 
@@ -281,8 +318,22 @@ static void cli_prefill_progress_cb(void *ud, const char *event, int current, in
 }
 
 static bool is_rendered_chat_prompt(const char *prompt) {
-    const char *bos = "<｜begin▁of▁sentence｜>";
-    return prompt && strncmp(prompt, bos, strlen(bos)) == 0;
+    static const char *prefixes[] = {
+        "<｜begin▁of▁sentence｜>",
+        "<｜User｜>",
+        "[gMASK]",
+        "<sop>",
+        "<|system|>",
+        "<|user|>",
+        "<|assistant|>",
+        "<|observation|>",
+    };
+    if (!prompt) return false;
+    for (size_t i = 0; i < sizeof(prefixes) / sizeof(prefixes[0]); i++) {
+        size_t n = strlen(prefixes[i]);
+        if (strncmp(prompt, prefixes[i], n) == 0) return true;
+    }
+    return false;
 }
 
 typedef struct {
@@ -410,7 +461,9 @@ static void print_generated_token(void *ud, int token) {
 }
 
 static void build_prompt(ds4_engine *engine, const cli_generation_options *gen, ds4_tokens *out) {
-    if (is_rendered_chat_prompt(gen->prompt)) {
+    if (gen->raw_prompt) {
+        ds4_tokenize_text(engine, gen->prompt ? gen->prompt : "", out);
+    } else if (is_rendered_chat_prompt(gen->prompt)) {
         ds4_tokenize_rendered_chat(engine, gen->prompt, out);
     } else {
         ds4_encode_chat_prompt(engine, gen->system, gen->prompt,
@@ -418,6 +471,16 @@ static void build_prompt(ds4_engine *engine, const cli_generation_options *gen,
     }
 }
 
+static void cli_apply_model_sampling_defaults(
+        ds4_engine             *engine,
+        cli_generation_options *gen) {
+    if (!engine || !gen || !ds4_engine_is_glm_dsa(engine)) return;
+
+    if (!gen->temperature_set) gen->temperature = 1.0f;
+    if (!gen->top_p_set) gen->top_p = 0.95f;
+    if (!gen->min_p_set) gen->min_p = 0.0f;
+}
+
 static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, const ds4_tokens *prompt) {
     ds4_session *session = NULL;
     if (ds4_session_create(&session, engine, cfg->gen.ctx_size) != 0) {
@@ -476,7 +539,7 @@ static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, con
     while (generated < max_tokens && !cli_interrupt_requested()) {
         int token = ds4_session_sample(session, cfg->gen.temperature, 0,
                                        cfg->gen.top_p, cfg->gen.min_p, &rng);
-        if (token == ds4_token_eos(engine)) break;
+        if (ds4_token_is_stop_for_think_mode(engine, token, think_mode)) break;
 
         int toks[17];
         int ntok = 0;
@@ -498,6 +561,16 @@ static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, con
                 return 1;
             }
         } else {
+            size_t piece_len = 0;
+            char *piece = ds4_token_text(engine, token, &piece_len);
+            token_printer_write_text(&printer, piece, piece_len);
+            fflush(stdout);
+            free(piece);
+            generated++;
+            if (generated >= max_tokens || cli_interrupt_requested()) {
+                continue;
+            }
+
             cli_dist_busy_set(cfg, true);
             int eval_rc = ds4_session_eval(session, token, err, sizeof(err));
             cli_dist_busy_set(cfg, false);
@@ -506,13 +579,12 @@ static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, con
                 ds4_session_free(session);
                 return 1;
             }
-            toks[0] = token;
-            ntok = 1;
+            continue;
         }
 
         bool stop = false;
         for (int j = 0; j < ntok; j++) {
-            if (toks[j] == ds4_token_eos(engine)) {
+            if (ds4_token_is_stop_for_think_mode(engine, toks[j], think_mode)) {
                 stop = true;
                 break;
             }
@@ -766,7 +838,7 @@ static int run_logprob_dump(ds4_engine *engine, const cli_config *cfg, const ds4
         }
         fputs("]}", fp);
 
-        if (token == ds4_token_eos(engine)) break;
+        if (ds4_token_is_stop(engine, token)) break;
         if (ds4_session_eval(session, token, err, sizeof(err)) != 0) {
             fprintf(stderr, "ds4: decode failed while dumping logprobs: %s\n", err);
             free(scores);
@@ -787,6 +859,176 @@ static int run_logprob_dump(ds4_engine *engine, const cli_config *cfg, const ds4
     return 0;
 }
 
+static void print_diag_token(FILE *fp, ds4_engine *engine, int token) {
+    fputs("{\"token\":", fp);
+    json_write_token(fp, engine, token);
+    fputc('}', fp);
+}
+
+static void print_diag_top(FILE *fp, ds4_engine *engine, const char *label,
+                           const ds4_token_score *scores, int n) {
+    fprintf(fp, "%s:", label);
+    for (int i = 0; i < n && scores[i].id >= 0; i++) {
+        fputc(' ', fp);
+        json_write_token(fp, engine, scores[i].id);
+        fprintf(fp, "@%.6g", scores[i].logit);
+    }
+    fputc('\n', fp);
+}
+
+static int run_decode_consistency(ds4_engine *engine, const cli_config *cfg,
+                                  const ds4_tokens *prompt) {
+    if (cfg->dist && cfg->dist->role != DS4_DISTRIBUTED_NONE) {
+        fprintf(stderr, "ds4: --decode-consistency is local-session only\n");
+        return 1;
+    }
+    if (cfg->gen.decode_consistency_tokens < 0) {
+        fprintf(stderr, "ds4: --decode-consistency requires a non-negative token count\n");
+        return 1;
+    }
+
+    const int vocab = ds4_engine_vocab_size(engine);
+    float *live_logits = malloc((size_t)vocab * sizeof(live_logits[0]));
+    float *fresh_logits = malloc((size_t)vocab * sizeof(fresh_logits[0]));
+    if (!live_logits || !fresh_logits) {
+        free(live_logits);
+        free(fresh_logits);
+        fprintf(stderr, "ds4: out of memory allocating diagnostic logits\n");
+        return 1;
+    }
+
+    ds4_tokens prefix = {0};
+    ds4_tokens_copy(&prefix, prompt);
+
+    ds4_session *live = NULL;
+    char err[160];
+    if (ds4_session_create(&live, engine, cfg->gen.ctx_size) != 0) {
+        fprintf(stderr, "ds4: --decode-consistency requires a graph session backend\n");
+        ds4_tokens_free(&prefix);
+        free(live_logits);
+        free(fresh_logits);
+        return 1;
+    }
+
+    cli_prefill_progress progress = {
+        .base_tokens = 0,
+        .input_tokens = prompt->len,
+        .use_color = ds4_log_is_tty(stderr),
+    };
+    ds4_session_set_progress(live, cli_prefill_progress_cb, &progress);
+    ds4_session_set_display_progress(live,
+                                     progress.use_color ? cli_prefill_progress_cb : NULL,
+                                     progress.use_color ? &progress : NULL);
+    if (ds4_session_sync(live, prompt, err, sizeof(err)) != 0) {
+        ds4_session_set_progress(live, NULL, NULL);
+        ds4_session_set_display_progress(live, NULL, NULL);
+        fprintf(stderr, "ds4: prompt processing failed: %s\n", err);
+        ds4_session_free(live);
+        ds4_tokens_free(&prefix);
+        free(live_logits);
+        free(fresh_logits);
+        return 1;
+    }
+    ds4_session_set_progress(live, NULL, NULL);
+    ds4_session_set_display_progress(live, NULL, NULL);
+
+    fprintf(stderr, "ds4: decode-consistency prompt_tokens=%d check_after=%d\n",
+            prompt->len, cfg->gen.decode_consistency_tokens);
+    for (int i = 0; i < cfg->gen.decode_consistency_tokens; i++) {
+        int token = ds4_session_argmax(live);
+        fprintf(stderr, "ds4: decode-consistency selected[%d]=", i);
+        print_diag_token(stderr, engine, token);
+        fputc('\n', stderr);
+        ds4_tokens_push(&prefix, token);
+        if (ds4_session_eval(live, token, err, sizeof(err)) != 0) {
+            fprintf(stderr, "ds4: decode failed during consistency check: %s\n", err);
+            ds4_session_free(live);
+            ds4_tokens_free(&prefix);
+            free(live_logits);
+            free(fresh_logits);
+            return 1;
+        }
+        if (ds4_token_is_stop(engine, token)) break;
+    }
+
+    ds4_token_score live_top[10];
+    int live_top_n = ds4_session_top_logprobs(live, live_top, 10);
+    if (ds4_session_copy_logits(live, live_logits, vocab) != vocab) {
+        fprintf(stderr, "ds4: failed to copy live logits\n");
+        ds4_session_free(live);
+        ds4_tokens_free(&prefix);
+        free(live_logits);
+        free(fresh_logits);
+        return 1;
+    }
+    ds4_session_free(live);
+    live = NULL;
+
+    ds4_session *fresh = NULL;
+    if (ds4_session_create(&fresh, engine, cfg->gen.ctx_size) != 0) {
+        fprintf(stderr, "ds4: failed to create fresh diagnostic session\n");
+        ds4_tokens_free(&prefix);
+        free(live_logits);
+        free(fresh_logits);
+        return 1;
+    }
+    progress.input_tokens = prefix.len;
+    ds4_session_set_progress(fresh, cli_prefill_progress_cb, &progress);
+    ds4_session_set_display_progress(fresh,
+                                     progress.use_color ? cli_prefill_progress_cb : NULL,
+                                     progress.use_color ? &progress : NULL);
+    if (ds4_session_sync(fresh, &prefix, err, sizeof(err)) != 0) {
+        ds4_session_set_progress(fresh, NULL, NULL);
+        ds4_session_set_display_progress(fresh, NULL, NULL);
+        fprintf(stderr, "ds4: fresh prompt processing failed: %s\n", err);
+        ds4_session_free(fresh);
+        ds4_tokens_free(&prefix);
+        free(live_logits);
+        free(fresh_logits);
+        return 1;
+    }
+    ds4_session_set_progress(fresh, NULL, NULL);
+    ds4_session_set_display_progress(fresh, NULL, NULL);
+
+    ds4_token_score fresh_top[10];
+    int fresh_top_n = ds4_session_top_logprobs(fresh, fresh_top, 10);
+    if (ds4_session_copy_logits(fresh, fresh_logits, vocab) != vocab) {
+        fprintf(stderr, "ds4: failed to copy fresh logits\n");
+        ds4_session_free(fresh);
+        ds4_tokens_free(&prefix);
+        free(live_logits);
+        free(fresh_logits);
+        return 1;
+    }
+    ds4_session_free(fresh);
+
+    double ss = 0.0;
+    float max_abs = 0.0f;
+    int max_i = 0;
+    for (int i = 0; i < vocab; i++) {
+        float d = fabsf(live_logits[i] - fresh_logits[i]);
+        if (d > max_abs) {
+            max_abs = d;
+            max_i = i;
+        }
+        ss += (double)d * (double)d;
+    }
+
+    fprintf(stderr,
+            "ds4: decode-consistency compared prefix_tokens=%d vocab=%d "
+            "max_abs=%.9g at token=%d live=%.9g fresh=%.9g rms=%.9g\n",
+            prefix.len, vocab, max_abs, max_i,
+            live_logits[max_i], fresh_logits[max_i],
+            sqrt(ss / (double)vocab));
+    print_diag_top(stderr, engine, "ds4: live_top", live_top, live_top_n);
+    print_diag_top(stderr, engine, "ds4: fresh_top", fresh_top, fresh_top_n);
+
+    ds4_tokens_free(&prefix);
+    free(live_logits);
+    free(fresh_logits);
+    return live_top_n > 0 && fresh_top_n > 0 && live_top[0].id == fresh_top[0].id ? 0 : 1;
+}
+
 static int run_perplexity_file(ds4_engine *engine, const cli_config *cfg) {
     char *text = read_prompt_file(cfg->gen.perplexity_file_path, true);
     ds4_tokens tokens = {0};
@@ -900,6 +1142,11 @@ static int run_generation(ds4_engine *engine, const cli_config *cfg) {
         ds4_tokens_free(&prompt);
         return rc;
     }
+    if (cfg->gen.decode_consistency_tokens > 0) {
+        rc = run_decode_consistency(engine, cfg, &prompt);
+        ds4_tokens_free(&prompt);
+        return rc;
+    }
 
     const bool diagnostic = cfg->gen.dump_tokens ||
                             cfg->gen.head_test ||
@@ -919,31 +1166,33 @@ static int run_generation(ds4_engine *engine, const cli_config *cfg) {
             fprintf(stderr, "ds4: diagnostic run completed on the native %s path.\n",
                     ds4_backend_name(cfg->engine.backend));
         }
-    } else if (cfg->engine.distributed.role == DS4_DISTRIBUTED_COORDINATOR ||
-               cfg->gen.temperature > 0.0f ||
-               ds4_engine_mtp_draft_tokens(engine) > 1) {
-        rc = run_sampled_generation(engine, cfg, &prompt);
     } else {
-        token_printer printer = {
-            .engine = engine,
-            .fp = stdout,
-            .format_thinking = ds4_think_mode_enabled(cli_effective_think_mode(&cfg->gen)),
-            .in_think = ds4_think_mode_enabled(cli_effective_think_mode(&cfg->gen)),
-            .use_color = isatty(fileno(stdout)) != 0,
-            .last_output_newline = true,
-        };
-        cli_prefill_progress progress = {
-            .base_tokens = 0,
-            .input_tokens = prompt.len,
-            .use_color = ds4_log_is_tty(stderr),
-        };
-        rc = ds4_engine_generate_argmax(engine, &prompt, cfg->gen.n_predict,
-                                        cfg->gen.ctx_size,
-                                        print_generated_token,
-                                        generation_done,
-                                        &printer,
-                                        cli_prefill_progress_cb,
-                                        &progress);
+        if (cfg->engine.distributed.role == DS4_DISTRIBUTED_COORDINATOR ||
+            cfg->gen.temperature > 0.0f ||
+            ds4_engine_mtp_draft_tokens(engine) > 1) {
+            rc = run_sampled_generation(engine, cfg, &prompt);
+        } else {
+            token_printer printer = {
+                .engine = engine,
+                .fp = stdout,
+                .format_thinking = ds4_think_mode_enabled(cli_effective_think_mode(&cfg->gen)),
+                .in_think = ds4_think_mode_enabled(cli_effective_think_mode(&cfg->gen)),
+                .use_color = isatty(fileno(stdout)) != 0,
+                .last_output_newline = true,
+            };
+            cli_prefill_progress progress = {
+                .base_tokens = 0,
+                .input_tokens = prompt.len,
+                .use_color = ds4_log_is_tty(stderr),
+            };
+            rc = ds4_engine_generate_argmax(engine, &prompt, cfg->gen.n_predict,
+                                            cfg->gen.ctx_size,
+                                            print_generated_token,
+                                            generation_done,
+                                            &printer,
+                                            cli_prefill_progress_cb,
+                                            &progress);
+        }
     }
 
     ds4_tokens_free(&prompt);
@@ -989,7 +1238,8 @@ typedef struct {
     ds4_session *session;
     ds4_tokens transcript;
     int ctx_size;
-    int max_prefix_tokens;
+    int think_prefix_pos;
+    int think_prefix_tokens;
 } repl_chat;
 
 static void tokens_insert(ds4_tokens *dst, int pos, const ds4_tokens *src) {
@@ -1019,23 +1269,49 @@ static void tokens_remove(ds4_tokens *dst, int pos, int n) {
     dst->len -= n;
 }
 
-/* Insert/remove the Think Max prefix inside the existing transcript.  The
- * prefix lives after BOS, before any system/developer text, which mirrors the
- * API rendering path.  Changing it invalidates the session because every later
- * token position would otherwise refer to the wrong prefix. */
-static void repl_chat_apply_max_prefix(ds4_engine *engine, repl_chat *chat, bool enable) {
-    if (enable && chat->max_prefix_tokens == 0) {
-        ds4_tokens prefix = {0};
-        ds4_chat_append_max_effort_prefix(engine, &prefix);
-        tokens_insert(&chat->transcript, 1, &prefix);
-        chat->max_prefix_tokens = prefix.len;
-        ds4_tokens_free(&prefix);
-        if (chat->session) ds4_session_invalidate(chat->session);
-    } else if (!enable && chat->max_prefix_tokens > 0) {
-        tokens_remove(&chat->transcript, 1, chat->max_prefix_tokens);
-        chat->max_prefix_tokens = 0;
+static const char *repl_glm_reasoning_effort_text(ds4_think_mode mode) {
+    switch (mode) {
+    case DS4_THINK_HIGH: return "Reasoning Effort: High";
+    case DS4_THINK_MAX:  return "Reasoning Effort: Max";
+    case DS4_THINK_NONE: return NULL;
+    }
+    return NULL;
+}
+
+static void repl_chat_build_think_prefix(ds4_engine *engine,
+                                         ds4_think_mode mode,
+                                         ds4_tokens *prefix) {
+    if (ds4_engine_is_glm_dsa(engine)) {
+        const char *effort = repl_glm_reasoning_effort_text(mode);
+        if (effort) ds4_chat_append_message(engine, prefix, "system", effort);
+    } else if (mode == DS4_THINK_MAX) {
+        ds4_chat_append_max_effort_prefix(engine, prefix);
+    }
+}
+
+/* Insert/replace the model-family thinking prefix inside the existing
+ * transcript.  It lives immediately after the BOS sequence and before any
+ * user/system text, matching the GGUF chat templates. */
+static void repl_chat_apply_think_prefix(ds4_engine *engine,
+                                         repl_chat *chat,
+                                         ds4_think_mode mode) {
+    ds4_tokens prefix = {0};
+    repl_chat_build_think_prefix(engine, mode, &prefix);
+
+    bool same = chat->think_prefix_tokens == prefix.len;
+    if (same && prefix.len > 0) {
+        same = !memcmp(chat->transcript.v + chat->think_prefix_pos,
+                       prefix.v,
+                       (size_t)prefix.len * sizeof(prefix.v[0]));
+    }
+    if (!same) {
+        tokens_remove(&chat->transcript, chat->think_prefix_pos,
+                      chat->think_prefix_tokens);
+        tokens_insert(&chat->transcript, chat->think_prefix_pos, &prefix);
+        chat->think_prefix_tokens = prefix.len;
         if (chat->session) ds4_session_invalidate(chat->session);
     }
+    ds4_tokens_free(&prefix);
 }
 
 static int repl_chat_create_session(ds4_engine *engine, repl_chat *chat, int ctx_size) {
@@ -1053,8 +1329,8 @@ static int repl_chat_create_session(ds4_engine *engine, repl_chat *chat, int ctx
 static int repl_chat_init(ds4_engine *engine, repl_chat *chat, const cli_config *cfg) {
     memset(chat, 0, sizeof(*chat));
     ds4_chat_begin(engine, &chat->transcript);
-    repl_chat_apply_max_prefix(engine, chat,
-                               cli_effective_think_mode(&cfg->gen) == DS4_THINK_MAX);
+    chat->think_prefix_pos = chat->transcript.len;
+    repl_chat_apply_think_prefix(engine, chat, cli_effective_think_mode(&cfg->gen));
     if (cfg->gen.system && cfg->gen.system[0]) {
         ds4_chat_append_message(engine, &chat->transcript, "system", cfg->gen.system);
     }
@@ -1075,6 +1351,10 @@ static int repl_chat_set_ctx(ds4_engine *engine, repl_chat *chat, int ctx_size)
     return repl_chat_create_session(engine, chat, ctx_size);
 }
 
+static bool repl_chat_assistant_turn_uses_eos(ds4_engine *engine) {
+    return !ds4_engine_is_glm_dsa(engine);
+}
+
 /* Run one interactive turn.  The transcript is tentatively extended with user
  * and assistant markers, then ds4_session_sync() decides whether this is a KV
  * continuation.  If prompt processing fails, the transcript rolls back before
@@ -1087,7 +1367,7 @@ static int run_chat_turn(ds4_engine *engine, cli_config *cfg, repl_chat *chat, c
 
     ds4_think_mode think_mode = ds4_think_mode_for_context(cfg->gen.think_mode,
                                                            chat->ctx_size);
-    repl_chat_apply_max_prefix(engine, chat, think_mode == DS4_THINK_MAX);
+    repl_chat_apply_think_prefix(engine, chat, think_mode);
     const int rollback_len = chat->transcript.len;
     ds4_chat_append_message(engine, &chat->transcript, "user", user_text);
     ds4_chat_append_assistant_prefix(engine, &chat->transcript, think_mode);
@@ -1147,7 +1427,7 @@ static int run_chat_turn(ds4_engine *engine, cli_config *cfg, repl_chat *chat, c
                                        cfg->gen.top_p,
                                        cfg->gen.min_p,
                                        &rng);
-        if (token == ds4_token_eos(engine)) break;
+        if (ds4_token_is_stop_for_think_mode(engine, token, think_mode)) break;
 
         int toks[17];
         int ntok = 0;
@@ -1168,20 +1448,29 @@ static int run_chat_turn(ds4_engine *engine, cli_config *cfg, repl_chat *chat, c
                 return 1;
             }
         } else {
+            size_t piece_len = 0;
+            char *piece = ds4_token_text(engine, token, &piece_len);
+            ds4_tokens_push(&chat->transcript, token);
+            token_printer_write_text(&printer, piece, piece_len);
+            fflush(stdout);
+            free(piece);
+            generated++;
+
             cli_dist_busy_set(cfg, true);
             int eval_rc = ds4_session_eval(chat->session, token, err, sizeof(err));
             cli_dist_busy_set(cfg, false);
             if (eval_rc != 0) {
                 fprintf(stderr, "ds4: decode failed: %s\n", err);
+                ds4_session_invalidate(chat->session);
                 return 1;
             }
-            toks[0] = token;
-            ntok = 1;
+            if (generated >= max_tokens) break;
+            continue;
         }
 
         bool stop = false;
         for (int j = 0; j < ntok; j++) {
-            if (toks[j] == ds4_token_eos(engine)) {
+            if (ds4_token_is_stop_for_think_mode(engine, toks[j], think_mode)) {
                 stop = true;
                 break;
             }
@@ -1203,7 +1492,7 @@ static int run_chat_turn(ds4_engine *engine, cli_config *cfg, repl_chat *chat, c
     if (interrupted && generated == 0) {
         chat->transcript.len = rollback_len;
         ds4_session_invalidate(chat->session);
-    } else {
+    } else if (repl_chat_assistant_turn_uses_eos(engine)) {
         ds4_tokens_push(&chat->transcript, ds4_token_eos(engine));
     }
 
@@ -1260,18 +1549,19 @@ static int run_repl(ds4_engine *engine, cli_config *cfg) {
             print_repl_help();
         } else if (!strcmp(cmd, "/think")) {
             cfg->gen.think_mode = DS4_THINK_HIGH;
-            repl_chat_apply_max_prefix(engine, &chat, false);
+            repl_chat_apply_think_prefix(engine, &chat, DS4_THINK_HIGH);
             puts("Thinking mode: high.");
         } else if (!strcmp(cmd, "/think-max")) {
             cfg->gen.think_mode = DS4_THINK_MAX;
             bool active = ds4_think_mode_for_context(cfg->gen.think_mode,
                                                      chat.ctx_size) == DS4_THINK_MAX;
-            repl_chat_apply_max_prefix(engine, &chat, active);
+            repl_chat_apply_think_prefix(engine, &chat,
+                                         active ? DS4_THINK_MAX : DS4_THINK_HIGH);
             cli_warn_think_max_downgraded(&cfg->gen, "/think-max");
             printf("Thinking mode: %s.\n", active ? "max" : "high (ctx below 393216)");
         } else if (!strcmp(cmd, "/nothink")) {
             cfg->gen.think_mode = DS4_THINK_NONE;
-            repl_chat_apply_max_prefix(engine, &chat, false);
+            repl_chat_apply_think_prefix(engine, &chat, DS4_THINK_NONE);
             puts("Thinking mode: none.");
         } else if (!strncmp(cmd, "/power", 6) && (cmd[6] == '\0' || isspace((unsigned char)cmd[6]))) {
             char *arg = trim_inplace(cmd + 6);
@@ -1296,15 +1586,16 @@ static int run_repl(ds4_engine *engine, cli_config *cfg) {
                 cfg->gen.ctx_size = parse_int(arg, "/ctx");
                 log_context_memory(cfg->engine.backend,
                                    cfg->gen.ctx_size,
-                                   cfg->engine.prefill_chunk);
+                                   cfg->engine.prefill_chunk,
+                                   cfg->engine.ssd_streaming);
                 rc = repl_chat_set_ctx(engine, &chat, cfg->gen.ctx_size);
                 if (rc != 0) {
                     linenoiseFree(line);
                     break;
                 }
-                bool active = ds4_think_mode_for_context(cfg->gen.think_mode,
-                                                         chat.ctx_size) == DS4_THINK_MAX;
-                repl_chat_apply_max_prefix(engine, &chat, active);
+                ds4_think_mode effective = ds4_think_mode_for_context(cfg->gen.think_mode,
+                                                                      chat.ctx_size);
+                repl_chat_apply_think_prefix(engine, &chat, effective);
                 cli_warn_think_max_downgraded(&cfg->gen, "/ctx");
             }
         } else if (!strcmp(cmd, "/quit") || !strcmp(cmd, "/exit")) {
@@ -1454,6 +1745,8 @@ static cli_config parse_options(int argc, char **argv) {
             c.gen.prompt = c.prompt_owned;
         } else if (!strcmp(arg, "-sys") || !strcmp(arg, "--system")) {
             c.gen.system = need_arg(&i, argc, argv, arg);
+        } else if (!strcmp(arg, "--raw") || !strcmp(arg, "--raw-prompt")) {
+            c.gen.raw_prompt = true;
         } else if (!strcmp(arg, "-m") || !strcmp(arg, "--model")) {
             c.engine.model_path = need_arg(&i, argc, argv, arg);
         } else if (!strcmp(arg, "--mtp")) {
@@ -1468,10 +1761,13 @@ static cli_config parse_options(int argc, char **argv) {
             c.gen.ctx_size = parse_int(need_arg(&i, argc, argv, arg), arg);
         } else if (!strcmp(arg, "--temp")) {
             c.gen.temperature = parse_float_range(need_arg(&i, argc, argv, arg), arg, 0.0f, 100.0f);
+            c.gen.temperature_set = true;
         } else if (!strcmp(arg, "--top-p")) {
             c.gen.top_p = parse_float_range(need_arg(&i, argc, argv, arg), arg, 0.0f, 1.0f);
+            c.gen.top_p_set = true;
         } else if (!strcmp(arg, "--min-p")) {
             c.gen.min_p = parse_float_range(need_arg(&i, argc, argv, arg), arg, 0.0f, 1.0f);
+            c.gen.min_p_set = true;
         } else if (!strcmp(arg, "--seed")) {
             c.gen.seed = parse_u64(need_arg(&i, argc, argv, arg), arg);
         } else if (!strcmp(arg, "--quality")) {
@@ -1491,6 +1787,10 @@ static cli_config parse_options(int argc, char **argv) {
             }
             c.engine.ssd_streaming_cache_experts = experts;
             c.engine.ssd_streaming_cache_bytes = bytes;
+        } else if (!strcmp(arg, "--ssd-streaming-full-layers")) {
+            int v = parse_nonnegative_int(need_arg(&i, argc, argv, arg), arg);
+            c.engine.ssd_streaming_full_layers = (uint32_t)v;
+            c.engine.ssd_streaming_full_layers_set = true;
         } else if (!strcmp(arg, "--ssd-streaming-preload-experts")) {
             int v = parse_int(need_arg(&i, argc, argv, arg), arg);
             if (v <= 0) {
@@ -1551,6 +1851,8 @@ static cli_config parse_options(int argc, char **argv) {
             c.gen.dump_logprobs_path = need_arg(&i, argc, argv, arg);
         } else if (!strcmp(arg, "--logprobs-top-k")) {
             c.gen.dump_logprobs_top_k = parse_int(need_arg(&i, argc, argv, arg), arg);
+        } else if (!strcmp(arg, "--decode-consistency")) {
+            c.gen.decode_consistency_tokens = parse_int(need_arg(&i, argc, argv, arg), arg);
         } else if (!strcmp(arg, "--perplexity-file")) {
             c.gen.perplexity_file_path = need_arg(&i, argc, argv, arg);
         } else if (!strcmp(arg, "--imatrix-dataset")) {
@@ -1650,12 +1952,16 @@ int main(int argc, char **argv) {
         return rc;
     }
     cfg.engine.inspect_only = cfg.inspect;
+    cfg.engine.first_token_test = cfg.gen.first_token_test;
+    cfg.engine.metal_graph_test = cfg.gen.metal_graph_test;
+    cfg.engine.context_size = cfg.gen.ctx_size;
     ds4_engine *engine = NULL;
     if (ds4_engine_open(&engine, &cfg.engine) != 0) {
         ds4_dist_options_free(cfg.dist);
         free(cfg.prompt_owned);
         return 1;
     }
+    cli_apply_model_sampling_defaults(engine, &cfg.gen);
     if (cfg.dist && cfg.dist->role == DS4_DISTRIBUTED_WORKER) {
         ds4_dist_generation_options dist_gen = {
             .prompt = cfg.gen.prompt,
@@ -1678,9 +1984,6 @@ int main(int argc, char **argv) {
         return rc;
     }
     if (!cfg.inspect) {
-        log_context_memory(cfg.engine.backend,
-                           cfg.gen.ctx_size,
-                           cfg.engine.prefill_chunk);
         cli_warn_think_max_downgraded(&cfg.gen, "--think-max");
     }
     int rc = 0;

ds4_cuda.cu

@@ -2478,6 +2478,7 @@ extern "C" int ds4_gpu_tensor_copy(ds4_gpu_tensor *dst, uint64_t dst_offset,
 
 extern "C" int ds4_gpu_begin_commands(void) { return 1; }
 extern "C" int ds4_gpu_flush_commands(void) { return cuda_ok(cudaDeviceSynchronize(), "flush"); }
+extern "C" int ds4_gpu_commands_active(void) { return 0; }
 extern "C" int ds4_gpu_signal_selected_readback_ready(uint64_t *event_value) {
     if (event_value) *event_value = 1;
     return cuda_ok(cudaDeviceSynchronize(), "selected readback signal");
@@ -2754,6 +2755,10 @@ extern "C" void ds4_gpu_set_quality(bool quality) {
     }
 }
 
+extern "C" void ds4_gpu_set_glm_model(bool enabled) {
+    (void)enabled;
+}
+
 extern "C" void ds4_gpu_set_ssd_streaming(bool enabled) {
     g_ssd_streaming_mode = enabled ? 1 : 0;
     g_stream_expert_runtime_cap = 0;
@@ -2766,6 +2771,10 @@ extern "C" void ds4_gpu_set_ssd_streaming(bool enabled) {
     }
 }
 
+extern "C" void ds4_gpu_set_glm_streaming_prefill_full_layer(bool enabled) {
+    (void)enabled;
+}
+
 extern "C" void ds4_gpu_set_streaming_expert_cache_budget(uint32_t experts) {
     g_stream_expert_budget_override = experts;
     g_stream_expert_runtime_cap = 0;
@@ -7779,6 +7788,20 @@ extern "C" int ds4_gpu_matmul_q8_0_tensor(ds4_gpu_tensor *out, const void *model
                                            in_dim, out_dim, x, n_tok, "q8_0");
 }
 
+extern "C" int ds4_gpu_matmul_q8_0_rows_scalar_tensor(
+        ds4_gpu_tensor *out,
+        const void *model_map,
+        uint64_t model_size,
+        uint64_t weight_offset,
+        uint64_t in_dim,
+        uint64_t out_dim,
+        const ds4_gpu_tensor *x,
+        uint64_t n_tok) {
+    (void)out; (void)model_map; (void)model_size; (void)weight_offset;
+    (void)in_dim; (void)out_dim; (void)x; (void)n_tok;
+    return 0;
+}
+
 extern "C" int ds4_gpu_matmul_q8_0_pair_tensor(
         ds4_gpu_tensor *out0,
         ds4_gpu_tensor *out1,
@@ -9497,6 +9520,26 @@ extern "C" int ds4_gpu_shared_gate_up_swiglu_q8_0_tensor(
                                         up_offset, in_dim, out_dim, x, 1) &&
            ds4_gpu_swiglu_tensor(mid, gate, up, (uint32_t)out_dim, clamp, 1.0f);
 }
+
+extern "C" int ds4_gpu_shared_gate_up_swiglu_q8_0_rows_scalar_tensor(
+        ds4_gpu_tensor       *gate,
+        ds4_gpu_tensor       *up,
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        uint64_t                n_tok,
+        float                   clamp) {
+    (void)gate; (void)up; (void)mid; (void)model_map; (void)model_size;
+    (void)gate_offset; (void)up_offset; (void)in_dim; (void)out_dim;
+    (void)x; (void)n_tok; (void)clamp;
+    return 0;
+}
+
 extern "C" int ds4_gpu_add_tensor(ds4_gpu_tensor *out, const ds4_gpu_tensor *a, const ds4_gpu_tensor *b, uint32_t n) {
     if (!out || !a || !b ||
         out->bytes < (uint64_t)n * sizeof(float) ||

ds4_distributed.c

@@ -4616,6 +4616,16 @@ static bool dist_kv_layer_tensor_bytes(
         uint32_t n_index_comp,
         uint64_t *out) {
     if (!layout || !out) return false;
+    if (ds4_engine_is_glm_dsa(engine)) {
+        return ds4_engine_glm_layer_payload_bytes(engine,
+                                                  layer,
+                                                  layout->raw_live,
+                                                  layout->head_dim,
+                                                  layout->indexer_head_dim,
+                                                  n_comp,
+                                                  n_index_comp,
+                                                  out);
+    }
     uint64_t bytes = 0;
     uint64_t tmp = 0;
     if (!dist_u64_mul(layout->raw_live, layout->head_dim, &tmp) ||
@@ -4668,12 +4678,15 @@ static bool dist_kv_layout_matches(
            a->raw_live == b->raw_live;
 }
 
-static bool dist_kv_raw_live_valid(const ds4_dist_kv_layout *layout) {
+static bool dist_kv_raw_live_valid(ds4_engine *engine, const ds4_dist_kv_layout *layout) {
     if (!layout || layout->raw_window == 0 || layout->raw_cap == 0) return false;
-    const uint32_t expected =
+    const uint32_t max_live =
         layout->token_count < layout->raw_window ? layout->token_count : layout->raw_window;
-    return layout->raw_live == expected &&
-           layout->raw_live <= layout->raw_cap;
+    if (ds4_engine_is_glm_dsa(engine)) {
+        return layout->raw_live <= max_live &&
+               layout->raw_live <= layout->raw_cap;
+    }
+    return layout->raw_live == max_live && layout->raw_live <= layout->raw_cap;
 }
 
 static int dist_kv_parse_layer_payload(
@@ -4725,7 +4738,7 @@ static int dist_kv_parse_layer_payload(
         return 1;
     }
     if (got.n_layers != (uint32_t)ds4_engine_layer_count(engine) ||
-        !dist_kv_raw_live_valid(&got)) {
+        !dist_kv_raw_live_valid(engine, &got)) {
         if (errlen) snprintf(err, errlen, "distributed KV shard layout is invalid");
         return 1;
     }
@@ -5251,7 +5264,7 @@ int ds4_dist_session_load_payload(
         layout.ctx > (uint32_t)ds4_session_ctx(owner) ||
         layout.token_count >= (uint32_t)ds4_session_ctx(owner) ||
         layout.vocab != (uint32_t)ds4_engine_vocab_size(d->state.engine) ||
-        !dist_kv_raw_live_valid(&layout)) {
+        !dist_kv_raw_live_valid(d->state.engine, &layout)) {
         if (errlen) snprintf(err, errlen, "DS4 KV payload does not match current distributed runtime");
         return 1;
     }
@@ -5648,10 +5661,6 @@ int ds4_dist_session_eval(
     }
     if (dist_session_ensure_route(d, err, errlen) != 0) return 1;
 
-    ds4_tokens transcript = {0};
-    ds4_tokens_copy(&transcript, checkpoint);
-    ds4_tokens_push(&transcript, token);
-
     int rc = dist_coordinator_eval_span(&d->state,
                                         owner,
                                         &d->plan,
@@ -5665,6 +5674,9 @@ int ds4_dist_session_eval(
                                         err,
                                         errlen);
     if (rc != 0) {
+        ds4_tokens transcript = {0};
+        ds4_tokens_copy(&transcript, checkpoint);
+        ds4_tokens_push(&transcript, token);
         if (dist_coordinator_rebuild_from_transcript(&d->state,
                                                      owner,
                                                      &d->plan,
@@ -5682,9 +5694,9 @@ int ds4_dist_session_eval(
             return 1;
         }
         d->plan_ready = true;
+        ds4_tokens_free(&transcript);
         rc = 0;
     }
-    ds4_tokens_free(&transcript);
     return rc;
 }
 

ds4_eval.c

@@ -1210,6 +1210,7 @@ typedef struct {
     uint32_t prefill_chunk;
     uint32_t ssd_streaming_cache_experts;
     uint64_t ssd_streaming_cache_bytes;
+    uint32_t ssd_streaming_full_layers;
     uint32_t ssd_streaming_preload_experts;
     uint64_t simulate_used_memory_bytes;
     int soft_limit_reply_budget;
@@ -1222,6 +1223,7 @@ typedef struct {
     bool quality;
     bool ssd_streaming;
     bool ssd_streaming_cold;
+    bool ssd_streaming_full_layers_set;
     bool self_test_extractors;
 } eval_config;
 
@@ -1440,6 +1442,16 @@ static int parse_int_arg(const char *s, const char *opt) {
     return (int)v;
 }
 
+static int parse_nonnegative_int_arg(const char *s, const char *opt) {
+    char *end = NULL;
+    long v = strtol(s, &end, 10);
+    if (s[0] == '\0' || *end != '\0' || v < 0 || v > INT_MAX) {
+        fprintf(stderr, "ds4-eval: invalid value for %s: %s\n", opt, s);
+        exit(2);
+    }
+    return (int)v;
+}
+
 static uint64_t parse_u64_arg(const char *s, const char *opt) {
     char *end = NULL;
     unsigned long long v = strtoull(s, &end, 10);
@@ -1602,6 +1614,10 @@ static eval_config parse_options(int argc, char **argv) {
             }
             c.ssd_streaming_cache_experts = experts;
             c.ssd_streaming_cache_bytes = bytes;
+        } else if (!strcmp(arg, "--ssd-streaming-full-layers")) {
+            int v = parse_nonnegative_int_arg(need_arg(&i, argc, argv, arg), arg);
+            c.ssd_streaming_full_layers = (uint32_t)v;
+            c.ssd_streaming_full_layers_set = true;
         } else if (!strcmp(arg, "--ssd-streaming-preload-experts")) {
             int v = parse_int_arg(need_arg(&i, argc, argv, arg), arg);
             if (v <= 0) {
@@ -3756,7 +3772,6 @@ static eval_run_result run_one_case(ds4_engine *engine, ds4_session *session,
     if (!tty && plain_in_think) plain_set_thinking_color(use_plain_color);
     tui_refresh(ui, "thinking");
 
-    const int eos = ds4_token_eos(engine);
     double t0 = ui->phase_start_sec;
     int forced_close_pos = -1;
     for (int i = 0; i < generation_limit; i++) {
@@ -3858,7 +3873,7 @@ static eval_run_result run_one_case(ds4_engine *engine, ds4_session *session,
         if (token < 0)
             token = ds4_session_sample(session, cfg->temperature, 0,
                                        cfg->top_p, cfg->min_p, rng);
-        if (token == eos) break;
+        if (ds4_token_is_stop(engine, token)) break;
         if (close_kind != EVAL_THINK_CLOSE_NONE &&
             think_close.kind == EVAL_THINK_CLOSE_NONE) {
             think_close.kind = close_kind;
@@ -4016,11 +4031,13 @@ static int parse_case_sequence(const char *arg, int ncases, int **seq_out, int *
 
 static void log_context_memory(ds4_backend backend,
                                int         ctx_size,
-                               uint32_t    prefill_chunk) {
+                               uint32_t    prefill_chunk,
+                               bool        ssd_streaming) {
     ds4_context_memory m =
-        ds4_context_memory_estimate_with_prefill(backend,
-                                                 ctx_size,
-                                                 prefill_chunk);
+        ds4_context_memory_estimate_with_prefill_mode(backend,
+                                                      ctx_size,
+                                                      prefill_chunk,
+                                                      ssd_streaming);
     fprintf(stderr,
             "ds4-eval: context buffers %.2f MiB (ctx=%d, backend=%s, prefill_chunk=%u, raw_kv_rows=%u, compressed_kv_rows=%u)\n",
             (double)m.total_bytes / (1024.0 * 1024.0),
@@ -4139,18 +4156,21 @@ int main(int argc, char **argv) {
         .mtp_path = cfg.mtp_path,
         .backend = cfg.backend,
         .n_threads = cfg.threads,
+        .context_size = cfg.ctx_size > 0 ? cfg.ctx_size : 0,
         .mtp_draft_tokens = 1,
         .mtp_margin = 3.0f,
         .power_percent = cfg.power_percent,
         .prefill_chunk = cfg.prefill_chunk,
         .ssd_streaming_cache_experts = cfg.ssd_streaming_cache_experts,
         .ssd_streaming_cache_bytes = cfg.ssd_streaming_cache_bytes,
+        .ssd_streaming_full_layers = cfg.ssd_streaming_full_layers,
         .ssd_streaming_preload_experts = cfg.ssd_streaming_preload_experts,
         .simulate_used_memory_bytes = cfg.simulate_used_memory_bytes,
         .warm_weights = cfg.warm_weights,
         .quality = cfg.quality,
         .ssd_streaming = cfg.ssd_streaming,
         .ssd_streaming_cold = cfg.ssd_streaming_cold,
+        .ssd_streaming_full_layers_set = cfg.ssd_streaming_full_layers_set,
         .distributed = cfg.dist,
     };
     char dist_err[256];
@@ -4190,7 +4210,10 @@ int main(int argc, char **argv) {
     fprintf(stderr, "ds4-eval: model shape %s\n", ds4_engine_model_name(engine));
     eval_warn_think_max_downgraded(&cfg);
     trace_write_header(trace, &cfg, ds4_engine_model_name(engine), ncases, max_prompt_tokens);
-    log_context_memory(cfg.backend, cfg.ctx_size, cfg.prefill_chunk);
+    log_context_memory(cfg.backend,
+                       cfg.ctx_size,
+                       cfg.prefill_chunk,
+                       cfg.ssd_streaming);
 
     ds4_session *session = NULL;
     if (ds4_session_create(&session, engine, cfg.ctx_size) != 0) {

ds4_gpu.h

@@ -39,7 +39,9 @@ int ds4_gpu_tensor_copy_f32_to_f16(ds4_gpu_tensor *dst, uint64_t dst_offset,
                                    uint64_t count);
 
 int ds4_gpu_begin_commands(void);
+int ds4_gpu_flush_encoder(void);
 int ds4_gpu_flush_commands(void);
+int ds4_gpu_commands_active(void);
 int ds4_gpu_signal_selected_readback_ready(uint64_t *event_value);
 int ds4_gpu_commit_and_wait_selected_readback(uint64_t event_value, const char *label);
 int ds4_gpu_wait_selected_readback_ready(uint64_t event_value, const char *label);
@@ -71,7 +73,9 @@ int ds4_gpu_preload_q4_expert_tables(const void *model_map, uint64_t model_size,
                                      uint32_t n_total_expert);
 int ds4_gpu_should_use_managed_kv_cache(uint64_t kv_cache_bytes, uint64_t context_bytes);
 void ds4_gpu_set_quality(bool quality);
+void ds4_gpu_set_glm_model(bool enabled);
 void ds4_gpu_set_ssd_streaming(bool enabled);
+void ds4_gpu_set_glm_streaming_prefill_full_layer(bool enabled);
 void ds4_gpu_set_streaming_expert_cache_budget(uint32_t experts);
 void ds4_gpu_set_streaming_expert_cache_expert_bytes(uint64_t bytes);
 uint64_t ds4_gpu_recommended_working_set_size(void);
@@ -103,6 +107,10 @@ int ds4_gpu_stream_expert_cache_begin_selected_load(
         const ds4_gpu_stream_expert_table *table,
         const int32_t                     *selected_ids,
         uint32_t                           n_selected);
+int ds4_gpu_glm_stream_expert_cache_begin_selected_load_tensor(
+        const ds4_gpu_stream_expert_table *table,
+        const ds4_gpu_tensor              *selected,
+        uint32_t                           n_selected);
 #if defined(DS4_ROCM_BUILD) || (!defined(DS4_NO_GPU) && !defined(__APPLE__))
 int ds4_gpu_stream_expert_cache_prepare_selected_batch(
         const ds4_gpu_stream_expert_table *table,
@@ -157,6 +165,25 @@ int ds4_gpu_embed_tokens_hc_tensor(
         uint32_t                n_embd,
         uint32_t                n_hc);
 
+int ds4_gpu_embed_token_q8_0_tensor(
+        ds4_gpu_tensor *out,
+        const void       *model_map,
+        uint64_t          model_size,
+        uint64_t          weight_offset,
+        uint32_t          n_vocab,
+        uint32_t          token,
+        uint32_t          n_embd);
+
+int ds4_gpu_embed_tokens_q8_0_tensor(
+        ds4_gpu_tensor       *out,
+        const ds4_gpu_tensor *tokens,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                weight_offset,
+        uint32_t                n_vocab,
+        uint32_t                n_tokens,
+        uint32_t                n_embd);
+
 int ds4_gpu_indexer_score_one_tensor(
         ds4_gpu_tensor       *scores,
         const ds4_gpu_tensor *q,
@@ -231,6 +258,36 @@ int ds4_gpu_matmul_q8_0_tensor(
         const ds4_gpu_tensor *x,
         uint64_t                n_tok);
 
+int ds4_gpu_matmul_q8_0_decode_mpp_tensor(
+        ds4_gpu_tensor       *out,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                weight_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        uint64_t                n_tok);
+
+int ds4_gpu_matmul_q8_0_decode_mpp_model_view_tensor(
+        ds4_gpu_tensor       *out,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                weight_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        uint64_t                n_tok);
+
+int ds4_gpu_matmul_q8_0_rows_scalar_tensor(
+        ds4_gpu_tensor       *out,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                weight_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        uint64_t                n_tok);
+
 /* Optional fused GPU operations.
  *
  * These are acceleration hooks, not required backend primitives.  A backend
@@ -275,6 +332,58 @@ int ds4_gpu_shared_gate_up_swiglu_q8_0_tensor(
         const ds4_gpu_tensor *x,
         float                   clamp);
 
+int ds4_gpu_shared_mid_swiglu_q8_0_tensor(
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        float                   clamp);
+
+int ds4_gpu_shared_gate_up_swiglu_q8_0_model_view_tensor(
+        ds4_gpu_tensor       *gate,
+        ds4_gpu_tensor       *up,
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        float                   clamp);
+
+int ds4_gpu_shared_gate_up_swiglu_q8_0_rows_tensor(
+        ds4_gpu_tensor       *gate,
+        ds4_gpu_tensor       *up,
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        uint64_t                n_tok,
+        float                   clamp);
+
+int ds4_gpu_shared_gate_up_swiglu_q8_0_rows_scalar_tensor(
+        ds4_gpu_tensor       *gate,
+        ds4_gpu_tensor       *up,
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                in_dim,
+        uint64_t                out_dim,
+        const ds4_gpu_tensor *x,
+        uint64_t                n_tok,
+        float                   clamp);
+
 int ds4_gpu_matmul_f16_tensor(
         ds4_gpu_tensor       *out,
         const void             *model_map,
@@ -345,6 +454,17 @@ int ds4_gpu_rms_norm_weight_rows_tensor(
         uint32_t                rows,
         float                   eps);
 
+int ds4_gpu_add_rms_norm_weight_tensor(
+        ds4_gpu_tensor       *norm_out,
+        ds4_gpu_tensor       *sum_out,
+        const ds4_gpu_tensor *a,
+        const ds4_gpu_tensor *b,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                weight_offset,
+        uint32_t                n,
+        float                   eps);
+
 int ds4_gpu_dsv4_qkv_rms_norm_rows_tensor(
         ds4_gpu_tensor       *q_out,
         const ds4_gpu_tensor *q,
@@ -434,6 +554,431 @@ int ds4_gpu_rope_tail_tensor(
         float             beta_fast,
         float             beta_slow);
 
+int ds4_gpu_glm_rope_tail_tensor(
+        ds4_gpu_tensor *x,
+        uint32_t        n_tokens,
+        uint32_t        n_head,
+        uint32_t        head_dim,
+        uint32_t        rot_dim,
+        uint32_t        pos0,
+        uint32_t        n_ctx_orig,
+        float           freq_base,
+        float           freq_scale,
+        float           ext_factor,
+        float           attn_factor,
+        float           beta_fast,
+        float           beta_slow);
+
+int ds4_gpu_glm_kv_lora_rms_norm_tensor(
+        ds4_gpu_tensor       *out,
+        const ds4_gpu_tensor *kv_raw,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              weight_offset,
+        uint32_t              n_tokens,
+        uint32_t              kv_raw_dim,
+        uint32_t              kv_lora_dim,
+        float                 eps);
+
+int ds4_gpu_glm_k_b_project_tensor(
+        ds4_gpu_tensor       *out,
+        const ds4_gpu_tensor *kv_norm,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              weight_offset,
+        uint32_t              n_tokens,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              n_head);
+
+int ds4_gpu_glm_store_compact_kv_tensor(
+        ds4_gpu_tensor       *kv_lora_cache,
+        ds4_gpu_tensor       *k_rope_cache,
+        const ds4_gpu_tensor *kv_norm,
+        const ds4_gpu_tensor *kv_raw,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_cap,
+        uint32_t              kv_raw_dim,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_rope,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_qkv_norm_store_compact_kv_tensor(
+        ds4_gpu_tensor       *q_out,
+        const ds4_gpu_tensor *q,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              q_weight_offset,
+        uint32_t              q_n,
+        ds4_gpu_tensor       *kv_lora_cache,
+        ds4_gpu_tensor       *k_rope_cache,
+        const ds4_gpu_tensor *kv_raw,
+        uint64_t              kv_weight_offset,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_cap,
+        uint32_t              kv_raw_dim,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_rope,
+        bool                  cache_f16,
+        float                 eps);
+
+int ds4_gpu_glm_store_indexer_k_tensor(
+        ds4_gpu_tensor       *indexer_key_cache,
+        const ds4_gpu_tensor *raw_k,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              weight_offset,
+        uint64_t              bias_offset,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_cap,
+        uint32_t              head_dim,
+        uint32_t              rot_dim,
+        uint32_t              n_ctx_orig,
+        float                 eps,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_build_kv_cache_tensor(
+        ds4_gpu_tensor       *key_cache,
+        ds4_gpu_tensor       *value_cache,
+        const ds4_gpu_tensor *kv_raw,
+        const ds4_gpu_tensor *k_nope,
+        const ds4_gpu_tensor *value,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_cap,
+        uint32_t              n_head,
+        uint32_t              kv_raw_dim,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              value_dim,
+        uint32_t              n_ctx_orig,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_build_kv_cache_flash_tensor(
+        ds4_gpu_tensor       *key_cache,
+        ds4_gpu_tensor       *value_cache,
+        const ds4_gpu_tensor *kv_raw,
+        const ds4_gpu_tensor *k_nope,
+        const ds4_gpu_tensor *value,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_cap,
+        uint32_t              n_head,
+        uint32_t              kv_raw_dim,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              value_dim,
+        uint32_t              n_ctx_orig,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_attention_full_tensor(
+        ds4_gpu_tensor       *heads,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *key_cache,
+        const ds4_gpu_tensor *value_cache,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_len,
+        uint32_t              cache_cap,
+        uint32_t              n_head,
+        uint32_t              qk_dim,
+        uint32_t              value_dim,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_fill_selected_range_tensor(
+        ds4_gpu_tensor *selected,
+        uint32_t        n_selected);
+
+int ds4_gpu_glm_fill_selected_range_batch_tensor(
+        ds4_gpu_tensor *selected,
+        uint32_t        n_tokens,
+        uint32_t        pos0,
+        uint32_t        n_selected,
+        uint32_t        pad_row);
+
+int ds4_gpu_glm_indexer_rope_tail_tensor(
+        ds4_gpu_tensor *x,
+        uint32_t        n_tokens,
+        uint32_t        n_head,
+        uint32_t        head_dim,
+        uint32_t        rot_dim,
+        uint32_t        pos0,
+        uint32_t        n_ctx_orig,
+        float           freq_base,
+        float           freq_scale,
+        float           ext_factor,
+        float           attn_factor,
+        float           beta_fast,
+        float           beta_slow);
+
+int ds4_gpu_glm_indexer_score_one_tensor(
+        ds4_gpu_tensor       *scores,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *weights,
+        const ds4_gpu_tensor *indexer_key_cache,
+        uint32_t              n_rows,
+        uint32_t              n_head,
+        uint32_t              head_dim,
+        float                 scale,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_indexer_scores_batch_tensor(
+        ds4_gpu_tensor       *scores,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *weights,
+        const ds4_gpu_tensor *indexer_key_cache,
+        uint32_t              n_rows,
+        uint32_t              n_tokens,
+        uint32_t              pos0,
+        uint32_t              n_head,
+        uint32_t              head_dim,
+        float                 scale,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_qk_lowrank_q8_0_tensor(
+        ds4_gpu_tensor       *qk_low,
+        const ds4_gpu_tensor *q,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              weight_offset,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_dim);
+
+int ds4_gpu_glm_qk_lowrank_q8_0_batch_tensor(
+        ds4_gpu_tensor       *qk_low,
+        const ds4_gpu_tensor *q,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              weight_offset,
+        uint32_t              n_tokens,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_dim);
+
+int ds4_gpu_glm_value_project_q8_0_batch_heads_tensor(
+        ds4_gpu_tensor       *heads,
+        const ds4_gpu_tensor *lora,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              weight_offset,
+        uint32_t              n_tokens,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              value_dim);
+
+int ds4_gpu_glm_attention_indexed_decode_tensor(
+        ds4_gpu_tensor       *heads,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *qk_low,
+        const ds4_gpu_tensor *kv_lora_cache,
+        const ds4_gpu_tensor *k_rope_cache,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              value_weight_offset,
+        const ds4_gpu_tensor *selected,
+        uint32_t              n_selected,
+        uint32_t              cache_cap,
+        bool                  cache_f16,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              value_dim,
+        uint32_t              n_ctx_orig,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow);
+
+int ds4_gpu_glm_attention_indexed_decode_split_group8_tensor(
+        ds4_gpu_tensor       *heads,
+        ds4_gpu_tensor       *partial_lora,
+        ds4_gpu_tensor       *partial_ms,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *qk_low,
+        const ds4_gpu_tensor *kv_lora_cache,
+        const ds4_gpu_tensor *k_rope_cache,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              value_weight_offset,
+        const ds4_gpu_tensor *selected,
+        uint32_t              n_selected,
+        bool                  selected_rows_valid,
+        uint32_t              cache_cap,
+        bool                  cache_f16,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              value_dim,
+        uint32_t              n_ctx_orig,
+        uint32_t              block_rows,
+        uint32_t              n_blocks,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow);
+
+int ds4_gpu_glm_attention_indexed_batch_tensor(
+        ds4_gpu_tensor       *heads,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *qk_low,
+        const ds4_gpu_tensor *kv_lora_cache,
+        const ds4_gpu_tensor *k_rope_cache,
+        const void           *model_map,
+        uint64_t              model_size,
+        uint64_t              value_weight_offset,
+        const ds4_gpu_tensor *selected,
+        uint32_t              n_tokens,
+        uint32_t              n_selected,
+        uint32_t              cache_cap,
+        bool                  cache_f16,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              value_dim,
+        uint32_t              n_ctx_orig,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow);
+
+int ds4_gpu_sort_i32_rows_asc_tensor(
+        ds4_gpu_tensor       *dst,
+        const ds4_gpu_tensor *src,
+        uint32_t              row_width,
+        uint32_t              n_rows);
+
+int ds4_gpu_glm_attention_indexed_batch_lora_tensor(
+        ds4_gpu_tensor       *lora_out,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *qk_low,
+        const ds4_gpu_tensor *kv_lora_cache,
+        const ds4_gpu_tensor *k_rope_cache,
+        const ds4_gpu_tensor *selected,
+        uint32_t              n_tokens,
+        uint32_t              n_selected,
+        uint32_t              cache_cap,
+        bool                  cache_f16,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              n_ctx_orig,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow);
+
+int ds4_gpu_glm_attention_indexed_batch_lora_causal_tensor(
+        ds4_gpu_tensor       *lora_out,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *qk_low,
+        const ds4_gpu_tensor *kv_lora_cache,
+        const ds4_gpu_tensor *k_rope_cache,
+        uint32_t              n_tokens,
+        uint32_t              pos0,
+        uint32_t              n_selected,
+        uint32_t              cache_cap,
+        bool                  cache_f16,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              n_ctx_orig,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow);
+
+int ds4_gpu_glm_attention_indexed_batch_lora_valid_tensor(
+        ds4_gpu_tensor       *lora_out,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *qk_low,
+        const ds4_gpu_tensor *kv_lora_cache,
+        const ds4_gpu_tensor *k_rope_cache,
+        const ds4_gpu_tensor *selected,
+        uint32_t              n_tokens,
+        uint32_t              n_selected,
+        uint32_t              cache_cap,
+        bool                  cache_f16,
+        uint32_t              n_head,
+        uint32_t              kv_lora_dim,
+        uint32_t              qk_nope,
+        uint32_t              qk_rope,
+        uint32_t              n_ctx_orig,
+        float                 freq_base,
+        float                 freq_scale,
+        float                 ext_factor,
+        float                 attn_factor,
+        float                 beta_fast,
+        float                 beta_slow);
+
+int ds4_gpu_glm_attention_flash_staged_tensor(
+        ds4_gpu_tensor       *heads,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *key_cache,
+        const ds4_gpu_tensor *value_cache,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_len,
+        uint32_t              cache_cap,
+        uint32_t              n_head,
+        uint32_t              qk_dim,
+        uint32_t              value_dim,
+        bool                  cache_f16);
+
+int ds4_gpu_glm_attention_flash_tensor(
+        ds4_gpu_tensor       *heads,
+        const ds4_gpu_tensor *q,
+        const ds4_gpu_tensor *key_cache,
+        const ds4_gpu_tensor *value_cache,
+        uint32_t              pos0,
+        uint32_t              n_tokens,
+        uint32_t              cache_len,
+        uint32_t              cache_cap,
+        uint32_t              n_head,
+        uint32_t              qk_dim,
+        uint32_t              value_dim,
+        bool                  cache_f16);
+
 /* Release decode fused KV finalizer: after the standalone RoPE kernel, this
  * performs DS4's FP8 non-RoPE KV round trip and writes the F16-rounded raw
  * attention cache row in one dispatch. */
@@ -762,6 +1307,13 @@ int ds4_gpu_add_tensor(
         const ds4_gpu_tensor *b,
         uint32_t                n);
 
+int ds4_gpu_add3_tensor(
+        ds4_gpu_tensor       *out,
+        const ds4_gpu_tensor *a,
+        const ds4_gpu_tensor *b,
+        const ds4_gpu_tensor *c,
+        uint32_t                n);
+
 int ds4_gpu_directional_steering_project_tensor(
         ds4_gpu_tensor       *x,
         const ds4_gpu_tensor *directions,
@@ -809,6 +1361,117 @@ int ds4_gpu_router_select_batch_tensor(
         float                   expert_weight_scale,
         uint32_t                n_tokens);
 
+int ds4_gpu_glm_router_select_tensor(
+        ds4_gpu_tensor       *selected,
+        ds4_gpu_tensor       *weights,
+        ds4_gpu_tensor       *probs,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                bias_offset,
+        const ds4_gpu_tensor *logits,
+        uint32_t                n_expert,
+        uint32_t                n_expert_used,
+        float                   expert_weight_scale);
+
+int ds4_gpu_glm_router_select_batch_tensor(
+        ds4_gpu_tensor       *selected,
+        ds4_gpu_tensor       *weights,
+        ds4_gpu_tensor       *probs,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                bias_offset,
+        const ds4_gpu_tensor *logits,
+        uint32_t                n_expert,
+        uint32_t                n_expert_used,
+        float                   expert_weight_scale,
+        uint32_t                n_tokens);
+
+int ds4_gpu_glm_routed_moe_one_tensor(
+        ds4_gpu_tensor       *out,
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                down_offset,
+        uint32_t                gate_type,
+        uint32_t                up_type,
+        uint32_t                down_type,
+        uint64_t                gate_expert_bytes,
+        uint64_t                gate_row_bytes,
+        uint64_t                up_expert_bytes,
+        uint64_t                up_row_bytes,
+        uint64_t                down_expert_bytes,
+        uint64_t                down_row_bytes,
+        uint32_t                expert_in_dim,
+        uint32_t                expert_mid_dim,
+        uint32_t                out_dim,
+        const ds4_gpu_tensor *selected,
+        const ds4_gpu_tensor *weights,
+        uint32_t                n_total_expert,
+        uint32_t                n_expert,
+        uint32_t                layer_index,
+        const ds4_gpu_tensor *x,
+        bool                    force_resident);
+
+int ds4_gpu_glm_routed_moe_batch_tensor(
+        ds4_gpu_tensor       *out,
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                down_offset,
+        uint32_t                gate_type,
+        uint32_t                up_type,
+        uint32_t                down_type,
+        uint64_t                gate_expert_bytes,
+        uint64_t                gate_row_bytes,
+        uint64_t                up_expert_bytes,
+        uint64_t                up_row_bytes,
+        uint64_t                down_expert_bytes,
+        uint64_t                down_row_bytes,
+        uint32_t                expert_in_dim,
+        uint32_t                expert_mid_dim,
+        uint32_t                out_dim,
+        const ds4_gpu_tensor *selected,
+        const ds4_gpu_tensor *weights,
+        uint32_t                n_total_expert,
+        uint32_t                n_expert,
+        uint32_t                layer_index,
+        const ds4_gpu_tensor *x,
+        uint32_t                n_tokens,
+        uint32_t                mid_token_stride);
+
+int ds4_gpu_glm_routed_moe_batch_direct_scalar_q4_tensor(
+        ds4_gpu_tensor       *out,
+        ds4_gpu_tensor       *mid,
+        const void             *model_map,
+        uint64_t                model_size,
+        uint64_t                gate_offset,
+        uint64_t                up_offset,
+        uint64_t                down_offset,
+        uint32_t                gate_type,
+        uint32_t                up_type,
+        uint32_t                down_type,
+        uint64_t                gate_expert_bytes,
+        uint64_t                gate_row_bytes,
+        uint64_t                up_expert_bytes,
+        uint64_t                up_row_bytes,
+        uint64_t                down_expert_bytes,
+        uint64_t                down_row_bytes,
+        uint32_t                expert_in_dim,
+        uint32_t                expert_mid_dim,
+        uint32_t                out_dim,
+        const ds4_gpu_tensor *selected,
+        const ds4_gpu_tensor *weights,
+        uint32_t                n_total_expert,
+        uint32_t                n_expert,
+        uint32_t                layer_index,
+        const ds4_gpu_tensor *x,
+        uint32_t                n_tokens,
+        uint32_t                mid_token_stride);
+
 int ds4_gpu_routed_moe_set_selected_override(const int32_t *selected, uint32_t n_selected);
 
 int ds4_gpu_routed_moe_one_tensor(
@@ -837,7 +1500,8 @@ int ds4_gpu_routed_moe_one_tensor(
         uint32_t                n_expert,
         float                   clamp,
         const ds4_gpu_tensor *x,
-        uint32_t                layer_index);
+        uint32_t                layer_index,
+        bool                    force_resident);
 
 int ds4_gpu_routed_moe_batch_tensor(
         ds4_gpu_tensor       *out,

ds4_help.c

@@ -164,8 +164,9 @@ static void print_model_runtime(FILE *fp, const help_colors *c,
     opt(fp, c, "--power N", "GPU duty-cycle target, 1..100. Default: 100");
     opt(fp, c, "--ssd-streaming", "Metal/CUDA/ROCm: opt in to SSD-backed model streaming instead of full residency.");
     opt(fp, c, "--ssd-streaming-cold", "SSD streaming: skip default popularity-based expert-cache preload.");
-    opt(fp, c, "--ssd-streaming-cache-experts N|NGB", "SSD streaming: routed expert cache as expert count or GiB, e.g. 32GB. Metal/ROCm default: 80% working set minus non-routed weights; CUDA default: backend fixed cache.");
-    opt(fp, c, "--ssd-streaming-preload-experts N", "SSD streaming: upfront popularity preload count. Default: auto hot seed capped at 4096; use --ssd-streaming-cold to skip.");
+    opt(fp, c, "--ssd-streaming-cache-experts N|NGB", "SSD streaming: N is an exact dynamic expert count; NGB is a routed memory budget that also reserves two full prefill layers. Auto: 80% working set minus non-routed weights; GLM Metal caps lower.");
+    opt(fp, c, "--ssd-streaming-full-layers N", "GLM Metal streaming: keep the first N routed layers fully resident. Default: auto from NGB expert budget; use 0 to disable.");
+    opt(fp, c, "--ssd-streaming-preload-experts N", "SSD streaming: upfront popularity preload count. DeepSeek auto-seeds by default; GLM demand-fills unless N is explicit.");
     opt(fp, c, "--simulate-used-memory NGB", "Diagnostic: lock N GiB before model load to simulate a smaller-memory machine.");
     opt(fp, c, "--prefill-chunk N", "Metal graph prefill chunk size. Default: auto (PRO long prompts use 8192; others use 4096).");
     if (full) {
@@ -192,6 +193,7 @@ static void print_sampling(FILE *fp, const help_colors *c, bool full) {
     opt(fp, c, "--top-p F", "Nucleus sampling probability.");
     opt(fp, c, "--min-p F", "Keep tokens scoring at least F times the top token.");
     opt(fp, c, "--seed N", "Sampling seed for reproducible non-greedy runs.");
+    para(fp, c, "GLM CLI and agent runs default to temperature 1.0, top-p 0.95, and min-p 0 unless those options are set explicitly.");
     opt(fp, c, "--think", "Use normal thinking mode.");
     opt(fp, c, "--think-max", "Use Think Max when context is large enough.");
     opt(fp, c, "--nothink", "Disable thinking and ask for direct replies.");
@@ -199,6 +201,7 @@ static void print_sampling(FILE *fp, const help_colors *c, bool full) {
         opt(fp, c, "-sys, --system TEXT", "System prompt. Empty string disables the default where supported.");
         opt(fp, c, "-p, --prompt TEXT", "One-shot prompt text.");
         opt(fp, c, "--prompt-file FILE", "Read one-shot prompt text from FILE.");
+        opt(fp, c, "--raw-prompt", "Tokenize the one-shot prompt without chat markers.");
     }
     fputc('\n', fp);
 }
@@ -248,6 +251,7 @@ static void print_cli_diagnostics(FILE *fp, const help_colors *c) {
     opt(fp, c, "--dump-logits FILE", "Write full next-token logits as JSON.");
     opt(fp, c, "--dump-logprobs FILE", "Write greedy continuation top-logprobs as JSON.");
     opt(fp, c, "--logprobs-top-k N", "Alternatives stored by --dump-logprobs. Default: 20");
+    opt(fp, c, "--decode-consistency N", "Compare N-token decode logits with a fresh full prefill.");
     opt(fp, c, "--expert-profile FILE", "Metal-only: write routed expert locality/cache simulation JSON.");
     opt(fp, c, "--perplexity-file FILE", "Score raw text with teacher-forced NLL.");
     opt(fp, c, "--imatrix-dataset FILE", "Rendered prompt dataset for imatrix collection.");
@@ -278,6 +282,7 @@ static void print_agent_specific(FILE *fp, const help_colors *c) {
     title(fp, c, "Agent Options");
     opt(fp, c, "-p, --prompt TEXT", "Submit an initial prompt after startup.");
     opt(fp, c, "--non-interactive", "Run without TUI. With -p: one turn; without -p: repeated stdin prompts.");
+    opt(fp, c, "--raw-prompt", "Non-interactive -p only: tokenize prompt without agent chat/tool text.");
     opt(fp, c, "-sys, --system TEXT", "Extra system prompt. Empty disables extra text.");
     opt(fp, c, "--trace FILE", "Write prompt, token, and DSML debug trace.");
     opt(fp, c, "--chdir DIR", "Change working directory before loading runtime assets.");
@@ -513,7 +518,7 @@ static void print_topic(FILE *fp, const help_colors *c, ds4_help_tool tool, cons
     else if (tool == DS4_HELP_AGENT && streq(topic, "tools")) {
         title(fp, c, "Agent Tool System");
         para(fp, c, "The agent can read, search, write, edit, run bash, and browse through Chrome-backed web tools.");
-        para(fp, c, "Tool calls are emitted by the model as DSML and rendered live in the terminal.");
+        para(fp, c, "DeepSeek-family models emit DSML tool calls; GLM models use native <tool_call> syntax. Both are rendered live in the terminal.");
         para(fp, c, "Edit uses exact old/new replacement; [upto] can bridge a unique head and tail for large anchored edits.");
         fputc('\n', fp);
     } else if (tool == DS4_HELP_BENCH && streq(topic, "benchmark")) print_bench_specific(fp, c);

gguf-tools/.gitignore

@@ -1,3 +1,11 @@
 deepseek4-quantize
 quality-testing/score_official
-quality-testing/data/
+quality-testing/score_llama
+quality-testing/data/*
+!quality-testing/data/README.md
+!quality-testing/data/glm52-openrouter-100/
+!quality-testing/data/glm52-openrouter-100/**
+!quality-testing/data/pro/
+!quality-testing/data/pro/**
+!quality-testing/data/flash/
+!quality-testing/data/flash/**

gguf-tools/Makefile

@@ -1,12 +1,21 @@
 CC ?= cc
+CXX ?= c++
 UNAME_S := $(shell uname -s)
+LLAMA_CPP_DIR ?= ../../deepseek-v4-quants/llama.cpp
+LLAMA_CPP_BUILD ?= $(LLAMA_CPP_DIR)/build-remote
+LLAMA_CPP_BIN ?= $(LLAMA_CPP_BUILD)/bin
+LLAMA_CPP_CXXFLAGS ?= -O3 -Wall -Wextra -std=c++17 \
+	-I$(LLAMA_CPP_DIR)/include -I$(LLAMA_CPP_DIR)/ggml/include
 
 ifeq ($(UNAME_S),Darwin)
 NATIVE_CPU_FLAG ?= -mcpu=native
 QUALITY_LDLIBS := -lm -pthread -framework Foundation -framework Metal
-QUALITY_TARGETS := ds4.o ds4_metal.o
-QUALITY_OBJS := ../ds4.o ../ds4_metal.o
+QUALITY_TARGETS := ds4.o ds4_metal.o ds4_distributed.o ds4_ssd.o rax.o
+QUALITY_OBJS := ../ds4.o ../ds4_metal.o ../ds4_distributed.o ../ds4_ssd.o ../rax.o
 QUALITY_LINK := $(CC) $(CFLAGS) -I.. -o
+LLAMA_CPP_LDLIBS ?= -L$(LLAMA_CPP_BIN) -Wl,-rpath,$(abspath $(LLAMA_CPP_BIN)) \
+	-lllama -lggml -lggml-base -lggml-cpu -lggml-blas -lggml-metal \
+	-framework Foundation -framework Metal -framework Accelerate
 else
 NATIVE_CPU_FLAG ?= -march=native
 CUDA_HOME ?= /usr/local/cuda
@@ -17,15 +26,17 @@ NVCC_ARCH_FLAGS := -arch=$(CUDA_ARCH)
 endif
 NVCCFLAGS ?= -O3 --use_fast_math $(NVCC_ARCH_FLAGS) -Xcompiler $(NATIVE_CPU_FLAG) -Xcompiler -pthread
 QUALITY_LDLIBS ?= -lm -Xcompiler -pthread -L$(CUDA_HOME)/targets/sbsa-linux/lib -L$(CUDA_HOME)/lib64 -lcudart -lcublas
-QUALITY_TARGETS := ds4.o ds4_cuda.o
-QUALITY_OBJS := ../ds4.o ../ds4_cuda.o
+QUALITY_TARGETS := ds4.o ds4_cuda.o ds4_distributed.o ds4_ssd.o rax.o
+QUALITY_OBJS := ../ds4.o ../ds4_cuda.o ../ds4_distributed.o ../ds4_ssd.o ../rax.o
 QUALITY_LINK := $(NVCC) $(NVCCFLAGS) -I.. -o
+LLAMA_CPP_LDLIBS ?= -L$(LLAMA_CPP_BIN) -Wl,-rpath,$(abspath $(LLAMA_CPP_BIN)) \
+	-lllama -lggml -lggml-base -lggml-cpu -ldl -lpthread
 endif
 
 CFLAGS ?= -O3 -Wall -Wextra -std=c11 $(NATIVE_CPU_FLAG)
 CPPFLAGS ?= -D_GNU_SOURCE
 
-.PHONY: all clean quality-score
+.PHONY: all clean quality-score quality-llama-score
 
 all: deepseek4-quantize
 
@@ -36,5 +47,8 @@ quality-score:
 	$(MAKE) -C .. $(QUALITY_TARGETS)
 	$(QUALITY_LINK) quality-testing/score_official quality-testing/score_official.c $(QUALITY_OBJS) $(QUALITY_LDLIBS)
 
+quality-llama-score:
+	$(CXX) $(LLAMA_CPP_CXXFLAGS) -o quality-testing/score_llama quality-testing/score_llama.cpp $(LLAMA_CPP_LDLIBS)
+
 clean:
-	rm -f deepseek4-quantize quality-testing/score_official
+	rm -f deepseek4-quantize quality-testing/score_official quality-testing/score_llama

gguf-tools/quality-testing/README.md

@@ -1,14 +1,33 @@
 # Official-Continuation Quality Testing
 
-This directory contains the 100 prompts and scripts used to compare local GGUF
-variants against official DeepSeek V4 continuations.
+This directory contains the prompts, tracked official fixtures, and scripts used
+to compare local GGUF variants against hosted-model continuations.
 
 The metric is target-token negative log likelihood: collect a deterministic
 official continuation, then ask each local GGUF how much probability it assigns
 to that exact continuation token by token.  This avoids judging quality from one
 sampled answer.
 
-## 1. Collect Official Continuations
+## 1. Tracked Fixture Sets
+
+Curated fixtures are kept in the repository so release QA can run without
+calling hosted APIs:
+
+- `data/glm52-openrouter-100`: 100 GLM 5.2 continuations collected through
+  OpenRouter `z-ai/glm-5.2` with `top_logprobs=20`.
+- `data/flash`: 100 DeepSeek V4 Flash continuations collected from the official
+  DeepSeek API with `top_logprobs=20`.
+- `data/pro`: 100 DeepSeek V4 PRO continuations collected from the official
+  DeepSeek API with `top_logprobs=20`.
+
+DeepSeek V4 Flash also has tracked official smoke vectors in
+`tests/test-vectors/`.  Those vectors drive `./ds4_test --logprob-vectors` and
+include short prompts plus long-prompt attention cases.
+
+The hosted APIs expose output-token logprobs and top-logprob alternatives, not
+full vocabulary logits.
+
+## 2. Collect Official Continuations
 
 ```sh
 export DEEPSEEK_API_KEY=...
@@ -19,6 +38,26 @@ python3 gguf-tools/quality-testing/collect_official.py \
   --max-tokens 24
 ```
 
+For GLM 5.2 through OpenRouter:
+
+```sh
+export OPENROUTER_API_KEY=...
+python3 gguf-tools/quality-testing/collect_official.py \
+  --model z-ai/glm-5.2 \
+  --endpoint https://openrouter.ai/api/v1/chat/completions \
+  --api-key-env OPENROUTER_API_KEY \
+  --prompts gguf-tools/quality-testing/prompts.jsonl \
+  --out gguf-tools/quality-testing/data/glm52-openrouter \
+  --count 100 \
+  --max-tokens 24 \
+  --top-logprobs 20 \
+  --token-limit-field max_tokens \
+  --provider-order parasail/fp8 \
+  --require-parameters \
+  --thinking omit \
+  --reasoning-effort none
+```
+
 Use one output directory per official model.  The default model is Flash, so
 `data/flash` is the recommended path for Flash continuations.  For PRO:
 
@@ -39,10 +78,11 @@ The script writes:
 - `data/<model>/responses/case_*.json`
 - `data/<model>/manifest.tsv`
 
-The prompt list is tracked in `prompts.jsonl`; the official responses are not
-tracked because they are derived from an external API.
+The prompt list is tracked in `prompts.jsonl`.  Curated fixture directories are
+also tracked after review; ad-hoc API collection directories should stay
+untracked until they are intentionally promoted into the release QA set.
 
-## 2. Build The Local Scorer
+## 3. Build The Local Scorer
 
 ```sh
 make -C gguf-tools quality-score
@@ -50,27 +90,41 @@ make -C gguf-tools quality-score
 
 The scorer links against the DS4 runtime and uses Metal by default.
 
-## 3. Score GGUF Variants
+## 4. Score GGUF Variants
 
 ```sh
 gguf-tools/quality-testing/score_official \
   ../deepseek-v4-quants/gguf/OLD.gguf \
-  gguf-tools/quality-testing/data/flash/manifest.tsv \
+  gguf-tools/quality-testing/data/pro/manifest.tsv \
   /tmp/old.tsv \
   4096
 
 gguf-tools/quality-testing/score_official \
   ../deepseek-v4-quants/gguf/NEW.gguf \
-  gguf-tools/quality-testing/data/flash/manifest.tsv \
+  gguf-tools/quality-testing/data/pro/manifest.tsv \
   /tmp/new.tsv \
   4096
 ```
 
-Use `data/pro/manifest.tsv` for PRO GGUFs.  The scorer and comparator do not
-care which model produced the manifest; the manifest path selects the
-continuation set.
+Use `data/flash/manifest.tsv` for Flash GGUFs,
+`data/glm52-openrouter-100/manifest.tsv` for GLM 5.2 GGUFs, and
+`data/pro/manifest.tsv` for PRO GGUFs.  The scorer and comparator do not care
+which model produced the manifest; the manifest path selects the continuation
+set.
 
-## 4. Compare
+For a full-residency vs SSD-streaming comparison, score the same model twice and
+add the streaming flags to one run:
+
+```sh
+gguf-tools/quality-testing/score_official \
+  /path/to/model.gguf \
+  gguf-tools/quality-testing/data/glm52-openrouter-100/manifest.tsv \
+  /tmp/streaming.tsv \
+  4096 \
+  --ssd-streaming
+```
+
+## 5. Compare
 
 ```sh
 python3 gguf-tools/quality-testing/compare_scores.py /tmp/old.tsv /tmp/new.tsv
@@ -86,3 +140,13 @@ Output fields:
   official first token.
 - `avg_greedy_lcp`: average greedy longest common prefix against the official
   continuation.
+- `api_target_mae`: when the manifest includes `response_file`, absolute
+  local-vs-API logprob delta for aligned official output tokens.
+- `api_top_coverage`: fraction of API top-logprob alternatives that map exactly
+  to one local tokenizer token.
+- `api_top1_rate`: how often the API top alternative equals the local greedy
+  token.
+- `api_topn_recall`: fraction of mapped API top-N alternatives found in the
+  local top-N for the same position.
+- `api_top_mae`: local-vs-API logprob MAE over mapped API top alternatives.
+- `api_pair_rate`: pairwise ordering agreement among mapped API alternatives.

gguf-tools/quality-testing/collect_official.py

@@ -1,5 +1,5 @@
 #!/usr/bin/env python3
-"""Collect official DeepSeek V4 continuations for local quant scoring."""
+"""Collect hosted-model continuations for local quant scoring."""
 
 from __future__ import annotations
 
@@ -129,18 +129,31 @@ def request_one(
     max_tokens: int,
     top_logprobs: int,
     thinking: str,
+    reasoning_effort: str,
+    token_limit_field: str,
+    provider_order: list[str],
+    provider_allow_fallbacks: bool,
+    provider_require_parameters: bool,
 ) -> dict:
     payload = {
         "model": model,
         "messages": [{"role": "user", "content": prompt}],
         "temperature": 0,
-        "max_tokens": max_tokens,
         "logprobs": True,
         "top_logprobs": top_logprobs,
         "stream": False,
     }
+    payload[token_limit_field] = max_tokens
     if thinking != "omit":
         payload["thinking"] = {"type": thinking}
+    if reasoning_effort != "omit":
+        payload["reasoning"] = {"effort": reasoning_effort}
+    if provider_order or provider_require_parameters:
+        provider = {"require_parameters": provider_require_parameters}
+        if provider_order:
+            provider["order"] = provider_order
+            provider["allow_fallbacks"] = provider_allow_fallbacks
+        payload["provider"] = provider
     req = urllib.request.Request(
         endpoint,
         data=json.dumps(payload).encode("utf-8"),
@@ -162,11 +175,29 @@ def fetch_with_retry(
     max_tokens: int,
     top_logprobs: int,
     thinking: str,
+    reasoning_effort: str,
+    token_limit_field: str,
+    provider_order: list[str],
+    provider_allow_fallbacks: bool,
+    provider_require_parameters: bool,
 ) -> dict:
     delay = 1.0
     for attempt in range(6):
         try:
-            return request_one(api_key, endpoint, model, prompt, max_tokens, top_logprobs, thinking)
+            return request_one(
+                api_key,
+                endpoint,
+                model,
+                prompt,
+                max_tokens,
+                top_logprobs,
+                thinking,
+                reasoning_effort,
+                token_limit_field,
+                provider_order,
+                provider_allow_fallbacks,
+                provider_require_parameters,
+            )
         except urllib.error.HTTPError as e:
             body = e.read().decode("utf-8", "replace")
             if e.code < 500 and e.code != 429:
@@ -187,17 +218,44 @@ def main() -> int:
     ap.add_argument("--prompts", default="gguf-tools/quality-testing/prompts.jsonl")
     ap.add_argument("--model", default=MODEL)
     ap.add_argument("--endpoint", default=ENDPOINT)
+    ap.add_argument("--api-key-env", default=None)
     ap.add_argument("--count", type=int, default=100)
     ap.add_argument("--max-tokens", type=int, default=24)
     ap.add_argument("--top-logprobs", type=int, default=5)
     ap.add_argument("--thinking", choices=("disabled", "enabled", "omit"), default="disabled")
+    ap.add_argument("--reasoning-effort",
+                    choices=("xhigh", "high", "medium", "low", "minimal", "none", "omit"),
+                    default=None)
+    ap.add_argument("--token-limit-field",
+                    choices=("auto", "max_tokens", "max_completion_tokens"),
+                    default="auto")
+    ap.add_argument("--provider-order",
+                    help="comma-separated OpenRouter provider slugs, for example ambient/fp8")
+    ap.add_argument("--allow-provider-fallbacks", action="store_true")
+    ap.add_argument("--require-parameters", action="store_true",
+                    help="for OpenRouter, route only to endpoints advertising the requested parameters")
     args = ap.parse_args()
     if args.top_logprobs < 0 or args.top_logprobs > 20:
         raise SystemExit("--top-logprobs must be between 0 and 20")
 
-    api_key = os.environ.get("DEEPSEEK_API_KEY")
+    openrouter = "openrouter.ai" in args.endpoint
+    api_key_env = args.api_key_env or ("OPENROUTER_API_KEY" if openrouter else "DEEPSEEK_API_KEY")
+    api_key = os.environ.get(api_key_env)
     if not api_key:
-        raise SystemExit("DEEPSEEK_API_KEY is not set")
+        raise SystemExit(f"{api_key_env} is not set")
+    thinking = args.thinking
+    reasoning_effort = args.reasoning_effort
+    if reasoning_effort is None:
+        reasoning_effort = "none" if openrouter else "omit"
+    if openrouter and args.thinking == "disabled":
+        thinking = "omit"
+    token_limit_field = args.token_limit_field
+    if token_limit_field == "auto":
+        token_limit_field = "max_tokens"
+    provider_order = []
+    if args.provider_order:
+        provider_order = [item.strip() for item in args.provider_order.split(",") if item.strip()]
+    provider_require_parameters = args.require_parameters
     prompts = load_prompts(Path(args.prompts))
 
     out = Path(args.out)
@@ -209,6 +267,12 @@ def main() -> int:
     rows = []
     total = min(args.count, len(prompts))
     print(f"model={args.model} endpoint={args.endpoint}", file=sys.stderr)
+    print(f"key_env={api_key_env} token_field={token_limit_field} thinking={thinking} reasoning={reasoning_effort}",
+          file=sys.stderr)
+    if provider_order or provider_require_parameters:
+        print(f"provider_order={','.join(provider_order) or '-'} "
+              f"provider_fallbacks={args.allow_provider_fallbacks} "
+              f"require_parameters={provider_require_parameters}", file=sys.stderr)
     for i, prompt in enumerate(prompts[: args.count]):
         case_id = f"case_{i:03d}"
         print(f"official {i + 1}/{total}: {case_id}", file=sys.stderr, flush=True)
@@ -219,12 +283,20 @@ def main() -> int:
             prompt,
             args.max_tokens,
             args.top_logprobs,
-            args.thinking,
+            thinking,
+            reasoning_effort,
+            token_limit_field,
+            provider_order,
+            args.allow_provider_fallbacks,
+            provider_require_parameters,
         )
         choice = response["choices"][0]
         content = choice.get("message", {}).get("content", "")
         if not content:
             print(f"warning: empty continuation for {case_id}", file=sys.stderr)
+        logprob_items = (choice.get("logprobs") or {}).get("content", []) or []
+        if args.top_logprobs > 0 and not logprob_items:
+            raise RuntimeError(f"{case_id}: response did not include output-token logprobs")
 
         prompt_path = out / "prompts" / f"{case_id}.txt"
         cont_path = out / "continuations" / f"{case_id}.txt"

gguf-tools/quality-testing/compare_scores.py

@@ -8,20 +8,89 @@ import sys
 from pathlib import Path
 
 
-def load(path: Path) -> dict[str, dict[str, float]]:
+def parse_cell(value: str) -> float | int:
+    v = float(value)
+    if v.is_integer():
+        return int(v)
+    return v
+
+
+def load(path: Path) -> dict[str, dict[str, float | int]]:
     with path.open(newline="", encoding="utf-8") as fp:
         rows = {}
         for row in csv.DictReader(fp, delimiter="\t"):
-            rows[row["id"]] = {
+            parsed = {
                 "target_tokens": int(row["target_tokens"]),
                 "nll": float(row["nll"]),
                 "avg_nll": float(row["avg_nll"]),
                 "first_match": int(row["first_match"]),
                 "greedy_lcp": int(row["greedy_lcp"]),
             }
+            for key, value in row.items():
+                if key is None or key == "id" or key in parsed or value is None or value == "":
+                    continue
+                parsed[key] = parse_cell(value)
+            rows[row["id"]] = parsed
         return rows
 
 
+def has_api(rows: dict[str, dict[str, float | int]], ids: list[str]) -> bool:
+    return all(
+        "api_target_tokens" in rows[case_id] and "api_top_items" in rows[case_id]
+        for case_id in ids
+    )
+
+
+def sum_key(rows: dict[str, dict[str, float | int]], ids: list[str], key: str) -> float:
+    return sum(float(rows[case_id].get(key, 0.0)) for case_id in ids)
+
+
+def weighted_avg(
+    rows: dict[str, dict[str, float | int]], ids: list[str], value_key: str, count_key: str
+) -> float:
+    count = sum_key(rows, ids, count_key)
+    if count == 0:
+        return 0.0
+    total = sum(
+        float(rows[case_id].get(value_key, 0.0)) * float(rows[case_id].get(count_key, 0.0))
+        for case_id in ids
+    )
+    return total / count
+
+
+def ratio(num: float, den: float) -> float:
+    return num / den if den else 0.0
+
+
+def print_api_summary(label: str, rows: dict[str, dict[str, float | int]], ids: list[str]) -> None:
+    target_tokens = sum_key(rows, ids, "api_target_tokens")
+    top_items = sum_key(rows, ids, "api_top_items")
+    top_mapped = sum_key(rows, ids, "api_top_mapped")
+    top1_count = sum_key(rows, ids, "api_top1_count")
+    top1_match = sum_key(rows, ids, "api_top1_match")
+    topn_ref = sum_key(rows, ids, "api_topn_ref")
+    topn_hit = sum_key(rows, ids, "api_topn_hit")
+    top_logprob_count = sum_key(rows, ids, "api_top_logprob_count")
+    pair_total = sum_key(rows, ids, "api_pair_total")
+    pair_agree = sum_key(rows, ids, "api_pair_agree")
+
+    print(f"{label}_api_target_tokens\t{int(target_tokens)}")
+    print(f"{label}_api_target_mae\t{weighted_avg(rows, ids, 'api_target_mae', 'api_target_tokens'):.9f}")
+    print(
+        f"{label}_api_target_mean_delta\t"
+        f"{weighted_avg(rows, ids, 'api_target_mean_delta', 'api_target_tokens'):.9f}"
+    )
+    print(f"{label}_api_top_coverage\t{ratio(top_mapped, top_items):.9f}")
+    print(f"{label}_api_top1_rate\t{ratio(top1_match, top1_count):.9f}")
+    print(f"{label}_api_topn_recall\t{ratio(topn_hit, topn_ref):.9f}")
+    print(f"{label}_api_top_mae\t{weighted_avg(rows, ids, 'api_top_mae', 'api_top_logprob_count'):.9f}")
+    print(
+        f"{label}_api_top_mean_delta\t"
+        f"{weighted_avg(rows, ids, 'api_top_mean_delta', 'api_top_logprob_count'):.9f}"
+    )
+    print(f"{label}_api_pair_rate\t{ratio(pair_agree, pair_total):.9f}")
+
+
 def main() -> int:
     if len(sys.argv) != 3:
         print(f"usage: {sys.argv[0]} OLD.tsv NEW.tsv", file=sys.stderr)
@@ -73,6 +142,9 @@ def main() -> int:
     print(f"case_wins_new_old_ties\t{new_case_wins}\t{old_case_wins}\t{ties}")
     print(f"first_token_matches_old_new\t{old_first}\t{new_first}")
     print(f"avg_greedy_lcp_old_new\t{old_lcp / len(ids):.3f}\t{new_lcp / len(ids):.3f}")
+    if has_api(old, ids) and has_api(new, ids):
+        print_api_summary("old", old, ids)
+        print_api_summary("new", new, ids)
 
     print("\nnew best cases:")
     for delta, case_id, t, old_avg, new_avg in sorted(deltas)[:8]:

gguf-tools/quality-testing/data/README.md

@@ -0,0 +1,21 @@
+# Official Quality Fixtures
+
+This directory contains curated hosted-model continuation fixtures that are
+safe to commit and use in release QA.
+
+- `glm52-openrouter-100`: 100 GLM 5.2 OpenRouter continuations with API
+  top-logprob slices.
+- `flash`: 100 DeepSeek V4 Flash official continuations with API top-logprob
+  slices.
+- `pro`: 100 DeepSeek V4 PRO official continuations with API top-logprob
+  slices.
+
+Each fixture directory contains:
+
+- `prompts/case_*.txt`: exact user prompts.
+- `continuations/case_*.txt`: deterministic hosted-model continuations.
+- `responses/case_*.json`: raw hosted responses, including logprob slices.
+- `manifest.tsv`: paths consumed by `score_official`.
+
+DeepSeek V4 Flash smoke vectors are also tracked in `tests/test-vectors/` and
+are run by `./ds4_test --logprob-vectors`.

gguf-tools/quality-testing/data/flash/continuations/case_000.txt

@@ -0,0 +1 @@
+Here is a detailed explanation of B-tree insertion, covering the standard algorithm, node splits, and the special case of splitting

gguf-tools/quality-testing/data/flash/continuations/case_001.txt

@@ -0,0 +1,5 @@
+Here's a concise design for a TCP echo server that handles slow clients:
+
+## Architecture
+
+**Event-driven, non-block

gguf-tools/quality-testing/data/flash/continuations/case_002.txt

@@ -0,0 +1 @@
+Here's a comprehensive comparison of memory-mapped (mmap) model weights versus copying all weights into private buffers on macOS

gguf-tools/quality-testing/data/flash/continuations/case_003.txt

@@ -0,0 +1 @@
+This is an excellent question because it gets to the heart of why **RMSNorm** (Root Mean Square Normalization)

gguf-tools/quality-testing/data/flash/continuations/case_004.txt

@@ -0,0 +1 @@
+This is an excellent question that gets at a core optimization challenge in modern LLM deployment. Let's break down why **

gguf-tools/quality-testing/data/flash/continuations/case_005.txt

@@ -0,0 +1,5 @@
+Here's pseudocode for a binary heap push (insert) operation:
+
+```
+function push(heap, value):
+    //

gguf-tools/quality-testing/data/flash/continuations/case_006.txt

@@ -0,0 +1 @@
+This is a great question that gets at the core of how modern LLMs generate text. Both **temperature** and **

gguf-tools/quality-testing/data/flash/continuations/case_007.txt

@@ -0,0 +1 @@
+Here is a summary of how speculative decoding preserves exact greedy output (i.e., the same tokens that would be generated by

gguf-tools/quality-testing/data/flash/continuations/case_008.txt

@@ -0,0 +1 @@
+This is an excellent question that gets to the heart of how Large Language Models (LLMs) are served efficiently. The

gguf-tools/quality-testing/data/flash/continuations/case_009.txt

@@ -0,0 +1,5 @@
+Here are three invariants for a special-token table used in a tokenizer:
+
+---
+
+### 1. **Un

gguf-tools/quality-testing/data/flash/continuations/case_010.txt

@@ -0,0 +1 @@
+Ecco una spiegazione dettagliata del funzionamento dell'inserimento in un **B-tree**, con

gguf-tools/quality-testing/data/flash/continuations/case_011.txt

@@ -0,0 +1,3 @@
+Ecco un progetto conciso per un server TCP echo con gestione di client lenti:
+
+## Architettura

gguf-tools/quality-testing/data/flash/continuations/case_012.txt

@@ -0,0 +1 @@
+Ecco un confronto dettagliato tra **mmap** (memory-mapped file) e **copia

gguf-tools/quality-testing/data/flash/continuations/case_013.txt

@@ -0,0 +1 @@
+Certamente. Ecco una derivazione dettagliata del perché **RMSNorm** (Root Mean Square Normalization

gguf-tools/quality-testing/data/flash/continuations/case_014.txt

@@ -0,0 +1 @@
+La cache KV (Key-Value) su disco è particolarmente utile nelle **sessioni agentiche** (d

gguf-tools/quality-testing/data/flash/continuations/case_015.txt

@@ -0,0 +1 @@
+Ecco lo pseudocodice per inserire un elemento in un heap binario (min-heap o max-he

gguf-tools/quality-testing/data/flash/continuations/case_016.txt

@@ -0,0 +1 @@
+Questa è un'ottima domanda che tocca il cuore del controllo della generazione nei modelli linguistici

gguf-tools/quality-testing/data/flash/continuations/case_017.txt

@@ -0,0 +1 @@
+La decodifica speculativa mantiene l'output esatto perché utilizza un **meccanismo di ver

gguf-tools/quality-testing/data/flash/continuations/case_018.txt

@@ -0,0 +1 @@
+Questa è un'ottima domanda che tocca un aspetto fondamentale dell'ottimizzazione dell'inferenza

gguf-tools/quality-testing/data/flash/continuations/case_019.txt

@@ -0,0 +1 @@
+Ecco tre invarianti per una tabella di token speciali (ad esempio in un parser o in un sistema di

gguf-tools/quality-testing/data/flash/continuations/case_020.txt

@@ -0,0 +1,3 @@
+## Binary Search for Insertion Point
+
+Binary search finds the insertion point for a target value in a sorted array by maintaining

gguf-tools/quality-testing/data/flash/continuations/case_021.txt

@@ -0,0 +1,4 @@
+Here's a short C function that clamps an integer to a range:
+
+```c
+int clamp(int value, int

gguf-tools/quality-testing/data/flash/continuations/case_022.txt

@@ -0,0 +1 @@
+This is a fundamental question in concurrent programming. The core difference lies in **what they protect** and **how they coordinate

gguf-tools/quality-testing/data/flash/continuations/case_023.txt

@@ -0,0 +1 @@
+In the context of a network server, **backpressure** is a mechanism for **flow control**. It refers to the

gguf-tools/quality-testing/data/flash/continuations/case_024.txt

@@ -0,0 +1 @@
+Here is a safe and robust format for storing a model checkpoint header, designed for clarity, extensibility, and error prevention

gguf-tools/quality-testing/data/flash/continuations/case_025.txt

@@ -0,0 +1 @@
+A **ring buffer** (also called a circular buffer) is a fixed-size data structure that uses a contiguous block of

gguf-tools/quality-testing/data/flash/continuations/case_026.txt

@@ -0,0 +1,4 @@
+Here is a brief, structured plan for testing long-context prompt chunking.
+
+### Goal
+Validate that chunking a

gguf-tools/quality-testing/data/flash/continuations/case_027.txt

@@ -0,0 +1 @@
+This is an excellent question that gets to the heart of why modern low-bit quantization (like 4-bit and 3

gguf-tools/quality-testing/data/flash/continuations/case_028.txt

@@ -0,0 +1 @@
+This is a great question that gets at a key performance optimization in modern large language model (LLM) inference, particularly

gguf-tools/quality-testing/data/flash/continuations/case_029.txt

@@ -0,0 +1 @@
+This is an excellent question because comparing quantizations is rarely a simple "A is better than B" scenario. The "

gguf-tools/quality-testing/data/flash/continuations/case_030.txt

@@ -0,0 +1,3 @@
+## Ricerca Binaria del Punto di Inserimento
+
+La ricerca binaria del punto di inserimento è una

gguf-tools/quality-testing/data/flash/continuations/case_031.txt

@@ -0,0 +1,4 @@
+Ecco una breve funzione C che limita un intero a un intervallo specificato:
+
+```c
+int clamp

gguf-tools/quality-testing/data/flash/continuations/case_032.txt

@@ -0,0 +1 @@
+Certamente. La differenza fondamentale tra un **mutex** (mutual exclusion) e un **contatore atom

gguf-tools/quality-testing/data/flash/continuations/case_033.txt

@@ -0,0 +1 @@
+In un server di rete, **backpressure** (contro-pressione o contropressione) è un me

gguf-tools/quality-testing/data/flash/continuations/case_034.txt

@@ -0,0 +1 @@
+Ecco un formato sicuro per salvare l'header di un checkpoint modello, progettato per garantire integ

gguf-tools/quality-testing/data/flash/continuations/case_035.txt

@@ -0,0 +1 @@
+Ti spiegherò il funzionamento di un **ring buffer** (o buffer circolare) e le tec

gguf-tools/quality-testing/data/flash/continuations/case_036.txt

@@ -0,0 +1 @@
+Ecco un piano breve e strutturato per testare il **chunking** (frammentazione) di

gguf-tools/quality-testing/data/flash/continuations/case_037.txt

@@ -0,0 +1 @@
+Nella quantizzazione low-bit (a bassa precisione) dei modelli di deep learning, una **matrice

gguf-tools/quality-testing/data/flash/continuations/case_038.txt

@@ -0,0 +1 @@
+Questa è un'ottima domanda che tocca un aspetto cruciale dell'ottimizzazione dei modelli

gguf-tools/quality-testing/data/flash/continuations/case_039.txt

@@ -0,0 +1 @@
+Per confrontare due quantizzazioni (ad esempio, due modelli di compressione di un segnale, due

gguf-tools/quality-testing/data/flash/continuations/case_040.txt

@@ -0,0 +1,3 @@
+Here are the first five measurements to take when diagnosing generation slowdowns at long context:
+
+1.  **Time-to-F

gguf-tools/quality-testing/data/flash/continuations/case_041.txt

@@ -0,0 +1 @@
+This is an excellent and fundamental question about how language models actually "think." The short answer is: **because the model

gguf-tools/quality-testing/data/flash/continuations/case_042.txt

@@ -0,0 +1,3 @@
+Let's break down **online softmax** for attention using simple variables.
+
+## The Problem: Standard Softmax in Attention

gguf-tools/quality-testing/data/flash/continuations/case_043.txt

@@ -0,0 +1,5 @@
+Here's a minimal JSON schema for a tool call with name and arguments:
+
+```json
+{
+  "type": "

gguf-tools/quality-testing/data/flash/continuations/case_044.txt

@@ -0,0 +1 @@
+To test that a file-backed mmap is not repeatedly remapped, you need to verify that the kernel is reusing

gguf-tools/quality-testing/data/flash/continuations/case_045.txt

@@ -0,0 +1 @@
+This is a fascinating question that gets to the heart of how modern graphics APIs like Metal, Vulkan, and DirectX

gguf-tools/quality-testing/data/flash/continuations/case_046.txt

@@ -0,0 +1 @@
+In a Mixture-of-Experts (MoE) layer, the **router** (also called a **gating

gguf-tools/quality-testing/data/flash/continuations/case_047.txt

@@ -0,0 +1 @@
+Here is the difference between **raw KV rows** and **compressed KV rows** in the context of storage engines (

gguf-tools/quality-testing/data/flash/continuations/case_048.txt

@@ -0,0 +1 @@
+Here is a comprehensive checklist for validating that a quantized model still follows instructions effectively. This checklist is designed to catch both

gguf-tools/quality-testing/data/flash/continuations/case_049.txt

@@ -0,0 +1 @@
+This is an excellent question that gets to the heart of how to properly test and understand a language model's capabilities. Evaluating

gguf-tools/quality-testing/data/flash/continuations/case_050.txt

@@ -0,0 +1 @@
+Ecco cinque misure iniziali per affrontare un problema di decoding lento in contesti lunghi (long

gguf-tools/quality-testing/data/flash/continuations/case_051.txt

@@ -0,0 +1 @@
+Questa è un'ottima domanda che tocca il cuore di come funzionano i modelli linguistici

gguf-tools/quality-testing/data/flash/continuations/case_052.txt

@@ -0,0 +1 @@
+Certamente! Spiegherò la **softmax** nel contesto del meccanismo di **attenzione**

gguf-tools/quality-testing/data/flash/continuations/case_053.txt

@@ -0,0 +1,5 @@
+Ecco uno schema JSON minimo per una tool call:
+
+```json
+{
+  "name": "nome_della

gguf-tools/quality-testing/data/flash/continuations/case_054.txt

@@ -0,0 +1 @@
+Per testare che un `mmap` su file non venga rimappato ripetutamente, puoi

gguf-tools/quality-testing/data/flash/continuations/case_055.txt

@@ -0,0 +1 @@
+Questa è una domanda molto interessante che tocca il cuore dell'ottimizzazione della grafica e della gest

gguf-tools/quality-testing/data/flash/continuations/case_056.txt

@@ -0,0 +1 @@
+In un **layer mixture-of-experts (MoE)**, il **router** (o gate network) svolge

gguf-tools/quality-testing/data/flash/continuations/case_057.txt

@@ -0,0 +1 @@
+Certamente. Ecco una descrizione chiara della differenza tra righe KV (Key-Value) raw e

gguf-tools/quality-testing/data/flash/continuations/case_058.txt

@@ -0,0 +1 @@
+Ecco una checklist strutturata per validare che un modello quantizzato (es. GPTQ, AWQ,

gguf-tools/quality-testing/data/flash/continuations/case_059.txt

@@ -0,0 +1 @@
+Valutare 50 prompt invece di un singolo prompt preferito è più utile per diversi motivi, specialmente

gguf-tools/quality-testing/data/flash/continuations/case_060.txt

@@ -0,0 +1,4 @@
+Here's a tiny Python function that returns the median of three numbers:
+
+```python
+def median_of_three(a,

gguf-tools/quality-testing/data/flash/continuations/case_061.txt

@@ -0,0 +1,3 @@
+This is a great question because it gets at the core of why we organize data.
+
+The reason sorted arrays make membership tests

gguf-tools/quality-testing/data/flash/continuations/case_062.txt

@@ -0,0 +1 @@
+**Eventual consistency** is a consistency model used in distributed databases (and other distributed systems) where, given enough time

gguf-tools/quality-testing/data/flash/continuations/case_063.txt

@@ -0,0 +1,3 @@
+Here's a classic example of a race condition in C where two threads increment a shared counter without synchronization:
+
+```c

gguf-tools/quality-testing/data/flash/continuations/case_064.txt

@@ -0,0 +1,2 @@
+This is a fundamental concept in computer science, networking, and system design. The simplest analogy is a **water pipe**.
+

gguf-tools/quality-testing/data/flash/continuations/case_065.txt

@@ -0,0 +1 @@
+A **trie** (also called a prefix tree) represents a set of strings by organizing them based on their **

gguf-tools/quality-testing/data/flash/continuations/case_066.txt

@@ -0,0 +1 @@
+Testing a command-line parser thoroughly requires covering both normal usage and edge cases that could break assumptions. Here's a systematic approach

gguf-tools/quality-testing/data/flash/continuations/case_067.txt

@@ -0,0 +1 @@
+This is an excellent question that gets to the heart of how virtual memory works. The short answer is that **`mm

gguf-tools/quality-testing/data/flash/continuations/case_068.txt

@@ -0,0 +1 @@
+This is a nuanced and insightful question that touches on the intersection of model compression (quantization) and the dynamics of Mi

gguf-tools/quality-testing/data/flash/continuations/case_069.txt

@@ -0,0 +1 @@
+Here is a short, focused checklist for reviewing a PR that involves memory lifetimes (e.g., Rust, C++, or

gguf-tools/quality-testing/data/flash/continuations/case_070.txt

@@ -0,0 +1,4 @@
+Ecco una semplice funzione Python per calcolare la mediana di tre numeri:
+
+```python
+def mediana(a

gguf-tools/quality-testing/data/flash/continuations/case_071.txt

@@ -0,0 +1 @@
+Questa è una domanda fondamentale in informatica. La risposta breve è: **perché permettono di us

gguf-tools/quality-testing/data/flash/continuations/case_072.txt

@@ -0,0 +1 @@
+In un database distribuito, la **consistenza eventuale** (o *eventual consistency*) è un modello

gguf-tools/quality-testing/data/flash/continuations/case_073.txt

@@ -0,0 +1,3 @@
+Ecco un breve esempio di race condition in C con due thread che condividono una variabile:
+
+```c

gguf-tools/quality-testing/data/flash/continuations/case_074.txt

@@ -0,0 +1 @@
+Certamente! Latenza e throughput sono due metriche fondamentali per misurare le prestazioni di un sistema (

gguf-tools/quality-testing/data/flash/continuations/case_075.txt

@@ -0,0 +1 @@
+Un **trie** (chiamato anche albero digitale o albero di prefissi) rappresenta un