Add GLM 5.2 inference and quality fixtures
This commit is contained in:
1 parent
63d9874c54
commit
005afedcd9
967 files changed
+1267733
-4988
No files matched your search
@@ -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
|
||||
@@ -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)
|
||||
@@ -73,6 +74,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)
|
||||
@@ -136,6 +140,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)
|
||||
@@ -277,4 +284,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
|
||||
+49
-11
@@ -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.
|
||||
|
||||
@@ -99,6 +135,7 @@ tiny routed-MoE verifier kernels, or shared `--mtp` support-model code changes:
|
||||
record the selected-expert footprint or stage timing in the DSpark log.
|
||||
|
||||
## 4. Metal PRO Path
|
||||
## 5. Metal PRO Path
|
||||
|
||||
PRO support is experimental, but release builds must not break it silently.
|
||||
|
||||
@@ -109,7 +146,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.
|
||||
|
||||
@@ -128,7 +165,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
|
||||
@@ -146,7 +183,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
|
||||
@@ -167,7 +204,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
|
||||
@@ -185,7 +222,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.
|
||||
|
||||
@@ -199,7 +236,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.
|
||||
@@ -212,7 +249,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.
|
||||
|
||||
@@ -244,7 +281,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.
|
||||
@@ -253,7 +290,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.
|
||||
@@ -262,11 +299,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
|
||||
|
||||
@@ -163,6 +163,19 @@ target-only continuation path; set `DS4_DSPARK_STRICT=1` for the same behavior
|
||||
without enabling all quality-mode kernels. Plain sampled and non-speculative
|
||||
session eval also skips DSpark draft preparation; DSpark speculation is
|
||||
currently a greedy argmax-only path.
|
||||
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:
|
||||
|
||||
@@ -231,14 +244,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
|
||||
|
||||
@@ -270,15 +288,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 \
|
||||
@@ -509,6 +523,77 @@ encryption or authentication, and is not release-stable yet; coordinator and
|
||||
workers should be built from the same commit and used on trusted machines and
|
||||
trusted networks.
|
||||
|
||||
## Tensor Parallelism
|
||||
|
||||
Tensor parallelism runs a single decode across two Macs connected with a
|
||||
Thunderbolt 5 cable, splitting the heavy per-layer work between the two
|
||||
GPUs and exchanging 16-24KB partial sums at synchronization gates inside the
|
||||
graph (RDMA over Thunderbolt when available, a dedicated TCP socket
|
||||
otherwise). Unlike the pipelined distributed mode above, both
|
||||
machines work on the *same token at the same time*, so it reduces
|
||||
per-token latency instead of just fitting a bigger model.
|
||||
|
||||
Each machine keeps one contiguous half of the routed experts resident. Dense,
|
||||
attention, shared-expert, embedding, and output weights remain replicated.
|
||||
This lets a model whose routed experts do not fit on one machine run fully
|
||||
resident across the pair; routed kernels never touch the peer's expert half.
|
||||
|
||||
### Running GLM 5.2 across two 128 GB MacBooks
|
||||
|
||||
One-time setup per boot, on **both** machines:
|
||||
|
||||
```
|
||||
# Let the GPU wire ~117 GB (default cap is ~75% of RAM; the resident
|
||||
# expert shard needs ~97.5 GiB plus KV/scratch).
|
||||
sudo sysctl iogpu.wired_limit_mb=120000
|
||||
|
||||
# RDMA over Thunderbolt needs an IPv4 address directly on the cabled
|
||||
# member interface (the bridge IP does not count). Use the interface
|
||||
# that is 'active' in ifconfig, e.g. en1 on one side and en6 on the
|
||||
# other. Skip this if you are fine with the TCP fallback.
|
||||
sudo ifconfig en1 inet 10.99.0.2/30 alias # machine A
|
||||
sudo ifconfig en6 inet 10.99.0.1/30 alias # machine B
|
||||
```
|
||||
|
||||
Both machines need the same tree, the same commit, and the same GGUF
|
||||
path. Start the worker first (machine B, dialing the coordinator's
|
||||
Thunderbolt bridge address), then the coordinator (machine A):
|
||||
|
||||
```
|
||||
# Machine B: worker. Computes its expert half in lockstep; never
|
||||
# tokenizes or samples.
|
||||
./ds4 -m models/GLM-5.2-UD-IQ2_XXS_RoutedIQ2XXS_blk78Q2K-DenseQ4.gguf \
|
||||
--tp-worker 172.31.255.2 9911
|
||||
|
||||
# Machine A: coordinator. Owns the prompt, sampling, and output.
|
||||
./ds4 -m models/GLM-5.2-UD-IQ2_XXS_RoutedIQ2XXS_blk78Q2K-DenseQ4.gguf \
|
||||
-c 8192 --tp-coordinator 9911 -p "Tell me something about the sea."
|
||||
```
|
||||
|
||||
Startup takes about 9 seconds per machine: each rank pre-faults its
|
||||
~100 GiB shard from SSD and pins it through a Metal residency set.
|
||||
DeepSeek V4 Flash works the same way with its own GGUF on both machines.
|
||||
DeepSeek gate vectors are 16 KB and ride as one RDMA message. GLM's
|
||||
6144-wide 24 KB vectors are split into two ordered RDMA messages.
|
||||
|
||||
Measured on two M5 Max 128 GB MacBooks (GLM 5.2, IQ2_XXS, 188 GiB):
|
||||
|
||||
| | two Macs, tensor parallel | one Mac, SSD streaming |
|
||||
|---|---|---|
|
||||
| decode | ~16.8 t/s (15.4 at 4k context) | ~4.8 t/s |
|
||||
| prefill (4096 tokens) | ~94 t/s | ~3-5 t/s |
|
||||
| residency | fully memory-resident | streams experts from SSD |
|
||||
|
||||
Notes: the coordinator mirrors every prompt sync and eval to the worker, so
|
||||
both KV caches stay in lockstep; prompt processing splits both the
|
||||
routed-expert GEMMs (by expert ownership) and the attention heads (a
|
||||
contiguous half per machine) with one bulk partial-sum exchange per
|
||||
layer per stage (`DS4_GLM_TP_TOKEN_PREFILL=1` selects a slower
|
||||
token-by-token prefill that exactly matches the single-machine arithmetic).
|
||||
The split graph is deterministic, but its changed floating-point reduction
|
||||
order is not generally byte-identical to single-machine execution. Design
|
||||
notes and the full bring-up log live in `misc/METAL_TENSOR_PARALLELISM.md`.
|
||||
|
||||
## Reducing heat, power usage and fan noise
|
||||
|
||||
Long local inference runs can keep the GPU busy for extended periods. If you
|
||||
|
||||
+83
-6
@@ -3,6 +3,10 @@ set -e
|
||||
|
||||
REPO="antirez/deepseek-v4-gguf"
|
||||
DSPARK_REPO="deepseek-ai/DeepSeek-V4-Flash-DSpark"
|
||||
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"
|
||||
@@ -11,6 +15,12 @@ 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"
|
||||
DSPARK_SUPPORT_FILE="DeepSeek-V4-Flash-DSpark-support.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"}
|
||||
@@ -27,7 +37,7 @@ TOKEN=${HF_TOKEN:-}
|
||||
|
||||
usage() {
|
||||
cat <<EOF
|
||||
DeepSeek V4 GGUF downloader
|
||||
DwarfStar GGUF downloader
|
||||
|
||||
Usage:
|
||||
./download_model.sh q2-imatrix [--token TOKEN]
|
||||
@@ -41,6 +51,10 @@ Usage:
|
||||
./download_model.sh dspark-source [--token TOKEN]
|
||||
./download_model.sh dspark-support-dry-run
|
||||
./download_model.sh dspark-support
|
||||
./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:
|
||||
|
||||
@@ -93,6 +107,21 @@ Targets:
|
||||
Converts the local DSpark source checkpoint into a standalone support
|
||||
GGUF for ds4 --mtp. Output is written under DS4_GGUF_DIR and is about
|
||||
6 GB. Requires the full DSpark source payloads.
|
||||
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
|
||||
@@ -120,6 +149,9 @@ After building DSpark support, enable it explicitly in greedy mode:
|
||||
|
||||
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
|
||||
}
|
||||
|
||||
@@ -135,6 +167,8 @@ LINK_MODEL=1
|
||||
DSPARK_SOURCE=0
|
||||
DSPARK_SUPPORT_DRY_RUN=0
|
||||
DSPARK_SUPPORT=0
|
||||
FORCE_HF_DOWNLOAD=0
|
||||
FLATTEN_DOWNLOADS=0
|
||||
|
||||
case "$MODEL" in
|
||||
q2-imatrix) MODEL_FILE=$Q2_IMATRIX_FILE ;;
|
||||
@@ -159,6 +193,30 @@ case "$MODEL" in
|
||||
dspark-support)
|
||||
DSPARK_SUPPORT=1
|
||||
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
|
||||
@@ -195,6 +253,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
|
||||
@@ -219,12 +280,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"
|
||||
@@ -240,7 +311,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
|
||||
@@ -257,6 +328,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
|
||||
@@ -376,7 +452,8 @@ build_dspark_support() {
|
||||
|
||||
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"
|
||||
@@ -386,7 +463,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
|
||||
|
||||
@@ -91,11 +91,12 @@ typedef struct {
|
||||
bool debug;
|
||||
} ds4_distributed_options;
|
||||
|
||||
/* Tensor parallelism: two identical machines run the full model in lockstep
|
||||
* and split the heavy per-layer matvecs, exchanging 16KB partial sums at two
|
||||
* gates per layer (see misc/METAL_TENSOR_PARALLELISM.md). The leader owns
|
||||
* prompt/sampling/output and listens; the worker dials in and mirrors every
|
||||
* session sync/eval. */
|
||||
/* Tensor parallelism: two identical machines run the model in lockstep and
|
||||
* split the heavy per-layer matvecs, exchanging partial sums at gates inside
|
||||
* the graph (see misc/METAL_TENSOR_PARALLELISM.md). Each rank keeps one
|
||||
* contiguous half of the routed experts resident; dense and shared weights
|
||||
* remain replicated. The leader owns prompt/sampling and listens; the worker
|
||||
* dials in and mirrors every session sync/eval. */
|
||||
typedef enum {
|
||||
DS4_TP_NONE = 0,
|
||||
DS4_TP_LEADER,
|
||||
@@ -123,6 +124,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;
|
||||
@@ -133,13 +135,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;
|
||||
@@ -180,18 +186,34 @@ int ds4_engine_power(ds4_engine *e);
|
||||
int ds4_engine_set_power(ds4_engine *e, int power_percent);
|
||||
const char *ds4_engine_model_name(ds4_engine *e);
|
||||
int ds4_engine_layer_count(ds4_engine *e);
|
||||
/* Decode gate schedule for the TP transport; see ds4_tp_identity. */
|
||||
void ds4_engine_tp_gate_schedule(ds4_engine *e,
|
||||
uint32_t *start,
|
||||
uint32_t *step,
|
||||
uint32_t *per_token);
|
||||
uint32_t ds4_engine_layer_compress_ratio(ds4_engine *e, uint32_t layer);
|
||||
uint64_t ds4_engine_hidden_f32_values(ds4_engine *e);
|
||||
int ds4_engine_embd_dim(ds4_engine *e);
|
||||
uint64_t ds4_engine_model_bytes(ds4_engine *e);
|
||||
int ds4_engine_tp_vocab_split(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
|
||||
@@ -201,6 +223,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,
|
||||
@@ -244,6 +271,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);
|
||||
|
||||
@@ -298,13 +330,16 @@ int ds4_session_top_logprobs(ds4_session *s, ds4_token_score *out, int k);
|
||||
int ds4_session_token_logprob(ds4_session *s, int token, ds4_token_score *out);
|
||||
int ds4_session_copy_logits(ds4_session *s, float *out, int cap);
|
||||
int ds4_session_set_logits(ds4_session *s, const float *logits, int n);
|
||||
/* Pay the one-time first-submission GPU cost outside any measured window;
|
||||
* used by the TP worker right after session create (no-op on CPU/GLM). */
|
||||
void ds4_session_gpu_warmup(ds4_session *s);
|
||||
int ds4_session_eval(ds4_session *s, int token, char *err, size_t errlen);
|
||||
int ds4_session_eval_speculative_argmax(ds4_session *s, int first_token,
|
||||
int max_tokens, int eos_token,
|
||||
int *accepted, int accepted_cap,
|
||||
char *err, size_t errlen);
|
||||
/* TP worker side of a mirrored speculative-verify block: run the replicated
|
||||
* batch verify for its KV side effects, then obey the leader's commit frame
|
||||
/* TP worker side of a mirrored speculative-verify block: run its half of the
|
||||
* batch verify for KV side effects, then obey the leader's commit frame
|
||||
* (keep, or roll back and replay). Only called from ds4_tp_worker_run. */
|
||||
int ds4_session_tp_spec_cycle(ds4_session *s, const int *drafts, int draft_n,
|
||||
char *err, size_t errlen);
|
||||
|
||||
+978
-146
File diff suppressed because it is too large.
Load diff
+95
-23
@@ -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;
|
||||
|
||||
@@ -6,7 +6,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
|
||||
@@ -29,16 +29,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;
|
||||
@@ -143,6 +148,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);
|
||||
@@ -192,19 +207,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) {
|
||||
@@ -245,7 +281,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;
|
||||
|
||||
@@ -282,8 +319,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 {
|
||||
@@ -411,7 +462,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,
|
||||
@@ -419,6 +472,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) {
|
||||
@@ -429,6 +492,9 @@ static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, con
|
||||
ds4_session_free(session);
|
||||
return 1;
|
||||
}
|
||||
/* Pay the one-time first-submission GPU cost before the prefill timer
|
||||
* starts (matches the TP worker's startup warmup). */
|
||||
ds4_session_gpu_warmup(session);
|
||||
|
||||
char err[160];
|
||||
ds4_think_mode think_mode = cli_effective_think_mode(&cfg->gen);
|
||||
@@ -477,7 +543,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;
|
||||
@@ -499,6 +565,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);
|
||||
@@ -507,13 +583,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;
|
||||
}
|
||||
@@ -767,7 +842,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);
|
||||
@@ -788,6 +863,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};
|
||||
@@ -901,6 +1146,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 ||
|
||||
@@ -920,36 +1170,39 @@ 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->engine.tp.role == DS4_TP_LEADER ||
|
||||
getenv("DS4_CLI_FORCE_SESSION") != NULL ||
|
||||
cfg->gen.temperature > 0.0f ||
|
||||
ds4_engine_mtp_draft_tokens(engine) > 1) {
|
||||
/* TP leaders always drive the session path: the sync/eval mirroring
|
||||
* that keeps the worker in lockstep lives there. The env override
|
||||
* exists so TP-vs-single-node validation compares like with like. */
|
||||
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->engine.tp.role == DS4_TP_LEADER ||
|
||||
getenv("DS4_CLI_FORCE_SESSION") != NULL ||
|
||||
cfg->gen.temperature > 0.0f ||
|
||||
ds4_engine_mtp_draft_tokens(engine) > 1) {
|
||||
/* TP leaders always drive the session path: the sync/eval
|
||||
* mirroring that keeps the worker in lockstep lives there.
|
||||
* The env override exists so TP-vs-single-node validation
|
||||
* compares like with like. */
|
||||
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);
|
||||
@@ -995,7 +1248,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) {
|
||||
@@ -1025,23 +1279,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) {
|
||||
@@ -1059,8 +1339,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);
|
||||
}
|
||||
@@ -1081,6 +1361,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
|
||||
@@ -1093,7 +1377,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);
|
||||
@@ -1153,7 +1437,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;
|
||||
@@ -1174,20 +1458,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;
|
||||
}
|
||||
@@ -1209,7 +1502,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));
|
||||
}
|
||||
|
||||
@@ -1266,18 +1559,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);
|
||||
@@ -1302,15 +1596,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")) {
|
||||
@@ -1474,6 +1769,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")) {
|
||||
@@ -1488,10 +1785,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")) {
|
||||
@@ -1511,6 +1811,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) {
|
||||
@@ -1571,6 +1875,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")) {
|
||||
@@ -1675,12 +1981,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.engine.tp.role == DS4_TP_WORKER) {
|
||||
int rc = ds4_tp_worker_run(engine, &cfg.engine.tp);
|
||||
ds4_engine_close(engine);
|
||||
@@ -1700,6 +2010,10 @@ int main(int argc, char **argv) {
|
||||
.quant_bits = (uint32_t)ds4_engine_routed_quant_bits(engine),
|
||||
.ctx_size = (uint32_t)cfg.gen.ctx_size,
|
||||
};
|
||||
ds4_engine_tp_gate_schedule(engine,
|
||||
&tp_id.gate_slot_start,
|
||||
&tp_id.gate_slot_step,
|
||||
&tp_id.gates_per_token);
|
||||
if (!ds4_tp_create(&tp_leader, &cfg.engine.tp, &tp_id, tp_err, sizeof(tp_err)) ||
|
||||
!ds4_engine_tp_bind(engine, tp_leader, tp_err, sizeof(tp_err))) {
|
||||
fprintf(stderr, "ds4: %s\n", tp_err);
|
||||
@@ -1732,9 +2046,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;
|
||||
|
||||
+89
@@ -2515,6 +2515,7 @@ extern "C" int ds4_gpu_pack_slot_rows_f32_tensor(
|
||||
|
||||
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");
|
||||
@@ -2791,6 +2792,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;
|
||||
@@ -2803,6 +2808,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;
|
||||
@@ -7838,6 +7847,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,
|
||||
@@ -8878,6 +8901,11 @@ extern "C" int ds4_gpu_attention_prefill_raw_heads_tensor(ds4_gpu_tensor *heads,
|
||||
n_tokens, window, n_head, head_dim);
|
||||
return cuda_ok(cudaGetLastError(), "attention_prefill_raw launch");
|
||||
}
|
||||
extern "C" int ds4_gpu_attention_prefill_raw_heads_range_tensor(ds4_gpu_tensor *heads, const void *model_map, uint64_t model_size, uint64_t sinks_offset, const ds4_gpu_tensor *q, const ds4_gpu_tensor *raw_kv, uint32_t q_row0, uint32_t n_q, uint32_t n_kv, uint32_t window, uint32_t n_head, uint32_t head_dim) {
|
||||
/* Rectangular problems only arise from the Metal-only TP row split. */
|
||||
if (q_row0 != 0 || n_q != n_kv) return 0;
|
||||
return ds4_gpu_attention_prefill_raw_heads_tensor(heads, model_map, model_size, sinks_offset, q, raw_kv, n_q, window, n_head, head_dim);
|
||||
}
|
||||
static int attention_decode_batch_launch(
|
||||
ds4_gpu_tensor *heads,
|
||||
const void *model_map,
|
||||
@@ -9298,6 +9326,11 @@ extern "C" int ds4_gpu_attention_prefill_static_mixed_heads_tensor(
|
||||
q, raw_kv, comp_kv, NULL, 0, n_tokens,
|
||||
n_comp, window, ratio, n_head, head_dim);
|
||||
}
|
||||
extern "C" int ds4_gpu_attention_prefill_static_mixed_heads_range_tensor(ds4_gpu_tensor *heads, const void *model_map, uint64_t model_size, uint64_t sinks_offset, const ds4_gpu_tensor *q, const ds4_gpu_tensor *raw_kv, const ds4_gpu_tensor *comp_kv, uint32_t comp_kv_f16, uint32_t q_row0, uint32_t n_q, uint32_t n_tokens, uint32_t n_comp, uint32_t window, uint32_t ratio, uint32_t n_head, uint32_t head_dim) {
|
||||
/* Rectangular problems only arise from the Metal-only TP row split. */
|
||||
if (q_row0 != 0 || n_q != n_tokens) return 0;
|
||||
return ds4_gpu_attention_prefill_static_mixed_heads_tensor(heads, model_map, model_size, sinks_offset, q, raw_kv, comp_kv, comp_kv_f16, n_tokens, n_comp, window, ratio, n_head, head_dim);
|
||||
}
|
||||
|
||||
extern "C" int ds4_gpu_attention_prefill_masked_mixed_heads_tensor(
|
||||
ds4_gpu_tensor *heads,
|
||||
@@ -9574,6 +9607,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) ||
|
||||
@@ -13349,10 +13402,46 @@ extern "C" void ds4_gpu_tp_set_batch_exchange(ds4_gpu_tp_batch_exchange_fn fn) {
|
||||
(void)fn;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_suspend_expert_sharding(int suspend) {
|
||||
(void)suspend;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_keepalive_pause(int paused) {
|
||||
(void)paused;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_set_attn_head_split(int enabled) {
|
||||
(void)enabled;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_model_residency_skip(int skip) {
|
||||
(void)skip;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_set_big_exchange(ds4_gpu_tp_big_exchange_fn fn) {
|
||||
(void)fn;
|
||||
}
|
||||
|
||||
extern "C" int ds4_gpu_tp_big_gate_encode(uint32_t layer, uint32_t rows,
|
||||
const ds4_gpu_tensor *out_t,
|
||||
ds4_gpu_tensor *in_t,
|
||||
uint64_t bytes) {
|
||||
(void)layer; (void)rows; (void)out_t; (void)in_t; (void)bytes;
|
||||
return 0;
|
||||
}
|
||||
extern "C" uint64_t ds4_gpu_tp_big_gate_kick(uint32_t layer, uint32_t rows,
|
||||
const ds4_gpu_tensor *out_t,
|
||||
ds4_gpu_tensor *in_t,
|
||||
uint64_t bytes) {
|
||||
(void)layer; (void)rows; (void)out_t; (void)in_t; (void)bytes;
|
||||
return 0;
|
||||
}
|
||||
|
||||
extern "C" int ds4_gpu_tp_big_gate_wait(uint64_t seq) {
|
||||
(void)seq;
|
||||
return 0;
|
||||
}
|
||||
|
||||
extern "C" int ds4_gpu_tp_batch_gate_encode(uint32_t layer, uint32_t rows) {
|
||||
(void)layer; (void)rows;
|
||||
fprintf(stderr, "ds4: tensor parallelism is Metal-only\n");
|
||||
|
||||
+23
-11
@@ -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;
|
||||
}
|
||||
|
||||
|
||||
+30
-7
@@ -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) {
|
||||
|
||||
+18
-3
@@ -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);
|
||||
}
|
||||
@@ -226,6 +229,16 @@ static void print_distributed(FILE *fp, const help_colors *c) {
|
||||
opt(fp, c, "--dist-replay-check", "Diagnostic: reset and replay prompt, then compare logits.");
|
||||
opt(fp, c, "--debug", "Print coordinator route/debug logs.");
|
||||
fputc('\n', fp);
|
||||
title(fp, c, "Tensor Parallelism");
|
||||
fputc('\n', fp);
|
||||
para(fp, c, "Tensor parallelism runs one token at a time across two Macs. Each rank keeps one contiguous half of the routed experts resident while dense and shared weights remain replicated. Start the worker first, then the coordinator.");
|
||||
fputc('\n', fp);
|
||||
opt(fp, c, "--tp-coordinator PORT", "Coordinate a TP pair and listen for one worker.");
|
||||
opt(fp, c, "--tp-coordinator-host HOST", "Coordinator listen address. Default: 0.0.0.0");
|
||||
opt(fp, c, "--tp-worker HOST PORT", "Dial the coordinator and mirror its session.");
|
||||
opt(fp, c, "--tp-transport auto|rdma|tcp", "Gate transport. Default: auto");
|
||||
opt(fp, c, "--tp-debug-hash N", "Cross-check hidden state every N tokens.");
|
||||
fputc('\n', fp);
|
||||
}
|
||||
|
||||
static void print_cli_diagnostics(FILE *fp, const help_colors *c);
|
||||
@@ -248,6 +261,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 +292,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 +528,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);
|
||||
|
||||
+10559
-789
File diff suppressed because it is too large.
Load diff
+24
@@ -142,10 +142,34 @@ extern "C" void ds4_gpu_tp_set_batch_exchange(ds4_gpu_tp_batch_exchange_fn fn) {
|
||||
(void)fn;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_suspend_expert_sharding(int suspend) {
|
||||
(void)suspend;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_keepalive_pause(int paused) {
|
||||
(void)paused;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_set_attn_head_split(int enabled) {
|
||||
(void)enabled;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_model_residency_skip(int skip) {
|
||||
(void)skip;
|
||||
}
|
||||
|
||||
extern "C" void ds4_gpu_tp_set_big_exchange(ds4_gpu_tp_big_exchange_fn fn) {
|
||||
(void)fn;
|
||||
}
|
||||
|
||||
extern "C" int ds4_gpu_tp_big_gate_encode(uint32_t layer, uint32_t rows,
|
||||
const ds4_gpu_tensor *out_t,
|
||||
ds4_gpu_tensor *in_t,
|
||||
uint64_t bytes) {
|
||||
(void)layer; (void)rows; (void)out_t; (void)in_t; (void)bytes;
|
||||
return 0;
|
||||
}
|
||||
|
||||
extern "C" int ds4_gpu_tp_batch_gate_encode(uint32_t layer, uint32_t rows) {
|
||||
(void)layer; (void)rows;
|
||||
fprintf(stderr, DS4_GPU_LOG_PREFIX "tensor parallelism is Metal-only\n");
|
||||
|
||||
+954
-103
File diff suppressed because it is too large.
Load diff
File diff suppressed because it is too large.
Load diff
@@ -38,7 +38,7 @@
|
||||
|
||||
#define DS4_TP_MAGIC UINT32_C(0x44533454) /* "DS4T" */
|
||||
#define DS4_TP_BATCH_MAGIC UINT32_C(0x44533442) /* "DS4B" */
|
||||
#define DS4_TP_PROTOCOL_VERSION 1u
|
||||
#define DS4_TP_PROTOCOL_VERSION 6u
|
||||
|
||||
/* Default gate timeout is generous: the first gate after a sync waits for
|
||||
* the peer's whole (possibly cold page cache) prefill. */
|
||||
@@ -55,8 +55,6 @@ typedef struct {
|
||||
uint32_t version;
|
||||
uint32_t role;
|
||||
uint32_t rdma_ok; /* this side has a usable verbs device */
|
||||
uint32_t null_split; /* DS4_TP_NULL_SPLIT bring-up mode, must match */
|
||||
uint32_t pad;
|
||||
uint64_t gguf_bytes;
|
||||
uint32_t model_id;
|
||||
uint32_t n_layer;
|
||||
@@ -64,6 +62,10 @@ typedef struct {
|
||||
uint32_t n_vocab;
|
||||
uint32_t quant_bits;
|
||||
uint32_t ctx_size;
|
||||
uint32_t gate_slot_start;
|
||||
uint32_t gate_slot_step;
|
||||
uint32_t gates_per_token;
|
||||
uint32_t pad;
|
||||
} ds4_tp_hello_fixed;
|
||||
|
||||
typedef struct {
|
||||
@@ -119,11 +121,13 @@ typedef struct {
|
||||
* RTR requires GRH addressing with the IPv4-mapped GID that appears only
|
||||
* once the Thunderbolt member interface has an IPv4 address of its own.
|
||||
* UC delivery is in-order and the gate sequence is globally deterministic
|
||||
* (86 gates per token, fixed order), so receives are pre-posted purely by
|
||||
* sequence number: recv for seq s lands in the slab in-slot (s-1) % slots
|
||||
* and its completion IS the arrival signal. */
|
||||
* (86 gates per token, fixed order). After any initial bulk prefill, decode
|
||||
* keeps a receive window posted by sequence number: recv for seq s lands in
|
||||
* the slab in-slot (s-1) % slots and its completion IS the arrival signal. */
|
||||
#define DS4_TP_RDMA_MAX_MSG 16384
|
||||
#define DS4_TP_RDMA_RECV_WINDOW 16
|
||||
#define DS4_TP_RDMA_BULK_SLOTS 64
|
||||
#define DS4_TP_RDMA_BULK_WR_TAG (UINT64_C(1) << 63)
|
||||
|
||||
typedef struct {
|
||||
ds4_tp_verbs_api api;
|
||||
@@ -139,7 +143,8 @@ typedef struct {
|
||||
ds4_tp_rdma_info peer;
|
||||
uint32_t send_outstanding; /* signaled sends not yet reaped */
|
||||
uint64_t recv_done; /* highest gate seq whose recv completed */
|
||||
uint64_t recv_posted; /* highest gate seq with a posted recv */
|
||||
uint64_t last_gate_seq; /* last real decode receive consumed */
|
||||
bool recv_window_active; /* decode recvs are queued ahead */
|
||||
pthread_mutex_t post_lock;
|
||||
} ds4_tp_rdma;
|
||||
#endif
|
||||
@@ -148,13 +153,17 @@ struct ds4_tp {
|
||||
ds4_tp_options opt;
|
||||
int rank; /* 0 leader, 1 worker */
|
||||
int control_fd;
|
||||
int data_fd; /* TCP gate socket (batch gates under RDMA too) */
|
||||
int data_fd; /* TCP fallback, headers, and verify gates */
|
||||
bool rdma_active;
|
||||
uint32_t peer_ctx;
|
||||
uint32_t n_layer;
|
||||
uint32_t n_embd;
|
||||
uint64_t vec_bytes;
|
||||
uint32_t n_slots;
|
||||
/* Decode gate schedule (see ds4_tp_identity). */
|
||||
uint32_t gate_slot_start;
|
||||
uint32_t gate_slot_step;
|
||||
uint32_t gates_per_token;
|
||||
uint8_t *slab;
|
||||
uint64_t slab_bytes;
|
||||
/* Slab regions, see ds4_tp.h layout comment. */
|
||||
@@ -321,11 +330,11 @@ bool ds4_tp_enabled(const ds4_tp_options *opt) {
|
||||
|
||||
void ds4_tp_usage(FILE *fp) {
|
||||
fprintf(fp,
|
||||
"Tensor parallelism (two identical machines, full model on both):\n"
|
||||
" --tp-lead <port> Lead a TP pair: listen for the worker,\n"
|
||||
" then run the normal CLI flow.\n"
|
||||
" --tp-lead-host <host> Leader listen address (default 0.0.0.0).\n"
|
||||
" --tp-worker <host> <port> Dial the leader and mirror its session.\n"
|
||||
"Tensor parallelism (two identical machines):\n"
|
||||
" --tp-coordinator <port> Coordinate a TP pair: listen for the\n"
|
||||
" worker, then run the normal CLI flow.\n"
|
||||
" --tp-coordinator-host <h> Coordinator listen address (default 0.0.0.0).\n"
|
||||
" --tp-worker <host> <port> Dial the coordinator and mirror its session.\n"
|
||||
" --tp-transport <auto|rdma|tcp> Gate transport (default auto).\n"
|
||||
" --tp-debug-hash <n> Cross-check hidden state every n tokens.\n");
|
||||
}
|
||||
@@ -340,16 +349,19 @@ int ds4_tp_parse_cli_arg(
|
||||
size_t errlen)
|
||||
{
|
||||
int i = *index;
|
||||
if (!strcmp(arg, "--tp-lead")) {
|
||||
if (!strcmp(arg, "--tp-coordinator") || !strcmp(arg, "--tp-lead")) {
|
||||
/* --tp-lead is the legacy spelling; --tp-coordinator matches the
|
||||
* pipelined distributed mode's coordinator/worker vocabulary. */
|
||||
if (i + 1 >= argc) goto missing;
|
||||
opt->role = DS4_TP_LEADER;
|
||||
opt->listen_port = atoi(argv[++i]);
|
||||
if (opt->listen_port <= 0 || opt->listen_port > 65535) {
|
||||
tp_set_err(err, errlen, "invalid --tp-lead port");
|
||||
tp_set_err(err, errlen, "invalid %s port", arg);
|
||||
return DS4_TP_CLI_ERROR;
|
||||
}
|
||||
if (!opt->listen_host) opt->listen_host = "0.0.0.0";
|
||||
} else if (!strcmp(arg, "--tp-lead-host")) {
|
||||
} else if (!strcmp(arg, "--tp-coordinator-host") ||
|
||||
!strcmp(arg, "--tp-lead-host")) {
|
||||
if (i + 1 >= argc) goto missing;
|
||||
opt->listen_host = argv[++i];
|
||||
} else if (!strcmp(arg, "--tp-worker")) {
|
||||
@@ -406,7 +418,7 @@ int ds4_tp_validate_engine_options(
|
||||
* verify block is mirrored to the worker via DS4_TP_FRAME_VERIFY and
|
||||
* the legacy MTP path falls back to per-token decode under TP. */
|
||||
if (opt->load_slice) {
|
||||
tp_set_err(err, errlen, "tensor parallelism loads the full model on both ranks");
|
||||
tp_set_err(err, errlen, "tensor parallelism does not use distributed layer slices");
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
@@ -616,6 +628,7 @@ static int tp_rdma_open(ds4_tp *tp, char *err, size_t errlen) {
|
||||
return 0;
|
||||
}
|
||||
r->max_inline = qia.cap.max_inline_data;
|
||||
|
||||
pthread_mutex_init(&r->post_lock, NULL);
|
||||
return 1;
|
||||
}
|
||||
@@ -690,21 +703,21 @@ static int tp_rdma_register_and_exchange(ds4_tp *tp, char *err, size_t errlen) {
|
||||
tp_set_err(err, errlen, "tp rdma: modify RTS: %s", strerror(errno));
|
||||
return 0;
|
||||
}
|
||||
if (tp->vec_bytes > DS4_TP_RDMA_MAX_MSG) {
|
||||
if (tp->vec_bytes > 2ull * DS4_TP_RDMA_MAX_MSG) {
|
||||
tp_set_err(err, errlen,
|
||||
"tp rdma: gate vector %llu bytes exceeds the driver's %u "
|
||||
"message limit (needs chunking)",
|
||||
"tp rdma: gate vector %llu bytes exceeds twice the driver's "
|
||||
"%u message limit",
|
||||
(unsigned long long)tp->vec_bytes, DS4_TP_RDMA_MAX_MSG);
|
||||
return 0;
|
||||
}
|
||||
/* Arm the initial recv window, then barrier on the control socket so
|
||||
* neither side can send a gate before the peer's recvs exist. */
|
||||
for (uint64_t s = 1; s <= DS4_TP_RDMA_RECV_WINDOW; s++) {
|
||||
if (!tp_rdma_post_gate_recv(tp, s)) {
|
||||
tp_set_err(err, errlen, "tp rdma: initial recv post failed");
|
||||
return 0;
|
||||
}
|
||||
if (tp->vec_bytes > DS4_TP_RDMA_MAX_MSG) {
|
||||
fprintf(stderr,
|
||||
"ds4-tp: rdma gate vectors ride as 2 chunked messages "
|
||||
"(%llu bytes > %u limit)\n",
|
||||
(unsigned long long)tp->vec_bytes, DS4_TP_RDMA_MAX_MSG);
|
||||
}
|
||||
/* Leave the receive queue empty for an initial bulk prefill. The first
|
||||
* decode gate arms the normal lookahead window after prefill finishes. */
|
||||
if (!tp_send_frame(tp->control_fd, DS4_TP_FRAME_RDMA_READY, NULL, 0)) {
|
||||
tp_set_err(err, errlen, "tp rdma: ready send failed");
|
||||
return 0;
|
||||
@@ -726,6 +739,16 @@ static const char *tp_wc_status_str(int status) {
|
||||
return buf;
|
||||
}
|
||||
|
||||
/* Slab slot a given gate seq lands in. DS4 fires every slot in order
|
||||
* (identity mapping); GLM's schedule from the hello skips dense layers
|
||||
* and the ATTN slots. */
|
||||
static uint32_t tp_gate_slot(const ds4_tp *tp, uint64_t seq) {
|
||||
if (tp->gates_per_token == 0)
|
||||
return (uint32_t)((seq - 1) % tp->n_slots);
|
||||
return tp->gate_slot_start +
|
||||
(uint32_t)((seq - 1) % tp->gates_per_token) * tp->gate_slot_step;
|
||||
}
|
||||
|
||||
/* Reap completions: send CQEs free send-queue slots, recv CQEs advance the
|
||||
* arrival watermark (UC is in-order, so gate seq recv completions arrive
|
||||
* monotonically). Returns 0 on any completion error. */
|
||||
@@ -755,50 +778,84 @@ static int tp_rdma_drain_cq(ds4_tp *tp) {
|
||||
* kernel reads. */
|
||||
static int tp_rdma_post_gate_recv(ds4_tp *tp, uint64_t seq) {
|
||||
ds4_tp_rdma *r = &tp->rdma;
|
||||
const uint32_t slot = (uint32_t)((seq - 1) % tp->n_slots);
|
||||
struct ibv_sge sge;
|
||||
struct ibv_recv_wr wr, *bad = NULL;
|
||||
memset(&wr, 0, sizeof(wr));
|
||||
sge.addr = (uintptr_t)(tp->slab + tp->in_off + (uint64_t)slot * tp->vec_bytes);
|
||||
sge.length = (uint32_t)tp->vec_bytes;
|
||||
sge.lkey = r->mr->lkey;
|
||||
wr.wr_id = seq;
|
||||
wr.sg_list = &sge;
|
||||
wr.num_sge = 1;
|
||||
if (ibv_post_recv(r->qp, &wr, &bad) != 0) {
|
||||
fprintf(stderr, "ds4-tp: rdma post_recv(seq %llu): %s\n",
|
||||
(unsigned long long)seq, strerror(errno));
|
||||
return 0;
|
||||
const uint32_t slot = tp_gate_slot(tp, seq);
|
||||
const uintptr_t base =
|
||||
(uintptr_t)(tp->slab + tp->in_off + (uint64_t)slot * tp->vec_bytes);
|
||||
/* Vectors above the driver's 16KB message cap ride as two chunks
|
||||
* landing contiguously in the slot. UC delivery is in-order and both
|
||||
* sides post/send strictly in seq order, so the k'th send always
|
||||
* matches the k'th recv; only the FINAL chunk carries the seq as
|
||||
* wr_id, so the arrival watermark advances when the slot is whole. */
|
||||
uint64_t off = 0;
|
||||
while (off < tp->vec_bytes) {
|
||||
const uint64_t len = tp->vec_bytes - off > DS4_TP_RDMA_MAX_MSG ?
|
||||
DS4_TP_RDMA_MAX_MSG : tp->vec_bytes - off;
|
||||
const int last = off + len == tp->vec_bytes;
|
||||
struct ibv_sge sge;
|
||||
struct ibv_recv_wr wr, *bad = NULL;
|
||||
memset(&wr, 0, sizeof(wr));
|
||||
sge.addr = base + off;
|
||||
sge.length = (uint32_t)len;
|
||||
sge.lkey = r->mr->lkey;
|
||||
wr.wr_id = last ? seq : 0;
|
||||
wr.sg_list = &sge;
|
||||
wr.num_sge = 1;
|
||||
if (ibv_post_recv(r->qp, &wr, &bad) != 0) {
|
||||
fprintf(stderr, "ds4-tp: rdma post_recv(seq %llu off %llu): %s\n",
|
||||
(unsigned long long)seq, (unsigned long long)off,
|
||||
strerror(errno));
|
||||
return 0;
|
||||
}
|
||||
off += len;
|
||||
}
|
||||
if (seq > r->recv_posted) r->recv_posted = seq;
|
||||
return 1;
|
||||
}
|
||||
|
||||
/* One gate: send our partial, wait for the peer's recv completion, keep the
|
||||
* recv window armed. */
|
||||
/* One decode gate: ensure the receive window is armed, send our partial,
|
||||
* wait for the peer's receive completion, and advance the window. */
|
||||
static int tp_rdma_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t gate, uint64_t seq) {
|
||||
ds4_tp_rdma *r = &tp->rdma;
|
||||
const uint32_t slot = layer * DS4_TP_GATES_PER_LAYER + gate;
|
||||
if (slot != (uint32_t)((seq - 1) % tp->n_slots)) {
|
||||
if (getenv("DS4_TP_GATE_TRACE")) {
|
||||
fprintf(stderr, "ds4-tp: gate trace l=%u g=%u seq=%llu want_slot=%u\n",
|
||||
layer, gate, (unsigned long long)seq, tp_gate_slot(tp, seq));
|
||||
}
|
||||
if (slot != tp_gate_slot(tp, seq)) {
|
||||
fprintf(stderr, "ds4-tp: gate order broke: layer %u gate %u vs seq %llu\n",
|
||||
layer, gate, (unsigned long long)seq);
|
||||
return 0;
|
||||
}
|
||||
struct ibv_sge sge;
|
||||
struct ibv_send_wr wr, *bad = NULL;
|
||||
memset(&wr, 0, sizeof(wr));
|
||||
sge.addr = (uintptr_t)(tp->slab + tp->out_off + (uint64_t)slot * tp->vec_bytes);
|
||||
sge.length = (uint32_t)tp->vec_bytes;
|
||||
sge.lkey = r->mr->lkey;
|
||||
wr.wr_id = seq;
|
||||
wr.sg_list = &sge;
|
||||
wr.num_sge = 1;
|
||||
wr.opcode = IBV_WR_SEND;
|
||||
wr.send_flags = IBV_SEND_SIGNALED;
|
||||
const uintptr_t send_base =
|
||||
(uintptr_t)(tp->slab + tp->out_off + (uint64_t)slot * tp->vec_bytes);
|
||||
pthread_mutex_lock(&r->post_lock);
|
||||
int ok = ibv_post_send(r->qp, &wr, &bad) == 0;
|
||||
if (ok) r->send_outstanding++;
|
||||
else fprintf(stderr, "ds4-tp: rdma post_send: %s\n", strerror(errno));
|
||||
int ok = 1;
|
||||
if (!r->recv_window_active) {
|
||||
for (uint64_t s = seq; ok && s < seq + DS4_TP_RDMA_RECV_WINDOW; s++)
|
||||
ok = tp_rdma_post_gate_recv(tp, s);
|
||||
if (ok) r->recv_window_active = true;
|
||||
}
|
||||
for (uint64_t off = 0; ok && off < tp->vec_bytes; ) {
|
||||
const uint64_t len = tp->vec_bytes - off > DS4_TP_RDMA_MAX_MSG ?
|
||||
DS4_TP_RDMA_MAX_MSG : tp->vec_bytes - off;
|
||||
struct ibv_sge sge;
|
||||
struct ibv_send_wr wr, *bad = NULL;
|
||||
memset(&wr, 0, sizeof(wr));
|
||||
sge.addr = send_base + off;
|
||||
sge.length = (uint32_t)len;
|
||||
sge.lkey = r->mr->lkey;
|
||||
wr.wr_id = seq;
|
||||
wr.sg_list = &sge;
|
||||
wr.num_sge = 1;
|
||||
wr.opcode = IBV_WR_SEND;
|
||||
wr.send_flags = IBV_SEND_SIGNALED;
|
||||
ok = ibv_post_send(r->qp, &wr, &bad) == 0;
|
||||
if (!ok) {
|
||||
fprintf(stderr, "ds4-tp: rdma post_send: %s\n", strerror(errno));
|
||||
} else {
|
||||
r->send_outstanding++;
|
||||
}
|
||||
off += len;
|
||||
}
|
||||
|
||||
double deadline = 0.0;
|
||||
while (ok && r->recv_done < seq) {
|
||||
@@ -811,10 +868,250 @@ static int tp_rdma_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t gate, uint
|
||||
}
|
||||
}
|
||||
if (ok) ok = tp_rdma_post_gate_recv(tp, seq + DS4_TP_RDMA_RECV_WINDOW);
|
||||
if (ok) r->last_gate_seq = seq;
|
||||
pthread_mutex_unlock(&r->post_lock);
|
||||
return ok;
|
||||
}
|
||||
|
||||
static int tp_rdma_big_gate_capable(const ds4_tp *tp) {
|
||||
const uint64_t stage_bytes =
|
||||
(uint64_t)DS4_TP_RDMA_BULK_SLOTS * DS4_TP_RDMA_MAX_MSG;
|
||||
const uint64_t batch_region_bytes =
|
||||
(uint64_t)tp->n_layer * DS4_TP_BATCH_MAX_ROWS * tp->vec_bytes;
|
||||
return tp->rdma.qp && tp->rdma.mr && batch_region_bytes >= stage_bytes;
|
||||
}
|
||||
|
||||
/* Decode keeps a lookahead window of receives on the latency QP. Before a
|
||||
* later prompt can reuse that QP for bulk rows, consume those receives with
|
||||
* dummy sends on both ranks. The TCP big-gate header exchange is the barrier
|
||||
* that guarantees both sides have reached this transition. */
|
||||
static int tp_rdma_drain_decode_window(ds4_tp *tp) {
|
||||
ds4_tp_rdma *r = &tp->rdma;
|
||||
if (!r->recv_window_active) return 1;
|
||||
|
||||
const uint32_t chunks_per_gate =
|
||||
(uint32_t)((tp->vec_bytes + DS4_TP_RDMA_MAX_MSG - 1u) /
|
||||
DS4_TP_RDMA_MAX_MSG);
|
||||
const uint32_t nwr = DS4_TP_RDMA_RECV_WINDOW * chunks_per_gate;
|
||||
struct ibv_sge sge[DS4_TP_RDMA_RECV_WINDOW * 2u];
|
||||
struct ibv_send_wr wr[DS4_TP_RDMA_RECV_WINDOW * 2u];
|
||||
memset(wr, 0, sizeof(wr));
|
||||
uint8_t *scratch = tp->slab + tp->batch_out_off;
|
||||
uint32_t wi = 0;
|
||||
for (uint32_t gate = 0; gate < DS4_TP_RDMA_RECV_WINDOW; gate++) {
|
||||
for (uint64_t off = 0; off < tp->vec_bytes; ) {
|
||||
const uint64_t len = tp->vec_bytes - off > DS4_TP_RDMA_MAX_MSG ?
|
||||
DS4_TP_RDMA_MAX_MSG : tp->vec_bytes - off;
|
||||
sge[wi] = (struct ibv_sge) {
|
||||
.addr = (uintptr_t)(scratch + off),
|
||||
.length = (uint32_t)len,
|
||||
.lkey = r->mr->lkey,
|
||||
};
|
||||
wr[wi].wr_id = DS4_TP_RDMA_BULK_WR_TAG | ((uint64_t)wi + 1u);
|
||||
wr[wi].sg_list = &sge[wi];
|
||||
wr[wi].num_sge = 1;
|
||||
wr[wi].opcode = IBV_WR_SEND;
|
||||
wr[wi].send_flags = wi + 1u == nwr ? IBV_SEND_SIGNALED : 0;
|
||||
if (wi > 0) wr[wi - 1u].next = &wr[wi];
|
||||
wi++;
|
||||
off += len;
|
||||
}
|
||||
}
|
||||
|
||||
pthread_mutex_lock(&r->post_lock);
|
||||
struct ibv_send_wr *bad = NULL;
|
||||
if (ibv_post_send(r->qp, wr, &bad) != 0) {
|
||||
fprintf(stderr, "ds4-tp: rdma receive-window drain post failed: %s\n",
|
||||
strerror(errno));
|
||||
pthread_mutex_unlock(&r->post_lock);
|
||||
return 0;
|
||||
}
|
||||
|
||||
uint32_t recv_done = 0;
|
||||
int send_done = 0;
|
||||
const double deadline = tp_now_sec() + (double)tp->timeout_sec;
|
||||
while (recv_done < nwr || !send_done) {
|
||||
struct ibv_wc wc[DS4_TP_RDMA_RECV_WINDOW * 2u + 1u];
|
||||
int n = ibv_poll_cq(r->cq,
|
||||
(int)(DS4_TP_RDMA_RECV_WINDOW * 2u + 1u), wc);
|
||||
if (n < 0) {
|
||||
pthread_mutex_unlock(&r->post_lock);
|
||||
return 0;
|
||||
}
|
||||
for (int i = 0; i < n; i++) {
|
||||
if (wc[i].status != IBV_WC_SUCCESS) {
|
||||
fprintf(stderr, "ds4-tp: rdma receive-window drain: %s\n",
|
||||
tp_wc_status_str(wc[i].status));
|
||||
pthread_mutex_unlock(&r->post_lock);
|
||||
return 0;
|
||||
}
|
||||
if (wc[i].opcode & IBV_WC_RECV) {
|
||||
recv_done++;
|
||||
} else if (wc[i].wr_id & DS4_TP_RDMA_BULK_WR_TAG) {
|
||||
send_done = 1;
|
||||
} else if (r->send_outstanding > 0) {
|
||||
r->send_outstanding--;
|
||||
}
|
||||
}
|
||||
if (tp_now_sec() > deadline) {
|
||||
fprintf(stderr,
|
||||
"ds4-tp: timeout draining RDMA receive window (%u/%u)\n",
|
||||
recv_done, nwr);
|
||||
pthread_mutex_unlock(&r->post_lock);
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
r->recv_done = r->last_gate_seq;
|
||||
r->recv_window_active = false;
|
||||
pthread_mutex_unlock(&r->post_lock);
|
||||
return 1;
|
||||
}
|
||||
|
||||
/* Large prefill row swaps share the latency QP. No future decode receives
|
||||
* are queued, so each round can post its 1 MiB receive window before sending
|
||||
* the matching 16 KiB messages. Verify scratch provides already-registered
|
||||
* staging memory and is idle during normal prefill. */
|
||||
static int tp_rdma_big_gate_exchange(ds4_tp *tp,
|
||||
const void *out,
|
||||
void *in,
|
||||
uint64_t bytes) {
|
||||
ds4_tp_rdma *r = &tp->rdma;
|
||||
if (!tp_rdma_big_gate_capable(tp) || r->recv_window_active) return 0;
|
||||
|
||||
/* Payloads already inside the registered slab (verify batches) can ride
|
||||
* directly. Ordinary prefill tensors use the idle verify regions as
|
||||
* registered staging because their standalone MTLBuffers are not in the
|
||||
* NIC memory region. */
|
||||
const uintptr_t slab_lo = (uintptr_t)tp->slab;
|
||||
const uintptr_t slab_hi = slab_lo + tp->slab_bytes;
|
||||
const uintptr_t out_lo = (uintptr_t)out;
|
||||
const uintptr_t in_lo = (uintptr_t)in;
|
||||
const bool direct =
|
||||
out_lo >= slab_lo && out_lo <= slab_hi && bytes <= slab_hi - out_lo &&
|
||||
in_lo >= slab_lo && in_lo <= slab_hi && bytes <= slab_hi - in_lo;
|
||||
uint8_t *stage_send = tp->slab + tp->batch_out_off;
|
||||
uint8_t *stage_recv = tp->slab + tp->batch_in_off;
|
||||
uint64_t off = 0;
|
||||
while (off < bytes) {
|
||||
const uint64_t remaining = bytes - off;
|
||||
uint32_t chunks = (uint32_t)((remaining + DS4_TP_RDMA_MAX_MSG - 1u) /
|
||||
DS4_TP_RDMA_MAX_MSG);
|
||||
if (chunks > DS4_TP_RDMA_BULK_SLOTS)
|
||||
chunks = DS4_TP_RDMA_BULK_SLOTS;
|
||||
|
||||
uint32_t lens[DS4_TP_RDMA_BULK_SLOTS];
|
||||
uint64_t chunk_off[DS4_TP_RDMA_BULK_SLOTS];
|
||||
uint64_t round_bytes = 0;
|
||||
for (uint32_t i = 0; i < chunks; i++) {
|
||||
const uint64_t left = remaining - round_bytes;
|
||||
lens[i] = (uint32_t)(left > DS4_TP_RDMA_MAX_MSG ?
|
||||
DS4_TP_RDMA_MAX_MSG : left);
|
||||
chunk_off[i] = direct ? round_bytes :
|
||||
(uint64_t)i * DS4_TP_RDMA_MAX_MSG;
|
||||
if (!direct) {
|
||||
memcpy(stage_send + chunk_off[i],
|
||||
(const uint8_t *)out + off + round_bytes, lens[i]);
|
||||
}
|
||||
round_bytes += lens[i];
|
||||
}
|
||||
|
||||
struct ibv_sge recv_sge[DS4_TP_RDMA_BULK_SLOTS];
|
||||
struct ibv_recv_wr recv_wr[DS4_TP_RDMA_BULK_SLOTS];
|
||||
memset(recv_wr, 0, sizeof(recv_wr));
|
||||
for (uint32_t i = 0; i < chunks; i++) {
|
||||
recv_sge[i] = (struct ibv_sge) {
|
||||
.addr = direct ? in_lo + off + chunk_off[i] :
|
||||
(uintptr_t)(stage_recv + chunk_off[i]),
|
||||
.length = lens[i],
|
||||
.lkey = r->mr->lkey,
|
||||
};
|
||||
recv_wr[i].wr_id = DS4_TP_RDMA_BULK_WR_TAG | ((uint64_t)i + 1u);
|
||||
recv_wr[i].sg_list = &recv_sge[i];
|
||||
recv_wr[i].num_sge = 1;
|
||||
recv_wr[i].next = i + 1u < chunks ? &recv_wr[i + 1u] : NULL;
|
||||
}
|
||||
struct ibv_recv_wr *bad_recv = NULL;
|
||||
if (ibv_post_recv(r->qp, recv_wr, &bad_recv) != 0) {
|
||||
fprintf(stderr, "ds4-tp: bulk rdma post_recv: %s\n",
|
||||
strerror(errno));
|
||||
return 0;
|
||||
}
|
||||
atomic_thread_fence(memory_order_release);
|
||||
struct ibv_sge send_sge[DS4_TP_RDMA_BULK_SLOTS];
|
||||
struct ibv_send_wr send_wr[DS4_TP_RDMA_BULK_SLOTS];
|
||||
memset(send_wr, 0, sizeof(send_wr));
|
||||
for (uint32_t i = 0; i < chunks; i++) {
|
||||
send_sge[i] = (struct ibv_sge) {
|
||||
.addr = direct ? out_lo + off + chunk_off[i] :
|
||||
(uintptr_t)(stage_send + chunk_off[i]),
|
||||
.length = lens[i],
|
||||
.lkey = r->mr->lkey,
|
||||
};
|
||||
send_wr[i].wr_id = DS4_TP_RDMA_BULK_WR_TAG | ((uint64_t)i + 1u);
|
||||
send_wr[i].sg_list = &send_sge[i];
|
||||
send_wr[i].num_sge = 1;
|
||||
send_wr[i].opcode = IBV_WR_SEND;
|
||||
send_wr[i].send_flags = i + 1u == chunks ? IBV_SEND_SIGNALED : 0;
|
||||
send_wr[i].next = i + 1u < chunks ? &send_wr[i + 1u] : NULL;
|
||||
}
|
||||
struct ibv_send_wr *bad_send = NULL;
|
||||
if (ibv_post_send(r->qp, send_wr, &bad_send) != 0) {
|
||||
fprintf(stderr, "ds4-tp: bulk rdma post_send: %s\n",
|
||||
strerror(errno));
|
||||
return 0;
|
||||
}
|
||||
|
||||
uint32_t recv_done = 0;
|
||||
int send_done = 0;
|
||||
const double deadline = tp_now_sec() + (double)tp->timeout_sec;
|
||||
while (recv_done < chunks || !send_done) {
|
||||
struct ibv_wc wc[DS4_TP_RDMA_BULK_SLOTS + 1u];
|
||||
int n = ibv_poll_cq(r->cq,
|
||||
(int)(DS4_TP_RDMA_BULK_SLOTS + 1u), wc);
|
||||
if (n < 0) return 0;
|
||||
for (int i = 0; i < n; i++) {
|
||||
if (wc[i].status != IBV_WC_SUCCESS) {
|
||||
fprintf(stderr,
|
||||
"ds4-tp: bulk rdma completion error: %s\n",
|
||||
tp_wc_status_str(wc[i].status));
|
||||
return 0;
|
||||
}
|
||||
if ((wc[i].wr_id & DS4_TP_RDMA_BULK_WR_TAG) == 0) {
|
||||
/* A final latency-QP send completion can remain queued
|
||||
* when a later prompt starts a bulk gate. */
|
||||
if (wc[i].opcode & IBV_WC_RECV) {
|
||||
if (wc[i].wr_id > r->recv_done)
|
||||
r->recv_done = wc[i].wr_id;
|
||||
} else if (r->send_outstanding > 0) {
|
||||
r->send_outstanding--;
|
||||
}
|
||||
continue;
|
||||
}
|
||||
if (wc[i].opcode & IBV_WC_RECV) recv_done++;
|
||||
else send_done = 1;
|
||||
}
|
||||
if (tp_now_sec() > deadline) {
|
||||
fprintf(stderr,
|
||||
"ds4-tp: timeout waiting for bulk RDMA round "
|
||||
"(%u/%u recvs, send=%d)\n",
|
||||
recv_done, chunks, send_done);
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
atomic_thread_fence(memory_order_acquire);
|
||||
if (!direct) {
|
||||
round_bytes = 0;
|
||||
for (uint32_t i = 0; i < chunks; i++) {
|
||||
memcpy((uint8_t *)in + off + round_bytes,
|
||||
stage_recv + chunk_off[i], lens[i]);
|
||||
round_bytes += lens[i];
|
||||
}
|
||||
}
|
||||
off += round_bytes;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static void tp_rdma_close(ds4_tp *tp) {
|
||||
ds4_tp_rdma *r = &tp->rdma;
|
||||
if (r->qp) r->api.destroy_qp(r->qp);
|
||||
@@ -838,7 +1135,6 @@ static int tp_hello_exchange(ds4_tp *tp, const ds4_tp_identity *id, int rdma_ok,
|
||||
.version = DS4_TP_PROTOCOL_VERSION,
|
||||
.role = (uint32_t)tp->opt.role,
|
||||
.rdma_ok = (uint32_t)rdma_ok,
|
||||
.null_split = getenv("DS4_TP_NULL_SPLIT") != NULL,
|
||||
.gguf_bytes = id->gguf_bytes,
|
||||
.model_id = id->model_id,
|
||||
.n_layer = id->n_layer,
|
||||
@@ -846,6 +1142,9 @@ static int tp_hello_exchange(ds4_tp *tp, const ds4_tp_identity *id, int rdma_ok,
|
||||
.n_vocab = id->n_vocab,
|
||||
.quant_bits = id->quant_bits,
|
||||
.ctx_size = id->ctx_size,
|
||||
.gate_slot_start = id->gate_slot_start,
|
||||
.gate_slot_step = id->gate_slot_step,
|
||||
.gates_per_token = id->gates_per_token,
|
||||
};
|
||||
ds4_tp_hello_fixed theirs;
|
||||
if (!tp_write_full(tp->control_fd, &mine, sizeof(mine)) ||
|
||||
@@ -866,13 +1165,12 @@ static int tp_hello_exchange(ds4_tp *tp, const ds4_tp_identity *id, int rdma_ok,
|
||||
tp_set_err(err, errlen, "tp hello: both sides claim role %u", mine.role);
|
||||
return 0;
|
||||
}
|
||||
if (theirs.null_split != mine.null_split) {
|
||||
tp_set_err(err, errlen, "tp hello: DS4_TP_NULL_SPLIT set on one side only");
|
||||
return 0;
|
||||
}
|
||||
if (theirs.gguf_bytes != mine.gguf_bytes || theirs.model_id != mine.model_id ||
|
||||
theirs.n_layer != mine.n_layer || theirs.n_embd != mine.n_embd ||
|
||||
theirs.n_vocab != mine.n_vocab || theirs.quant_bits != mine.quant_bits) {
|
||||
theirs.n_vocab != mine.n_vocab || theirs.quant_bits != mine.quant_bits ||
|
||||
theirs.gate_slot_start != mine.gate_slot_start ||
|
||||
theirs.gate_slot_step != mine.gate_slot_step ||
|
||||
theirs.gates_per_token != mine.gates_per_token) {
|
||||
tp_set_err(err, errlen,
|
||||
"tp hello: model mismatch (peer gguf=%llu id=%u layers=%u embd=%u "
|
||||
"vocab=%u qbits=%u)",
|
||||
@@ -885,6 +1183,9 @@ static int tp_hello_exchange(ds4_tp *tp, const ds4_tp_identity *id, int rdma_ok,
|
||||
tp->n_embd = id->n_embd;
|
||||
tp->vec_bytes = (uint64_t)id->n_embd * sizeof(float);
|
||||
tp->n_slots = id->n_layer * DS4_TP_GATES_PER_LAYER;
|
||||
tp->gate_slot_start = id->gate_slot_start;
|
||||
tp->gate_slot_step = id->gate_slot_step;
|
||||
tp->gates_per_token = id->gates_per_token;
|
||||
tp_slab_layout(tp);
|
||||
/* Transport decision: RDMA only when both sides can. */
|
||||
int want_rdma = tp->opt.transport != DS4_TP_TRANSPORT_TCP;
|
||||
@@ -920,7 +1221,8 @@ int ds4_tp_create(
|
||||
|
||||
int rdma_ok = 0;
|
||||
#ifdef DS4_TP_HAVE_VERBS
|
||||
if (opt->transport != DS4_TP_TRANSPORT_TCP)
|
||||
if (opt->transport != DS4_TP_TRANSPORT_TCP &&
|
||||
(uint64_t)id->n_embd * sizeof(float) <= 2ull * DS4_TP_RDMA_MAX_MSG)
|
||||
rdma_ok = tp_rdma_probe(&tp->rdma.api);
|
||||
#endif
|
||||
|
||||
@@ -951,8 +1253,8 @@ int ds4_tp_create(
|
||||
#endif
|
||||
{
|
||||
/* Second socket dedicated to gate traffic so control frames never
|
||||
* interleave with 16KB gate payloads. Created under RDMA too: the
|
||||
* verify-block batch gates always run over TCP. */
|
||||
* interleave with gate payloads. Created under RDMA too for
|
||||
* headers, verify-block gates, and transport fallback. */
|
||||
if (tp->rank == 0) {
|
||||
tp->data_fd = accept(listener, NULL, NULL);
|
||||
if (tp->data_fd < 0) {
|
||||
@@ -1051,17 +1353,36 @@ int ds4_tp_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t gate, uint64_t seq
|
||||
}
|
||||
|
||||
/* Verify-block batch gate: one exchange per layer moving all block rows at
|
||||
* once. Runs over the TCP data socket even when the row gates use RDMA —
|
||||
* it is off the per-token latency path (43 exchanges per verify block) and
|
||||
* the RDMA engine's 16KB/seq accounting stays untouched. Symmetric
|
||||
* write-then-read like the TCP row gate; 4MB socket buffers absorb the
|
||||
* <=128KB payloads without deadlock. */
|
||||
* once. The payload lives in the registered slab, so RDMA sends it directly;
|
||||
* TCP remains the symmetric write-then-read fallback. */
|
||||
int ds4_tp_batch_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t rows,
|
||||
uint64_t seq) {
|
||||
if (tp->data_fd < 0 || rows == 0 || rows > DS4_TP_BATCH_MAX_ROWS) return 0;
|
||||
const uint64_t bytes = (uint64_t)rows * tp->vec_bytes;
|
||||
ds4_tp_gate_header h = { DS4_TP_BATCH_MAGIC, (uint16_t)layer,
|
||||
(uint16_t)rows, seq };
|
||||
#ifdef DS4_TP_HAVE_VERBS
|
||||
if (tp->rdma_active && tp_rdma_big_gate_capable(tp)) {
|
||||
if (!tp_write_full(tp->data_fd, &h, sizeof(h))) return 0;
|
||||
ds4_tp_gate_header ph;
|
||||
if (!tp_read_full(tp->data_fd, &ph, sizeof(ph))) return 0;
|
||||
if (ph.magic != DS4_TP_BATCH_MAGIC || ph.layer != layer ||
|
||||
ph.gate != rows || ph.seq != seq) {
|
||||
fprintf(stderr,
|
||||
"ds4-tp: batch gate desync: got l=%u rows=%u seq=%llu, "
|
||||
"want l=%u rows=%u seq=%llu\n",
|
||||
ph.layer, ph.gate, (unsigned long long)ph.seq,
|
||||
layer, rows, (unsigned long long)seq);
|
||||
return 0;
|
||||
}
|
||||
if (!tp_rdma_drain_decode_window(tp)) return 0;
|
||||
return tp_rdma_big_gate_exchange(
|
||||
tp,
|
||||
tp->slab + ds4_tp_slab_batch_out_offset(tp, layer),
|
||||
tp->slab + ds4_tp_slab_batch_in_offset(tp, layer),
|
||||
bytes);
|
||||
}
|
||||
#endif
|
||||
struct iovec iov[2] = {
|
||||
{ &h, sizeof(h) },
|
||||
{ tp->slab + ds4_tp_slab_batch_out_offset(tp, layer), bytes },
|
||||
@@ -1099,6 +1420,51 @@ int ds4_tp_batch_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t rows,
|
||||
bytes);
|
||||
}
|
||||
|
||||
/* Prefill batch gate: RDMA uses the pipelined registered-slab path above.
|
||||
* The fallback alternates 2MB TCP write/read rounds in the same order, so
|
||||
* neither side can fill its send buffer while the peer is also only writing
|
||||
* (the 4MB socket buffers absorb one round). */
|
||||
#define DS4_TP_BIG_CHUNK (2ull * 1024ull * 1024ull)
|
||||
|
||||
int ds4_tp_big_gate_exchange(ds4_tp *tp, uint32_t layer, uint64_t seq,
|
||||
const void *out, void *in, uint64_t bytes) {
|
||||
if (tp->data_fd < 0 || !out || !in || bytes == 0) return 0;
|
||||
ds4_tp_gate_header h = { DS4_TP_BATCH_MAGIC, (uint16_t)layer, 0xB16u, seq };
|
||||
if (!tp_write_full(tp->data_fd, &h, sizeof(h))) return 0;
|
||||
ds4_tp_gate_header ph;
|
||||
if (!tp_read_full(tp->data_fd, &ph, sizeof(ph))) return 0;
|
||||
if (ph.magic != DS4_TP_BATCH_MAGIC || ph.layer != layer ||
|
||||
ph.gate != 0xB16u || ph.seq != seq) {
|
||||
fprintf(stderr,
|
||||
"ds4-tp: big gate desync: got l=%u tag=%x seq=%llu, want l=%u seq=%llu\n",
|
||||
ph.layer, ph.gate, (unsigned long long)ph.seq,
|
||||
layer, (unsigned long long)seq);
|
||||
return 0;
|
||||
}
|
||||
#ifdef DS4_TP_HAVE_VERBS
|
||||
if (tp->rdma_active && tp_rdma_big_gate_capable(tp)) {
|
||||
if (!tp_rdma_drain_decode_window(tp)) return 0;
|
||||
return tp_rdma_big_gate_exchange(tp, out, in, bytes);
|
||||
}
|
||||
#endif
|
||||
uint64_t off = 0;
|
||||
while (off < bytes) {
|
||||
const uint64_t n = bytes - off > DS4_TP_BIG_CHUNK ?
|
||||
DS4_TP_BIG_CHUNK : bytes - off;
|
||||
if (!tp_write_full(tp->data_fd, (const char *)out + off, n)) return 0;
|
||||
if (!tp_read_full(tp->data_fd, (char *)in + off, n)) return 0;
|
||||
off += n;
|
||||
}
|
||||
if (getenv("DS4_GLM_TP_DEBUG")) {
|
||||
const float *o = (const float *)out, *i = (const float *)in;
|
||||
fprintf(stderr,
|
||||
"ds4-tp: big gate l=%u seq=%llu out[0..3]=%g %g %g %g in[0..3]=%g %g %g %g\n",
|
||||
layer, (unsigned long long)seq,
|
||||
o[0], o[1], o[2], o[3], i[0], i[1], i[2], i[3]);
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
/* ------------------------------------------------------------------------
|
||||
* Lockstep control plane.
|
||||
* --------------------------------------------------------------------- */
|
||||
@@ -1297,6 +1663,10 @@ int ds4_tp_worker_run(ds4_engine *engine, const ds4_tp_options *opt) {
|
||||
.quant_bits = (uint32_t)ds4_engine_routed_quant_bits(engine),
|
||||
.ctx_size = 0, /* adopt the leader's */
|
||||
};
|
||||
ds4_engine_tp_gate_schedule(engine,
|
||||
&id.gate_slot_start,
|
||||
&id.gate_slot_step,
|
||||
&id.gates_per_token);
|
||||
|
||||
ds4_tp *tp = NULL;
|
||||
if (!ds4_tp_create(&tp, opt, &id, err, sizeof(err))) {
|
||||
@@ -1316,6 +1686,10 @@ int ds4_tp_worker_run(ds4_engine *engine, const ds4_tp_options *opt) {
|
||||
ds4_tp_free(tp);
|
||||
return 1;
|
||||
}
|
||||
/* The worker encodes no main-queue GPU work until the first mirrored
|
||||
* sync arrives, so its one-time first-submission cost would otherwise
|
||||
* land inside the leader-timed prefill. Pay it here instead. */
|
||||
ds4_session_gpu_warmup(session);
|
||||
ds4_log(stderr, DS4_LOG_OK, "tp worker ready (ctx %d)", ctx);
|
||||
|
||||
int rc = 0;
|
||||
@@ -1349,7 +1723,7 @@ int ds4_tp_worker_run(ds4_engine *engine, const ds4_tp_options *opt) {
|
||||
rc = 1;
|
||||
break;
|
||||
}
|
||||
if (getenv("DS4_TP_NULL_SPLIT") == NULL) {
|
||||
if (ds4_engine_tp_vocab_split(engine)) {
|
||||
const int vocab = ds4_engine_vocab_size(engine);
|
||||
const uint32_t vhalf = (uint32_t)vocab / 2u;
|
||||
float *lg = malloc((size_t)vocab * sizeof(float));
|
||||
|
||||
@@ -10,13 +10,14 @@
|
||||
|
||||
/* Tensor-parallel transport and lockstep protocol.
|
||||
*
|
||||
* Two ranks run the same full model. Rank 0 (leader) is a normal frontend
|
||||
* session that mirrors every ds4_session_sync()/ds4_session_eval() call to
|
||||
* rank 1 (worker) over a TCP control socket, so both engines execute the
|
||||
* identical graph sequence. Inside each decoded token, partial block
|
||||
* outputs are exchanged at two gates per layer through a registered memory
|
||||
* slab: RDMA WRITE + sequence flag when RDMA over Thunderbolt is available,
|
||||
* or a full-duplex TCP exchange as fallback.
|
||||
* Two ranks run the same logical model, each with one contiguous half of the
|
||||
* routed experts resident. Rank 0 (leader) is a normal frontend session that
|
||||
* mirrors every ds4_session_sync()/ds4_session_eval() call to rank 1 (worker)
|
||||
* over a TCP control socket, so both engines execute the identical graph
|
||||
* sequence.
|
||||
* Inside each decoded token, partial block outputs are exchanged through a
|
||||
* registered memory slab: two-sided RDMA SEND/RECV when RDMA over
|
||||
* Thunderbolt is available, or a full-duplex TCP exchange as fallback.
|
||||
*
|
||||
* Layering: ds4.c calls the session-mirroring and slab entry points;
|
||||
* ds4_metal.m only ever sees ds4_tp_gate_exchange() through a callback
|
||||
@@ -43,6 +44,15 @@ typedef struct {
|
||||
uint32_t n_vocab;
|
||||
uint32_t quant_bits;
|
||||
uint32_t ctx_size;
|
||||
/* Decode gate schedule, used to place RDMA recvs into the right slab
|
||||
* slot: slot(seq) = start + ((seq-1) % per_token) * step.
|
||||
* per_token 0 falls back to the identity mapping over all slots
|
||||
* (DS4: every layer fires ATTN then FFN). GLM fires one FFN gate per
|
||||
* sparse layer only, so its schedule skips the dense prefix and the
|
||||
* ATTN slots. Exchanged in the hello; both sides must agree. */
|
||||
uint32_t gate_slot_start;
|
||||
uint32_t gate_slot_step;
|
||||
uint32_t gates_per_token;
|
||||
} ds4_tp_identity;
|
||||
|
||||
bool ds4_tp_enabled(const ds4_tp_options *opt);
|
||||
@@ -112,12 +122,17 @@ int ds4_tp_attach_slab(ds4_tp *tp, void *base, char *err, size_t errlen);
|
||||
* Called from the GPU gate service thread. Returns 0 on failure. */
|
||||
int ds4_tp_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t gate, uint64_t seq);
|
||||
|
||||
/* Verify-block batch gate: exchange `rows` row partials for one layer in a
|
||||
* single symmetric transfer over the TCP data socket (used even when the
|
||||
* row gates run RDMA). Called from the GPU gate service thread. */
|
||||
/* Verify-block batch gate: exchange `rows` row partials for one layer in one
|
||||
* bulk RDMA transfer, with a symmetric TCP transfer as fallback. Called from
|
||||
* the GPU gate service thread. */
|
||||
int ds4_tp_batch_gate_exchange(ds4_tp *tp, uint32_t layer, uint32_t rows,
|
||||
uint64_t seq);
|
||||
|
||||
/* Prefill batch gate: arbitrary-size symmetric payload exchange over bulk
|
||||
* RDMA, with interleaved 2MB TCP rounds as fallback (see ds4_tp.c). */
|
||||
int ds4_tp_big_gate_exchange(ds4_tp *tp, uint32_t layer, uint64_t seq,
|
||||
const void *out, void *in, uint64_t bytes);
|
||||
|
||||
/* Lockstep mirroring (leader side) and worker loop primitives. */
|
||||
int ds4_tp_send_sync(ds4_tp *tp, const int *tokens, uint32_t n_tokens);
|
||||
int ds4_tp_wait_sync_ack(ds4_tp *tp, char *err, size_t errlen);
|
||||
@@ -166,11 +181,10 @@ int ds4_tp_send_logits_half(ds4_tp *tp, const float *half, uint32_t count);
|
||||
int ds4_tp_recv_logits_half(ds4_tp *tp, float *half, uint32_t count);
|
||||
|
||||
/* Speculative verify mirroring. The leader announces a draft block right
|
||||
* before running the (replicated, ungated) batch verify; the worker runs
|
||||
* the same verify for its KV side effects and then blocks on the commit
|
||||
* frame, which carries the leader's decision: full_accept keeps the pushed
|
||||
* rows, otherwise both sides roll back and replay replay_n tokens through
|
||||
* the gated single-token decode in lockstep. */
|
||||
* before both ranks run the expert-split batch verify; the worker then blocks
|
||||
* on the commit frame, which carries the leader's decision: full_accept keeps
|
||||
* the pushed rows, otherwise both sides roll back and replay replay_n tokens
|
||||
* through the gated single-token decode in lockstep. */
|
||||
int ds4_tp_send_verify(ds4_tp *tp, const int *drafts, uint32_t n);
|
||||
int ds4_tp_send_verify_commit(ds4_tp *tp, int32_t full_accept, int32_t replay_n);
|
||||
int ds4_tp_recv_verify_commit(ds4_tp *tp, int32_t *full_accept, int32_t *replay_n);
|
||||
|
||||
@@ -1,4 +1,12 @@
|
||||
deepseek4-quantize
|
||||
gguf-requantize-dense
|
||||
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/**
|
||||
+20
-6
@@ -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
|
||||
@@ -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.
|
||||
@@ -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"
|
||||
|
||||
@@ -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]:
|
||||
|
||||
@@ -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`.
|
||||
@@ -0,0 +1 @@
|
||||
Here is a detailed explanation of B-tree insertion, covering the standard algorithm, node splits, and the special case of splitting
|
||||
@@ -0,0 +1,5 @@
|
||||
Here's a concise design for a TCP echo server that handles slow clients:
|
||||
|
||||
## Architecture
|
||||
|
||||
**Event-driven, non-block
|
||||
@@ -0,0 +1 @@
|
||||
Here's a comprehensive comparison of memory-mapped (mmap) model weights versus copying all weights into private buffers on macOS
|
||||
@@ -0,0 +1 @@
|
||||
This is an excellent question because it gets to the heart of why **RMSNorm** (Root Mean Square Normalization)
|
||||
@@ -0,0 +1 @@
|
||||
This is an excellent question that gets at a core optimization challenge in modern LLM deployment. Let's break down why **
|
||||
@@ -0,0 +1,5 @@
|
||||
Here's pseudocode for a binary heap push (insert) operation:
|
||||
|
||||
```
|
||||
function push(heap, value):
|
||||
//
|
||||
@@ -0,0 +1 @@
|
||||
This is a great question that gets at the core of how modern LLMs generate text. Both **temperature** and **
|
||||
@@ -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
|
||||
@@ -0,0 +1 @@
|
||||
This is an excellent question that gets to the heart of how Large Language Models (LLMs) are served efficiently. The
|
||||
@@ -0,0 +1,5 @@
|
||||
Here are three invariants for a special-token table used in a tokenizer:
|
||||
|
||||
---
|
||||
|
||||
### 1. **Un
|
||||
@@ -0,0 +1 @@
|
||||
Ecco una spiegazione dettagliata del funzionamento dell'inserimento in un **B-tree**, con
|
||||
@@ -0,0 +1,3 @@
|
||||
Ecco un progetto conciso per un server TCP echo con gestione di client lenti:
|
||||
|
||||
## Architettura
|
||||
@@ -0,0 +1 @@
|
||||
Ecco un confronto dettagliato tra **mmap** (memory-mapped file) e **copia
|
||||
@@ -0,0 +1 @@
|
||||
Certamente. Ecco una derivazione dettagliata del perché **RMSNorm** (Root Mean Square Normalization
|
||||
@@ -0,0 +1 @@
|
||||
La cache KV (Key-Value) su disco è particolarmente utile nelle **sessioni agentiche** (d
|
||||
@@ -0,0 +1 @@
|
||||
Ecco lo pseudocodice per inserire un elemento in un heap binario (min-heap o max-he
|
||||
@@ -0,0 +1 @@
|
||||
Questa è un'ottima domanda che tocca il cuore del controllo della generazione nei modelli linguistici
|
||||
@@ -0,0 +1 @@
|
||||
La decodifica speculativa mantiene l'output esatto perché utilizza un **meccanismo di ver
|
||||
@@ -0,0 +1 @@
|
||||
Questa è un'ottima domanda che tocca un aspetto fondamentale dell'ottimizzazione dell'inferenza
|
||||
@@ -0,0 +1 @@
|
||||
Ecco tre invarianti per una tabella di token speciali (ad esempio in un parser o in un sistema di
|
||||
@@ -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
|
||||
@@ -0,0 +1,4 @@
|
||||
Here's a short C function that clamps an integer to a range:
|
||||
|
||||
```c
|
||||
int clamp(int value, int
|
||||
@@ -0,0 +1 @@
|
||||
This is a fundamental question in concurrent programming. The core difference lies in **what they protect** and **how they coordinate
|
||||
@@ -0,0 +1 @@
|
||||
In the context of a network server, **backpressure** is a mechanism for **flow control**. It refers to the
|
||||
@@ -0,0 +1 @@
|
||||
Here is a safe and robust format for storing a model checkpoint header, designed for clarity, extensibility, and error prevention
|
||||
@@ -0,0 +1 @@
|
||||
A **ring buffer** (also called a circular buffer) is a fixed-size data structure that uses a contiguous block of
|
||||
@@ -0,0 +1,4 @@
|
||||
Here is a brief, structured plan for testing long-context prompt chunking.
|
||||
|
||||
### Goal
|
||||
Validate that chunking a
|
||||
@@ -0,0 +1 @@
|
||||
This is an excellent question that gets to the heart of why modern low-bit quantization (like 4-bit and 3
|
||||
@@ -0,0 +1 @@
|
||||
This is a great question that gets at a key performance optimization in modern large language model (LLM) inference, particularly
|
||||
@@ -0,0 +1 @@
|
||||
This is an excellent question because comparing quantizations is rarely a simple "A is better than B" scenario. The "
|
||||
@@ -0,0 +1,3 @@
|
||||
## Ricerca Binaria del Punto di Inserimento
|
||||
|
||||
La ricerca binaria del punto di inserimento è una
|
||||
@@ -0,0 +1,4 @@
|
||||
Ecco una breve funzione C che limita un intero a un intervallo specificato:
|
||||
|
||||
```c
|
||||
int clamp
|
||||
@@ -0,0 +1 @@
|
||||
Certamente. La differenza fondamentale tra un **mutex** (mutual exclusion) e un **contatore atom
|
||||
@@ -0,0 +1 @@
|
||||
In un server di rete, **backpressure** (contro-pressione o contropressione) è un me
|
||||
@@ -0,0 +1 @@
|
||||
Ecco un formato sicuro per salvare l'header di un checkpoint modello, progettato per garantire integ
|
||||
@@ -0,0 +1 @@
|
||||
Ti spiegherò il funzionamento di un **ring buffer** (o buffer circolare) e le tec
|
||||
@@ -0,0 +1 @@
|
||||
Ecco un piano breve e strutturato per testare il **chunking** (frammentazione) di
|
||||
@@ -0,0 +1 @@
|
||||
Nella quantizzazione low-bit (a bassa precisione) dei modelli di deep learning, una **matrice
|
||||
@@ -0,0 +1 @@
|
||||
Questa è un'ottima domanda che tocca un aspetto cruciale dell'ottimizzazione dei modelli
|
||||
@@ -0,0 +1 @@
|
||||
Per confrontare due quantizzazioni (ad esempio, due modelli di compressione di un segnale, due
|
||||
@@ -0,0 +1,3 @@
|
||||
Here are the first five measurements to take when diagnosing generation slowdowns at long context:
|
||||
|
||||
1. **Time-to-F
|
||||
@@ -0,0 +1 @@
|
||||
This is an excellent and fundamental question about how language models actually "think." The short answer is: **because the model
|
||||
@@ -0,0 +1,3 @@
|
||||
Let's break down **online softmax** for attention using simple variables.
|
||||
|
||||
## The Problem: Standard Softmax in Attention
|
||||
@@ -0,0 +1,5 @@
|
||||
Here's a minimal JSON schema for a tool call with name and arguments:
|
||||
|
||||
```json
|
||||
{
|
||||
"type": "
|
||||
@@ -0,0 +1 @@
|
||||
To test that a file-backed mmap is not repeatedly remapped, you need to verify that the kernel is reusing
|
||||
@@ -0,0 +1 @@
|
||||
This is a fascinating question that gets to the heart of how modern graphics APIs like Metal, Vulkan, and DirectX
|
||||
@@ -0,0 +1 @@
|
||||
In a Mixture-of-Experts (MoE) layer, the **router** (also called a **gating
|
||||
@@ -0,0 +1 @@
|
||||
Here is the difference between **raw KV rows** and **compressed KV rows** in the context of storage engines (
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -0,0 +1 @@
|
||||
Ecco cinque misure iniziali per affrontare un problema di decoding lento in contesti lunghi (long
|
||||
@@ -0,0 +1 @@
|
||||
Questa è un'ottima domanda che tocca il cuore di come funzionano i modelli linguistici
|
||||
@@ -0,0 +1 @@
|
||||
Certamente! Spiegherò la **softmax** nel contesto del meccanismo di **attenzione**
|
||||
@@ -0,0 +1,5 @@
|
||||
Ecco uno schema JSON minimo per una tool call:
|
||||
|
||||
```json
|
||||
{
|
||||
"name": "nome_della
|
||||
@@ -0,0 +1 @@
|
||||
Per testare che un `mmap` su file non venga rimappato ripetutamente, puoi
|
||||
@@ -0,0 +1 @@
|
||||
Questa è una domanda molto interessante che tocca il cuore dell'ottimizzazione della grafica e della gest
|
||||
@@ -0,0 +1 @@
|
||||
In un **layer mixture-of-experts (MoE)**, il **router** (o gate network) svolge
|
||||
@@ -0,0 +1 @@
|
||||
Certamente. Ecco una descrizione chiara della differenza tra righe KV (Key-Value) raw e
|
||||
@@ -0,0 +1 @@
|
||||
Ecco una checklist strutturata per validare che un modello quantizzato (es. GPTQ, AWQ,
|
||||
@@ -0,0 +1 @@
|
||||
Valutare 50 prompt invece di un singolo prompt preferito è più utile per diversi motivi, specialmente
|
||||
@@ -0,0 +1,4 @@
|
||||
Here's a tiny Python function that returns the median of three numbers:
|
||||
|
||||
```python
|
||||
def median_of_three(a,
|
||||
@@ -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
|
||||
@@ -0,0 +1 @@
|
||||
**Eventual consistency** is a consistency model used in distributed databases (and other distributed systems) where, given enough time
|
||||
@@ -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
|
||||
@@ -0,0 +1,2 @@
|
||||
This is a fundamental concept in computer science, networking, and system design. The simplest analogy is a **water pipe**.
|
||||
|
||||
@@ -0,0 +1 @@
|
||||
A **trie** (also called a prefix tree) represents a set of strings by organizing them based on their **
|
||||
@@ -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
|
||||
@@ -0,0 +1 @@
|
||||
This is an excellent question that gets to the heart of how virtual memory works. The short answer is that **`mm
|
||||
@@ -0,0 +1 @@
|
||||
This is a nuanced and insightful question that touches on the intersection of model compression (quantization) and the dynamics of Mi
|
||||
@@ -0,0 +1 @@
|
||||
Here is a short, focused checklist for reviewing a PR that involves memory lifetimes (e.g., Rust, C++, or
|
||||
@@ -0,0 +1,4 @@
|
||||
Ecco una semplice funzione Python per calcolare la mediana di tre numeri:
|
||||
|
||||
```python
|
||||
def mediana(a
|
||||
@@ -0,0 +1 @@
|
||||
Questa è una domanda fondamentale in informatica. La risposta breve è: **perché permettono di us
|
||||
@@ -0,0 +1 @@
|
||||
In un database distribuito, la **consistenza eventuale** (o *eventual consistency*) è un modello
|
||||
Loaded 100 of 967 files, more files were not shown because too many files have changed in this diff.
Show more
Reference in new issue
Block a user