31 Commits
Author SHA1 Message Date
antirez 448d5695d1 Tune Laguna sampling defaults 2026-07-27 12:56:54 +02:00
antirez d8d9688584 Merge Laguna CUDA and cross-backend DFlash 2026-07-27 00:35:40 +02:00
Salvatore Sanfilippo 8f620f31e2 Tune DFlash confidence and exact verifier kernels 2026-07-26 15:07:20 -07:00
Salvatore Sanfilippo 5f3c78da2e Pipeline DFlash verification across GPU backends 2026-07-26 15:07:20 -07:00
Salvatore Sanfilippo 0ee6d45650 Support quantized DFlash draft models 2026-07-26 15:07:20 -07:00
antirez 31db3a87c5 Accelerate Laguna DFlash on CUDA 2026-07-26 22:46:49 +02:00
antirez 30f4e1785c Optimize Laguna CUDA generation on DGX Spark 2026-07-26 15:33:23 +02:00
antirez 6aec2354be Add fast Laguna CUDA support for DGX Spark 2026-07-26 14:49:31 +02:00
antirez 62b59a48e2 Speed up Laguna ROCm decoding and DFlash 2026-07-26 10:59:32 +02:00
Salvatore Sanfilippo bf605b17d8 Add Laguna DFlash speculative decoding on Metal
Load Poolside's standalone DFlash GGUF, draft and verify greedy Laguna tokens on Metal, expose CLI/agent/server controls, pause speculation when measured throughput loses, and sanitize/test non-finite target features.
2026-07-25 03:19:27 -07:00
Salvatore Sanfilippo 5c5f95a58d Harden tagged agent tool-call decoding 2026-07-25 03:10:44 -07:00
Salvatore Sanfilippo 6bd7fd37df Add Laguna ROCm inference
Support Laguna Q2/Q3 and Q4 routed experts on ROCm, accelerate decode projections and Q8 paths, enable the agent backend, and document the supported model variants.
2026-07-25 03:10:44 -07:00
Salvatore Sanfilippo 3b0ec5f1d8 Add Laguna Q2/Q3 download target 2026-07-24 03:49:37 -07:00
Salvatore Sanfilippo 938227a248 Support mixed Q3_K Laguna experts 2026-07-24 03:47:39 -07:00
Salvatore Sanfilippo 92347005d2 Support Q2_K routed Laguna experts 2026-07-24 02:59:13 -07:00
Salvatore Sanfilippo 7e3dbef7e3 Support revised Laguna Q4_K_M quant 2026-07-23 09:48:33 -07:00
Salvatore Sanfilippo 8a927009e6 Add native Laguna S 2.1 support
Implement Laguna S 2.1 model loading, chat and tool protocols, optimized Metal decode and layer-major prefill, long-context attention, CLI/agent/server integration, and backend compatibility guards. Add the official Q4_K_M downloader, usage documentation, regression coverage, and OpenRouter quality fixtures.
2026-07-23 04:53:53 -07:00
antirez efdadd41e2 README: update the AI assisted models list. 2026-07-20 16:50:40 +02:00
antirez 646db66656 README: remove title now that there is the logo. 2026-07-20 16:41:44 +02:00
antirez cb5103449b Add DwarfStar logo and credits 2026-07-20 16:40:44 +02:00
Ivan Fioravanti 427e281d24 Optimize Metal prefill and decode kernels 2026-07-20 11:09:49 +02:00
antirez e0824ddcaa Make release builds warning-free 2026-07-20 11:09:28 +02:00
antirez a185c36dc4 Stabilize CUDA and Metal session batching 2026-07-18 18:50:59 +02:00
antirez ef8d923ad7 Add ROCm GLM 5.2 support 2026-07-18 18:50:30 +02:00
antirez 66cbce5289 Unify distributed inference CLI and documentation 2026-07-18 18:50:24 +02:00
antirez 36ee8c1407 Add CUDA tensor parallelism and session batching 2026-07-18 18:50:03 +02:00
antirez 36cd0ca915 Add native Metal session batching 2026-07-18 18:49:00 +02:00
antirez 005afedcd9 Add GLM 5.2 inference and quality fixtures 2026-07-18 18:48:45 +02:00
antirez 63d9874c54 Add two-machine Metal tensor parallelism 2026-07-18 18:47:58 +02:00
antirez fc9efd1cc7 Add DSpark speculative decoding 2026-07-18 18:47:26 +02:00
antirez 519c4d8fe7 Fix PRO streaming and sampling correctness 2026-07-18 18:46:19 +02:00
1319 changed files with 1357062 additions and 7930 deletions

No files matched your search

+3
View File
@@ -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
+17
View File
@@ -4,11 +4,28 @@
/ds4-eval
/ds4-agent
/ds4_native
/ds4_agent_test
/ds4_server_test
/ds4_test
/ds4flash.gguf
/TODO.md
/gguf/
/core
/core.*
/tests/test_engine_mgpu_placement
/tests/test_engine_mgpu_runtime
/tests/test_engine_mgpu_refusal
/tests/test_engine_correctness
/tests/test_gpu_args
/tests/test_gpu_lookup_cache_strict
/tests/test_gpu_model_cache
/tests/test_gpu_xdev
/tests/test_cuda_mixed_batch
/tests/test_cuda_session_batch
/tests/test_layer_pack
/tests/test_metal_session_batch
/tests/test_q4k_dot
/tests/test_sampling
*.o
*.dSYM/
__pycache__/
+194 -32
View File
@@ -3,22 +3,29 @@ UNAME_S := $(shell uname -s)
ifeq ($(UNAME_S),Darwin)
NATIVE_CPU_FLAG ?= -mcpu=native
SAMPLING_TEST :=
else
NATIVE_CPU_FLAG ?= -march=native
SAMPLING_TEST := tests/test_sampling
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)
ROCM_SRCS := $(wildcard rocm/*.cuh)
DS4_TEST_MODEL ?= ds4flash.gguf
DS4_TEST_MTP ?= gguf/DeepSeek-V4-Flash-MTP-Q4K-Q8_0-F32.gguf
DS4_DSPARK_MODEL ?= $(DS4_TEST_MODEL)
DS4_DSPARK_SUPPORT ?= gguf/DeepSeek-V4-Flash-DSpark-support.gguf
ifeq ($(UNAME_S),Darwin)
METAL_LDLIBS := $(LDLIBS) -framework Foundation -framework Metal
CORE_OBJS = ds4.o ds4_distributed.o ds4_ssd.o ds4_metal.o
CPU_CORE_OBJS = ds4_cpu.o ds4_distributed.o ds4_ssd.o
CORE_OBJS = ds4.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_metal.o ds4_layer_pack.o
CPU_CORE_OBJS = ds4_cpu.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_layer_pack.o
else
CFLAGS += -D_GNU_SOURCE -fno-finite-math-only
CUDA_HOME ?= /usr/local/cuda
@@ -28,8 +35,8 @@ ifneq ($(strip $(CUDA_ARCH)),)
NVCC_ARCH_FLAGS := -arch=$(CUDA_ARCH)
endif
NVCCFLAGS ?= -O3 -g -lineinfo --use_fast_math $(NVCC_ARCH_FLAGS) -Xcompiler $(NATIVE_CPU_FLAG) -Xcompiler -pthread
CORE_OBJS = ds4.o ds4_distributed.o ds4_ssd.o ds4_cuda.o
CPU_CORE_OBJS = ds4_cpu.o ds4_distributed.o ds4_ssd.o
CORE_OBJS = ds4.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_cuda.o ds4_layer_pack.o
CPU_CORE_OBJS = ds4_cpu.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_layer_pack.o
CUDA_LDLIBS ?= -lm -Xcompiler -pthread -L$(CUDA_HOME)/targets/sbsa-linux/lib -L$(CUDA_HOME)/lib64 -lcudart -lcublas
HIPCC ?= $(shell command -v hipcc 2>/dev/null || echo /opt/rocm/bin/hipcc)
ROCM_ARCH ?= gfx1151
@@ -40,7 +47,7 @@ DS4_LINK_LIBS ?= $(CUDA_LDLIBS)
METAL_LDLIBS := $(LDLIBS)
endif
.PHONY: all help clean test cpu cuda cuda-spark cuda-generic cuda-regression strix-halo rocm
.PHONY: all help clean test test-metal-session-batch test-cuda-session-batch test-cuda-mixed-batch dspark-acceptance dspark-verify-depth mtp-verify-depth cpu cuda cuda-spark cuda-generic cuda-regression strix-halo rocm
ifeq ($(UNAME_S),Darwin)
all: ds4 ds4-server ds4-bench ds4-eval ds4-agent
@@ -50,29 +57,43 @@ help:
@echo " make Build Metal ./ds4, ./ds4-server, ./ds4-bench, ./ds4-eval, and ./ds4-agent"
@echo " make cpu Build CPU-only ./ds4, ./ds4-server, ./ds4-bench, ./ds4-eval, and ./ds4-agent"
@echo " make test Build and run tests"
@echo " make dspark-verify-depth Run DSpark speculative verification smoke if support GGUF is present"
@echo " make mtp-verify-depth Run legacy MTP speculative verification smoke if MTP GGUF is present"
@echo " make clean Remove build outputs"
ds4: ds4_cli.o ds4_help.o linenoise.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_cli.o ds4_help.o linenoise.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4: ds4_cli.o ds4_help.o linenoise.o ds4_gpu_args.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_cli.o ds4_help.o linenoise.o ds4_gpu_args.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4-server: ds4_server.o ds4_help.o ds4_kvstore.o rax.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_server.o ds4_help.o ds4_kvstore.o rax.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4-server: ds4_server.o ds4_help.o ds4_kvstore.o rax.o ds4_gpu_args.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_server.o ds4_help.o ds4_kvstore.o rax.o ds4_gpu_args.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4-bench: ds4_bench.o ds4_help.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_bench.o ds4_help.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4-bench: ds4_bench.o ds4_help.o ds4_gpu_args.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_bench.o ds4_help.o ds4_gpu_args.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4-eval: ds4_eval.o ds4_help.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_eval.o ds4_help.o $(CORE_OBJS) $(METAL_LDLIBS)
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)
ds4-agent: ds4_agent.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o ds4_gpu_args.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_agent.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o ds4_gpu_args.o $(CORE_OBJS) $(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)
$(CC) $(CFLAGS) -o ds4-bench ds4_bench_cpu.o ds4_help.o $(CPU_CORE_OBJS) $(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)
tests/test_metal_session_batch.o: tests/test_metal_session_batch.c ds4.h
$(CC) $(CFLAGS) -I. -c -o $@ tests/test_metal_session_batch.c
tests/test_metal_session_batch: tests/test_metal_session_batch.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ $^ $(METAL_LDLIBS)
test-metal-session-batch: tests/test_metal_session_batch
DS4_TEST_MODEL="$(DS4_TEST_MODEL)" ./tests/test_metal_session_batch
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 ds4_gpu_args_cpu.o $(CPU_CORE_OBJS)
$(CC) $(CFLAGS) -o ds4 ds4_cli_cpu.o ds4_help.o linenoise.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-server ds4_server_cpu.o ds4_help.o ds4_kvstore.o rax.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-bench ds4_bench_cpu.o ds4_help.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-eval ds4_eval_cpu.o ds4_help.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-agent ds4_agent_cpu.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-agent ds4_agent_cpu.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
cuda-regression:
@echo "cuda-regression requires a CUDA build"
@@ -88,6 +109,8 @@ help:
@echo " make rocm Alias for make strix-halo"
@echo " make cpu Build CPU-only ./ds4, ./ds4-server, ./ds4-bench, ./ds4-eval, and ./ds4-agent"
@echo " make test Build and run tests"
@echo " make dspark-verify-depth Run DSpark speculative verification smoke if support GGUF is present"
@echo " make mtp-verify-depth Run legacy MTP speculative verification smoke if MTP GGUF is present"
@echo " make clean Remove build outputs"
cuda-spark:
@@ -106,34 +129,37 @@ cuda:
strix-halo:
$(MAKE) -B ds4 ds4-server ds4-bench ds4-eval ds4-agent \
CORE_OBJS="ds4.o ds4_distributed.o ds4_ssd.o ds4_rocm.o" \
CORE_OBJS="ds4.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_rocm.o ds4_rocm_compat.o ds4_rocm_unavailable.o ds4_layer_pack.o" \
CFLAGS="$(CFLAGS) -DDS4_ROCM_BUILD" \
DS4_LINK="$(HIPCC) $(ROCM_CFLAGS)" \
DS4_LINK_LIBS="$(ROCM_LDLIBS)"
rocm: strix-halo
ds4: ds4_cli.o ds4_help.o linenoise.o $(CORE_OBJS)
ds4: ds4_cli.o ds4_help.o linenoise.o ds4_gpu_args.o $(CORE_OBJS)
$(DS4_LINK) -o $@ $^ $(DS4_LINK_LIBS)
ds4-server: ds4_server.o ds4_help.o ds4_kvstore.o rax.o $(CORE_OBJS)
ds4-server: ds4_server.o ds4_help.o ds4_kvstore.o rax.o ds4_gpu_args.o $(CORE_OBJS)
$(DS4_LINK) -o $@ $^ $(DS4_LINK_LIBS)
ds4-bench: ds4_bench.o ds4_help.o $(CORE_OBJS)
ds4-bench: ds4_bench.o ds4_help.o ds4_gpu_args.o $(CORE_OBJS)
$(DS4_LINK) -o $@ $^ $(DS4_LINK_LIBS)
ds4-eval: ds4_eval.o ds4_help.o $(CORE_OBJS)
$(DS4_LINK) -o $@ $^ $(DS4_LINK_LIBS)
ds4-agent: ds4_agent.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o $(CORE_OBJS)
ds4-agent: ds4_agent.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o ds4_gpu_args.o $(CORE_OBJS)
$(DS4_LINK) -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)
$(CC) $(CFLAGS) -o ds4-bench ds4_bench_cpu.o ds4_help.o $(CPU_CORE_OBJS) $(LDLIBS)
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 ds4_gpu_args_cpu.o $(CPU_CORE_OBJS)
$(CC) $(CFLAGS) -o ds4 ds4_cli_cpu.o ds4_help.o linenoise.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-server ds4_server_cpu.o ds4_help.o ds4_kvstore.o rax.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-bench ds4_bench_cpu.o ds4_help.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-eval ds4_eval_cpu.o ds4_help.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-agent ds4_agent_cpu.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o $(CPU_CORE_OBJS) $(LDLIBS)
$(CC) $(CFLAGS) -o ds4-agent ds4_agent_cpu.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o ds4_gpu_args_cpu.o $(CPU_CORE_OBJS) $(LDLIBS)
cuda-regression: tests/cuda_long_context_smoke
./tests/cuda_long_context_smoke
@@ -151,9 +177,15 @@ ds4_cli.o: ds4_cli.c ds4.h ds4_ssd.h ds4_distributed.h ds4_help.h linenoise.h
ds4_distributed.o: ds4_distributed.c ds4_distributed.h ds4.h ds4_ssd.h
$(CC) $(CFLAGS) -c -o $@ ds4_distributed.c
ds4_tp.o: ds4_tp.c ds4_tp.h ds4.h ds4_ssd.h
$(CC) $(CFLAGS) -c -o $@ ds4_tp.c
ds4_help.o: ds4_help.c ds4_help.h
$(CC) $(CFLAGS) -c -o $@ ds4_help.c
ds4_gpu_args.o: ds4_gpu_args.c ds4_gpu_args.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -c -o $@ ds4_gpu_args.c
ds4_server.o: ds4_server.c ds4.h ds4_ssd.h ds4_distributed.h ds4_help.h ds4_kvstore.h rax.h
$(CC) $(CFLAGS) -c -o $@ ds4_server.c
@@ -188,11 +220,14 @@ linenoise.o: linenoise.c linenoise.h
$(CC) $(CFLAGS) -c -o $@ linenoise.c
ds4_cpu.o: ds4.c ds4.h ds4_ssd.h ds4_distributed.h ds4_gpu.h
$(CC) $(CFLAGS) -DDS4_NO_GPU -c -o $@ ds4.c
$(CC) $(CFLAGS) -Wno-unused-function -DDS4_NO_GPU -c -o $@ ds4.c
ds4_cli_cpu.o: ds4_cli.c ds4.h ds4_ssd.h ds4_distributed.h ds4_help.h linenoise.h
$(CC) $(CFLAGS) -DDS4_NO_GPU -c -o $@ ds4_cli.c
ds4_gpu_args_cpu.o: ds4_gpu_args.c ds4_gpu_args.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -DDS4_NO_GPU -c -o $@ ds4_gpu_args.c
ds4_server_cpu.o: ds4_server.c ds4.h ds4_ssd.h ds4_distributed.h ds4_help.h ds4_kvstore.h rax.h
$(CC) $(CFLAGS) -DDS4_NO_GPU -c -o $@ ds4_server.c
@@ -208,15 +243,108 @@ ds4_agent_cpu.o: ds4_agent.c ds4.h ds4_ssd.h ds4_distributed.h ds4_help.h ds4_kv
ds4_metal.o: ds4_metal.m ds4_gpu.h $(METAL_SRCS)
$(CC) $(OBJCFLAGS) -c -o $@ ds4_metal.m
ds4_cuda.o: ds4_cuda.cu ds4_gpu.h ds4_iq2_tables_cuda.inc
ds4_cuda.o: ds4_cuda.cu ds4_gpu.h ds4_gpu_mgpu.h ds4_iq2_tables_cuda.inc \
rocm/ds4_rocm_laguna.cuh rocm/ds4_rocm_dflash.cuh
$(NVCC) $(NVCCFLAGS) -c -o $@ ds4_cuda.cu
ds4_rocm.o: ds4_rocm.cu ds4_gpu.h ds4_iq2_tables_cuda.inc $(ROCM_SRCS)
$(HIPCC) $(ROCM_CFLAGS) -c -o $@ ds4_rocm.cu
ds4_rocm_compat.o: ds4_rocm_compat.cu ds4_gpu.h ds4_gpu_mgpu.h ds4_gpu_args.h
$(HIPCC) $(ROCM_CFLAGS) -c -o $@ ds4_rocm_compat.cu
ds4_rocm_unavailable.o: ds4_rocm_unavailable.cu
$(HIPCC) $(ROCM_CFLAGS) -c -o $@ ds4_rocm_unavailable.cu
tests/cuda_long_context_smoke: tests/cuda_long_context_smoke.o ds4_cuda.o
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
tests/test_layer_pack.o: tests/test_layer_pack.c ds4_layer_pack.h
$(CC) $(CFLAGS) -I. -c -o $@ $<
tests/test_layer_pack: tests/test_layer_pack.o ds4_layer_pack.o
$(CC) $(CFLAGS) -o $@ $^ $(LDLIBS)
tests/test_gpu_args.o: tests/test_gpu_args.c ds4_gpu_args.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -I. -DDS4_NO_GPU -c -o $@ $<
tests/test_gpu_args: tests/test_gpu_args.o ds4_gpu_args_cpu.o
$(CC) $(CFLAGS) -o $@ $^ $(LDLIBS)
ds4_cpu_test_hooks.o: ds4.c ds4.h ds4_gpu.h ds4_gpu_mgpu.h ds4_layer_pack.h
$(CC) $(CFLAGS) -Wno-unused-function -DDS4_NO_GPU -DDS4_TEST_HOOKS -c -o $@ ds4.c
tests/test_engine_mgpu_placement.o: tests/test_engine_mgpu_placement.c ds4.h ds4_gpu_mgpu.h ds4_layer_pack.h
$(CC) $(CFLAGS) -I. -c -o $@ $<
tests/test_engine_mgpu_placement: tests/test_engine_mgpu_placement.o ds4_cpu_test_hooks.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_layer_pack.o
$(CC) $(CFLAGS) -o $@ $^ $(LDLIBS)
ifneq ($(UNAME_S),Darwin)
tests/test_gpu_xdev.o: tests/test_gpu_xdev.c ds4_gpu.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_gpu_xdev: tests/test_gpu_xdev.o ds4_cuda.o
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
tests/test_gpu_model_cache.o: tests/test_gpu_model_cache.c ds4_gpu.h
$(CC) $(CFLAGS) -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_gpu_model_cache: tests/test_gpu_model_cache.o ds4_cuda.o
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
tests/test_gpu_lookup_cache_strict.o: tests/test_gpu_lookup_cache_strict.c ds4_gpu.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_gpu_lookup_cache_strict: tests/test_gpu_lookup_cache_strict.o ds4_cuda.o
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
ds4_cuda_test_hooks.o: ds4.c ds4.h ds4_gpu.h ds4_gpu_mgpu.h ds4_layer_pack.h
$(CC) $(CFLAGS) -Wno-unused-function -DDS4_TEST_HOOKS -I$(CUDA_HOME)/include -c -o $@ ds4.c
tests/test_engine_mgpu_refusal.o: tests/test_engine_mgpu_refusal.c ds4.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_engine_mgpu_refusal: tests/test_engine_mgpu_refusal.o ds4_gpu_args.o ds4_kvstore.o rax.o $(CORE_OBJS)
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
tests/test_engine_mgpu_runtime.o: tests/test_engine_mgpu_runtime.c ds4.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -DDS4_TEST_HOOKS -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_engine_mgpu_runtime: tests/test_engine_mgpu_runtime.o ds4_cuda_test_hooks.o ds4_gpu_args.o ds4_kvstore.o rax.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_cuda.o ds4_layer_pack.o
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
tests/test_engine_correctness.o: tests/test_engine_correctness.c ds4.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_engine_correctness: tests/test_engine_correctness.o ds4_gpu_args.o ds4_kvstore.o rax.o $(CORE_OBJS)
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
tests/test_sampling.o: tests/test_sampling.c ds4.h
$(CC) $(CFLAGS) -DDS4_TEST_HOOKS -I. -c -o $@ $<
tests/test_sampling: tests/test_sampling.o ds4_cuda_test_hooks.o ds4_gpu_args.o ds4_kvstore.o rax.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_cuda.o ds4_layer_pack.o
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
tests/test_cuda_session_batch.o: tests/test_cuda_session_batch.c ds4.h ds4_gpu_args.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_cuda_session_batch: tests/test_cuda_session_batch.o ds4_gpu_args.o ds4_kvstore.o rax.o $(CORE_OBJS)
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
test-cuda-session-batch: tests/test_cuda_session_batch
DS4_TEST_MODEL="$(DS4_TEST_MODEL)" ./tests/test_cuda_session_batch
tests/test_cuda_mixed_batch.o: tests/test_cuda_mixed_batch.c ds4.h ds4_gpu_args.h ds4_gpu_mgpu.h
$(CC) $(CFLAGS) -DDS4_TEST_HOOKS -I. -I$(CUDA_HOME)/include -c -o $@ $<
tests/test_cuda_mixed_batch: tests/test_cuda_mixed_batch.o ds4_cuda_test_hooks.o ds4_gpu_args.o ds4_kvstore.o rax.o ds4_distributed.o ds4_tp.o ds4_ssd.o ds4_cuda.o ds4_layer_pack.o
$(NVCC) $(NVCCFLAGS) -o $@ $^ $(CUDA_LDLIBS)
test-cuda-mixed-batch: tests/test_cuda_mixed_batch
DS4_TEST_MODEL="$(DS4_TEST_MODEL)" ./tests/test_cuda_mixed_batch
endif
ds4_test: ds4_test.o ds4_help.o ds4_kvstore.o rax.o $(CORE_OBJS)
ifeq ($(UNAME_S),Darwin)
$(CC) $(CFLAGS) -o $@ ds4_test.o ds4_help.o ds4_kvstore.o rax.o $(CORE_OBJS) $(METAL_LDLIBS)
@@ -231,14 +359,48 @@ else
$(NVCC) $(NVCCFLAGS) -o $@ ds4_agent_test.o ds4_help.o ds4_web.o ds4_kvstore.o linenoise.o $(CORE_OBJS) $(CUDA_LDLIBS)
endif
test: ds4_test ds4_agent_test ds4-eval q4k-dot-test
test: ds4_test ds4_agent_test ds4-eval q4k-dot-test \
tests/test_layer_pack tests/test_engine_mgpu_placement tests/test_gpu_args \
$(SAMPLING_TEST) ds4 ds4-server ds4-bench ds4-agent
./ds4-eval --self-test-extractors
./ds4_agent_test
./ds4_test
./tests/test_layer_pack
./tests/test_engine_mgpu_placement
./tests/test_gpu_args
./tests/test_gpu_args_cli.sh
ifneq ($(UNAME_S),Darwin)
./tests/test_sampling
endif
dspark-acceptance: ds4
DS4_DSPARK_MODEL="$(DS4_DSPARK_MODEL)" \
DS4_DSPARK_SUPPORT="$(DS4_DSPARK_SUPPORT)" \
sh tests/dspark_acceptance_fixture.sh
dspark-verify-depth: ds4_test
@if [ ! -f "$(DS4_TEST_MODEL)" ]; then \
echo "dspark-verify-depth: skipped, missing model $(DS4_TEST_MODEL)"; \
elif [ ! -f "$(DS4_DSPARK_SUPPORT)" ]; then \
echo "dspark-verify-depth: skipped, missing DSpark support $(DS4_DSPARK_SUPPORT)"; \
echo "dspark-verify-depth: run make dspark-support or set DS4_DSPARK_SUPPORT=FILE"; \
else \
DS4_TEST_MODEL="$(DS4_TEST_MODEL)" DS4_TEST_DSPARK="$(DS4_DSPARK_SUPPORT)" ./ds4_test --dspark-verify-depth; \
fi
mtp-verify-depth: ds4_test
@if [ ! -f "$(DS4_TEST_MODEL)" ]; then \
echo "mtp-verify-depth: skipped, missing model $(DS4_TEST_MODEL)"; \
elif [ ! -f "$(DS4_TEST_MTP)" ]; then \
echo "mtp-verify-depth: skipped, missing MTP support $(DS4_TEST_MTP)"; \
echo "mtp-verify-depth: run ./download_model.sh mtp or set DS4_TEST_MTP=FILE"; \
else \
DS4_TEST_MODEL="$(DS4_TEST_MODEL)" DS4_TEST_MTP="$(DS4_TEST_MTP)" ./ds4_test --mtp-verify-depth; \
fi
q4k-dot-test: tests/test_q4k_dot.c
$(CC) -O2 -Wall -Wextra -std=c99 -o tests/test_q4k_dot tests/test_q4k_dot.c -lm -pthread
./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 tests/test_metal_session_batch tests/test_gpu_xdev tests/test_gpu_model_cache tests/test_gpu_lookup_cache_strict tests/test_engine_mgpu_refusal tests/test_engine_mgpu_runtime tests/test_engine_correctness tests/test_sampling tests/test_cuda_session_batch tests/test_cuda_mixed_batch tests/*.o *.o tests/cuda_long_context_smoke tests/cuda_long_context_smoke.o
+255 -14
View File
@@ -12,7 +12,7 @@ run.
Preferred release test hosts:
- CUDA / DGX Spark: `toor@192.168.0.180`.
- CUDA / DGX Spark: `toor@192.168.60.184`.
- Metal / distributed Mac testing: `mac-m5max-it` and `mac-m5max-us`.
- ROCm: The Strix Halo system at antirez@strixhalo (Framework Desktop).
@@ -36,6 +36,15 @@ in this system.
`make clean && make`.
- Build CPU-only binaries as a compile check only:
`make clean && make cpu`.
- Treat compiler warnings as build failures. Save each release and test build's
complete output and require no `warning:` or NVCC `warning #` lines. Fix the
source when possible; use a narrow target-specific suppression only when a
test deliberately compiles a partial translation unit.
- Repeat the warning-free build gate on the release hardware:
`make clean && make` on Metal,
`make clean && make cuda-spark` on DGX Spark,
`make clean && make cuda-generic CUDA_HOME=/usr` on the multi-GPU CUDA host,
`make clean && make strix-halo` on Strix Halo.
- Run whitespace checks before committing:
`git diff --check`.
- Confirm `./ds4 --help`, `./ds4-server --help`, and `./ds4-agent --help` render
@@ -45,6 +54,13 @@ in this system.
- Run the default suite:
`make test`.
- Run `tests/test_gpu_args_cli.sh` explicitly after changing executable option
parsing or multi-GPU placement. Invalid values and device/budget count
mismatches must reach the shared GPU parser in all four binaries; an
`unknown option` response from a binary that advertises the flag is a
release blocker. On CUDA, also start `ds4-server` once with
`--gpu-vram auto` and the intended `--gpu-devices` list and preserve the
resolved layout line.
- Run the vector checks explicitly after any tokenizer, template, KV, kernel,
quantization, or prompt-rendering change:
`./ds4_test --logprob-vectors`
@@ -54,7 +70,43 @@ in this system.
`./ds4_test --server`.
- Run `./ds4-eval --self-test-extractors`.
## 3. Metal Flash Path
## 3. Official Continuation Quality Gates
These tests are release-blocking after tokenizer, template, KV-cache, attention,
MoE routing, quantization, logit, or model-graph changes. They are
teacher-forced continuation checks against hosted-model output and API
top-logprob slices, so do not replace them with one sampled chat answer.
- Build the scorer:
`make -C gguf-tools quality-score`.
- Run the tracked DeepSeek V4 Flash smoke vectors:
`./ds4_test --logprob-vectors`.
This covers the official Flash API vectors in `tests/test-vectors/`, including
short prompts and long-prompt attention cases.
- Run the 100-case DeepSeek V4 Flash fixture for every released Flash GGUF:
`gguf-tools/quality-testing/score_official /path/to/deepseek-v4-flash.gguf gguf-tools/quality-testing/data/flash/manifest.tsv /tmp/flash.tsv 4096`.
- Run the 100-case GLM 5.2 OpenRouter fixture for every released GLM GGUF:
`gguf-tools/quality-testing/score_official models/GLM-5.2-UD-Q4_K_XL.gguf gguf-tools/quality-testing/data/glm52-openrouter-100/manifest.tsv /tmp/glm52-q4.tsv 4096`.
Current Q4 XL reference band: first-token match `95/100`, API top-1 agreement
about `0.942`, and API pair-order agreement about `0.880`.
- Run the same GLM fixture for reduced-precision GLM release files. The Q2
routed reference is lower quality but should stay near first-token match
`92/100`, API top-1 agreement about `0.890`, and API pair-order agreement
about `0.800` unless the quantization changed deliberately.
- Run the 100-case DeepSeek V4 PRO fixture for every released PRO GGUF:
`gguf-tools/quality-testing/score_official /path/to/deepseek-v4-pro.gguf gguf-tools/quality-testing/data/pro/manifest.tsv /tmp/pro.tsv 4096`.
- For SSD streaming, run the same official-continuation scorer once with full
residency and once with `--ssd-streaming` for the release model. The summary
and API agreement should stay in the same quality band.
- Compare any candidate against the previous release or last-known-good output:
`python3 gguf-tools/quality-testing/compare_scores.py /tmp/old.tsv /tmp/new.tsv`.
Treat a large first-token-match drop, a clear NLL regression, or a material
API top-1/pair-order regression as a blocker unless the release notes call out
an intentional quality tradeoff.
- Keep the raw `summary` and `api_summary` lines in the release notes or QA log.
Do not use stale manifests from `misc/` as release evidence.
## 4. Metal Flash Path
Use the normal Flash GGUF that 128 GB users run.
@@ -73,7 +125,78 @@ Use the normal Flash GGUF that 128 GB users run.
generation speed, and KV bytes with the last known good numbers for the same
machine.
## 4. Metal PRO Path
### DSpark / DeepSpec Runtime
DSpark is opt-in, but it mutates the verifier, target-hidden capture, support
model loading, and scheduler paths. Run these whenever DSpark support,
speculative verification, confidence/scheduler policy, target hidden capture,
tiny routed-MoE verifier kernels, or shared `--mtp` support-model code changes:
- Default greedy acceptance fixture:
`DS4_DSPARK_MODEL=/Users/antirez/ds4/ds4flash.gguf DS4_DSPARK_SUPPORT=/Users/antirez/ds4/gguf/DeepSeek-V4-Flash-DSpark-support.gguf make dspark-acceptance`.
- 64-token guardrail:
`DS4_DSPARK_FIXTURE_TOKENS=64 DS4_DSPARK_MODEL=/Users/antirez/ds4/ds4flash.gguf DS4_DSPARK_SUPPORT=/Users/antirez/ds4/gguf/DeepSeek-V4-Flash-DSpark-support.gguf make dspark-acceptance`.
- Fixed-block partial-accept fallback:
`DS4_DSPARK_FIXTURE_CONFIDENCE=0 DS4_DSPARK_FIXTURE_TOKENS=8 DS4_DSPARK_FIXTURE_REQUIRE_PARTIAL=1 DS4_DSPARK_MODEL=/Users/antirez/ds4/ds4flash.gguf DS4_DSPARK_SUPPORT=/Users/antirez/ds4/gguf/DeepSeek-V4-Flash-DSpark-support.gguf make dspark-acceptance`.
- DSpark verifier invariant smoke:
`DS4_TEST_MODEL=/Users/antirez/ds4/ds4flash.gguf DS4_DSPARK_SUPPORT=/Users/antirez/ds4/gguf/DeepSeek-V4-Flash-DSpark-support.gguf make dspark-verify-depth`.
- If shared support-model or verifier structures changed, also run legacy MTP:
`make mtp-verify-depth` with `DS4_TEST_MTP` set to a one-stage MTP support
GGUF, or confirm the target skips only because the optional file is missing.
- Record `c_add` `accepted_draft`, `errors=0`, `verify_layer`, and `net_saved`
for both 32-token and 64-token runs. A faster run with lower proposal
quality is a regression unless it was an intentional scheduler change.
- If verifier MoE kernels changed, run one diagnostic `c_add` profile with
`DS4_DSPARK_VERIFY_SELECTED_PROFILE=1` or the Metal MoE stage profiler and
record the selected-expert footprint or stage timing in the DSpark log.
### Session Microbatching And Metal TP
Run these gates whenever session scheduling, batched decode, mixed
prefill/decode, QKV projection, shared or routed experts, tensor parallelism,
or backend fallback selection changes.
- On a single Metal machine, run the full-vocabulary exact-logit oracle with
2, 4, 8, and 16 sessions:
`DS4_TEST_MODEL=/path/to/ds4flash.gguf DS4_TEST_SESSION_COUNT=N make test-metal-session-batch`.
Compatible resident Q8 runs must report `native_shared=1 native_qkv=1` at
every tested count. The 16-session run covers row counts above the old
artificial eight-row limit.
- Repeat the four-session oracle with
`DS4_METAL_SESSION_BATCH_SHARED=0` and with
`DS4_METAL_SESSION_BATCH_QKV=0`. The first run must use the complete fallback;
the second may batch the shared expert only. Both must remain bit-exact.
- The oracle must cover reversed row ordering, at least six decode steps, and a
mixed prefill/decode call. Any nonzero differing-logit count is a blocker;
argmax-only agreement is insufficient.
- Benchmark 1, 2, 4, 8, and 16 simultaneous resident sessions on the same host
and model. Record model-step latency and aggregate decode tokens/second, not
only request completion speed. The current Metal path batches QKV and part of
the shared expert, but still runs attention, routed experts, shared down, and
the output head per session. Treat flat aggregate scaling as unfinished
implementation work, not evidence that Metal cannot benefit from batching.
- On `mac-m5max-it` and `mac-m5max-us`, run the same oracle in physical TP mode
over explicit `tcp` and `rdma` transports. Set `DS4_TEST_TP_MODE=leader` on
the leader and `DS4_TEST_TP_MODE=worker DS4_TEST_TP_LEADER_HOST=HOST` on the
worker, with a unique `DS4_TEST_TP_PORT`. Run at least 2 and 4 sessions and
preserve both logs. TP currently uses the ordered per-session fallback, so
native single-machine row-grid flags must remain off in the TP logs.
- For the current TB5 MacBook link, US is `10.99.0.2` on `en1`/`rdma_en1` and
IT is `10.99.0.1` on `en6`/`rdma_en6`; both use GID index 1. Before testing,
require `rdma_ctl status` to report `enabled` and `ibv_devinfo -v` to show
`PORT_ACTIVE` plus the corresponding `::ffff:10.99.0.x` GID. Force the
device and GID with `--rdma-device NAME --rdma-gid-index 1` if automatic
selection is ambiguous. A working TB IP ping alone is not RDMA evidence.
- Kill the TP worker during one batch with `DS4_TEST_TP_DISCONNECT=1` on the
leader. The operation must fail cleanly, invalidate every affected session,
and return control without hanging.
- Verify unsupported combinations explicitly: GLM, DSpark/MTP support models,
SSD streaming, quality/reference modes, steering, resident Q4 expert
overlap, and CPU-router modes must use the established exact fallback or
reject the combination before evaluation. They must not partially activate
native batching.
## 5. Metal PRO Path
PRO support is experimental, but release builds must not break it silently.
@@ -84,7 +207,70 @@ 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. GLM 5.2
GLM has a different template, model shape, MTP block, attention layout,
tensor-parallel gate width, and streaming policy. Flash or PRO success does not
substitute for this matrix.
- On a 512 GB Metal machine, run short greedy prompts with both the Q4 XL and
reduced-precision Q2 release GGUFs. Cover thinking and no-thinking templates,
and verify the server reports the GLM model family rather than a DeepSeek
alias internally.
- Run the OpenRouter smoke vectors explicitly:
`DS4_TEST_MODEL=/path/to/glm.gguf
DS4_TEST_VECTOR_FILE=tests/test-vectors/glm-openrouter/official.vec
./ds4_test --logprob-vectors`.
Preserve the report as a diagnostic. The hosted vectors include very
low-probability top-20 tails whose membership is not stable after GLM routed
expert quantization, so an individual `official top token missing locally`
assertion is not by itself a release blocker. Selected-token mismatches must
remain consistent with the model's 100-case first-token band, and the
section 3 scorer is the release gate for aggregate GLM quality.
- Run the 100-case Q4 XL and Q2 official fixtures from section 3 and preserve
both `summary` and `api_summary` lines. Compare against the documented Q4 and
Q2 reference bands independently.
- Run `tests/glm_long_context_smoke.sh` with the release-advertised context on
the 512 GB Metal host. The generated continuation must begin with `>` and
contain none of the known token-corruption markers.
- Exercise integrated GLM MTP with `--glm-mtp-timing` on a deterministic
prompt. Compare the greedy text to
a non-MTP run, require clean speculative cycles, and record acceptance and
timing. Also run once with MTP disabled to prove ordinary decode remains the
default.
- Run the Metal session oracle with 2 and 4 GLM sessions. It must report
`family=glm native_shared=0 native_qkv=0` and remain exact, including mixed
prefill/decode; the DeepSeek-only row-grid kernels must not activate.
- Run resident and SSD-streaming GLM Q2 prompts with the same greedy input.
Compare first token and top-logprob sanity, and record the selected full-layer
prefix and dynamic expert-cache budget.
- Run physical two-machine GLM TP over TCP and RDMA with short and long prompts.
Record prefill/decode speed, transport, rank residency, and clean shutdown.
Repeat one run with `--tensor-parallel-token-prefill` as the exact-arithmetic
diagnostic. Use a GGUF whose routed-expert type has ownership-aware GLM TP
kernels. Also test a Q4-routed GLM file as a negative gate: until Q4 ownership
kernels are implemented, both ranks must reject it clearly before evaluation
rather than loading a partial split or hanging.
- On the eight-GPU CUDA host, run one resident GLM Q2 prompt, a long-context
prompt, integrated GLM MTP, and concurrent server requests. Use ordinary
eight-GPU layer placement for GLM; do not pass the Flash-specific
`--cuda-tensor-parallel` option. Multi-tier GLM prefill must
report progress through the tier-switching token-major path, and decode,
cache updates, and output-head/logit assembly must complete without CPU spill.
Auto-placement must reserve each layer's compact DSA/indexer cache and the
graph workspace before loading weights; a late graph-allocation failure is a
release blocker. Confirm the long-context layout stays within every device's
budget and uses additional tiers when the cache no longer fits on the earlier
ones.
The long-context harness can select this backend with
`DS4_GLM_BACKEND=cuda` and pass placement flags through
`DS4_GLM_EXTRA_ARGS="--gpu-vram auto --gpu-devices 0,2,4,6,1,3,5,7"`.
- Through `ds4-server`, exercise OpenAI chat, Responses, and Anthropic requests
against GLM, including thinking and SSE. DeepSeek compatibility endpoint
aliases may resolve to the loaded model, but rendered prompts and generated
text must use the GLM template.
## 7. SSD Streaming
SSD streaming is a capacity path, so test both correctness and user experience.
@@ -103,25 +289,57 @@ 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
## 8. 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
Use the DGX Spark / GB10 host `toor@192.168.60.184`. Do not claim CUDA is
release-ready without this pass.
- Fetch or push the exact release commit to the CUDA machine.
- Build:
`make clean && make cuda-spark`.
- Require both the DGX Spark build and the eight-GPU CUDA build to complete
without compiler warnings.
- Run:
`make cuda-regression`.
- Run a short CLI prompt with the Flash GGUF and record generation t/s.
- Run a longer prompt that exercises routed experts past a few thousand tokens.
- On the eight-GPU CUDA host, run the full-vocabulary decode oracle:
`DS4_TEST_MODEL=/path/to/flash.gguf make test-cuda-session-batch`.
Preserve the per-batch timing for 2, 4, and 8 rows and require
`nonexact_logits=0`. Run the released Q4 file and the reduced-precision Q2
file: Q4 exercises grouped routed/shared stages, while unsupported Q2 native
MoE shapes must retain the ordered exact fallback.
- With CUDA TP attention enabled, compatible Q4 runs must use grouped
attention-core, QKV, KV-store, and attention-post by default and remain
full-vocabulary exact against isolated decode. On the eight-L40S host, the
16-row decode step must remain above 110 aggregate tokens/s. Repeat once with
`DS4_CUDA_TP_ATTN=0` only as rollback coverage; it is not the production
configuration.
- Run native mixed prefill/decode at the default frontier and at compressed
context:
`DS4_TEST_MODEL=/path/to/flash.gguf make test-cuda-mixed-batch` and
`DS4_TEST_CONTEXT=4096 DS4_TEST_MIXED_INITIAL=2048 DS4_TEST_MIXED_ROUNDS=8
DS4_TEST_MODEL=/path/to/flash.gguf make test-cuda-mixed-batch`.
Every round must report exact logits and `mode=native`; a serialized fallback
is a failure for the eight-GPU TP/EP topology. Under CUDA TP attention, the
native mixed step must use the same exact grouped decode stages when their
capability checks pass; record correctness and speedup separately. Also
force an 800-row prefill quantum with
`DS4_TEST_ALLOW_FALLBACK=1`; it must report the serialized safety fallback.
- Start `ds4-server` with 8 and 16 batched sessions and issue at least that many
simultaneous requests with mixed prompt lengths. Verify no session mix-up,
deadlock, or starvation and record aggregate generation throughput.
- On DGX Spark, verify the same public batch API and server concurrency use the
single-GPU fallback without creating peer-only TP/EP state. The eight-GPU
native oracle is not a valid Spark test because its topology is intentionally
unavailable there.
- If CUDA Q4, distributed, streaming hooks, tensor span loading, or model cache
code changed, test the specific GGUF and split mode that uses that path.
- Verify that any CUDA-only warning fixes are also clean on macOS and do not
change Metal behavior.
## 7. ROCm / Strix Halo
## 9. 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
@@ -130,6 +348,7 @@ a substitute for CUDA or Metal release testing.
- Fetch or push the exact release commit to the Strix Halo machine.
- Build:
`make clean && make strix-halo`.
- Require the ROCm build to complete without compiler warnings.
- Use the q2 Flash imatrix GGUF for release smoke tests:
`DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix.gguf`.
- Do not use the mixed q2-q4 or Q4 Flash GGUFs for routine Strix Halo QA yet.
@@ -139,10 +358,21 @@ a substitute for CUDA or Metal release testing.
`./ds4 -m gguf/DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix.gguf --ctx 4096 --nothink -p "Reply with exactly: OK"`.
- Run one longer prompt if ROCm kernels, backend hooks, tensor loading, model
cache, KV cache, or graph prefill code changed.
- Run the GLM Q2 release model through ROCm SSD streaming with at least four
generated tokens:
`./ds4 --rocm -m gguf/GLM-5.2-UD-Q2_K_RoutedQ2K.gguf --ssd-streaming --ctx 4096 --nothink --tokens 4 -p "Reply with exactly: OK"`.
Startup must select a cache budget that passes the memory guard without an
override, and both compact indexed prefill and decode must complete.
- Run one longer GLM prompt with the release-advertised Strix context after
changes to GLM attention, typed quantized projections, streaming expert
caches, or memory budgeting. Record the context, cache split, and whether
the continuation stays free of token-corruption markers.
- Run the same GLM model with `--glm-mtp-timing --temp 0`. At least one draft
verification cycle must complete without a `glm mtp step failed` message.
- Record startup memory/cache messages, prefill speed, generation speed, and
whether the backend reports `ROCm backend initialized`.
## 8. Distributed Inference
## 10. Distributed Inference
Distributed code has regressed around route setup, KV snapshots, request IDs,
and split model loading. Test it whenever distributed, KV, session, or model
@@ -160,7 +390,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
## 11. Disk KV Cache
Disk KV cache bugs are high impact for server users.
@@ -174,7 +404,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
## 12. Server APIs
The server must keep compatibility across OpenAI, Responses, and Anthropic
clients.
@@ -184,10 +414,13 @@ clients.
- Test OpenAI chat completion, OpenAI Responses, and Anthropic messages.
- Test SSE streaming with thinking enabled and disabled.
- Test keepalive during long prefill and confirm clients do not time out.
- On the eight-L40S CUDA TP target, start `ds4-server` with the release TP
options and verify all 16 100k-context sessions allocate. Startup must report
a 2048-token prefill cap; a silent fallback to 4096 is an OOM regression.
- Test `--trace` and confirm rendered prompts, cache decisions, generated text,
and tool-parser events are useful without leaking unrelated state.
## 11. ds4-agent
## 13. ds4-agent
The agent is the most stateful component. Test it manually, not only by build.
@@ -219,7 +452,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
## 14. Download Script And Model Files
- Test `download_model.sh` in a temporary directory so local weights are not
overwritten.
@@ -228,7 +461,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
## 15. Performance And Power
- Run `ds4-bench` on the release machine and compare with tracked CSV baselines.
- Test `--power 100` is not throttled.
@@ -237,17 +470,25 @@ 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
## 16. Release Sign-off
Do not sign off until:
- macOS Metal Flash passed.
- GLM 5.2 Metal, official-quality, MTP, batching-fallback, and applicable TP or
CUDA gates 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
not validated.
- Metal, CUDA, ROCm, CPU-only, and test builds completed without compiler
warnings on every release target that was validated.
- Disk KV cache was exercised.
- Server API streaming was exercised.
- Agent interruption and tool loops were exercised manually.
- Speed is within expected variance for the same hardware and model.
- Metal 2/4/8/16-session exactness and forced fallback gates passed.
- Physical Metal TP batching and CUDA native decode/mixed batching passed when
those backends are part of the release.
- Any skipped item is written down with the reason.
+530 -87
View File
@@ -1,47 +1,47 @@
# DwarfStar
<p align="center">
<img src="logo.svg" alt="DwarfStar logo" width="220">
</p>
**DwarfStar** is a small native inference engine optimized first for
**DeepSeek V4 Flash**, with support for **DeepSeek V4 PRO** on very high-memory
machines. It is
intentionally narrow: not a generic GGUF runner, not a wrapper around another
runtime: it is completely self-contained. Other than running the model in a
correct and fast way, the project goal is to provide DeepSeek specific loading,
prompt rendering, tool calling, KV state handling (RAM and on-disk), server
API and integrated coding agent, all ready to work with coding agents or with
the provided CLI interface. There are also tools for GGUF and imatrix generation,
and for quality and speed testing.
**DeepSeek V4 Flash**. It also supports **GLM 5.2**, **Laguna S 2.1**, and, on
very high-memory machines, **DeepSeek V4 PRO**. It is self-contained and
deliberately narrow, not a general GGUF runner. Model loading, prompt rendering,
tool calls, KV state, the HTTP server, and the coding agent are built and tested
together.
The repository also includes tools and data for GGUF, imatrix, quality, and speed.
We support the following backends:
* **Metal** is our primary target. Starting from MacBooks with 96GB of RAM (or less, using SSD streaming).
* **NVIDIA CUDA / DGX Spark**, CUDA with special care for the DGX Spark.
* **Strix Halo (ROCm)**, systems like the Framework Desktop and other systems based on the same GPU and unified RAM design.
Supported backends:
* **Metal**, the primary target, on Macs with 96 GB or more. Smaller machines
can use SSD streaming.
* **NVIDIA CUDA**, including multi-GPU systems and DGX Spark.
* **ROCm** on Strix Halo systems such as the Framework Desktop.
This project would not exist without **llama.cpp and GGML**, make sure to read
the acknowledgements section, a big thank you to Georgi Gerganov and all the
other contributors.
**Note that DeepSeek v4** is not our only target. Right now Flash and PRO are the
perfect fit because of capabilities, size, KV cache efficiency. If tomorrow a
better open weight model is released for the 128GB size, we could switch, the same
for other important size classes like 512GB of RAM. The project is stictly
opportunistic depending on what open weight models exist in a given moment.
If a new model will be supported, the old one may be removed completely and
no longer supported, unless there is some kind of overlap of abilities.
Model support is intentionally opportunistic. The project follows the best open
weights for useful local machine sizes, especially 128 GB laptops and 512 GB
workstations. A model may be removed when a better replacement arrives.
# So, what can I do with this software?
* You can run a very capable models in your consumer hardware, a MacBook, a DGX Spark, or a Strix Halo for example. Even if you have not enough RAM, with SSD streaming, you can run it at a decent speed.
* Using the CUDA multi-GPU support and with ds4-server micro batching of decoding and generation, you can turn a server with old-ish CUDA cards (Ada Lovelace architecture), no longer supported for new models by vLLM, into a multi-user LLM server for your company. We tested this setup with 8xL40S NVIDIA cards and multiple sessions with very good results. 120 t/s aggreated generation, 2000 t/s prefill.
* Using two MacBook M5 Max / M3 Ultra RDMA, you can run 4 bit DeepSeek Flash or GLM 5.2 with tensor parallelism.
* You can also use pipeline paralellism to glue together multiple systems to sum their RAM and run larger models.
## Motivations
* Very capable open weight models finally exist. DeepSeek v4 Flash feels quasi-frontier. The PRO is even better. Both resist 2 bit quantization very well.
* Very capable computers like MacBooks, the DGX Spark now exist.
* DeepSeek v4 kv cache design makes it pratical to run very big contexts. Other vendors are using this approach.
* This few hundred billions models are strictly better than smaller (even if dense) models, regardless of what benchmarks say.
* Capable open-weight models now fit on high-end personal machines.
* DeepSeek V4 Flash and PRO, GLM 5.2, tolerate aggressive routed-expert quantization.
* Compressed KV caches and fast local SSDs make long contexts practical.
* The idea of an inference system specialized for a few models.
That said, a few important things about this project:
# AI full disclosure
* The local inference landscape contains many excellent projects, but new models are released continuously, and the attention immediately gets captured by the next model to implement. This project takes a deliberately narrow bet: one model at a time, official-vector validation (logits obtained with the official implementation), long-context tests, and enough agent integration to know if it really works. The exact model may change as the landscape evolves, but the constraint remains: local inference credible on high end personal machines or Mac Studios, starting from 96/128GB of memory.
* This software is developed with **strong assistance from GPT 5.5** and with humans leading the ideas, testing, and debugging. We say this openly because it shaped how the project was built. If you are not happy with AI-developed code, this software is not for you. The acknowledgement below is equally important: this would not exist without `llama.cpp` and GGML, largely written by hand.
* This implementation is based on the idea that compressed KV caches like the one of DeepSeek v4 and the fast SSD disks of modern MacBooks should change our idea that KV cache belongs to RAM. **The KV cache is actually a first-class disk citizen**. Fast SSD disks also changed the inference game from the point of view of "model needs to fit RAM": while having more RAM the the model size is still preferred, SSD streaming allows to turn the available amount of RAM from a hard cutoff (can I run this model or not?) to continuous spectrum of speed levels.
* Our vision is that local inference should be a set of three things working well together, out of the box: A) inference engine with HTTP API + B) GGUF specially crafted to run well under a given engine and given assumptions + C) testing and validation with coding agents implementations. D) Purpose built agents for specific models and execution environments. DwarfStar only runs with the GGUF files provided. It gets tested against officially obtained logits at different context sizes. This project exists because we wanted to make one local model feel finished end to end, not just runnable. However this is beta quality code, so probably we are not still there, especially since recently we introduced large new features: distributed inference, SSD streaming, and other minor improvements.
* The optimized graph path targets **Metal on macOS** and **CUDA on Linux**. The CPU path is only for correctness checks and model/tokenizer diagnostics. For CPU-only Linux builds, use `make cpu`; it builds the normal `./ds4` and `./ds4-server` binaries without CUDA or Metal. On macOS, **warning: current macOS versions have a bug in the virtual memory implementation that will crash the kernel** if you try to run the CPU code. Remember? Software sucks. It was not possible to fix the CPU inference to avoid crashing, since each time you have to restart the computer, which is not funny. Help us, if you have the guts.
* This software is developed with **strong assistance from GPT 5.5, 5.6, Claude Fable** and with humans leading the ideas, testing, and debugging. We say this openly because it shaped how the project was built. If you are not happy with AI-developed code, this software is not for you. The acknowledgement below is equally important: this would not exist without `llama.cpp` and GGML, largely written by hand.
## Acknowledgements to llama.cpp and GGML
@@ -58,14 +58,9 @@ notice in our `LICENSE` file.
## Status
The code and GGUF files are to be considered of **beta quality** because
inference and model serving is a complicated matter and all this exists
only for a few days. It will take months to reach a more stable form.
However, we try to keep the project in a usable state, and we are making
progress. If you have issues, make sure to use `--trace` to log the
sessions, and open issues including the full trace.
The `ds4-agent` is alpha quality, the project was later added.
The software is currently very fast changing. Consider it beta quality.
Before each release, a big QA run is executed, however instabilities
are definitely possible.
## More Documentation
@@ -75,6 +70,8 @@ next sections.
- [CONTRIBUTING.md](CONTRIBUTING.md): correctness and speed regression testing
guide for contributors. **Read this before sending a pull request**.
- [QA_BEFORE_RELEASES.md](QA_BEFORE_RELEASES.md): the complete release test
matrix, including the remote Metal, CUDA, and ROCm machines.
- [gguf-tools/README.md](gguf-tools/README.md): offline GGUF generation,
imatrix collection, quantization tooling, and quality checks.
- [gguf-tools/imatrix/README.md](gguf-tools/imatrix/README.md): how the
@@ -92,11 +89,13 @@ next sections.
## Model Weights
This implementation only works with the DeepSeek V4 Flash and PRO GGUFs published for
this project. It is not a general GGUF loader, and arbitrary DeepSeek/GGUF files
This implementation only works with the DeepSeek V4, GLM 5.2, and Laguna S 2.1
GGUFs listed below. It is not a general GGUF loader, and arbitrary GGUF files
will not have the tensor layout, quantization mix, metadata, or optional MTP
state expected by the engine. The 2 bit quantizations provided here are not
a joke: they behave well, work under coding agents, call tools in a reliable way.
state expected by the engine. The 2 bit DeepSeek and GLM quantizations provided
here are verified to be actually high quality: they behave well, work under
coding agents, and call tools reliably.
The 2 bit quants use a very asymmetrical quantization: only the routed MoE
experts are quantized, up/gate at `IQ2_XXS`, down at `Q2_K`. They are the
majority of all the model space: the other components (shared experts,
@@ -109,6 +108,8 @@ Download one main model. **Prefer the imatrix versions.**
./download_model.sh q2-q4-imatrix # 96/128 GB RAM machines, q2 with last 6 layers q4
./download_model.sh q4-imatrix # >= 256 GB RAM machines, imatrix-tuned q4
./download_model.sh pro-q2-imatrix # 512 GB RAM machines, PRO q2 imatrix quant
./download_model.sh laguna-q4 # >= 96 GB unified memory, official Poolside Q4_K_M
./download_model.sh laguna-q2-q3 # 64 GB class, mixed routed Q2_K/Q3_K
```
For the full PRO Q4 distributed run, download one half on each machine:
@@ -139,12 +140,155 @@ but must be enabled explicitly with `--mtp`. The current MTP/speculative
decoding path is still experimental: it is correctness-gated and currently
provides at most a slight speedup, not a meaningful generation-speed win.
GLM 5.2 support is 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-iq2xxs # antirez routed IQ2_XXS single-file GGUF
./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.
These formats do not all support the same execution modes. The Q4 files work
for normal Metal and CUDA inference. Two-Mac tensor parallelism currently
requires an ownership-aware IQ2_XXS or Q2_K routed layout; a routed Q4 GLM
must be rejected before evaluation.
GLM's MTP block is part of the main GGUF; it does not use the separate Flash
MTP file. Ordinary decode remains the default. `--glm-mtp` enables experimental
greedy speculation. `--glm-mtp-timing` also enables it and prints acceptance
and timing counters:
```sh
./ds4 -m gguf/GLM-5.2-UD-IQ2_XXS_RoutedIQ2XXS_blk78Q2K.gguf \
--glm-mtp-timing --temp 0
```
GLM inference uses the Metal, CUDA, or ROCm graph backend. Directional steering,
`--power` below 100, an explicit `--prefill-chunk`, and the external `--mtp`
file are not supported for GLM yet.
Laguna S 2.1 support targets Poolside's official imatrix-quantized Q4_K_M GGUF.
The current 63.56 GiB recipe uses Q4_K routed experts and Q8_0 signal-path
weights. It runs with full residency on Metal, CUDA, or ROCm. DwarfStar also accepts
Poolside's earlier 70.01 GiB recipe with F16 attention and mixed Q4_K/Q6_K
experts on Metal. These files fit comfortably on a 96 or 128 GiB machine.
For lower-memory systems, `laguna-q2-q3` downloads a 44.95 GiB mixed quant. It
retains the official file's dense weights, uses Q2_K routed experts in layers
1 through 20, and Q3_K routed experts in layers 21 through 47. This is the
smaller Laguna layout and is intended for full residency on 64 GB class
systems. The current Q8_0-signal Q4_K_M and mixed Q2_K/Q3_K layouts run on
Metal, CUDA, and ROCm. Linux validation covers NVIDIA GB10 in DGX Spark and the
Ryzen AI Max+ 395 / Radeon 8060S (`gfx1151`) in Strix Halo.
CLI, agent, and server use Laguna's native chat, interleaved reasoning, and
tagged tool-call formats:
```sh
./download_model.sh laguna-q4
./ds4 -m gguf/laguna-s-2.1-Q4_K_M.gguf -c 32768 -p "Explain this repository"
./ds4-agent -m gguf/laguna-s-2.1-Q4_K_M.gguf -c 32768
./ds4-server -m gguf/laguna-s-2.1-Q4_K_M.gguf -c 32768
```
Poolside's standalone DFlash model can accelerate greedy Laguna decoding on
Metal, CUDA, and ROCm without changing the generated tokens. The download target uses
a 1.04 GiB Q8_0 quant of Poolside's support model: it drafts faster and uses
half the memory of BF16, while the full Laguna model still verifies every
accepted token. Download it separately and pass it alongside any supported
Laguna GGUF:
```sh
./download_model.sh laguna-dflash
./ds4 -m gguf/laguna-s-2.1-Q4_K_M.gguf \
--dflash gguf/laguna-s-2.1-DFlash-Q8_0.gguf --temp 0
./ds4-agent -m gguf/laguna-s-2.1-Q4_K_M.gguf \
--dflash gguf/laguna-s-2.1-DFlash-Q8_0.gguf
./ds4-server -m gguf/laguna-s-2.1-Q4_K_M.gguf \
--dflash gguf/laguna-s-2.1-DFlash-Q8_0.gguf
```
The default explores up to 15 draft positions on CUDA and up to three on Metal
and ROCm. It stops before
verification when a proposal's probability is below 0.4. Use
`--dflash-draft N` to tune the 1 through 15 range and
`--dflash-p-min P` to tune the confidence cutoff; set the latter to `0` to
keep a fixed verifier width. A nonzero cutoff is faster, but varying the
verifier batch width can resolve nearly tied greedy logits differently because
floating-point reductions are not batch-invariant. Use `0` when fixed-width
reproducibility matters. DwarfStar automatically falls back to ordinary
Laguna decoding when sampling is stochastic or when speculation is slower for
the current turn. Set `DS4_DFLASH_TIMING=1` to print per-cycle draft and
verify timings.
For Strix Halo:
```sh
./download_model.sh laguna-q2-q3
make strix-halo
./ds4 --rocm -m gguf/laguna-s-2.1-RoutedQ2_K-Last27Q3_K.gguf -c 8192
./ds4-agent --rocm -m gguf/laguna-s-2.1-RoutedQ2_K-Last27Q3_K.gguf -c 8192
./ds4-server --rocm -m gguf/laguna-s-2.1-RoutedQ2_K-Last27Q3_K.gguf -c 8192
```
For DGX Spark:
```sh
./download_model.sh laguna-q2-q3
make cuda-spark
./ds4 --cuda -m gguf/laguna-s-2.1-RoutedQ2_K-Last27Q3_K.gguf -c 8192
./ds4-agent --cuda -m gguf/laguna-s-2.1-RoutedQ2_K-Last27Q3_K.gguf -c 8192
./ds4-server --cuda -m gguf/laguna-s-2.1-RoutedQ2_K-Last27Q3_K.gguf -c 8192
```
On the development DGX Spark with that mixed Q2_K/Q3_K GGUF, a 4,096-token
prefill runs at about 580 tokens/s and generation after that prompt at about
26.7 tokens/s. Long CUDA prefills compact routed tokens by expert, dequantize
only active expert matrices, and use variable-size FP16 tensor-core GEMMs with
FP32 projection outputs. Decode reuses each Laguna KV-cache row across grouped
query heads and uses denser Q2_K/Q3_K warp layouts; short prefills retain the
quantized matvec kernels. Set `DS4_CUDA_MOE_NO_TC_PREFILL=1` to select the
quantized long-prefill fallback. The decode optimizations can be isolated with
`DS4_CUDA_LAGUNA_NO_GQA_DECODE=1`, `DS4_CUDA_MOE_NO_Q2_HALFWARP=1`, and
`DS4_CUDA_MOE_NO_Q3_GROUP_DOWN=1`.
On the development Strix Halo with the official Q4_K_M GGUF, a 2,048-token
prefill runs at about 252 tokens/s and short-context generation at about 25.8
tokens/s; after an 8,192-token prompt prefill is about 194 tokens/s and
generation about 24.1 tokens/s. The smaller mixed Q2_K/Q3_K quant is faster
still. Laguna currently requires full model residency; SSD streaming,
distributed inference, and tensor parallelism are not implemented.
DFlash speculative decoding works on ROCm. With a fixed verifier width
(`--dflash-p-min 0`) it is token-exact against ordinary decoding in the
regression corpus, but on this GPU it is only a win on highly predictable
text. Each draft position routes to its own ten experts, so the verifier's
expert traffic grows with the draft length while a Strix Halo decode step is
already bandwidth-bound; break-even needs roughly 60% of draft tokens accepted.
DwarfStar measures throughput for the first cycles of every turn and drops back
to ordinary decoding when speculation is not paying, so leaving `--dflash`
enabled is safe.
The shipped GGUF is configured for a 262144-token context. Laguna defaults to
temperature 0.7, top-k 20, top-p 0.95, and min-p 0.05; explicit sampling
options always take precedence. Use `--nothink` or the
`laguna-s-2.1-chat` server alias for direct replies, and preserve reasoning
content between tool calls when building a client.
Then build:
```sh
make # macOS Metal
make cuda-spark # Linux CUDA, DGX Spark / GB10
make cuda-generic # Linux CUDA, other local CUDA GPUs
make strix-halo # Linux ROCm, AMD Strix Halo
make cpu # CPU-only diagnostics build
```
@@ -152,8 +296,56 @@ make cpu # CPU-only diagnostics build
select another supported GGUF from `./gguf/`. Run `./ds4 --help` and
`./ds4-server --help` for the full flag list.
## DSpark Speculative Decoding
DSpark is an auxiliary draft model released by DeepSeek for DeepSeek V4 Flash.
It reads hidden states from the main model and proposes up to five future
tokens. DwarfStar checks those proposals with the main Flash model and commits
only the accepted prefix. The main model remains authoritative; a rejected or
low-confidence suffix falls back to ordinary target decoding.
The possible gain is faster generation: when several proposed tokens are
accepted, one target verification pass advances the stream by several tokens.
It does not accelerate prefill, and the draft and verification work is not
free. Predictable continuations, especially code, tend to benefit most;
low-yield prompts can be no faster or even slower. DSpark is therefore still
experimental and explicitly opt-in.
The released DSpark checkpoint is packaged here as a separate support GGUF of
about 5.6 GiB. It is not a standalone model. Download it once:
```sh
./download_model.sh dspark-support
```
The same support file can be used with the Flash `q2-imatrix`,
`q2-q4-imatrix`, and `q4-imatrix` models listed above. For now **DeepSeek
V4 PRO** is not supported. On Metal, the main model may be resident or use
`--ssd-streaming`; the support model still adds its own weights and runtime
state to the memory requirement. DSpark replaces the legacy one-stage MTP
support model for that run rather than stacking with it.
Run it with greedy decoding:
```sh
./ds4 -m ds4flash.gguf \
--mtp gguf/DeepSeek-V4-Flash-DSpark-support.gguf \
--dspark --temp 0
```
`--mtp` supplies the support GGUF, while `--dspark` selects the DSpark runtime.
The default confidence threshold is `0.9`; it prunes suffixes that are unlikely
to repay their verification cost. `--dspark-confidence 0` forces fixed
five-token blocks and is intended for diagnostics. Sampled decoding does not
use DSpark proposals. `--quality` and `--dspark-strict` also keep target-only
decoding, which is useful for comparisons and correctness checks.
## Speed
*Warning: some of those numbers may no longer be updated, because of the optimization
efforts that improved the runtime speed without updating the benchmark
results.*
These are single-run Metal CLI numbers with `--ctx 32768`, `--nothink`, greedy
decoding, and `-n 256`. The short prompt is a normal small Italian story
prompt. The long prompts exercise chunked prefill plus long-context decode.
@@ -181,8 +373,8 @@ Q4 requires the larger-memory machine class, so M3 Max Q4 numbers are `N/A`.
The normal Metal path tries to make the model resident in GPU-addressable
memory. This is the fastest path and should remain your default when the model
fits. When it does not fit, DwarfStar also has a Metal-only **SSD streaming**
capacity mode. In this mode the non-routed model weights stay resident, while
fits. DwarfStar also has an **SSD streaming** capacity mode on Metal and for
GLM 5.2 on ROCm. In this mode the non-routed model weights stay resident, while
routed MoE experts are kept in an in-memory cache and loaded from the GGUF file
on cache misses.
@@ -206,14 +398,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 backend 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 backend's recommended working set, subtracts non-routed weights, then
applies the same routed-prefill headroom before sizing the dynamic cache. Leave
the hot expert preload enabled for normal use; use `--ssd-streaming-cold` and
`--ssd-streaming-preload-experts N` only for measurements.
### Practical SSD streaming examples
@@ -245,15 +442,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 \
@@ -264,18 +457,40 @@ is stable, re-enable thinking with a conservative generation limit:
--tokens 1500
```
GLM 5.2 uses the same option. Its streaming path keeps the largest full-layer
prefix that fits resident, then uses the remaining budget for a dynamic expert
cache. Start with the automatic budget:
```sh
./ds4 \
-m gguf/GLM-5.2-UD-IQ2_XXS_RoutedIQ2XXS_blk78Q2K.gguf \
--ssd-streaming \
--ctx 32768
```
The important startup line is the cache report. Start conservative, then
increase the cache if the machine has headroom.
## Distributed Inference
On a 128GB Strix Halo, use the routed Q2_K model and a 4096-token context as the
starting point. The automatic cache budget leaves room for the GLM graph and KV
state:
Distributed inference lets DwarfStar **run a model that is too large for one machine** by
```sh
./download_model.sh glm-antirez-q2
make strix-halo
./ds4 --rocm -m gguf/GLM-5.2-UD-Q2_K_RoutedQ2K.gguf \
--ssd-streaming --ctx 4096
```
## Distributed inference with pipeline parallelism
Pipeline parallelism lets DwarfStar **run a model that is too large for one machine** by
splitting transformer layers across multiple machines. The main example is the
full 4-bit Flash quant across two 128 GB MacBooks: each process maps only its
own layer slice, activations are sent over TCP, and the coordinator keeps normal
CLI/API behavior.
Distributed inference also allows to **speed up prefill** by
Pipeline parallelism can also **speed up prefill** by
using multiple GPUs at the same time to process different micro-batches at
different layers, like in an assembly line. Only prefill can be accelerated this
way. Generation is purely autoregressive: each token must finish across the
@@ -484,6 +699,171 @@ 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 over RDMA
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:
```sh
# 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
```
Check the verbs device before loading the model:
```sh
rdma_ctl status
ibv_devinfo -v
```
The device must be active and expose the IPv4-mapped GID for the address above,
for example `::ffff:10.99.0.2`. A working IP ping does not prove that RDMA is
active.
Both machines need the same tree, commit, and GGUF path. Tensor parallelism is
always a 50/50 split with one worker, so do not pass `--layers`. Start the worker
first; it retries while the coordinator loads. The worker must dial the address
on the Thunderbolt member interface, not the bridge address:
```sh
MODEL=gguf/GLM-5.2-UD-IQ2_XXS_RoutedIQ2XXS_blk78Q2K.gguf
# Machine B: worker.
./ds4 -m "$MODEL" --tensor-parallel --role worker \
--coordinator 10.99.0.2 9911 --transport rdma
# Machine A: coordinator.
./ds4 -m "$MODEL" --tensor-parallel --role coordinator \
--listen 10.99.0.2 9911 --transport rdma -c 8192 \
-p "Tell me something about the sea."
```
The active verbs device and IPv4-mapped GID are selected automatically. If that
is ambiguous, add `--rdma-device rdma_en6 --rdma-gid-index 1` on the worker and
the matching `rdma_en1` flags on the coordinator. Use `--transport tcp` on both
sides to force TCP. Tensor parallel roles are currently exposed by the `ds4`
CLI, not by `ds4-server` or `ds4-agent`.
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 (`--tensor-parallel-token-prefill` 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.
## Tensor Parallelism across CUDA GPUs
On a single CUDA server, `--cuda-tensor-parallel` splits DeepSeek V4 Flash
tensor and routed-expert work across an even number of GPUs. This is separate
from the Mac-to-Mac mode above: it does not use `--role`, RDMA, or the
distributed layer pipeline. GPU placement and memory budgets are selected with
the normal `--gpu-devices` and `--gpu-vram` options.
The device order is significant. With `N` devices, the first `N/2` logical
tiers are contiguous layer-pipeline homes and the second `N/2` tiers are their
tensor-parallel partners. Specify all homes first and then all partners, with
the closest P2P pair at matching positions. For example, the tested L40S host
uses physical pairs `(0,1)`, `(2,3)`, `(4,5)`, and `(6,7)`, expressed as
`0,2,4,6,1,3,5,7`. Each pair stores a 50/50 split of the routed experts, and
the vocabulary head is row-sharded across the participating output tiers.
Those large tensors are not duplicated. Dense attention, router, and shared
expert weights are replicated within each pair.
For maximum throughput on eight 48 GB L40S cards, use the imatrix Q4 model.
Its routed `Q4_K` layout has the native grouped multi-session kernels; the Q2
model is the lower-memory choice (including tested four-card runs), but its
unsupported grouped routed shapes use the exact fallback and have lower
aggregate serving throughput. Download and build the L40S target with:
```sh
./download_model.sh q4-imatrix
make cuda CUDA_ARCH=sm_89
```
This is the interactive-agent setup used on the eight-L40S server:
```sh
MODEL=gguf/DeepSeek-V4-Flash-Q4KExperts-F16HC-F16Compressor-F16Indexer-Q8Attn-Q8Shared-Q8Out-chat-v2-imatrix.gguf
./ds4-agent --cuda --cuda-tensor-parallel \
--gpu-vram auto \
--gpu-devices 0,2,4,6,1,3,5,7 \
--model "$MODEL" \
--ctx 100000
```
For serving, keep multiple KV sessions resident so decode rows can be grouped
across requests. The tested host is configured for up to 16 resident sessions:
```sh
./ds4-server --cuda --cuda-tensor-parallel \
--gpu-vram auto \
--gpu-devices 0,2,4,6,1,3,5,7 \
--model "$MODEL" \
--ctx 100000 \
--batched-session 16 \
--host 0.0.0.0
```
The equivalent local launchers are `./run-nvidia-tp-agent.sh` and
`./run-nvidia-tp-server.sh`. The server launcher also enables the on-disk KV
cache. Reduce the session count or context size if the requested resident KV
caches do not fit after model loading. CUDA TP, half-resident expert ownership,
output sharding, pipelined prefill, and compatible grouped decode are selected
by `--cuda-tensor-parallel`; no `DS4_CUDA_*` environment tuning is required.
Without an explicit `--prefill-chunk`, this mode uses 2048-token chunks so the
tested 16-session, 100k-context layout retains enough VRAM for resident KV
caches. An explicit `--prefill-chunk` remains an override for other topologies.
Any even card count that can hold the selected model and graph scratch is a
valid topology. On this class of 48 GB card, the useful measured endpoints are
Q2 on four cards (two pipeline stages) and Q4 on eight cards (four stages).
For a four-card PIX-paired subset such as physical GPUs `0,1,4,5`, the ordered
list is `0,4,1,5`. Two cards do not have enough memory for these Flash models.
This mode currently requires DeepSeek V4 Flash and an even multi-GPU
placement. GLM 5.2 instead uses normal layer placement across the selected
CUDA devices. DGX Spark is a single-GPU target and must not be started with
`--cuda-tensor-parallel`.
## Reducing heat, power usage and fan noise
Long local inference runs can keep the GPU busy for extended periods. If you
@@ -497,8 +877,8 @@ and inserting small sleeps between work units: during prefill it sleeps between
layers, and during generation it sleeps between decoded tokens. This reduces
sustained load without changing model output.
The option is available on the CLI, server, agent, eval, and benchmark tools,
for example:
The option is available on the CLI, server, agent, eval, and benchmark tools
for DeepSeek models. GLM 5.2 currently accepts only `--power 100`. For example:
```sh
./ds4 --power 50
@@ -566,12 +946,10 @@ exponential sweeps. Output is CSV with one row per frontier: latest prefill
interval tokens/sec, generation tokens/sec at that frontier, and
`kvcache_bytes`.
Sessions prefill long prompts in 4096-token chunks by default. Set
`DS4_METAL_PREFILL_CHUNK=N` to compare another chunk size, for example `2048`
to match the strict official-vector checkpoint path, or
`DS4_METAL_PREFILL_CHUNK=0` to prefill a prompt as one whole batch when memory
allows. Changing the chunk changes the KV checkpoint/logit path, so compare it
as an explicit run configuration.
Sessions prefill long prompts in 4096-token chunks by default. Use
`--prefill-chunk 2048`, for example, to match the strict official-vector
checkpoint path. Changing the chunk changes the KV checkpoint/logit path, so
compare it as an explicit run configuration.
Chunked Metal prefill reuses the same range-capable layer-major graph for each
chunk, preserving absolute compressor/indexer boundaries while avoiding the old
per-layer chunk dispatch path.
@@ -705,14 +1083,34 @@ Start a local OpenAI/Anthropic-compatible server:
Use `--chdir /path/to/ds4` when launching `ds4-server` from another directory,
so relative runtime files such as `metal/*.metal` resolve from the project tree.
The server keeps one mutable backend/KV checkpoint in memory,
so stateless clients that resend a longer version of the same prompt can reuse
the shared prefix instead of pre-filling from token zero.
By default the server keeps one mutable backend/KV checkpoint in memory, so
stateless clients that resend a longer version of the same prompt can reuse the
shared prefix instead of pre-filling from token zero.
Request parsing and sockets run in client threads, but inference itself is
serialized through one graph worker. The current server does not batch multiple
independent requests together; concurrent requests wait their turn on the single
live graph/session.
`--batched-session N` preallocates `N` independent resident KV sessions. Ready
decode steps are evaluated together, while long prefills alternate in bounded
chunks so one request does not block every decoder. Requests beyond `N` wait
for a resident slot. If disk KV caching is enabled, an idle slot is persisted
before reuse and can be restored when that conversation returns; an active
request is never evicted. Choose `N` and `--ctx` so all resident KV allocations
fit in GPU memory. Without this option, inference retains the original
single-session behavior.
Batching is exact: when a native batched kernel is unavailable, DwarfStar runs
the affected rows in a fixed order and returns the same full logits as separate
session evaluations. The current backend behavior is:
| Backend and model | Session execution |
| --- | --- |
| Metal, resident DeepSeek Flash | Native shared-expert and QKV batching from two rows upward when supported; ordered fallback otherwise. |
| Metal, GLM 5.2 | Ordered exact fallback. |
| CUDA, DeepSeek Flash on a supported multi-GPU TP/EP layout | Native decode and mixed prefill/decode, with exact fallbacks for unsupported kernel shapes. |
| CUDA single GPU, including DGX Spark | Ordered exact fallback. |
`N` resident sessions allocate `N` KV states, so a context size that fits once
may not fit eight times. Native batching can improve aggregate throughput; an
ordered fallback provides concurrency and fairness, but not the same speedup.
MTP speculative decoding is disabled while native session batching is active.
Supported endpoints:
@@ -747,8 +1145,11 @@ controls. Tool uses are returned as Anthropic `tool_use` blocks.
Default sampled API generation uses `temperature=1`, `top_p=1`, and
`min_p=0.05`, so the default filter is relative probability rather than
nucleus mass. In thinking mode DwarfStar uses those fixed sampling defaults and
ignores client sampling knobs, matching DeepSeek's fixed-thinking API behavior.
nucleus mass. In thinking mode DwarfStar applies those fixed sampling defaults
to any knob the request omits, matching DeepSeek's fixed-thinking API behavior,
but sampling parameters set explicitly in the request always win: a
`temperature=0` request is greedy through the whole reasoning phase, so
benchmark harnesses get deterministic thinking-mode output.
The chat, Responses, and Anthropic endpoints support SSE streaming. In thinking
mode, reasoning is streamed in the native API shape instead of being mixed into
@@ -951,8 +1352,8 @@ matches the local `~/bin/claude-ds4` setup:
#!/bin/sh
unset ANTHROPIC_API_KEY
export ANTHROPIC_BASE_URL="${DS4_ANTHROPIC_BASE_URL:-http://127.0.0.1:8000}"
export ANTHROPIC_AUTH_TOKEN="${DS4_API_KEY:-dsv4-local}"
export ANTHROPIC_BASE_URL="http://127.0.0.1:8000"
export ANTHROPIC_AUTH_TOKEN="dsv4-local"
export ANTHROPIC_MODEL="deepseek-v4-flash"
export ANTHROPIC_CUSTOM_MODEL_OPTION="deepseek-v4-flash"
@@ -1191,6 +1592,25 @@ make cuda CUDA_ARCH=sm_120
make cuda CUDA_ARCH=native
```
CUDA builds accept `--gpu-vram N[,N,...]` and `--gpu-devices N[,N,...]` in the
CLI, server, agent, and benchmark. VRAM values are per-device GiB budgets;
`--gpu-vram auto` uses the free memory reported by CUDA. The device list controls
the placement order and must have the same number of entries as an explicit
budget list. Placement reserves graph and KV memory for the requested context
and refuses to start if model layers would spill to the CPU.
Without `--cuda-tensor-parallel`, CUDA uses normal layer placement across the
listed devices. This is also the supported multi-GPU layout for GLM 5.2:
```sh
./ds4 -m gguf/GLM-5.2-UD-Q2_K_RoutedQ2K.gguf \
--gpu-vram auto --gpu-devices 0,2,4,6,1,3,5,7 \
--ctx 32768 -p "Hello"
```
For the DeepSeek Flash tensor/expert-parallel layout, see
"Tensor Parallelism across CUDA GPUs" above.
There is also a CPU reference/debug path:
```sh
@@ -1224,10 +1644,9 @@ captured from the official DeepSeek V4 Flash API. The requests use
`top_logprobs` slice exposed by the API. Local vectors are generated with
`./ds4 --dump-logprobs` and compared by token bytes, so tokenizer/template or
attention regressions show up before they become long generation failures. The
C runner pins `DS4_METAL_PREFILL_CHUNK=2048` for this strict API-vector
comparison.
C runner pins a 2048-token prefill chunk for this strict API-vector comparison.
All project tests are driven by the C runner, with a small `ds4-eval`
The core local tests are driven by the C runner, with a small `ds4-eval`
extractor self-test run first:
```sh
@@ -1236,6 +1655,24 @@ make test # ./ds4-eval --self-test-extractors && ./ds4_test --a
./ds4_test --server
```
The batching tests are model-backed and must run on the matching GPU backend:
```sh
# Metal, with DS4_TEST_SESSION_COUNT set to 2, 4, 8, and 16.
DS4_TEST_MODEL=/path/to/model.gguf DS4_TEST_SESSION_COUNT=4 \
make test-metal-session-batch
# CUDA multi-GPU Flash.
DS4_TEST_MODEL=/path/to/model.gguf make test-cuda-session-batch
DS4_TEST_MODEL=/path/to/model.gguf make test-cuda-mixed-batch
```
For GLM, run the same Metal session test with a GLM GGUF and run
`tests/glm_long_context_smoke.sh /path/to/model.gguf`. The official 100-case
quality scorers, two-Mac TCP/RDMA tests, CUDA matrix, and manual agent checks are
release gates rather than quick local tests; follow
[QA_BEFORE_RELEASES.md](QA_BEFORE_RELEASES.md).
## Debugging Notes
When a generation looks wrong, three small tools are usually enough to get a
@@ -1257,3 +1694,9 @@ first answer:
logit/model issues.
- `ds4-server --trace` writes the rendered prompts, cache decisions, generated
text, and tool-parser events for a whole agent session.
## Logo
The DwarfStar logo was designed by hand by Salvatore Sanfilippo, made more
graphical with AI, and manually reworked by Ben Gnomino, whose human touch made
it rock.
+159 -10
View File
@@ -1,6 +1,11 @@
#!/bin/sh
set -e
GLM_UNSLOTH_REPO="unsloth/GLM-5.2-GGUF"
GLM_ANTIREZ_REPO="antirez/GLM-5.2-GGUF"
LAGUNA_REPO="poolside/Laguna-S-2.1-GGUF"
LAGUNA_ANTIREZ_REPO="antirez/Laguna-S-2.1-GGUF"
LAGUNA_REVISION="706fa69799926b6afde1af9e24ca2a4923f110a1"
REPO="antirez/deepseek-v4-gguf"
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"
@@ -9,6 +14,16 @@ PRO_Q2_IMATRIX_FILE="DeepSeek-V4-Pro-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-Instruct-
PRO_Q4_LAYERS00_30_FILE="DeepSeek-V4-Pro-Q4K-Layers00-30.gguf"
PRO_Q4_LAYERS31_OUTPUT_FILE="DeepSeek-V4-Pro-Q4K-Layers-31-output.gguf"
MTP_FILE="DeepSeek-V4-Flash-MTP-Q4K-Q8_0-F32.gguf"
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"
LAGUNA_Q4_FILE="laguna-s-2.1-Q4_K_M.gguf"
LAGUNA_Q2_Q3_FILE="laguna-s-2.1-RoutedQ2_K-Last27Q3_K.gguf"
LAGUNA_DFLASH_FILE="laguna-s-2.1-DFlash-Q8_0.gguf"
ROOT=$(CDPATH= cd -- "$(dirname -- "$0")" && pwd)
OUT_DIR=${DS4_GGUF_DIR:-"$ROOT/gguf"}
@@ -17,10 +32,11 @@ case "$OUT_DIR" in
*) OUT_DIR="$ROOT/$OUT_DIR" ;;
esac
TOKEN=${HF_TOKEN:-}
HF_REVISION=
usage() {
cat <<EOF
DeepSeek V4 GGUF downloader
DwarfStar GGUF downloader
Usage:
./download_model.sh q2-imatrix [--token TOKEN]
@@ -31,6 +47,14 @@ Usage:
./download_model.sh pro-q4-layers31-output [--token TOKEN]
./download_model.sh pro-q4-split [--token TOKEN]
./download_model.sh mtp [--token TOKEN]
./download_model.sh dspark-support [--token TOKEN]
./download_model.sh glm-unsloth-q4 [--token TOKEN]
./download_model.sh glm-antirez-iq2xxs [--token TOKEN]
./download_model.sh glm-antirez-q2 [--token TOKEN]
./download_model.sh glm-antirez-q4 [--token TOKEN]
./download_model.sh laguna-q4 [--token TOKEN]
./download_model.sh laguna-q2-q3 [--token TOKEN]
./download_model.sh laguna-dflash [--token TOKEN]
Targets:
@@ -69,6 +93,42 @@ Targets:
It is useful with q2-imatrix, q2-q4-imatrix, and q4-imatrix, but must be
enabled explicitly with --mtp when running ds4 or ds4-server.
dspark-support
Optional DSpark speculative decoding support GGUF, about 6 GB. Enable it
with --dspark and --mtp when running ds4 or ds4-server.
glm-unsloth-q4
GLM 5.2 Unsloth UD-Q4_K_XL quant from unsloth/GLM-5.2-GGUF.
Downloads all 11 shards and links ./ds4flash.gguf to the first shard.
glm-antirez-iq2xxs
GLM 5.2 antirez routed IQ2_XXS GGUF from antirez/GLM-5.2-GGUF.
Includes Q2_K block 78 and is intended for reduced-memory testing.
glm-antirez-q2
GLM 5.2 antirez routed Q2_K GGUF from antirez/GLM-5.2-GGUF.
About 262 GB on disk.
glm-antirez-q4
GLM 5.2 antirez routed Q4_K GGUF from antirez/GLM-5.2-GGUF.
About 434 GB on disk.
laguna-q4
Official imatrix-quantized Laguna S 2.1 Q4_K_M GGUF from Poolside.
About 68 GB on disk; supported by Metal, CUDA, and ROCm with full model
residency.
laguna-q2-q3
Mixed Laguna S 2.1 routed-expert quant for 64 GB systems. Routed
layers 1..20 use Q2_K and layers 21..47 use Q3_K; all other tensors
retain the official Q4_K_M layout. 44.95 GiB on disk; supported by
Metal, CUDA, and ROCm with full model residency.
laguna-dflash
Q8_0 Laguna S 2.1 DFlash speculative support GGUF, quantized from
Poolside's official support model. About 1.04 GiB on disk. This is an
optional support model and does not replace or relink the main model.
Options:
--token TOKEN Hugging Face token. Otherwise HF_TOKEN or the local HF token
cache is used if present.
@@ -87,8 +147,12 @@ Then the default commands work:
After downloading mtp, enable it explicitly, for example:
./ds4 --mtp <download directory>/$MTP_FILE --mtp-draft 2
PRO files are downloaded with the official Hugging Face downloader because
they are too large for the curl path used by the smaller GGUF files.
After downloading DSpark support, enable it explicitly in greedy mode:
./ds4 --dspark --mtp <download directory>/$DSPARK_SUPPORT_FILE --temp 0
Large PRO, GLM, and Laguna files use the official Hugging Face downloader
because they are too large, sharded, or nested for the curl path used by the
smaller DeepSeek Flash GGUF files.
EOF
}
@@ -101,6 +165,8 @@ MODEL=$1
shift
MODEL_FILES=
LINK_MODEL=1
FORCE_HF_DOWNLOAD=0
FLATTEN_DOWNLOADS=0
case "$MODEL" in
q2-imatrix) MODEL_FILE=$Q2_IMATRIX_FILE ;;
@@ -114,6 +180,49 @@ case "$MODEL" in
LINK_MODEL=0
;;
mtp) MODEL_FILE=$MTP_FILE; LINK_MODEL=0 ;;
dspark-support) MODEL_FILE=$DSPARK_SUPPORT_FILE; LINK_MODEL=0 ;;
glm-unsloth-q4)
REPO=$GLM_UNSLOTH_REPO
MODEL_FILE=$GLM_UNSLOTH_Q4_FIRST_FILE
MODEL_FILES=
for part in 00001 00002 00003 00004 00005 00006 00007 00008 00009 00010 00011; do
MODEL_FILES="$MODEL_FILES $GLM_UNSLOTH_Q4_REMOTE_BASE-${part}-of-00011.gguf"
done
FORCE_HF_DOWNLOAD=1
FLATTEN_DOWNLOADS=1
;;
glm-antirez-q2)
REPO=$GLM_ANTIREZ_REPO
MODEL_FILE=$GLM_ANTIREZ_Q2_FILE
FORCE_HF_DOWNLOAD=1
;;
glm-antirez-iq2xxs)
REPO=$GLM_ANTIREZ_REPO
MODEL_FILE=$GLM_ANTIREZ_IQ2XXS_FILE
FORCE_HF_DOWNLOAD=1
;;
glm-antirez-q4)
REPO=$GLM_ANTIREZ_REPO
MODEL_FILE=$GLM_ANTIREZ_Q4_FILE
FORCE_HF_DOWNLOAD=1
;;
laguna-q4)
REPO=$LAGUNA_REPO
MODEL_FILE=$LAGUNA_Q4_FILE
FORCE_HF_DOWNLOAD=1
HF_REVISION=$LAGUNA_REVISION
;;
laguna-q2-q3)
REPO=$LAGUNA_ANTIREZ_REPO
MODEL_FILE=$LAGUNA_Q2_Q3_FILE
FORCE_HF_DOWNLOAD=1
;;
laguna-dflash)
REPO=$LAGUNA_ANTIREZ_REPO
MODEL_FILE=$LAGUNA_DFLASH_FILE
FORCE_HF_DOWNLOAD=1
LINK_MODEL=0
;;
-h|--help|help)
usage
exit 0
@@ -149,6 +258,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
@@ -164,15 +276,31 @@ find_hf_command() {
printf '%s\n' hf
return 0
fi
for dir in "$HOME"/Library/Python/*/bin "$HOME"/.local/bin; do
if [ -x "$dir/hf" ]; then
printf '%s\n' "$dir/hf"
return 0
fi
done
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"
@@ -182,13 +310,13 @@ download_one_hf() {
if [ -e "$part" ]; then
echo "Found curl partial download: $part" >&2
echo "The Hugging Face downloader cannot resume curl .part files." >&2
echo "Move or remove that partial download before retrying this PRO target." >&2
echo "Move or remove that partial download before retrying this target." >&2
exit 1
fi
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
@@ -196,15 +324,27 @@ download_one_hf() {
echo "Downloading $file"
echo "from https://huggingface.co/$REPO"
if [ -n "$HF_REVISION" ]; then
echo "revision $HF_REVISION"
fi
echo "using $HF_CMD download"
echo "If the download stops, run the same command again to resume it."
if [ -n "$TOKEN" ]; then
if [ -n "$TOKEN" ] && [ -n "$HF_REVISION" ]; then
"$HF_CMD" download "$REPO" "$file" --revision "$HF_REVISION" --repo-type model --local-dir "$OUT_DIR" --token "$TOKEN"
elif [ -n "$TOKEN" ]; then
"$HF_CMD" download "$REPO" "$file" --repo-type model --local-dir "$OUT_DIR" --token "$TOKEN"
elif [ -n "$HF_REVISION" ]; then
"$HF_CMD" download "$REPO" "$file" --revision "$HF_REVISION" --repo-type model --local-dir "$OUT_DIR"
else
"$HF_CMD" download "$REPO" "$file" --repo-type model --local-dir "$OUT_DIR"
fi
if [ "$hf_out" != "$out" ] && [ -s "$hf_out" ]; then
mv "$hf_out" "$out"
rmdir "$(dirname "$hf_out")" 2>/dev/null || true
fi
if [ ! -s "$out" ]; then
echo "Hugging Face download finished but expected file is missing: $out" >&2
exit 1
@@ -213,7 +353,8 @@ download_one_hf() {
download_one() {
file=$1
out="$OUT_DIR/$file"
local_file=$(local_download_name "$file")
out="$OUT_DIR/$local_file"
part="$out.part"
aria2_part="$out.aria2"
url="https://huggingface.co/$REPO/resolve/main/$file"
@@ -223,7 +364,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
@@ -262,6 +403,14 @@ if [ "$MODEL" = "mtp" ]; then
echo "MTP is an optional component for q2-imatrix, q2-q4-imatrix, and q4-imatrix."
echo "Enable it explicitly, for example:"
echo " ./ds4 --mtp $OUT_DIR/$MTP_FILE --mtp-draft 2"
elif [ "$MODEL" = "dspark-support" ]; then
echo
echo "DSpark support downloaded. Enable it explicitly in greedy mode:"
echo " ./ds4 --dspark -m ./ds4flash.gguf --mtp $OUT_DIR/$DSPARK_SUPPORT_FILE --temp 0"
elif [ "$MODEL" = "laguna-dflash" ]; then
echo
echo "Laguna DFlash support downloaded. Enable it explicitly in greedy mode:"
echo " ./ds4 -m <laguna-model.gguf> --dflash $OUT_DIR/$LAGUNA_DFLASH_FILE --temp 0"
elif [ "$MODEL" = "pro-q4-layers00-30" ] || [ "$MODEL" = "pro-q4-layers31-output" ] || [ "$MODEL" = "pro-q4-split" ]; then
echo
echo "Downloaded PRO Q4 distributed split file(s). Use them with --layers,"
+43777 -1897
View File
File diff suppressed because it is too large. Load diff
+154
View File
@@ -91,14 +91,52 @@ typedef struct {
bool debug;
} ds4_distributed_options;
/* 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,
DS4_TP_WORKER,
} ds4_tp_role;
typedef enum {
DS4_TP_TRANSPORT_AUTO = 0,
DS4_TP_TRANSPORT_RDMA,
DS4_TP_TRANSPORT_TCP,
} ds4_tp_transport;
typedef struct {
ds4_tp_role role;
bool requested; /* --tensor-parallel with shared role options */
const char *listen_host; /* leader listens here for the worker */
int listen_port;
const char *leader_host; /* worker dials the leader */
int leader_port;
ds4_tp_transport transport;
const char *rdma_device;
int rdma_gid_index;
bool rdma_gid_index_set;
bool glm_token_prefill;
int debug_hash; /* cross-check hidden state every N tokens */
} ds4_tp_options;
typedef struct {
const char *model_path;
const char *mtp_path;
const char *dflash_path;
ds4_backend backend;
int n_threads;
int context_size;
uint32_t prefill_chunk;
int mtp_draft_tokens;
int dflash_draft_tokens;
float dflash_p_min;
float mtp_margin;
float dspark_confidence_threshold;
const char *directional_steering_file;
const char *expert_profile_path;
float directional_steering_attn;
@@ -106,18 +144,34 @@ 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 glm_mtp;
bool glm_mtp_timing;
bool dspark;
bool dspark_strict;
bool dflash_p_min_set;
bool dspark_confidence_threshold_set;
bool cuda_tensor_parallel;
bool ssd_streaming;
bool ssd_streaming_cold;
bool ssd_streaming_full_layers_set;
bool inspect_only;
/* Multi-GPU placement uses this to price per-layer KV storage. */
int placement_ctx_hint;
/* Server batch mode serializes execution and can share prefill scratch. */
bool share_session_prefill_workspace;
bool first_token_test;
bool metal_graph_test;
bool load_slice;
uint32_t load_layer_start;
uint32_t load_layer_end;
bool load_output;
ds4_distributed_options distributed;
ds4_tp_options tp;
} ds4_engine_options;
typedef void (*ds4_token_emit_fn)(void *ud, int token);
@@ -145,23 +199,67 @@ typedef struct {
} ds4_session_payload_file;
int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt);
/* Multi-GPU pipeline-parallel entry point (wave 2).
*
* Accepts an optional ds4_gpu_config (defined in ds4_gpu_mgpu.h) that
* lets callers describe a multi-GPU placement target. Passing NULL is
* back-compatible with ds4_engine_open and produces identical engine
* state — bit-equivalent execution at runtime.
*
* When a non-NULL config is supplied AND the computed placement spans
* more than one tier (either multiple GPUs or any CPU-spill), this
* wave-2 implementation prints the layout and refuses to open: full
* multi-tier execution wiring lands in a follow-up task
* (mgpu-graph-session-execution). Callers receive a non-zero return
* and a documented stderr notice. */
/* ds4_gpu_config is declared in ds4_gpu_mgpu.h, which callers should
* include separately. We forward-declare it here so this header can be
* used as-is (callers passing NULL don't need the struct definition). */
struct ds4_gpu_config;
int ds4_engine_create_with_gpu_config(ds4_engine **out,
const ds4_engine_options *opt,
const struct ds4_gpu_config *gpu_cfg);
void ds4_engine_close(ds4_engine *e);
void ds4_engine_summary(ds4_engine *e);
int ds4_engine_vocab_size(ds4_engine *e);
uint32_t ds4_engine_prefill_chunk(ds4_engine *e);
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);
bool ds4_engine_is_laguna(ds4_engine *e);
const char *ds4_engine_default_system_prompt(ds4_engine *e);
void ds4_engine_sampling_defaults(ds4_engine *e, float *temperature,
int *top_k, float *top_p, float *min_p);
const char *ds4_backend_name(ds4_backend backend);
bool ds4_think_mode_enabled(ds4_think_mode mode);
const char *ds4_think_mode_name(ds4_think_mode mode);
const char *ds4_think_max_prefix(void);
const char *ds4_glm_reasoning_effort_text(ds4_think_mode mode);
uint32_t ds4_think_max_min_context(void);
ds4_think_mode ds4_think_mode_for_context(ds4_think_mode mode, int ctx_size);
/* Uses the active model shape selected by ds4_engine_open(); call after opening
@@ -171,6 +269,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,
@@ -211,12 +314,25 @@ void ds4_encode_chat_prompt(
void ds4_chat_append_max_effort_prefix(ds4_engine *e, ds4_tokens *tokens);
void ds4_chat_append_message(ds4_engine *e, ds4_tokens *tokens, const char *role, const char *content);
void ds4_chat_append_assistant_prefix(ds4_engine *e, ds4_tokens *tokens, ds4_think_mode think_mode);
void ds4_chat_append_assistant_end(ds4_engine *e, ds4_tokens *tokens);
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);
/* Tensor-parallel binding: allocates the GPU gate slab, registers it with
* the transport and arms the per-layer gate machinery. Call once, after
* ds4_tp_create() and before any session work. Transport lifecycle stays
* with the caller. */
struct ds4_tp;
int ds4_engine_tp_bind(ds4_engine *e, struct ds4_tp *tp, char *err, size_t errlen);
int ds4_session_create(ds4_session **out, ds4_engine *e, int ctx_size);
void ds4_session_free(ds4_session *s);
int ds4_session_power(ds4_session *s);
@@ -230,6 +346,10 @@ void ds4_session_set_display_progress(ds4_session *s, ds4_session_progress_fn fn
* safe boundaries where the live checkpoint is either unchanged or represents a
* valid token prefix, and returns DS4_SESSION_SYNC_INTERRUPTED when it stops. */
void ds4_session_set_cancel(ds4_session *s, ds4_session_cancel_fn fn, void *ud);
/* Enable optional speculative support-model state for the next sync/decode
* sequence. Frontends disable it for sampling modes that cannot use a greedy
* verifier, avoiding support-model work on every ordinary decode token. */
void ds4_session_set_speculative_enabled(ds4_session *s, bool enabled);
void ds4_session_report_progress(ds4_session *s, const char *event, int current, int total);
/* Distributed coordinator sessions return 1 when the full layer route is
* available, 0 when it is still incomplete, and -1 for a local API error. */
@@ -246,6 +366,7 @@ typedef enum {
/* Synchronize the live session to a full prompt token prefix. If the current
* checkpoint is a prefix, only the suffix is evaluated; otherwise the backend
* state is refilled from scratch. */
#define DS4_SESSION_SYNC_INTERRUPTED 2
int ds4_session_sync(ds4_session *s, const ds4_tokens *prompt, char *err, size_t errlen);
bool ds4_session_rewrite_requires_rebuild(int live_len, int canonical_len, int common);
ds4_session_rewrite_result ds4_session_rewrite_from_common(
@@ -257,15 +378,48 @@ int ds4_session_argmax_excluding(ds4_session *s, int excluded_id);
int ds4_sample_logits(const float *logits, int n_vocab, float temperature,
int top_k, float top_p, float min_p, uint64_t *rng);
int ds4_session_sample(ds4_session *s, float temperature, int top_k, float top_p, float min_p, uint64_t *rng);
#ifdef DS4_TEST_HOOKS
int ds4_test_sample_logits(const float *logits, uint32_t n_vocab,
float temperature, int top_k,
float top_p, float min_p, uint64_t *rng,
float *prob_scratch);
uint64_t ds4_test_mixed_native_count(void);
#endif
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);
typedef struct {
ds4_session *session;
int token;
} ds4_decode_item;
/* Advance independent sessions by one token each. Batch size one is exactly
* ds4_session_eval(). Backends without native batching use a correctness-first
* sequential fallback. */
int ds4_sessions_eval_batch(ds4_decode_item *items, int count,
char *err, size_t errlen);
/* Advance one resumed prefill suffix and an independent decode batch as one
* scheduling step. Unsupported combinations use the ordinary serialized
* session operations. */
int ds4_sessions_eval_batch_with_prefill(
ds4_decode_item *items, int count,
ds4_session *prefill_session, const ds4_tokens *prefill_prompt,
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 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);
void ds4_session_invalidate(ds4_session *s);
void ds4_session_rewind(ds4_session *s, int pos);
int ds4_session_pos(ds4_session *s);
+1419 -124
View File
File diff suppressed because it is too large. Load diff
+163 -24
View File
@@ -1,15 +1,16 @@
#include "ds4.h"
#include "ds4_distributed.h"
#include "ds4_gpu_args.h"
#include "ds4_help.h"
/* Purpose-built throughput benchmark.
*
* 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 +23,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;
@@ -29,6 +32,8 @@ typedef struct {
const char *system;
const char *csv_path;
const char *expert_profile_path;
const char *gpu_vram_arg;
const char *gpu_devices_arg;
ds4_backend backend;
int threads;
int ctx_start;
@@ -40,6 +45,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 +55,9 @@ typedef struct {
bool quality;
bool ssd_streaming;
bool ssd_streaming_cold;
bool ssd_streaming_full_layers_set;
bool cuda_tensor_parallel;
bool show_output;
} bench_config;
static double bench_now_sec(void) {
@@ -57,6 +66,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);
}
@@ -241,6 +273,12 @@ static bench_config parse_options(int argc, char **argv) {
} else if (!strcmp(arg, "--cuda")) {
c.backend = DS4_BACKEND_CUDA;
#endif
} else if (!strcmp(arg, "--gpu-vram")) {
c.gpu_vram_arg = need_arg(&i, argc, argv, arg);
} else if (!strcmp(arg, "--gpu-devices")) {
c.gpu_devices_arg = need_arg(&i, argc, argv, arg);
} else if (!strcmp(arg, "--cuda-tensor-parallel")) {
c.cuda_tensor_parallel = true;
} else if (!strcmp(arg, "--cpu")) {
c.backend = DS4_BACKEND_CPU;
} else if (!strcmp(arg, "--quality")) {
@@ -260,6 +298,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) {
@@ -284,6 +326,8 @@ static bench_config parse_options(int argc, char **argv) {
}
} else if (!strcmp(arg, "--warm-weights")) {
c.warm_weights = true;
} else if (!strcmp(arg, "--show-output")) {
c.show_output = true;
} else {
fprintf(stderr, "ds4-bench: unknown option: %s\n", arg);
usage(stderr, NULL);
@@ -441,11 +485,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),
@@ -508,20 +554,44 @@ static void maybe_warn_distributed_step_shape(const bench_config *cfg, ds4_sessi
int main(int argc, char **argv) {
bench_config cfg = parse_options(argc, argv);
/* Hint the packer at the largest ctx this bench run will exercise
* so per-layer KV bytes are priced for the real session size, not
* a stale 4096 default. Single-tier and CPU paths ignore this. */
int placement_ctx_hint = cfg.ctx_max;
if (cfg.ctx_alloc > placement_ctx_hint) placement_ctx_hint = cfg.ctx_alloc;
ds4_gpu_config gpu_cfg = {0};
bool skip_cuda = false;
const bool have_gpu_config = cfg.gpu_vram_arg || cfg.gpu_devices_arg;
if (have_gpu_config) {
char gpu_err[256];
if (parse_gpu_vram_arg(cfg.gpu_vram_arg, cfg.gpu_devices_arg,
&gpu_cfg, &skip_cuda,
gpu_err, sizeof(gpu_err)) != 0) {
fprintf(stderr, "ds4-bench: %s\n", gpu_err);
return 2;
}
cfg.backend = skip_cuda ? DS4_BACKEND_CPU : DS4_BACKEND_CUDA;
}
ds4_engine_options opt = {
.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,
.warm_weights = cfg.warm_weights,
.quality = cfg.quality,
.cuda_tensor_parallel = cfg.cuda_tensor_parallel,
.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,
};
@@ -531,8 +601,25 @@ int main(int argc, char **argv) {
return 2;
}
ds4_engine *engine = NULL;
if (ds4_engine_open(&engine, &opt) != 0) return 1;
log_context_memory(cfg.backend, cfg.ctx_alloc, cfg.prefill_chunk);
if (have_gpu_config && !skip_cuda) {
const bool was_auto =
(cfg.gpu_vram_arg && !strcmp(cfg.gpu_vram_arg, "auto")) ||
(!cfg.gpu_vram_arg && cfg.gpu_devices_arg);
char layout[256];
if (format_gpu_layout_line(&gpu_cfg, was_auto,
layout, sizeof(layout)) > 0) {
fprintf(stdout, "%s\n", layout);
fflush(stdout);
}
if (ds4_engine_create_with_gpu_config(
&engine, &opt, &gpu_cfg) != 0) return 1;
} else if (ds4_engine_open(&engine, &opt) != 0) {
return 1;
}
log_context_memory(opt.backend,
cfg.ctx_alloc,
ds4_engine_prefill_chunk(engine),
cfg.ssd_streaming);
char *text = read_file(cfg.prompt_path ? cfg.prompt_path : cfg.chat_prompt_path);
ds4_tokens prompt = {0};
@@ -581,12 +668,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 +702,41 @@ 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;
int *gen_token_buf = cfg.show_output && cfg.gen_tokens > 0
? malloc((size_t)cfg.gen_tokens * sizeof(gen_token_buf[0]))
: NULL;
int gen_token_count = 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 +749,40 @@ 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;
if (gen_token_buf) gen_token_buf[gen_token_count++] = token;
gen_done++;
}
const double gen_t1 = bench_now_sec();
if (cfg.show_output && gen_token_buf && gen_token_count > 0) {
fprintf(stderr, "ds4-bench: gen[ctx=%d] decoded text: \"", frontier);
for (int i = 0; i < gen_token_count; i++) {
size_t tlen = 0;
char *txt = ds4_token_text(engine, gen_token_buf[i], &tlen);
if (txt) {
fwrite(txt, 1, tlen, stderr);
free(txt);
}
}
fprintf(stderr, "\"\n");
fflush(stderr);
}
free(gen_token_buf);
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 +795,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;
+606 -79
View File
@@ -1,11 +1,13 @@
#include "ds4.h"
#include "ds4_distributed.h"
#include "ds4_gpu_args.h"
#include "ds4_tp.h"
#include "ds4_help.h"
#include "linenoise.h"
/* 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
@@ -25,19 +27,58 @@
#include <time.h>
#include <unistd.h>
static bool cli_env_flag_enabled(const char *name, bool defval) {
const char *v = getenv(name);
if (!v || !v[0]) return defval;
return strcmp(v, "0") != 0;
}
static bool cli_splitkv_spec_requested(void) {
if (cli_env_flag_enabled("DS4_CUDA_NO_SPLITKV_SPEC", false)) return false;
return cli_env_flag_enabled("DS4_CUDA_SPLITKV_SPEC", false);
}
static bool cli_greedy_fast_attention_requested(void) {
if (!cli_env_flag_enabled("DS4_CUDA_NO_GREEDY_SPLITKV", false) &&
cli_env_flag_enabled("DS4_CUDA_GREEDY_SPLITKV", false))
{
return true;
}
if (!cli_env_flag_enabled("DS4_CUDA_NO_GREEDY_VEC4", false) &&
cli_env_flag_enabled("DS4_CUDA_GREEDY_VEC4", false))
{
return true;
}
return false;
}
static bool cli_greedy_argmax_requested(bool speculative_requested) {
if (cli_greedy_fast_attention_requested()) return true;
if (speculative_requested) return false;
return cli_env_flag_enabled("DS4_CUDA_GREEDY_TOP1", true);
}
typedef struct {
const char *prompt;
const char *system;
bool system_set;
bool raw_prompt;
int n_predict;
int ctx_size;
float temperature;
int top_k;
float top_p;
float min_p;
bool temperature_set;
bool top_k_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;
@@ -57,6 +98,10 @@ typedef struct {
cli_generation_options gen;
char *prompt_owned;
bool inspect;
/* CLI flag wiring: raw argv values for --gpu-vram and --gpu-devices.
* Resolved post-parse via parse_gpu_vram_arg(). */
const char *gpu_vram_arg;
const char *gpu_devices_arg;
} cli_config;
static volatile sig_atomic_t cli_interrupted;
@@ -142,6 +187,16 @@ static int parse_int(const char *s, const char *opt) {
return (int)v;
}
static int parse_nonnegative_int(const char *s, const char *opt) {
char *end = NULL;
long v = strtol(s, &end, 10);
if (s[0] == '\0' || *end != '\0' || v < 0 || v > INT32_MAX) {
fprintf(stderr, "ds4: invalid value for %s: %s\n", opt, s);
exit(2);
}
return (int)v;
}
static uint64_t parse_u64(const char *s, const char *opt) {
char *end = NULL;
unsigned long long v = strtoull(s, &end, 10);
@@ -191,19 +246,40 @@ static ds4_backend default_backend(void) {
static void log_context_memory(ds4_backend backend,
int ctx_size,
uint32_t prefill_chunk) {
uint32_t prefill_chunk,
bool ssd_streaming) {
ds4_context_memory m =
ds4_context_memory_estimate_with_prefill(backend,
ctx_size,
prefill_chunk);
ds4_context_memory_estimate_with_prefill_mode(backend,
ctx_size,
prefill_chunk,
ssd_streaming);
const uint64_t kv_bytes = m.raw_bytes + m.compressed_bytes;
const uint64_t total_bytes = kv_bytes + m.scratch_bytes;
const bool color = ds4_log_is_tty(stderr);
const char *green = color ? "\x1b[32m" : "";
const char *bright_green = color ? "\x1b[1;32m" : "";
const char *reset = color ? "\x1b[0m" : "";
fprintf(stderr,
"ds4: context buffers %.2f MiB (ctx=%d, backend=%s, prefill_chunk=%u, raw_kv_rows=%u, compressed_kv_rows=%u)\n",
(double)m.total_bytes / (1024.0 * 1024.0),
"%sds4: memory: KV %.2f GiB (raw %.2f + compressed %.2f) "
"+ buffers %.2f GiB = %s%.2f GiB context%s\n",
green,
(double)kv_bytes / 1073741824.0,
(double)m.raw_bytes / 1073741824.0,
(double)m.compressed_bytes / 1073741824.0,
(double)m.scratch_bytes / 1073741824.0,
bright_green,
(double)total_bytes / 1073741824.0,
reset);
fprintf(stderr,
"%sds4: memory detail: ctx=%d prefill_cap=%u raw_kv_rows=%u "
"compressed_kv_rows=%u backend=%s%s\n",
green,
ctx_size,
ds4_backend_name(backend),
m.prefill_cap,
m.raw_cap,
m.comp_cap);
m.comp_cap,
ds4_backend_name(backend),
reset);
}
static ds4_think_mode cli_effective_think_mode(const cli_generation_options *gen) {
@@ -244,7 +320,8 @@ static void cli_prefill_progress_cb(void *ud, const char *event, int current, in
(void)total;
cli_prefill_progress *p = ud;
if (!p || !event || p->input_tokens <= 0) return;
const bool is_display = strcmp(event, "prefill_display") == 0;
const bool is_display =
strcmp(event, "prefill_display") == 0;
if (strcmp(event, "prefill_chunk") && !is_display) return;
if (is_display && !p->use_color) return;
@@ -281,8 +358,23 @@ 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[] = {
"〈|EOS|〉",
"<|begin▁of▁sentence|>",
"<|User|>",
"[gMASK]",
"<sop>",
"<|system|>",
"<|user|>",
"<|assistant|>",
"<|observation|>",
};
if (!prompt) return false;
for (size_t i = 0; i < sizeof(prefixes) / sizeof(prefixes[0]); i++) {
size_t n = strlen(prefixes[i]);
if (strncmp(prompt, prefixes[i], n) == 0) return true;
}
return false;
}
typedef struct {
@@ -410,7 +502,9 @@ static void print_generated_token(void *ud, int token) {
}
static void build_prompt(ds4_engine *engine, const cli_generation_options *gen, ds4_tokens *out) {
if (is_rendered_chat_prompt(gen->prompt)) {
if (gen->raw_prompt) {
ds4_tokenize_text(engine, gen->prompt ? gen->prompt : "", out);
} else if (is_rendered_chat_prompt(gen->prompt)) {
ds4_tokenize_rendered_chat(engine, gen->prompt, out);
} else {
ds4_encode_chat_prompt(engine, gen->system, gen->prompt,
@@ -418,6 +512,19 @@ 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) return;
float temperature, top_p, min_p;
int top_k;
ds4_engine_sampling_defaults(engine, &temperature, &top_k, &top_p, &min_p);
if (!gen->temperature_set) gen->temperature = temperature;
if (!gen->top_k_set) gen->top_k = top_k;
if (!gen->top_p_set) gen->top_p = top_p;
if (!gen->min_p_set) gen->min_p = min_p;
}
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) {
@@ -428,6 +535,13 @@ 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);
ds4_session_set_speculative_enabled(
session,
cfg->gen.temperature <= 0.0f &&
getenv("DS4_MTP_SPEC_DISABLE") == NULL);
char err[160];
ds4_think_mode think_mode = cli_effective_think_mode(&cfg->gen);
@@ -472,11 +586,25 @@ static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, con
uint64_t rng = cfg->gen.seed ? cfg->gen.seed :
((uint64_t)time(NULL) ^ ((uint64_t)getpid() << 32) ^ (uint64_t)clock());
int generated = 0;
const bool speculative_argmax = cfg->gen.temperature <= 0.0f &&
((ds4_engine_mtp_draft_tokens(engine) > 1 &&
getenv("DS4_MTP_SPEC_DISABLE") == NULL) ||
cli_splitkv_spec_requested());
const bool greedy_argmax = cfg->gen.temperature <= 0.0f &&
cli_greedy_argmax_requested(speculative_argmax);
bool have_greedy_next = false;
int greedy_next = -1;
const double t_decode0 = cli_now_sec();
while (generated < max_tokens && !cli_interrupt_requested()) {
int token = ds4_session_sample(session, cfg->gen.temperature, 0,
int token;
if (greedy_argmax && have_greedy_next) {
token = greedy_next;
have_greedy_next = false;
} else {
token = ds4_session_sample(session, cfg->gen.temperature, cfg->gen.top_k,
cfg->gen.top_p, cfg->gen.min_p, &rng);
if (token == ds4_token_eos(engine)) break;
}
if (ds4_token_is_stop_for_think_mode(engine, token, think_mode)) break;
int toks[17];
int ntok = 0;
@@ -498,6 +626,16 @@ static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, con
return 1;
}
} else {
size_t piece_len = 0;
char *piece = ds4_token_text(engine, token, &piece_len);
token_printer_write_text(&printer, piece, piece_len);
fflush(stdout);
free(piece);
generated++;
if (generated >= max_tokens || cli_interrupt_requested()) {
continue;
}
cli_dist_busy_set(cfg, true);
int eval_rc = ds4_session_eval(session, token, err, sizeof(err));
cli_dist_busy_set(cfg, false);
@@ -506,13 +644,12 @@ static int run_sampled_generation(ds4_engine *engine, const cli_config *cfg, con
ds4_session_free(session);
return 1;
}
toks[0] = token;
ntok = 1;
continue;
}
bool stop = false;
for (int j = 0; j < ntok; j++) {
if (toks[j] == ds4_token_eos(engine)) {
if (ds4_token_is_stop_for_think_mode(engine, toks[j], think_mode)) {
stop = true;
break;
}
@@ -766,7 +903,7 @@ static int run_logprob_dump(ds4_engine *engine, const cli_config *cfg, const ds4
}
fputs("]}", fp);
if (token == ds4_token_eos(engine)) break;
if (ds4_token_is_stop(engine, token)) break;
if (ds4_session_eval(session, token, err, sizeof(err)) != 0) {
fprintf(stderr, "ds4: decode failed while dumping logprobs: %s\n", err);
free(scores);
@@ -787,6 +924,176 @@ static int run_logprob_dump(ds4_engine *engine, const cli_config *cfg, const ds4
return 0;
}
static void print_diag_token(FILE *fp, ds4_engine *engine, int token) {
fputs("{\"token\":", fp);
json_write_token(fp, engine, token);
fputc('}', fp);
}
static void print_diag_top(FILE *fp, ds4_engine *engine, const char *label,
const ds4_token_score *scores, int n) {
fprintf(fp, "%s:", label);
for (int i = 0; i < n && scores[i].id >= 0; i++) {
fputc(' ', fp);
json_write_token(fp, engine, scores[i].id);
fprintf(fp, "@%.6g", scores[i].logit);
}
fputc('\n', fp);
}
static int run_decode_consistency(ds4_engine *engine, const cli_config *cfg,
const ds4_tokens *prompt) {
if (cfg->dist && cfg->dist->role != DS4_DISTRIBUTED_NONE) {
fprintf(stderr, "ds4: --decode-consistency is local-session only\n");
return 1;
}
if (cfg->gen.decode_consistency_tokens < 0) {
fprintf(stderr, "ds4: --decode-consistency requires a non-negative token count\n");
return 1;
}
const int vocab = ds4_engine_vocab_size(engine);
float *live_logits = malloc((size_t)vocab * sizeof(live_logits[0]));
float *fresh_logits = malloc((size_t)vocab * sizeof(fresh_logits[0]));
if (!live_logits || !fresh_logits) {
free(live_logits);
free(fresh_logits);
fprintf(stderr, "ds4: out of memory allocating diagnostic logits\n");
return 1;
}
ds4_tokens prefix = {0};
ds4_tokens_copy(&prefix, prompt);
ds4_session *live = NULL;
char err[160];
if (ds4_session_create(&live, engine, cfg->gen.ctx_size) != 0) {
fprintf(stderr, "ds4: --decode-consistency requires a graph session backend\n");
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return 1;
}
cli_prefill_progress progress = {
.base_tokens = 0,
.input_tokens = prompt->len,
.use_color = ds4_log_is_tty(stderr),
};
ds4_session_set_progress(live, cli_prefill_progress_cb, &progress);
ds4_session_set_display_progress(live,
progress.use_color ? cli_prefill_progress_cb : NULL,
progress.use_color ? &progress : NULL);
if (ds4_session_sync(live, prompt, err, sizeof(err)) != 0) {
ds4_session_set_progress(live, NULL, NULL);
ds4_session_set_display_progress(live, NULL, NULL);
fprintf(stderr, "ds4: prompt processing failed: %s\n", err);
ds4_session_free(live);
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return 1;
}
ds4_session_set_progress(live, NULL, NULL);
ds4_session_set_display_progress(live, NULL, NULL);
fprintf(stderr, "ds4: decode-consistency prompt_tokens=%d check_after=%d\n",
prompt->len, cfg->gen.decode_consistency_tokens);
for (int i = 0; i < cfg->gen.decode_consistency_tokens; i++) {
int token = ds4_session_argmax(live);
fprintf(stderr, "ds4: decode-consistency selected[%d]=", i);
print_diag_token(stderr, engine, token);
fputc('\n', stderr);
ds4_tokens_push(&prefix, token);
if (ds4_session_eval(live, token, err, sizeof(err)) != 0) {
fprintf(stderr, "ds4: decode failed during consistency check: %s\n", err);
ds4_session_free(live);
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return 1;
}
if (ds4_token_is_stop(engine, token)) break;
}
ds4_token_score live_top[10];
int live_top_n = ds4_session_top_logprobs(live, live_top, 10);
if (ds4_session_copy_logits(live, live_logits, vocab) != vocab) {
fprintf(stderr, "ds4: failed to copy live logits\n");
ds4_session_free(live);
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return 1;
}
ds4_session_free(live);
live = NULL;
ds4_session *fresh = NULL;
if (ds4_session_create(&fresh, engine, cfg->gen.ctx_size) != 0) {
fprintf(stderr, "ds4: failed to create fresh diagnostic session\n");
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return 1;
}
progress.input_tokens = prefix.len;
ds4_session_set_progress(fresh, cli_prefill_progress_cb, &progress);
ds4_session_set_display_progress(fresh,
progress.use_color ? cli_prefill_progress_cb : NULL,
progress.use_color ? &progress : NULL);
if (ds4_session_sync(fresh, &prefix, err, sizeof(err)) != 0) {
ds4_session_set_progress(fresh, NULL, NULL);
ds4_session_set_display_progress(fresh, NULL, NULL);
fprintf(stderr, "ds4: fresh prompt processing failed: %s\n", err);
ds4_session_free(fresh);
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return 1;
}
ds4_session_set_progress(fresh, NULL, NULL);
ds4_session_set_display_progress(fresh, NULL, NULL);
ds4_token_score fresh_top[10];
int fresh_top_n = ds4_session_top_logprobs(fresh, fresh_top, 10);
if (ds4_session_copy_logits(fresh, fresh_logits, vocab) != vocab) {
fprintf(stderr, "ds4: failed to copy fresh logits\n");
ds4_session_free(fresh);
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return 1;
}
ds4_session_free(fresh);
double ss = 0.0;
float max_abs = 0.0f;
int max_i = 0;
for (int i = 0; i < vocab; i++) {
float d = fabsf(live_logits[i] - fresh_logits[i]);
if (d > max_abs) {
max_abs = d;
max_i = i;
}
ss += (double)d * (double)d;
}
fprintf(stderr,
"ds4: decode-consistency compared prefix_tokens=%d vocab=%d "
"max_abs=%.9g at token=%d live=%.9g fresh=%.9g rms=%.9g\n",
prefix.len, vocab, max_abs, max_i,
live_logits[max_i], fresh_logits[max_i],
sqrt(ss / (double)vocab));
print_diag_top(stderr, engine, "ds4: live_top", live_top, live_top_n);
print_diag_top(stderr, engine, "ds4: fresh_top", fresh_top, fresh_top_n);
ds4_tokens_free(&prefix);
free(live_logits);
free(fresh_logits);
return live_top_n > 0 && fresh_top_n > 0 && live_top[0].id == fresh_top[0].id ? 0 : 1;
}
static int run_perplexity_file(ds4_engine *engine, const cli_config *cfg) {
char *text = read_prompt_file(cfg->gen.perplexity_file_path, true);
ds4_tokens tokens = {0};
@@ -900,6 +1207,11 @@ static int run_generation(ds4_engine *engine, const cli_config *cfg) {
ds4_tokens_free(&prompt);
return rc;
}
if (cfg->gen.decode_consistency_tokens > 0) {
rc = run_decode_consistency(engine, cfg, &prompt);
ds4_tokens_free(&prompt);
return rc;
}
const bool diagnostic = cfg->gen.dump_tokens ||
cfg->gen.head_test ||
@@ -919,31 +1231,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->gen.temperature > 0.0f ||
ds4_engine_mtp_draft_tokens(engine) > 1) {
rc = run_sampled_generation(engine, cfg, &prompt);
} else {
token_printer printer = {
.engine = engine,
.fp = stdout,
.format_thinking = ds4_think_mode_enabled(cli_effective_think_mode(&cfg->gen)),
.in_think = ds4_think_mode_enabled(cli_effective_think_mode(&cfg->gen)),
.use_color = isatty(fileno(stdout)) != 0,
.last_output_newline = true,
};
cli_prefill_progress progress = {
.base_tokens = 0,
.input_tokens = prompt.len,
.use_color = ds4_log_is_tty(stderr),
};
rc = ds4_engine_generate_argmax(engine, &prompt, cfg->gen.n_predict,
cfg->gen.ctx_size,
print_generated_token,
generation_done,
&printer,
cli_prefill_progress_cb,
&progress);
if (cfg->engine.distributed.role == DS4_DISTRIBUTED_COORDINATOR ||
cfg->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);
@@ -989,7 +1309,8 @@ typedef struct {
ds4_session *session;
ds4_tokens transcript;
int ctx_size;
int max_prefix_tokens;
int think_prefix_pos;
int think_prefix_tokens;
} repl_chat;
static void tokens_insert(ds4_tokens *dst, int pos, const ds4_tokens *src) {
@@ -1019,23 +1340,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 (!ds4_engine_is_laguna(engine) && mode == DS4_THINK_MAX) {
ds4_chat_append_max_effort_prefix(engine, prefix);
}
}
/* Insert/replace the model-family thinking prefix inside the existing
* transcript. It lives immediately after the BOS sequence and before any
* user/system text, matching the GGUF chat templates. */
static void repl_chat_apply_think_prefix(ds4_engine *engine,
repl_chat *chat,
ds4_think_mode mode) {
ds4_tokens prefix = {0};
repl_chat_build_think_prefix(engine, mode, &prefix);
bool same = chat->think_prefix_tokens == prefix.len;
if (same && prefix.len > 0) {
same = !memcmp(chat->transcript.v + chat->think_prefix_pos,
prefix.v,
(size_t)prefix.len * sizeof(prefix.v[0]));
}
if (!same) {
tokens_remove(&chat->transcript, chat->think_prefix_pos,
chat->think_prefix_tokens);
tokens_insert(&chat->transcript, chat->think_prefix_pos, &prefix);
chat->think_prefix_tokens = prefix.len;
if (chat->session) ds4_session_invalidate(chat->session);
}
ds4_tokens_free(&prefix);
}
static int repl_chat_create_session(ds4_engine *engine, repl_chat *chat, int ctx_size) {
@@ -1053,8 +1400,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);
}
@@ -1085,9 +1432,13 @@ static int run_chat_turn(ds4_engine *engine, cli_config *cfg, repl_chat *chat, c
return 1;
}
ds4_session_set_speculative_enabled(
chat->session,
cfg->gen.temperature <= 0.0f &&
getenv("DS4_MTP_SPEC_DISABLE") == NULL);
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);
@@ -1139,15 +1490,29 @@ static int run_chat_turn(ds4_engine *engine, cli_config *cfg, repl_chat *chat, c
uint64_t rng = cfg->gen.seed ? cfg->gen.seed :
((uint64_t)time(NULL) ^ ((uint64_t)getpid() << 32) ^ (uint64_t)clock());
int generated = 0;
const bool speculative_argmax = cfg->gen.temperature <= 0.0f &&
((ds4_engine_mtp_draft_tokens(engine) > 1 &&
getenv("DS4_MTP_SPEC_DISABLE") == NULL) ||
cli_splitkv_spec_requested());
const bool greedy_argmax = cfg->gen.temperature <= 0.0f &&
cli_greedy_argmax_requested(speculative_argmax);
bool have_greedy_next = false;
int greedy_next = -1;
const double t_decode0 = cli_now_sec();
while (generated < max_tokens && !cli_interrupt_requested()) {
int token = ds4_session_sample(chat->session,
int token;
if (greedy_argmax && have_greedy_next) {
token = greedy_next;
have_greedy_next = false;
} else {
token = ds4_session_sample(chat->session,
cfg->gen.temperature,
0,
cfg->gen.top_k,
cfg->gen.top_p,
cfg->gen.min_p,
&rng);
if (token == ds4_token_eos(engine)) break;
}
if (ds4_token_is_stop_for_think_mode(engine, token, think_mode)) break;
int toks[17];
int ntok = 0;
@@ -1168,20 +1533,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;
}
@@ -1204,7 +1578,7 @@ static int run_chat_turn(ds4_engine *engine, cli_config *cfg, repl_chat *chat, c
chat->transcript.len = rollback_len;
ds4_session_invalidate(chat->session);
} else {
ds4_tokens_push(&chat->transcript, ds4_token_eos(engine));
ds4_chat_append_assistant_end(engine, &chat->transcript);
}
const double prefill_s = t_prefill1 - t_prefill0;
@@ -1260,18 +1634,19 @@ static int run_repl(ds4_engine *engine, cli_config *cfg) {
print_repl_help();
} else if (!strcmp(cmd, "/think")) {
cfg->gen.think_mode = DS4_THINK_HIGH;
repl_chat_apply_max_prefix(engine, &chat, false);
repl_chat_apply_think_prefix(engine, &chat, DS4_THINK_HIGH);
puts("Thinking mode: high.");
} else if (!strcmp(cmd, "/think-max")) {
cfg->gen.think_mode = DS4_THINK_MAX;
bool active = ds4_think_mode_for_context(cfg->gen.think_mode,
chat.ctx_size) == DS4_THINK_MAX;
repl_chat_apply_max_prefix(engine, &chat, active);
repl_chat_apply_think_prefix(engine, &chat,
active ? DS4_THINK_MAX : DS4_THINK_HIGH);
cli_warn_think_max_downgraded(&cfg->gen, "/think-max");
printf("Thinking mode: %s.\n", active ? "max" : "high (ctx below 393216)");
} else if (!strcmp(cmd, "/nothink")) {
cfg->gen.think_mode = DS4_THINK_NONE;
repl_chat_apply_max_prefix(engine, &chat, false);
repl_chat_apply_think_prefix(engine, &chat, DS4_THINK_NONE);
puts("Thinking mode: none.");
} else if (!strncmp(cmd, "/power", 6) && (cmd[6] == '\0' || isspace((unsigned char)cmd[6]))) {
char *arg = trim_inplace(cmd + 6);
@@ -1296,15 +1671,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);
ds4_engine_prefill_chunk(engine),
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")) {
@@ -1395,11 +1771,12 @@ static cli_config parse_options(int argc, char **argv) {
.model_path = "ds4flash.gguf",
.backend = default_backend(),
.mtp_draft_tokens = 1,
.dflash_draft_tokens = 0,
.mtp_margin = 3.0f,
},
.gen = {
.prompt = NULL,
.system = "You are a helpful assistant",
.system = NULL,
.n_predict = 50000,
.ctx_size = 32768,
.temperature = DS4_DEFAULT_TEMPERATURE,
@@ -1439,6 +1816,20 @@ static cli_config parse_options(int argc, char **argv) {
}
if (dist_parse == DS4_DIST_CLI_MATCHED) continue;
char tp_parse_err[256] = {0};
ds4_tp_cli_parse_result tp_parse = ds4_tp_parse_cli_arg(arg,
&i,
argc,
argv,
&c.engine.tp,
tp_parse_err,
sizeof(tp_parse_err));
if (tp_parse == DS4_TP_CLI_ERROR) {
fprintf(stderr, "ds4: %s\n", tp_parse_err[0] ? tp_parse_err : "invalid tensor-parallel option");
exit(2);
}
if (tp_parse == DS4_TP_CLI_MATCHED) continue;
if (!strcmp(arg, "-p") || !strcmp(arg, "--prompt")) {
if (c.gen.prompt) {
fprintf(stderr, "ds4: specify only one prompt source\n");
@@ -1454,24 +1845,58 @@ 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);
c.gen.system_set = true;
} 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")) {
c.engine.mtp_path = need_arg(&i, argc, argv, arg);
} else if (!strcmp(arg, "--dflash")) {
c.engine.dflash_path = need_arg(&i, argc, argv, arg);
} else if (!strcmp(arg, "--dflash-draft")) {
c.engine.dflash_draft_tokens =
parse_int(need_arg(&i, argc, argv, arg), arg);
} else if (!strcmp(arg, "--dflash-p-min")) {
c.engine.dflash_p_min =
parse_float_range(
need_arg(&i, argc, argv, arg), arg, 0.0f, 1.0f);
c.engine.dflash_p_min_set = true;
} else if (!strcmp(arg, "--mtp-draft")) {
c.engine.mtp_draft_tokens = parse_int(need_arg(&i, argc, argv, arg), arg);
} else if (!strcmp(arg, "--mtp-margin")) {
c.engine.mtp_margin = parse_float_range(need_arg(&i, argc, argv, arg), arg, 0.0f, 1000.0f);
} else if (!strcmp(arg, "--glm-mtp")) {
c.engine.glm_mtp = true;
} else if (!strcmp(arg, "--glm-mtp-timing")) {
c.engine.glm_mtp = true;
c.engine.glm_mtp_timing = true;
} else if (!strcmp(arg, "--dspark")) {
c.engine.dspark = true;
} else if (!strcmp(arg, "--dspark-confidence")) {
c.engine.dspark = true;
c.engine.dspark_confidence_threshold =
parse_float_range(need_arg(&i, argc, argv, arg), arg, 0.0f, 1.0f);
c.engine.dspark_confidence_threshold_set = true;
} else if (!strcmp(arg, "--dspark-strict")) {
c.engine.dspark = true;
c.engine.dspark_strict = true;
} else if (!strcmp(arg, "-n") || !strcmp(arg, "--tokens")) {
c.gen.n_predict = parse_int(need_arg(&i, argc, argv, arg), arg);
} else if (!strcmp(arg, "-c") || !strcmp(arg, "--ctx")) {
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-k")) {
c.gen.top_k = parse_nonnegative_int(need_arg(&i, argc, argv, arg), arg);
c.gen.top_k_set = true;
} else if (!strcmp(arg, "--top-p")) {
c.gen.top_p = parse_float_range(need_arg(&i, argc, argv, arg), arg, 0.0f, 1.0f);
c.gen.top_p_set = true;
} else if (!strcmp(arg, "--min-p")) {
c.gen.min_p = parse_float_range(need_arg(&i, argc, argv, arg), arg, 0.0f, 1.0f);
c.gen.min_p_set = true;
} else if (!strcmp(arg, "--seed")) {
c.gen.seed = parse_u64(need_arg(&i, argc, argv, arg), arg);
} else if (!strcmp(arg, "--quality")) {
@@ -1491,6 +1916,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) {
@@ -1543,6 +1972,12 @@ static cli_config parse_options(int argc, char **argv) {
} else if (!strcmp(arg, "--cuda")) {
c.engine.backend = DS4_BACKEND_CUDA;
#endif
} else if (!strcmp(arg, "--gpu-vram")) {
c.gpu_vram_arg = need_arg(&i, argc, argv, arg);
} else if (!strcmp(arg, "--gpu-devices")) {
c.gpu_devices_arg = need_arg(&i, argc, argv, arg);
} else if (!strcmp(arg, "--cuda-tensor-parallel")) {
c.engine.cuda_tensor_parallel = true;
} else if (!strcmp(arg, "--dump-tokens")) {
c.gen.dump_tokens = true;
} else if (!strcmp(arg, "--dump-logits")) {
@@ -1551,6 +1986,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")) {
@@ -1625,11 +2062,21 @@ static cli_config parse_options(int argc, char **argv) {
fprintf(stderr, "ds4: --perplexity-file does not use -p/--prompt-file\n");
exit(2);
}
char tp_err[256];
if (!ds4_tp_adopt_distributed_options(&c.engine.tp, c.dist,
tp_err, sizeof(tp_err))) {
fprintf(stderr, "ds4: %s\n", tp_err);
exit(2);
}
char dist_err[256];
if (ds4_dist_prepare_engine_options(c.dist, &c.engine, dist_err, sizeof(dist_err)) != 0) {
fprintf(stderr, "ds4: %s\n", dist_err);
exit(2);
}
if (!ds4_tp_validate_engine_options(&c.engine, tp_err, sizeof(tp_err))) {
fprintf(stderr, "ds4: %s\n", tp_err);
exit(2);
}
return c;
}
@@ -1650,12 +2097,89 @@ 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;
cfg.engine.placement_ctx_hint = cfg.gen.ctx_size;
ds4_engine *engine = NULL;
if (ds4_engine_open(&engine, &cfg.engine) != 0) {
if (cfg.gpu_vram_arg || cfg.gpu_devices_arg) {
ds4_gpu_config gpu_cfg = {0};
bool skip_cuda = false;
char errbuf[256];
if (parse_gpu_vram_arg(cfg.gpu_vram_arg, cfg.gpu_devices_arg,
&gpu_cfg, &skip_cuda,
errbuf, sizeof(errbuf)) != 0) {
fprintf(stderr, "ds4: %s\n", errbuf);
ds4_dist_options_free(cfg.dist);
free(cfg.prompt_owned);
return 2;
}
cfg.engine.backend = skip_cuda ? DS4_BACKEND_CPU : DS4_BACKEND_CUDA;
if (skip_cuda) {
if (ds4_engine_open(&engine, &cfg.engine) != 0) {
ds4_dist_options_free(cfg.dist);
free(cfg.prompt_owned);
return 1;
}
} else {
const bool was_auto =
(cfg.gpu_vram_arg && !strcmp(cfg.gpu_vram_arg, "auto")) ||
(!cfg.gpu_vram_arg && cfg.gpu_devices_arg);
char layout[256];
if (format_gpu_layout_line(&gpu_cfg, was_auto,
layout, sizeof(layout)) > 0) {
fprintf(stdout, "%s\n", layout);
fflush(stdout);
}
if (ds4_engine_create_with_gpu_config(&engine, &cfg.engine,
&gpu_cfg) != 0) {
ds4_dist_options_free(cfg.dist);
free(cfg.prompt_owned);
return 1;
}
}
} else 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.gen.system_set) {
cfg.gen.system = ds4_engine_default_system_prompt(engine);
}
if (cfg.engine.tp.role == DS4_TP_WORKER) {
int rc = ds4_tp_worker_run(engine, &cfg.engine.tp);
ds4_engine_close(engine);
ds4_dist_options_free(cfg.dist);
free(cfg.prompt_owned);
return rc;
}
ds4_tp *tp_leader = NULL;
if (cfg.engine.tp.role == DS4_TP_LEADER) {
char tp_err[256] = "";
ds4_tp_identity tp_id = {
.gguf_bytes = ds4_engine_model_bytes(engine),
.model_id = (uint32_t)ds4_engine_model_id(engine),
.n_layer = (uint32_t)ds4_engine_layer_count(engine),
.n_embd = (uint32_t)ds4_engine_embd_dim(engine),
.n_vocab = (uint32_t)ds4_engine_vocab_size(engine),
.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);
ds4_tp_free(tp_leader);
ds4_engine_close(engine);
ds4_dist_options_free(cfg.dist);
free(cfg.prompt_owned);
return 1;
}
}
if (cfg.dist && cfg.dist->role == DS4_DISTRIBUTED_WORKER) {
ds4_dist_generation_options dist_gen = {
.prompt = cfg.gen.prompt,
@@ -1680,7 +2204,8 @@ int main(int argc, char **argv) {
if (!cfg.inspect) {
log_context_memory(cfg.engine.backend,
cfg.gen.ctx_size,
cfg.engine.prefill_chunk);
ds4_engine_prefill_chunk(engine),
cfg.engine.ssd_streaming);
cli_warn_think_max_downgraded(&cfg.gen, "--think-max");
}
int rc = 0;
@@ -1700,7 +2225,9 @@ int main(int argc, char **argv) {
} else {
rc = run_generation(engine, &cfg);
}
if (tp_leader) ds4_tp_send_stop(tp_leader);
ds4_engine_close(engine);
ds4_tp_free(tp_leader);
ds4_dist_options_free(cfg.dist);
free(cfg.prompt_owned);
return rc;
+19149 -2789
View File
File diff suppressed because it is too large. Load diff
+23 -11
View File
@@ -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
View File
@@ -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) {
+1873 -3
View File
File diff suppressed because it is too large. Load diff
+230
View File
@@ -0,0 +1,230 @@
/* ds4_gpu_args.c — shared parser for --gpu-vram and --gpu-devices.
*
* See ds4_gpu_args.h. The actual CUDA probing for the "auto" branch
* lives in ds4_cuda.cu so this .c file never includes <cuda_runtime.h>
* directly (mirroring the rest of the tree's separation between .c and
* .cu compilation units).
*/
#include "ds4_gpu_args.h"
#include <ctype.h>
#include <errno.h>
#include <stdarg.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#define DS4_GPU_ARGS_DEFAULT_SAFETY_MARGIN ((size_t)512 * 1024 * 1024)
/* Helper: copy at most outlen-1 bytes of msg into out and NUL-terminate. */
static void errbuf_set(char *out, size_t outlen, const char *msg) {
if (!out || outlen == 0) return;
size_t n = strlen(msg);
if (n >= outlen) n = outlen - 1;
memcpy(out, msg, n);
out[n] = '\0';
}
static void errbuf_setf(char *out, size_t outlen, const char *fmt, ...) {
if (!out || outlen == 0) return;
va_list ap;
va_start(ap, fmt);
vsnprintf(out, outlen, fmt, ap);
va_end(ap);
}
/* Maximum per-entry values. Generous, but small enough to keep all
* downstream arithmetic safely inside size_t/int. */
#define DS4_GPU_ARGS_MAX_VRAM_GB 16384 /* 16 TiB per device */
#define DS4_GPU_ARGS_MAX_DEVICE_INDEX 1023 /* CUDA visible-device ID cap */
/* Parse a comma-separated list of nonnegative integers into out_count
* entries of out[]. Returns 0 on success, nonzero on parse error.
* Empty entries (",,") or trailing/leading commas error. Per-entry
* upper bound enforced via max_value (use LONG_MAX for "no cap"). */
static int parse_csv_int_list(const char *s,
long *out,
int *out_count,
int max_count,
long max_value,
char *errbuf,
size_t errbuflen,
const char *what) {
*out_count = 0;
if (!s || !*s) {
errbuf_setf(errbuf, errbuflen, "%s: empty value", what);
return 1;
}
const char *p = s;
while (*p) {
/* Disallow leading whitespace inside the list. */
while (*p == ' ' || *p == '\t') p++;
char *end = NULL;
errno = 0;
long v = strtol(p, &end, 10);
if (end == p) {
errbuf_setf(errbuf, errbuflen, "%s: not a number: '%s'", what, p);
return 1;
}
if (errno != 0 || v < 0 || v > max_value) {
errbuf_setf(errbuf, errbuflen,
"%s: out-of-range value: '%s' (allowed 0..%ld)",
what, p, max_value);
return 1;
}
if (*out_count >= max_count) {
errbuf_setf(errbuf, errbuflen, "%s: too many entries (max %d)", what, max_count);
return 1;
}
out[*out_count] = v;
(*out_count)++;
while (*end == ' ' || *end == '\t') end++;
if (*end == ',') {
end++;
if (*end == '\0' || *end == ',') {
errbuf_setf(errbuf, errbuflen, "%s: empty entry near '%s'", what, end);
return 1;
}
} else if (*end != '\0') {
errbuf_setf(errbuf, errbuflen, "%s: junk after value: '%s'", what, end);
return 1;
}
p = end;
}
if (*out_count == 0) {
errbuf_setf(errbuf, errbuflen, "%s: no values parsed", what);
return 1;
}
return 0;
}
int parse_gpu_vram_arg(const char *vram_arg,
const char *devices_arg,
ds4_gpu_config *out,
bool *out_skip_cuda,
char *errbuf,
size_t errbuflen) {
if (!out || !out_skip_cuda) {
errbuf_set(errbuf, errbuflen, "internal: NULL out parameter");
return 1;
}
*out_skip_cuda = false;
/* Caller zeroed *out — just confirm. */
out->n_gpus = 0;
out->safety_margin_bytes = DS4_GPU_ARGS_DEFAULT_SAFETY_MARGIN;
/* Parse devices_arg first so we have the filter when handling "auto". */
long dev_list[DS4_MAX_GPUS] = {0};
int dev_count = 0;
if (devices_arg) {
if (parse_csv_int_list(devices_arg, dev_list, &dev_count,
DS4_MAX_GPUS,
DS4_GPU_ARGS_MAX_DEVICE_INDEX,
errbuf, errbuflen,
"--gpu-devices") != 0) {
return 1;
}
}
/* Case 1: vram_arg unset. */
if (vram_arg == NULL) {
if (devices_arg == NULL) {
errbuf_set(errbuf, errbuflen, "internal: parse_gpu_vram_arg called with both args NULL");
return 1;
}
/* --gpu-devices given without --gpu-vram: implicit auto with filter. */
vram_arg = "auto";
}
/* Case 2: --gpu-vram 0 — short-circuit. */
if (!strcmp(vram_arg, "0")) {
if (devices_arg) {
errbuf_set(errbuf, errbuflen,
"--gpu-vram 0 conflicts with --gpu-devices");
return 1;
}
*out_skip_cuda = true;
return 0;
}
/* Case 3: --gpu-vram auto. */
if (!strcmp(vram_arg, "auto")) {
#if !defined(DS4_NO_GPU) && !defined(__APPLE__)
int filter_buf[DS4_MAX_GPUS];
const int *filter_ptr = NULL;
if (dev_count > 0) {
for (int i = 0; i < dev_count; i++) filter_buf[i] = (int)dev_list[i];
filter_ptr = filter_buf;
}
if (ds4_gpu_args_probe_auto_cuda(filter_ptr, dev_count, out,
out->safety_margin_bytes,
errbuf, errbuflen) != 0) {
return 1;
}
return 0;
#else
(void)dev_list; (void)dev_count;
errbuf_set(errbuf, errbuflen,
"--gpu-vram auto requires a CUDA build");
return 1;
#endif
}
/* Case 4: explicit GB budgets — "N[,N...]" */
long vram_list[DS4_MAX_GPUS] = {0};
int vram_count = 0;
if (parse_csv_int_list(vram_arg, vram_list, &vram_count,
DS4_MAX_GPUS,
DS4_GPU_ARGS_MAX_VRAM_GB,
errbuf, errbuflen,
"--gpu-vram") != 0) {
return 1;
}
if (dev_count > 0 && dev_count != vram_count) {
errbuf_setf(errbuf, errbuflen,
"--gpu-devices count (%d) does not match --gpu-vram count (%d)",
dev_count, vram_count);
return 1;
}
out->n_gpus = vram_count;
for (int i = 0; i < vram_count; i++) {
int dev_id = (dev_count > 0) ? (int)dev_list[i] : i;
out->device_indices[i] = dev_id;
out->vram_bytes[i] = (size_t)vram_list[i] * (size_t)1024 * 1024 * 1024;
}
return 0;
}
int format_gpu_layout_line(const ds4_gpu_config *cfg,
bool was_auto,
char *out,
size_t outlen) {
if (!out || outlen == 0 || !cfg) return -1;
/* "ds4: GPU config: N devices [d0,d1,...] requested, budgets G0,G1,... GB; auto=<bool>" */
int written = snprintf(out, outlen,
"ds4: GPU config: %d device%s [",
cfg->n_gpus, cfg->n_gpus == 1 ? "" : "s");
if (written < 0 || (size_t)written >= outlen) return written;
for (int i = 0; i < cfg->n_gpus; i++) {
int n = snprintf(out + written, outlen - (size_t)written,
"%s%d", i ? "," : "", cfg->device_indices[i]);
if (n < 0 || (size_t)(written + n) >= outlen) return written + n;
written += n;
}
int n = snprintf(out + written, outlen - (size_t)written,
"] requested, budgets ");
if (n < 0 || (size_t)(written + n) >= outlen) return written + n;
written += n;
for (int i = 0; i < cfg->n_gpus; i++) {
double gb = (double)cfg->vram_bytes[i] / (1024.0 * 1024.0 * 1024.0);
int m = snprintf(out + written, outlen - (size_t)written,
"%s%.0f", i ? "," : "", gb);
if (m < 0 || (size_t)(written + m) >= outlen) return written + m;
written += m;
}
int t = snprintf(out + written, outlen - (size_t)written,
" GB; auto=%s", was_auto ? "true" : "false");
if (t < 0) return t;
written += t;
return written;
}
+84
View File
@@ -0,0 +1,84 @@
/* ds4_gpu_args.h — shared CLI argument parser for --gpu-vram and
* --gpu-devices flags. Used by ds4_cli, ds4_server, ds4_agent, and
* ds4_bench. See mgpu-cli-wiring task (wave 2) for design.
*
* The parser does not link CUDA headers directly. On CUDA builds the
* "auto" path is gated and dispatches to ds4_gpu_args_probe_auto_cuda,
* which is defined in ds4_cuda.cu. On Metal/CPU builds, "auto" returns
* an error.
*/
#ifndef DS4_GPU_ARGS_H
#define DS4_GPU_ARGS_H
#include <stdbool.h>
#include <stddef.h>
#include "ds4_gpu_mgpu.h"
#ifdef __cplusplus
extern "C" {
#endif
/* Single canonical parser. Populates *out and *out_skip_cuda on
* success; returns 0 on success, nonzero on parse error (errbuf
* populated). Caller zeroes out and out_skip_cuda before calling.
*
* vram_arg may be:
* NULL -- flag unset
* "0" -- skip CUDA entirely; sets *out_skip_cuda = true
* "auto" -- query free VRAM via CUDA (CUDA build only)
* "N[,N...]" -- explicit GB budgets per device
*
* devices_arg may be:
* NULL -- 0..n-1 of visible devices (or 0..vram_count-1)
* "N[,N...]" -- explicit CUDA device indices
*
* Constraints:
* - If both args are given and counts differ, error.
* - Negative or non-integer values error.
* - Counts > DS4_MAX_GPUS error.
*/
int parse_gpu_vram_arg(
const char *vram_arg,
const char *devices_arg,
ds4_gpu_config *out,
bool *out_skip_cuda,
char *errbuf,
size_t errbuflen);
/* Format the layout echo line written to stdout right before
* ds4_engine_create_with_gpu_config is called.
*
* Example output:
* ds4: GPU config: 2 devices [0,1] requested, budgets 40,12 GB; auto=true
*
* Returns the number of bytes written (excluding the trailing NUL), or
* -1 if out is NULL/outlen is 0.
*/
int format_gpu_layout_line(const ds4_gpu_config *cfg,
bool was_auto,
char *out,
size_t outlen);
/* CUDA-only auto probe. Defined in ds4_cuda.cu under the same gate.
* Returns 0 on success, nonzero on failure (errbuf populated).
* device_filter may be NULL (use all visible devices) or a list of
* length filter_len (use only those device indices).
*
* Declared here for the parser's use; not called by application code
* directly. The Mac/CPU builds skip this branch via the
* #if !defined(DS4_NO_GPU) && !defined(__APPLE__) gate in the parser. */
#if !defined(DS4_NO_GPU) && !defined(__APPLE__)
int ds4_gpu_args_probe_auto_cuda(const int *device_filter,
int filter_len,
ds4_gpu_config *out,
size_t safety_margin_bytes,
char *errbuf,
size_t errbuflen);
#endif
#ifdef __cplusplus
}
#endif
#endif /* DS4_GPU_ARGS_H */
+237
View File
@@ -0,0 +1,237 @@
/* ds4_gpu_mgpu.h — multi-GPU plumbing types and APIs (v0).
*
* This header carries the new multi-GPU additions for the multi-GPU plumbing PP work
* (device-aware CUDA). It is included from ds4_cuda.cu and from
* downstream tasks that need access to g_gpu[], g_n_gpus, g_gpu_peer_ok[],
* the ds4_gpu_config struct, and the new tensor APIs.
*
* Why not in ds4_gpu.h? The legacy ds4_gpu.h is included from C-only
* callers (ds4.c, ds4_cli.c, etc.) and from the Metal build, but is NOT
* included from ds4_cuda.cu historically. That asymmetry hid pre-existing
* signature mismatches between the legacy header and ds4_cuda.cu. We keep
* the legacy header opaque and put the new shared types here, so this
* file is the single source of truth for both ds4_cuda.cu and downstream
* multi-GPU tasks without disturbing the legacy contract.
*
* The struct definitions reference CUDA-specific handle types via void *
* placeholders so the header is safe to include from C builds, Metal
* builds, and the CUDA build (where ds4_cuda.cu casts the void * back
* to cudaStream_t / cublasHandle_t / cudaEvent_t internally).
*/
#ifndef DS4_GPU_MGPU_H
#define DS4_GPU_MGPU_H
#include <stddef.h>
#include <stdint.h>
#ifdef __cplusplus
extern "C" {
#endif
#define DS4_MAX_GPUS 16
/* Complete definition of the previously-opaque ds4_gpu_tensor, plus a
* typedef so the new API prototypes below can use the bare name
* `ds4_gpu_tensor *` in both C and C++ without forcing callers to
* include ds4_gpu.h first. Callers that include this header can
* stack-allocate or struct-embed tensors and pass to
* ds4_gpu_tensor_alloc_on. */
struct ds4_gpu_tensor {
void *ptr;
uint64_t bytes;
int owner;
int device_id; /* -1 means legacy/untagged → treat as device 0 */
};
#ifndef DS4_GPU_TENSOR_DEFINED
#define DS4_GPU_TENSOR_DEFINED
typedef struct ds4_gpu_tensor ds4_gpu_tensor;
#endif
#ifndef DS4_GPU_ATTENTION_DECODE_ROW_DEFINED
#define DS4_GPU_ATTENTION_DECODE_ROW_DEFINED
#define DS4_GPU_ATTENTION_DECODE_BATCH_MAX 32u
typedef struct {
uint64_t raw_kv;
uint64_t comp_kv;
uint64_t topk;
uint32_t pos;
uint32_t n_raw;
uint32_t raw_cap;
uint32_t raw_start;
uint32_t n_comp;
uint32_t top_k;
uint32_t window;
uint32_t ratio;
uint32_t indexed;
} ds4_gpu_attention_decode_row;
#endif
/* Tagged so headers (notably ds4.h) can forward-declare `struct
* ds4_gpu_config` without dragging in this entire header. */
typedef struct ds4_gpu_config {
int device_indices[DS4_MAX_GPUS]; /* CUDA device IDs to use */
/* Explicit per-device budget in bytes. The engine does NOT auto-fill
* missing budgets - a value of 0 means "zero bytes of budget for
* this slot" and (combined with reserves) will push placement to
* CPU spill. Auto-detection (e.g. mapping --gpu-vram auto to
* cudaMemGetInfo) is the caller's job; see CLI flag wiring for the
* canonical CLI path. The engine emits a clear stderr and refuses
* if n_gpus > 0 and every vram_bytes[] is 0 (almost certainly a
* caller bug from zero-initializing the struct). */
size_t vram_bytes[DS4_MAX_GPUS];
int n_gpus;
size_t safety_margin_bytes; /* per-device reserve */
} ds4_gpu_config;
typedef struct {
int device_id;
void *stream; /* cudaStream_t under CUDA */
void *cublas; /* cublasHandle_t under CUDA */
int cublas_ready;
void *scratch;
size_t scratch_bytes;
size_t budget_bytes;
size_t used_bytes;
void *boundary_event; /* cudaEvent_t under CUDA */
} ds4_gpu_ctx;
extern ds4_gpu_ctx g_gpu[DS4_MAX_GPUS];
extern int g_n_gpus;
extern int g_gpu_peer_ok[DS4_MAX_GPUS][DS4_MAX_GPUS];
/* Primary multi-device init. The existing ds4_gpu_init (declared in
* ds4_gpu.h) is a thin shim that builds a single-device config for
* device 0 and calls this. */
int ds4_gpu_init_multi(const ds4_gpu_config *cfg);
/* Caller-supplied struct alloc on a specific device. Returns 0 on
* success, nonzero on error. Pair with ds4_gpu_tensor_free_in_place. */
int ds4_gpu_tensor_alloc_on(ds4_gpu_tensor *t, int device_id, uint64_t bytes);
void ds4_gpu_tensor_free_in_place(ds4_gpu_tensor *t);
/* Heap-allocated tensor on a specific logical tier; mirrors the legacy
* ds4_gpu_tensor_alloc ABI (returns ds4_gpu_tensor *) but with a tier
* parameter. Returns NULL on failure. Used by the multi-tier graph
* allocations in ds4.c. Single-tier callers can continue using the
* legacy ds4_gpu_tensor_alloc(bytes) which is equivalent to
* ds4_gpu_tensor_alloc_ptr_on(0, bytes). */
ds4_gpu_tensor *ds4_gpu_tensor_alloc_ptr_on(int tier, uint64_t bytes);
int ds4_gpu_tensor_copy_async(ds4_gpu_tensor *dst, const ds4_gpu_tensor *src, uint64_t bytes);
void ds4_gpu_enable_q8_dequant_gemm(void);
/* Heap-allocated MANAGED-memory tensor on a specific logical tier. Used
* for the KV-cache pool. Stamps tier on the tensor so subsequent
* ds4_gpu_tensor_free runs under the correct device. Note managed memory
* pages on first-touch and is not strictly device-bound; the stamped
* tier records the home tier for accounting + free. */
ds4_gpu_tensor *ds4_gpu_tensor_alloc_managed_on(int tier, uint64_t bytes);
/* Cross-device tensor copy. Same-device → cudaMemcpyAsync; peer-capable
* cross-device → cudaMemcpyPeerAsync with event sync; non-peer → pinned
* host bounce (per src→dst pair). Honors DS4_FORCE_HOST_BOUNCE=1. */
int ds4_gpu_tensor_copy_xdev(ds4_gpu_tensor *dst,
const ds4_gpu_tensor *src,
uint64_t bytes);
int ds4_gpu_tensor_copy_xdev_default(ds4_gpu_tensor *dst,
const ds4_gpu_tensor *src,
uint64_t bytes);
/* Grouped default-stream handoff whose copies execute on the destination
* device. The source default stream records readiness; the destination waits,
* performs all three copies, then naturally orders its following kernels.
* This allows source-device work submitted afterward to overlap the handoff. */
int ds4_gpu_tensor_copy_xdev3_default_dst(
ds4_gpu_tensor *dst0,
const ds4_gpu_tensor *src0,
uint64_t bytes0,
ds4_gpu_tensor *dst1,
const ds4_gpu_tensor *src1,
uint64_t bytes1,
ds4_gpu_tensor *dst2,
const ds4_gpu_tensor *src2,
uint64_t bytes2);
/* Cross-device copy of three buffers between the same source/destination
* tiers, recording a single readiness event for the destination stream. This
* is meant for tiny grouped activation handoffs where three independent
* ds4_gpu_tensor_copy_xdev calls would spend more time in event plumbing than
* in the copies themselves. Falls back internally when a grouped peer copy is
* not applicable. */
int ds4_gpu_tensor_copy_xdev3(ds4_gpu_tensor *dst0,
const ds4_gpu_tensor *src0,
uint64_t bytes0,
ds4_gpu_tensor *dst1,
const ds4_gpu_tensor *src1,
uint64_t bytes1,
ds4_gpu_tensor *dst2,
const ds4_gpu_tensor *src2,
uint64_t bytes2);
/* Cross-device copy that also orders against prior work on the destination
* stream before overwriting dst. This is needed for pipelined prefill slots
* where the next producer reuses a destination buffer that the destination
* stage just consumed. */
int ds4_gpu_tensor_copy_xdev_ordered(ds4_gpu_tensor *dst,
const ds4_gpu_tensor *src,
uint64_t bytes);
/* Order a peer read without copying: records readiness on src's stream and
* makes dst_tier's stream wait for it. */
int ds4_gpu_tensor_wait_xdev(const ds4_gpu_tensor *src, int dst_tier);
int ds4_gpu_tensor_wait_xdev_default(const ds4_gpu_tensor *src, int dst_tier);
/* Cross-device float add used by CUDA tensor-parallel reductions.
* See ds4_gpu.h for the full contract. */
int ds4_gpu_add_xdev_tensor(ds4_gpu_tensor *out,
const ds4_gpu_tensor *local,
const ds4_gpu_tensor *remote,
ds4_gpu_tensor *remote_tmp,
uint32_t n);
/* Returns the device_id recorded on the tensor; -1 if untagged. */
int ds4_gpu_tensor_device(const ds4_gpu_tensor *t);
/* Set the current CUDA device by LOGICAL tier index (0..g_n_gpus-1).
*
* This is the canonical shim for per-layer device routing in the
* multi-tier execution path. The caller passes a logical tier index;
* the shim internally indexes g_gpu[tier].device_id and calls
* cudaSetDevice. Returns 0 on success, nonzero on error or if the
* tier index is out of range.
*
* Wave-2 multi-GPU placement scaffolding adds this shim but does not
* exercise the multi-tier execution path; multi-GPU execution
* (follow-up) is its first caller. */
int ds4_gpu_set_current_device(int logical_tier);
int ds4_gpu_set_current_device_fenced(int logical_tier);
/* Register the mmap'd host model pointer for selective-cache lookups
* WITHOUT triggering any device-side copy. This bypasses the
* DS4_CUDA_COPY_MODEL environment-variable branch that
* ds4_gpu_set_model_map normally honors, which is essential for
* multi-tier startup: we want only per-device selective tensor caches,
* never the whole-model copy. Returns 1 on success, 0 on error. */
int ds4_gpu_register_model_map_no_copy(const void *model_map, uint64_t model_size);
/* Strict per-device selective-cache lookup (no fallback).
*
* Returns 1 only if a covering entry exists whose device_id matches the
* caller-supplied expected_device (a PHYSICAL CUDA device id). Otherwise
* returns 0 — no host-pointer fallback, no different-device match. The
* caller is expected to have cudaSetDevice'd to expected_device before
* invoking; the returned pointer is valid to consume from that device's
* kernel.
*
* Used by multi-tier kernel-dispatch resolvers in
* multi-GPU execution (multi-GPU execution). Single-tier callers should
* keep using ds4_gpu_lookup_cache for back-compat behavior. */
int ds4_gpu_lookup_cache_strict(uint64_t source_offset,
uint64_t bytes,
int expected_device,
void **out_device_ptr);
#ifdef __cplusplus
} /* extern "C" */
#endif
#endif /* DS4_GPU_MGPU_H */
+36 -5
View File
@@ -153,6 +153,11 @@ static void print_model_runtime(FILE *fp, const help_colors *c,
#else
opt(fp, c, "--metal | --cuda | --cpu", "Select the backend explicitly.");
opt(fp, c, "--backend NAME", "Backend name: metal, cuda, or cpu.");
opt(fp, c, "--gpu-vram N[,N,...]|auto", "CUDA VRAM budgets per device, in GiB, or auto-detect free VRAM.");
opt(fp, c, "--gpu-devices N[,N,...]", "CUDA device indices used by multi-GPU placement.");
if (tool != DS4_HELP_EVAL) {
opt(fp, c, "--cuda-tensor-parallel", "Enable the paired DeepSeek tensor/expert path on an even multi-GPU CUDA placement.");
}
#endif
if (tool != DS4_HELP_BENCH) {
opt(fp, c, "-c, --ctx N", "Allocated context tokens.");
@@ -164,17 +169,26 @@ 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).");
opt(fp, c, "--prefill-chunk N", "Graph prefill chunk size. Default: CUDA TP 2048; PRO 8192; Laguna 16384; others 4096.");
if (full) {
if (tool != DS4_HELP_BENCH) {
opt(fp, c, "--mtp FILE", "Optional MTP support GGUF used for draft-token probes.");
}
if (tool == DS4_HELP_DS4 || tool == DS4_HELP_AGENT || tool == DS4_HELP_SERVER) {
opt(fp, c, "--dflash FILE", "Laguna DFlash support GGUF for greedy speculative decoding.");
opt(fp, c, "--dflash-draft N", "Maximum adaptive DFlash draft positions, 1..15. CUDA default: 15; other backends: 3");
opt(fp, c, "--dflash-p-min P", "Stop before proposals below probability P, 0..1. Default: 0.4; 0 keeps fixed verifier width");
opt(fp, c, "--mtp-draft N", "Maximum autoregressive MTP draft tokens. Default: 1");
opt(fp, c, "--mtp-margin F", "Verifier confidence margin for fast MTP acceptance. Default: 3");
opt(fp, c, "--glm-mtp", "Enable integrated greedy GLM MTP speculation.");
opt(fp, c, "--glm-mtp-timing", "Enable GLM MTP and print acceptance/timing counters.");
opt(fp, c, "--dspark", "Enable DSpark using the support GGUF passed with --mtp.");
opt(fp, c, "--dspark-confidence F", "Enable DSpark with confidence pruning threshold 0..1. Default: 0.9");
opt(fp, c, "--dspark-strict", "Load DSpark support but keep target-only decode.");
}
opt(fp, c, "--quality", "Prefer exact kernels where faster approximate paths exist.");
opt(fp, c, "--warm-weights", "Touch mapped tensor pages at startup to reduce first-use stalls.");
@@ -189,9 +203,11 @@ static void print_sampling(FILE *fp, const help_colors *c, bool full) {
title(fp, c, "Prompt And Sampling");
opt(fp, c, "-n, --tokens N", "Maximum generated tokens.");
opt(fp, c, "--temp F", "Sampling temperature. 0 is greedy/deterministic.");
opt(fp, c, "--top-k N", "Sample only from the N highest-scoring tokens. 0 disables.");
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 defaults to temperature 1.0, top-p 0.95, and min-p 0. Laguna defaults to temperature 0.7, top-k 20, top-p 0.95, and min-p 0.05. Explicit options always win.");
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 +215,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 +243,17 @@ 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 uses the same coordinator/worker addresses as distributed mode, but always runs one 50/50 worker. Add --tensor-parallel, omit --layers, start the worker, then start the coordinator.");
fputc('\n', fp);
opt(fp, c, "--tensor-parallel", "Switch --role/--listen/--coordinator to two-machine tensor parallelism.");
opt(fp, c, "--transport auto|rdma|tcp", "Tensor gate transport. Default: auto");
opt(fp, c, "--rdma-device NAME", "Select a verbs device when auto-detection is ambiguous.");
opt(fp, c, "--rdma-gid-index N", "Select the local verbs GID index.");
opt(fp, c, "--tensor-parallel-token-prefill", "GLM diagnostic: prefill one token at a time for exact arithmetic.");
opt(fp, c, "--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 +276,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 +307,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.");
@@ -305,8 +335,9 @@ static void print_server_api(FILE *fp, const help_colors *c) {
opt(fp, c, "--port N", "Bind port. Default: 8000");
opt(fp, c, "--cors", "Add Access-Control-Allow-* headers for browser JS clients.");
opt(fp, c, "--trace FILE", "Write prompts, cache decisions, output, and tool calls.");
opt(fp, c, "--batched-session N", "Keep N resident sessions and batch decode-ready requests.");
para(fp, c, "Endpoints: /v1/chat/completions, /v1/responses, /v1/completions, and /v1/messages.");
para(fp, c, "Model endpoint aliases include deepseek-v4-flash and deepseek-v4-pro; both serve the loaded GGUF.");
para(fp, c, "Model aliases are available for DeepSeek V4, GLM-5.2, and Laguna-S-2.1; every alias serves the loaded GGUF.");
fputc('\n', fp);
}
@@ -513,7 +544,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 and Laguna 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);
+149
View File
@@ -0,0 +1,149 @@
/* Monotonic-contiguous multi-tier layer placement packer.
*
* Pure C99: no CUDA, no platform-specific code. Used by both the CUDA and
* Metal/CPU builds. See ds4_layer_pack.h for the API contract and the
* design doc docs/superpowers/specs/2026-05-26-multi-gpu-pp-v0-design.md
* for the rationale. */
#include "ds4_layer_pack.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
int ds4_compute_layer_placement(const size_t *entry_bytes,
int n_entries,
const ds4_layer_pack_config *cfg,
int *device_for_entry) {
if (!entry_bytes || !cfg || !device_for_entry) return 1;
if (n_entries < 0) return 2;
if (cfg->n_gpus < 0 || cfg->n_gpus > DS4_LAYER_PACK_MAX_GPUS) return 3;
/* Work over a local copy of budgets so the caller's struct is untouched. */
size_t budget[DS4_LAYER_PACK_MAX_GPUS];
for (int d = 0; d < cfg->n_gpus; d++) {
budget[d] = cfg->gpu_budget_bytes[d];
}
int d = 0;
for (int e = 0; e < n_entries; e++) {
/* Advance until the current device can hold this entry, or we run
* out of devices. Strict greater-than: exact fits stay on d. */
while (d < cfg->n_gpus && entry_bytes[e] > budget[d]) {
d++;
}
if (d < cfg->n_gpus) {
device_for_entry[e] = d;
budget[d] -= entry_bytes[e];
} else {
device_for_entry[e] = DS4_LAYER_PACK_CPU;
/* d stays at cfg->n_gpus so every subsequent entry also lands
* on the CPU tier. */
}
}
return 0;
}
/* Internal helper: append " + embedding" or " + output head" to a buffer if
* the matching pseudo-layer lives on this tier. */
static void append_pseudo_layer_tags(char *buf, size_t buflen,
const int *device_for_entry,
int tier,
int n_entries,
int n_layers) {
/* entry 0 -> embedding, entry n_layers+1 -> output head */
if (n_entries > 0 && device_for_entry[0] == tier) {
strncat(buf, " + embedding", buflen - strlen(buf) - 1);
}
if (n_entries > 1 && n_layers + 1 < n_entries &&
device_for_entry[n_layers + 1] == tier) {
strncat(buf, " + output head", buflen - strlen(buf) - 1);
}
}
/* Emit "L-R" for a contiguous transformer-layer range assigned to a tier.
* Returns 1 if anything was printed (i.e. tier owned at least one layer). */
static int print_layer_range_for_tier(FILE *out, const int *device_for_entry,
int tier, int n_entries, int n_layers) {
/* Transformer layers live at indices 1..n_layers in device_for_entry.
* Their human-facing numbers are 0..n_layers-1 (i.e. entry index - 1). */
int first = -1;
int last = -1;
for (int i = 1; i <= n_layers && i < n_entries; i++) {
if (device_for_entry[i] == tier) {
if (first < 0) first = i - 1; /* human-facing index */
last = i - 1;
}
}
if (first < 0) {
fprintf(out, "(no transformer layers)");
return 0;
}
if (first == last) {
fprintf(out, "layer %d", first);
} else {
fprintf(out, "layers %d-%d", first, last);
}
return 1;
}
void ds4_layer_pack_print(FILE *out,
const int *device_for_entry,
int n_entries,
int n_layers,
const size_t *entry_bytes,
const size_t *gpu_used_bytes,
const size_t *gpu_budget_bytes,
int n_gpus) {
if (!out || !device_for_entry) return;
(void)entry_bytes; /* not needed for the textual layout */
fprintf(out, "multi-GPU layout:\n");
for (int d = 0; d < n_gpus; d++) {
fprintf(out, " GPU%d: ", d);
int have_layers = print_layer_range_for_tier(out, device_for_entry,
d, n_entries, n_layers);
/* Pseudo-layer tags. */
char tags[64];
tags[0] = '\0';
append_pseudo_layer_tags(tags, sizeof(tags), device_for_entry, d,
n_entries, n_layers);
if (tags[0]) {
if (!have_layers) {
/* Erase the "(no transformer layers)" stub and just write
* the tag with leading "+ " stripped. */
/* The stub was already written; appending " + embedding"
* is still informative. */
}
fprintf(out, "%s", tags);
}
/* Usage line. */
if (gpu_used_bytes && gpu_budget_bytes) {
const double used_gb = (double)gpu_used_bytes[d] / 1073741824.0;
const double budget_gb = (double)gpu_budget_bytes[d] / 1073741824.0;
/* Two-space pad before the parens matches the design doc. */
fprintf(out, " (%.1f / %.1f GB)", used_gb, budget_gb);
}
fputc('\n', out);
}
/* CPU tier: only print if at least one entry is on CPU. */
int cpu_present = 0;
for (int i = 0; i < n_entries; i++) {
if (device_for_entry[i] == DS4_LAYER_PACK_CPU) { cpu_present = 1; break; }
}
if (cpu_present) {
fprintf(out, " CPU : ");
int have_layers = print_layer_range_for_tier(out, device_for_entry,
DS4_LAYER_PACK_CPU,
n_entries, n_layers);
char tags[64];
tags[0] = '\0';
append_pseudo_layer_tags(tags, sizeof(tags), device_for_entry,
DS4_LAYER_PACK_CPU, n_entries, n_layers);
if (tags[0]) {
(void)have_layers;
fprintf(out, "%s", tags);
}
fputc('\n', out);
}
}
+77
View File
@@ -0,0 +1,77 @@
#ifndef DS4_LAYER_PACK_H
#define DS4_LAYER_PACK_H
#include <stddef.h>
#include <stdio.h>
#define DS4_LAYER_PACK_MAX_GPUS 16
#define DS4_LAYER_PACK_CPU (-1)
typedef struct {
/* Per-device budget in BYTES (already net of per-device overhead such as
* cuBLAS workspace, global scratch, logits buffer, MTP, steering, and the
* configured safety margin). The packer treats these as raw capacity. */
size_t gpu_budget_bytes[DS4_LAYER_PACK_MAX_GPUS];
int n_gpus;
} ds4_layer_pack_config;
#ifdef __cplusplus
extern "C" {
#endif
/* Compute a monotonic-contiguous layer placement.
*
* Inputs:
* entry_bytes[] — per-entry byte footprint, in FORWARD ORDER:
* entry 0 = embedding pseudo-layer
* entry 1..n_layers = transformer layers
* entry n_layers+1 = output-head pseudo-layer
* n_entries — total number of entries (typically n_layers + 2)
* cfg — per-device budgets
*
* Output:
* device_for_entry[] — filled with target device for each entry; CPU is
* DS4_LAYER_PACK_CPU. Caller-owned, must be
* large enough for n_entries int slots.
*
* Returns:
* 0 on success.
* Nonzero on configuration errors (null pointer, n_entries < 0,
* n_gpus out of range).
*
* Per the design doc, an entry that exceeds every GPU budget spills to CPU;
* by the monotonicity rule every entry after it also goes to CPU. There is
* no error path for "entry too large" — the CPU tier is always available. */
int ds4_compute_layer_placement(const size_t *entry_bytes,
int n_entries,
const ds4_layer_pack_config *cfg,
int *device_for_entry);
/* Print a human-readable layout summary matching the design doc's format:
*
* multi-GPU layout:
* GPU0: layers 0-21 + embedding (38.4 / 40.0 GB)
* GPU1: layers 22-31 (11.7 / 12.0 GB)
* CPU : layers 32-42 + output head
*
* Entry-naming convention:
* entry 0 -> "embedding"
* entry i in 1..n_layers -> transformer layer (numbered i-1)
* entry n_layers+1 -> "output head"
*
* gpu_used_bytes[] and gpu_budget_bytes[] may be NULL — in that case the
* "(used / budget)" line is omitted for the GPU lines. */
void ds4_layer_pack_print(FILE *out,
const int *device_for_entry,
int n_entries,
int n_layers,
const size_t *entry_bytes,
const size_t *gpu_used_bytes,
const size_t *gpu_budget_bytes,
int n_gpus);
#ifdef __cplusplus
}
#endif
#endif /* DS4_LAYER_PACK_H */
+16786 -1232
View File
File diff suppressed because it is too large. Load diff
+98
View File
@@ -69,6 +69,13 @@ typedef struct {
uint16_t dmin;
} cuda_block_q2_K;
typedef struct {
uint8_t hmask[CUDA_QK_K / 8];
uint8_t qs[CUDA_QK_K / 4];
uint8_t scales[12];
uint16_t d;
} cuda_block_q3_K;
typedef struct {
uint16_t d;
uint16_t dmin;
@@ -87,6 +94,8 @@ typedef struct {
uint16_t qs[CUDA_QK_K / 8];
} cuda_block_iq2_xxs;
static_assert(sizeof(cuda_block_q3_K) == 110, "Q3_K block layout mismatch");
#include "ds4_iq2_tables_cuda.inc"
#include "rocm/ds4_rocm_runtime.cuh"
@@ -97,6 +106,10 @@ typedef struct {
#include "rocm/ds4_rocm_norm_rope.cuh"
#include "rocm/ds4_rocm_laguna.cuh"
#include "rocm/ds4_rocm_dflash.cuh"
#include "rocm/ds4_rocm_fp8_kv.cuh"
#include "rocm/ds4_rocm_attention.cuh"
@@ -126,6 +139,91 @@ typedef struct {
#include "rocm/ds4_rocm_moe_launch.cuh"
#include "rocm/ds4_rocm_glm.cuh"
#include "rocm/ds4_rocm_hc_output_launch.cuh"
#include "rocm/ds4_rocm_current_api_compat.cuh"
/* Tensor-parallel gates are Metal-only; stubs keep shared graph code
* linkable (TP option validation rejects non-Metal backends). */
extern "C" int ds4_gpu_tp_gate_encode(uint32_t layer, uint32_t gate) {
(void)layer; (void)gate;
fprintf(stderr, DS4_GPU_LOG_PREFIX "tensor parallelism is Metal-only\n");
return 0;
}
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");
return 0;
}
extern "C" int ds4_gpu_matmul_q8_0_kslice_tensor(
ds4_gpu_tensor *out, const void *model_map, uint64_t model_size,
uint64_t weight_offset, uint64_t full_in_dim, uint64_t k_off,
uint64_t k_cnt, uint64_t out_dim, const ds4_gpu_tensor *x,
uint64_t x_elem_off) {
(void)out; (void)model_map; (void)model_size; (void)weight_offset;
(void)full_in_dim; (void)k_off; (void)k_cnt; (void)out_dim; (void)x;
(void)x_elem_off;
fprintf(stderr, DS4_GPU_LOG_PREFIX "tensor parallelism is Metal-only\n");
return 0;
}
extern "C" int ds4_gpu_attention_output_q8_tp_tensor(
ds4_gpu_tensor *out, ds4_gpu_tensor *low, const void *model_map,
uint64_t model_size, uint64_t out_a_offset, uint64_t out_b_offset,
uint64_t group_dim, uint64_t rank, uint32_t n_groups_total,
uint32_t group0, uint32_t group_cnt, uint64_t out_dim,
const ds4_gpu_tensor *heads) {
(void)out; (void)low; (void)model_map; (void)model_size;
(void)out_a_offset; (void)out_b_offset; (void)group_dim; (void)rank;
(void)n_groups_total; (void)group0; (void)group_cnt; (void)out_dim;
(void)heads;
fprintf(stderr, DS4_GPU_LOG_PREFIX "tensor parallelism is Metal-only\n");
return 0;
}
extern "C" int ds4_gpu_hc_expand_add_tensor(
ds4_gpu_tensor *out_hc, const ds4_gpu_tensor *block_out,
const ds4_gpu_tensor *block_add, const ds4_gpu_tensor *residual_hc,
const ds4_gpu_tensor *post, const ds4_gpu_tensor *comb,
uint32_t n_embd, uint32_t n_hc) {
(void)out_hc; (void)block_out; (void)block_add; (void)residual_hc;
(void)post; (void)comb; (void)n_embd; (void)n_hc;
fprintf(stderr, DS4_GPU_LOG_PREFIX "tensor parallelism is Metal-only\n");
return 0;
}
+423
View File
@@ -0,0 +1,423 @@
#include <hip/hip_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "ds4_gpu_mgpu.h"
#include "ds4_gpu.h"
#include "ds4_gpu_args.h"
ds4_gpu_ctx g_gpu[DS4_MAX_GPUS] = {};
int g_n_gpus = 1;
int g_gpu_peer_ok[DS4_MAX_GPUS][DS4_MAX_GPUS] = {{1}};
static int rocm_tier_valid(int tier) {
return tier == 0 && g_n_gpus == 1;
}
extern "C" int ds4_gpu_init_multi(const ds4_gpu_config *cfg) {
if (!cfg || cfg->n_gpus != 1) {
fprintf(stderr, "ds4: ROCm supports one GPU per process\n");
return 0;
}
g_gpu[0].device_id = cfg->device_indices[0];
if (hipSetDevice(g_gpu[0].device_id) != hipSuccess) return 0;
return ds4_gpu_init();
}
extern "C" int ds4_gpu_set_current_device(int tier) {
if (!rocm_tier_valid(tier)) return 1;
return hipSetDevice(g_gpu[0].device_id) == hipSuccess ? 0 : 1;
}
extern "C" int ds4_gpu_set_current_device_fenced(int tier) {
return ds4_gpu_set_current_device(tier);
}
extern "C" int ds4_gpu_tensor_alloc_on(ds4_gpu_tensor *t, int tier,
uint64_t bytes) {
if (!t) return 1;
if (!rocm_tier_valid(tier)) return 2;
if (bytes == 0) bytes = 1;
if (hipMalloc(&t->ptr, (size_t)bytes) != hipSuccess) return 3;
t->bytes = bytes;
t->owner = 1;
t->device_id = 0;
return 0;
}
extern "C" ds4_gpu_tensor *ds4_gpu_tensor_alloc_ptr_on(int tier,
uint64_t bytes) {
if (!rocm_tier_valid(tier)) return NULL;
return ds4_gpu_tensor_alloc(bytes);
}
extern "C" ds4_gpu_tensor *ds4_gpu_tensor_alloc_managed_on(int tier,
uint64_t bytes) {
if (!rocm_tier_valid(tier)) return NULL;
return ds4_gpu_tensor_alloc_managed(bytes);
}
extern "C" void ds4_gpu_tensor_free_in_place(ds4_gpu_tensor *t) {
if (!t) return;
if (t->owner && t->ptr) (void)hipFree(t->ptr);
memset(t, 0, sizeof(*t));
}
extern "C" int ds4_gpu_tensor_device(const ds4_gpu_tensor *t) {
return t ? 0 : -1;
}
extern "C" int ds4_gpu_tensor_copy_async(ds4_gpu_tensor *dst,
const ds4_gpu_tensor *src,
uint64_t bytes) {
if (!dst || !src || bytes > dst->bytes || bytes > src->bytes) return 0;
if (bytes == 0) return 1;
return hipMemcpyAsync(dst->ptr, src->ptr, (size_t)bytes,
hipMemcpyDeviceToDevice, 0) == hipSuccess;
}
extern "C" int ds4_gpu_tensor_copy_xdev(ds4_gpu_tensor *dst,
const ds4_gpu_tensor *src,
uint64_t bytes) {
return ds4_gpu_tensor_copy(dst, 0, src, 0, bytes);
}
extern "C" int ds4_gpu_tensor_copy_xdev_default(ds4_gpu_tensor *dst,
const ds4_gpu_tensor *src,
uint64_t bytes) {
return ds4_gpu_tensor_copy_xdev(dst, src, bytes);
}
extern "C" int ds4_gpu_tensor_copy_xdev_ordered(ds4_gpu_tensor *dst,
const ds4_gpu_tensor *src,
uint64_t bytes) {
return ds4_gpu_tensor_copy_xdev(dst, src, bytes);
}
extern "C" int ds4_gpu_tensor_copy_xdev3(
ds4_gpu_tensor *dst0, const ds4_gpu_tensor *src0, uint64_t bytes0,
ds4_gpu_tensor *dst1, const ds4_gpu_tensor *src1, uint64_t bytes1,
ds4_gpu_tensor *dst2, const ds4_gpu_tensor *src2, uint64_t bytes2) {
return (bytes0 == 0 || ds4_gpu_tensor_copy_xdev(dst0, src0, bytes0)) &&
(bytes1 == 0 || ds4_gpu_tensor_copy_xdev(dst1, src1, bytes1)) &&
(bytes2 == 0 || ds4_gpu_tensor_copy_xdev(dst2, src2, bytes2));
}
extern "C" int ds4_gpu_tensor_copy_xdev3_default_dst(
ds4_gpu_tensor *dst0, const ds4_gpu_tensor *src0, uint64_t bytes0,
ds4_gpu_tensor *dst1, const ds4_gpu_tensor *src1, uint64_t bytes1,
ds4_gpu_tensor *dst2, const ds4_gpu_tensor *src2, uint64_t bytes2) {
return ds4_gpu_tensor_copy_xdev3(dst0, src0, bytes0, dst1, src1, bytes1,
dst2, src2, bytes2);
}
extern "C" int ds4_gpu_tensor_wait_xdev(const ds4_gpu_tensor *src,
int dst_tier) {
return src && rocm_tier_valid(dst_tier);
}
extern "C" int ds4_gpu_tensor_wait_xdev_default(const ds4_gpu_tensor *src,
int dst_tier) {
return ds4_gpu_tensor_wait_xdev(src, dst_tier);
}
extern "C" uint64_t ds4_gpu_tier_free_vram(int tier) {
size_t free_bytes = 0;
size_t total_bytes = 0;
if (!rocm_tier_valid(tier) ||
hipMemGetInfo(&free_bytes, &total_bytes) != hipSuccess) {
return 0;
}
return (uint64_t)free_bytes;
}
extern "C" int ds4_gpu_args_probe_auto_cuda(
const int *device_filter, int filter_len, ds4_gpu_config *out,
size_t safety_margin_bytes, char *errbuf, size_t errbuflen) {
if (!out) {
if (errbuf && errbuflen) snprintf(errbuf, errbuflen, "internal: NULL out");
return 1;
}
int visible = 0;
hipError_t rc = hipGetDeviceCount(&visible);
if (rc != hipSuccess || visible <= 0) {
if (errbuf && errbuflen) {
snprintf(errbuf, errbuflen, "hipGetDeviceCount failed: %s",
rc == hipSuccess ? "no devices" : hipGetErrorString(rc));
}
return 1;
}
if (filter_len > 1 || (!device_filter && visible > 1)) {
if (errbuf && errbuflen) {
snprintf(errbuf, errbuflen,
"ROCm supports one GPU per process; select one device");
}
return 1;
}
const int device = device_filter && filter_len == 1 ? device_filter[0] : 0;
if (device < 0 || device >= visible || hipSetDevice(device) != hipSuccess) {
if (errbuf && errbuflen) {
snprintf(errbuf, errbuflen, "invalid ROCm device %d", device);
}
return 1;
}
size_t free_bytes = 0;
size_t total_bytes = 0;
rc = hipMemGetInfo(&free_bytes, &total_bytes);
if (rc != hipSuccess) {
if (errbuf && errbuflen) {
snprintf(errbuf, errbuflen, "hipMemGetInfo failed: %s",
hipGetErrorString(rc));
}
return 1;
}
const size_t reserve_floor = (size_t)2ull * 1024ull * 1024ull * 1024ull;
const size_t reserve_pct = free_bytes / 20u;
const size_t reserve = reserve_floor > reserve_pct ? reserve_floor : reserve_pct;
memset(out, 0, sizeof(*out));
out->device_indices[0] = device;
out->vram_bytes[0] = free_bytes > reserve ? free_bytes - reserve : 0;
out->n_gpus = 1;
out->safety_margin_bytes = safety_margin_bytes;
return 0;
}
extern "C" void ds4_gpu_enable_q8_dequant_gemm(void) {
}
static int g_rocm_q8_cache_suppressed = 0;
extern "C" int ds4_gpu_q8_cache_suppressed(void) {
return g_rocm_q8_cache_suppressed;
}
extern "C" void ds4_gpu_set_q8_cache_suppressed(int suppressed) {
g_rocm_q8_cache_suppressed = suppressed != 0;
}
extern "C" int ds4_gpu_set_decode_fast_attention(int enabled) {
(void)enabled;
return 0;
}
extern "C" int ds4_gpu_set_decode_score_vec4(int enabled) {
(void)enabled;
return 0;
}
extern "C" int ds4_gpu_matmul_f16_router_rows_exact_tensor(
ds4_gpu_tensor *out, const void *model_map, uint64_t model_size,
uint64_t weight_offset, const ds4_gpu_tensor *x, uint32_t n_rows) {
return ds4_gpu_matmul_f16_tensor(out, model_map, model_size, weight_offset,
4096u, 256u, x, n_rows);
}
extern "C" int ds4_gpu_dsv4_qkv_rms_norm_rows_kv_rope_tensor(
ds4_gpu_tensor *q_out, const ds4_gpu_tensor *q,
const void *model_map, uint64_t model_size,
uint64_t q_weight_offset, uint32_t q_n,
ds4_gpu_tensor *kv_out, const ds4_gpu_tensor *kv,
uint64_t kv_weight_offset, uint32_t kv_n, uint32_t rows,
uint32_t kv_n_head, uint32_t kv_head_dim, uint32_t n_rot,
uint32_t pos0, uint32_t n_ctx_orig, bool inverse,
float freq_base, float freq_scale, float ext_factor,
float attn_factor, float beta_fast, float beta_slow, float eps) {
return ds4_gpu_dsv4_qkv_rms_norm_rows_tensor(
q_out, q, model_map, model_size, q_weight_offset, q_n,
kv_out, kv, kv_weight_offset, kv_n, rows, eps) != 0 &&
ds4_gpu_rope_tail_tensor(
kv_out, rows, kv_n_head, kv_head_dim, n_rot, pos0,
n_ctx_orig, inverse, freq_base, freq_scale, ext_factor,
attn_factor, beta_fast, beta_slow) != 0;
}
extern "C" int ds4_gpu_embed_token_quant_tensor(
ds4_gpu_tensor *out, const void *model_map, uint64_t model_size,
uint64_t weight_offset, uint32_t weight_type, uint32_t n_vocab,
uint32_t token, uint32_t n_embd) {
if (weight_type != 8u) return 0;
return ds4_gpu_embed_token_q8_0_tensor(out, model_map, model_size,
weight_offset, n_vocab, token,
n_embd);
}
extern "C" int ds4_gpu_embed_tokens_quant_tensor(
ds4_gpu_tensor *out, const ds4_gpu_tensor *tokens,
const void *model_map, uint64_t model_size, uint64_t weight_offset,
uint32_t weight_type, uint32_t n_vocab, uint32_t n_tokens,
uint32_t n_embd) {
if (weight_type != 8u) return 0;
return ds4_gpu_embed_tokens_q8_0_tensor(out, tokens, model_map, model_size,
weight_offset, n_vocab, n_tokens,
n_embd);
}
extern "C" int ds4_gpu_matmul_quant_tensor(
ds4_gpu_tensor *out, const void *model_map, uint64_t model_size,
uint64_t weight_offset, uint32_t weight_type, uint64_t in_dim,
uint64_t out_dim, const ds4_gpu_tensor *x, uint64_t n_tok) {
if (weight_type == 8u) {
return ds4_gpu_matmul_q8_0_tensor(out, model_map, model_size,
weight_offset, in_dim, out_dim, x,
n_tok);
}
if (weight_type == 1u) {
return ds4_gpu_matmul_f16_tensor(out, model_map, model_size,
weight_offset, in_dim, out_dim, x,
n_tok);
}
return 0;
}
extern "C" int ds4_gpu_matmul_quant_decode_mpp_model_view_tensor(
ds4_gpu_tensor *out, const void *model_map, uint64_t model_size,
uint64_t weight_offset, uint32_t weight_type, uint64_t in_dim,
uint64_t out_dim, const ds4_gpu_tensor *x, uint64_t n_tok) {
if (weight_type == 8u) {
return ds4_gpu_matmul_q8_0_decode_mpp_model_view_tensor(
out, model_map, model_size, weight_offset, in_dim, out_dim,
x, n_tok);
}
return ds4_gpu_matmul_quant_tensor(out, model_map, model_size,
weight_offset, weight_type, in_dim,
out_dim, x, n_tok);
}
extern "C" int ds4_gpu_glm_k_b_project_typed_tensor(
ds4_gpu_tensor *out, const ds4_gpu_tensor *kv_norm,
const void *model_map, uint64_t model_size, uint64_t weight_offset,
uint32_t weight_type, uint32_t n_tokens, uint32_t kv_lora_dim,
uint32_t qk_nope, uint32_t n_head) {
if (weight_type != 8u) return 0;
return ds4_gpu_glm_k_b_project_tensor(out, kv_norm, model_map, model_size,
weight_offset, n_tokens,
kv_lora_dim, qk_nope, n_head);
}
extern "C" int ds4_gpu_glm_qk_lowrank_typed_tensor(
ds4_gpu_tensor *qk_low, const ds4_gpu_tensor *q,
const void *model_map, uint64_t model_size, uint64_t weight_offset,
uint32_t weight_type, uint32_t n_head, uint32_t kv_lora_dim,
uint32_t qk_nope, uint32_t qk_dim) {
if (weight_type != 8u) return 0;
return ds4_gpu_glm_qk_lowrank_q8_0_tensor(
qk_low, q, model_map, model_size, weight_offset, n_head,
kv_lora_dim, qk_nope, qk_dim);
}
extern "C" int ds4_gpu_glm_qk_lowrank_typed_batch_tensor(
ds4_gpu_tensor *qk_low, const ds4_gpu_tensor *q,
const void *model_map, uint64_t model_size, uint64_t weight_offset,
uint32_t weight_type, uint32_t n_tokens, uint32_t n_head,
uint32_t kv_lora_dim, uint32_t qk_nope, uint32_t qk_dim) {
if (weight_type != 8u) return 0;
return ds4_gpu_glm_qk_lowrank_q8_0_batch_tensor(
qk_low, q, model_map, model_size, weight_offset, n_tokens, n_head,
kv_lora_dim, qk_nope, qk_dim);
}
extern "C" int ds4_gpu_glm_value_project_typed_batch_heads_tensor(
ds4_gpu_tensor *heads, const ds4_gpu_tensor *lora,
const void *model_map, uint64_t model_size, uint64_t weight_offset,
uint32_t weight_type, uint32_t n_tokens, uint32_t n_head,
uint32_t kv_lora_dim, uint32_t value_dim) {
if (weight_type != 8u) return 0;
return ds4_gpu_glm_value_project_q8_0_batch_heads_tensor(
heads, lora, model_map, model_size, weight_offset, n_tokens,
n_head, kv_lora_dim, value_dim);
}
extern "C" int ds4_gpu_glm_attention_indexed_decode_typed_tensor(
ds4_gpu_tensor *heads, const ds4_gpu_tensor *q,
const ds4_gpu_tensor *qk_low, const ds4_gpu_tensor *kv_lora_cache,
const ds4_gpu_tensor *k_rope_cache, const void *model_map,
uint64_t model_size, uint64_t value_weight_offset,
uint32_t value_weight_type, const ds4_gpu_tensor *selected,
uint32_t n_selected, uint32_t cache_cap, bool cache_f16,
uint32_t n_head, uint32_t kv_lora_dim, uint32_t qk_nope,
uint32_t qk_rope, uint32_t value_dim, uint32_t n_ctx_orig,
float freq_base, float freq_scale, float ext_factor,
float attn_factor, float beta_fast, float beta_slow) {
if (value_weight_type != 8u) return 0;
return ds4_gpu_glm_attention_indexed_decode_tensor(
heads, q, qk_low, kv_lora_cache, k_rope_cache, model_map,
model_size, value_weight_offset, selected, n_selected, cache_cap,
cache_f16, n_head, kv_lora_dim, qk_nope, qk_rope, value_dim,
n_ctx_orig, freq_base, freq_scale, ext_factor, attn_factor,
beta_fast, beta_slow);
}
extern "C" int ds4_gpu_glm_attention_indexed_batch_typed_tensor(
ds4_gpu_tensor *heads, const ds4_gpu_tensor *q,
const ds4_gpu_tensor *qk_low, const ds4_gpu_tensor *kv_lora_cache,
const ds4_gpu_tensor *k_rope_cache, const void *model_map,
uint64_t model_size, uint64_t value_weight_offset,
uint32_t value_weight_type, const ds4_gpu_tensor *selected,
uint32_t n_tokens, uint32_t n_selected, uint32_t cache_cap,
bool cache_f16, uint32_t n_head, uint32_t kv_lora_dim,
uint32_t qk_nope, uint32_t qk_rope, uint32_t value_dim,
uint32_t n_ctx_orig, float freq_base, float freq_scale,
float ext_factor, float attn_factor, float beta_fast,
float beta_slow) {
if (value_weight_type != 8u) return 0;
return ds4_gpu_glm_attention_indexed_batch_tensor(
heads, q, qk_low, kv_lora_cache, k_rope_cache, model_map,
model_size, value_weight_offset, selected, n_tokens, n_selected,
cache_cap, cache_f16, n_head, kv_lora_dim, qk_nope, qk_rope,
value_dim, n_ctx_orig, freq_base, freq_scale, ext_factor,
attn_factor, beta_fast, beta_slow);
}
extern "C" int ds4_gpu_matmul_q6_K_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_laguna_routed_shared_moe_one_tensor(
ds4_gpu_tensor *routed_out, ds4_gpu_tensor *routed_mid,
ds4_gpu_tensor *shared_out, ds4_gpu_tensor *shared_mid,
const void *model_map, uint64_t model_size,
const ds4_gpu_laguna_moe_desc *routed,
const ds4_gpu_laguna_moe_desc *shared,
uint32_t expert_in_dim, uint32_t expert_mid_dim, uint32_t out_dim,
const ds4_gpu_tensor *selected, const ds4_gpu_tensor *weights,
uint32_t n_total_expert, uint32_t n_expert,
const ds4_gpu_tensor *shared_selected,
const ds4_gpu_tensor *shared_weight, const ds4_gpu_tensor *x) {
(void)routed_out; (void)routed_mid; (void)shared_out; (void)shared_mid;
(void)model_map; (void)model_size; (void)routed; (void)shared;
(void)expert_in_dim; (void)expert_mid_dim; (void)out_dim;
(void)selected; (void)weights; (void)n_total_expert; (void)n_expert;
(void)shared_selected; (void)shared_weight; (void)x;
return 0;
}
extern "C" int ds4_gpu_laguna_qkvg_f16_tensor(
ds4_gpu_tensor *q, ds4_gpu_tensor *k, ds4_gpu_tensor *v,
ds4_gpu_tensor *gate, const void *model_map, uint64_t model_size,
uint64_t q_weight_offset, uint64_t k_weight_offset,
uint64_t v_weight_offset, uint64_t gate_weight_offset,
uint32_t in_dim, uint32_t q_dim, uint32_t kv_dim,
uint32_t gate_dim, const ds4_gpu_tensor *x) {
(void)q; (void)k; (void)v; (void)gate; (void)model_map; (void)model_size;
(void)q_weight_offset; (void)k_weight_offset; (void)v_weight_offset;
(void)gate_weight_offset; (void)in_dim; (void)q_dim; (void)kv_dim;
(void)gate_dim; (void)x;
return 0;
}
extern "C" int ds4_gpu_laguna_attn_output_residual_f16_tensor(
ds4_gpu_tensor *out, const void *model_map, uint64_t model_size,
uint64_t weight_offset, uint32_t in_dim, uint32_t out_dim,
const ds4_gpu_tensor *x, const ds4_gpu_tensor *residual) {
(void)out; (void)model_map; (void)model_size; (void)weight_offset;
(void)in_dim; (void)out_dim; (void)x; (void)residual;
return 0;
}
+41
View File
@@ -0,0 +1,41 @@
/* Shared graph code references CUDA-only fused and multi-device hooks even
* when ROCm validation guarantees that those paths cannot be selected. Keep
* the unavailable surface in one place. These C-linkage stubs intentionally
* accept any argument list because they never inspect their arguments. */
#include <stdint.h>
#define ROCM_UNAVAILABLE_INT(name) extern "C" int name(...) { return 0; }
ROCM_UNAVAILABLE_INT(ds4_gpu_add_xdev_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_attention_decode_rows_rope_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_attention_noncausal_raw_batch_heads_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_attention_output_low_q4_K_slice_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_attention_output_low_q8_rows_exact_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_attention_output_q4_K_batch_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_attention_prefill_raw_heads_range_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_attention_prefill_static_mixed_heads_range_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_device_cache_support_tensors)
ROCM_UNAVAILABLE_INT(ds4_gpu_device_cache_tensors)
ROCM_UNAVAILABLE_INT(ds4_gpu_dspark_markov_argmax_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_indexer_top1_value_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_kv_fp8_store_raw_decode_rows_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_matmul_q8_0_kslice_hc_expand_add_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_matmul_q8_0_kslice_rows_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_matmul_q8_0_top1_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_matmul_quant_kslice_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_moe_handoff_pack_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_register_model_map_no_copy)
ROCM_UNAVAILABLE_INT(ds4_gpu_register_support_map)
ROCM_UNAVAILABLE_INT(ds4_gpu_rope_tail_decode_rows_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_routed_moe_batch_owned_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_routed_moe_one_owned_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_routed_moe_owned_packed_combine_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_routed_moe_owned_slots_combine_rows_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_routed_moe_owned_slots_combine_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_shared_down_hc_expand_add_q8_0_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_shared_down_hc_expand_owned_q8_0_tensor)
ROCM_UNAVAILABLE_INT(ds4_gpu_shared_mid_swiglu_q8_0_decode_exact_tensor)
extern "C" uint64_t ds4_gpu_tp_big_gate_kick(...) { return 0; }
ROCM_UNAVAILABLE_INT(ds4_gpu_tp_big_gate_wait)
+2425 -376
View File
File diff suppressed because it is too large. Load diff
+31 -2
View File
@@ -77,6 +77,34 @@ uint32_t ds4_ssd_cache_experts_for_byte_budget(uint64_t bytes,
return (uint32_t)experts;
}
static uint64_t ds4_ssd_auto_cache_percent(void) {
const char *env = getenv("DS4_SSD_AUTO_CACHE_PCT");
if (env && env[0]) {
errno = 0;
char *end = NULL;
unsigned long v = strtoul(env, &end, 10);
if (end != env && *end == '\0' && errno == 0 &&
v >= 50 && v <= 95) {
return (uint64_t)v;
}
fprintf(stderr,
"ds4: invalid DS4_SSD_AUTO_CACHE_PCT=%s (want 50..95); "
"using default\n",
env);
}
/*
* Decode on the ROCm streaming path is SSD-bandwidth bound: every routed
* expert miss is a random NVMe read, so a larger resident expert cache is
* the biggest decode-throughput lever. BUT the expert cache lives in the
* same physical RAM as the OS on a unified-memory APU, and transient
* spikes (pinned read/upload staging, the prefill headroom, cache growth)
* ride on top of the steady-state plan. Pushing the split too high runs
* the machine out of RAM and trips the Linux OOM killer. 80% was measured
* safe here; opt into more only deliberately via DS4_SSD_AUTO_CACHE_PCT.
*/
return 80;
}
bool ds4_ssd_auto_cache_plan(uint64_t recommended_bytes,
uint64_t non_routed_bytes,
uint64_t per_expert_bytes,
@@ -86,9 +114,10 @@ bool ds4_ssd_auto_cache_plan(uint64_t recommended_bytes,
memset(out, 0, sizeof(*out));
if (recommended_bytes == 0 || per_expert_bytes == 0) return false;
const uint64_t pct = ds4_ssd_auto_cache_percent();
out->model_target_bytes =
recommended_bytes > UINT64_MAX / 4ull ?
UINT64_MAX : (recommended_bytes * 4ull) / 5ull;
recommended_bytes > UINT64_MAX / pct ?
UINT64_MAX : (recommended_bytes * pct) / 100ull;
if (out->model_target_bytes > non_routed_bytes) {
out->cache_bytes = out->model_target_bytes - non_routed_bytes;
}
File diff suppressed because it is too large. Load diff
+2219
View File
File diff suppressed because it is too large. Load diff
+234
View File
@@ -0,0 +1,234 @@
#ifndef DS4_TP_H
#define DS4_TP_H
#include <stdbool.h>
#include <stddef.h>
#include <stdint.h>
#include <stdio.h>
#include "ds4.h"
/* Tensor-parallel transport and lockstep protocol.
*
* 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
* registered with the GPU gate machinery. Nothing here touches tensors.
*/
typedef struct ds4_tp ds4_tp;
enum {
DS4_TP_GATE_ATTN = 0,
DS4_TP_GATE_FFN = 1,
DS4_TP_GATES_PER_LAYER = 2,
/* Max rows in a verify-block batch gate (speculative blocks are <=5). */
DS4_TP_BATCH_MAX_ROWS = 8,
};
/* Engine identity exchanged in the hello so a mismatched pair aborts before
* any inference runs. */
typedef struct {
uint64_t gguf_bytes;
uint32_t model_id;
uint32_t n_layer;
uint32_t n_embd;
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);
typedef enum {
DS4_TP_CLI_ERROR = -1,
DS4_TP_CLI_NOT_MATCHED = 0,
DS4_TP_CLI_MATCHED = 1,
} ds4_tp_cli_parse_result;
/* CLI parsing, same contract as ds4_dist_parse_cli_arg(): returns 1 when the
* argument was consumed, 0 when not matched, -1 on error (err filled). */
int ds4_tp_parse_cli_arg(
const char *arg,
int *index,
int argc,
char **argv,
ds4_tp_options *opt,
char *err,
size_t errlen);
int ds4_tp_adopt_distributed_options(
ds4_tp_options *tp,
ds4_distributed_options *dist,
char *err,
size_t errlen);
void ds4_tp_usage(FILE *fp);
/* Validates option combinations that TP cannot run with (SSD streaming,
* distributed mode, MTP drafting, CPU backend). */
int ds4_tp_validate_engine_options(
const ds4_engine_options *opt,
char *err,
size_t errlen);
/* Connection bring-up. The leader listens and accepts one worker; the
* worker dials with retry. Both then exchange and validate identities.
* Blocking; call after the engine is loaded (identity needs the shape). */
int ds4_tp_create(
ds4_tp **out,
const ds4_tp_options *opt,
const ds4_tp_identity *id,
char *err,
size_t errlen);
void ds4_tp_free(ds4_tp *tp);
int ds4_tp_rank(const ds4_tp *tp);
bool ds4_tp_is_rdma(const ds4_tp *tp);
uint32_t ds4_tp_peer_ctx(const ds4_tp *tp);
bool ds4_tp_failed(const ds4_tp *tp);
void ds4_tp_mark_failed(ds4_tp *tp);
/* Gate slab. The engine allocates one shared GPU-visible block and hands
* its base VA here; ds4_tp registers it with the NIC (RDMA) and exchanges
* remote keys. Layout, all offsets from base, S = n_layer * 2 slots:
*
* out vectors S * vec_bytes written by local GPU kernels
* in vectors S * vec_bytes RDMA/TCP-written with the peer partials
* in seq flags S * 8 written strictly after each in vector
* token slot 16 {seq u64, token i32, pad} leader->worker
* (gpu flags, then batch out/in: n_layer * BATCH_MAX_ROWS * vec_bytes
* each, row partials for the speculative verify-block gates)
*
* vec_bytes = n_embd * 4 (f32 partials, never quantized on the wire). */
uint64_t ds4_tp_slab_bytes(uint32_t n_layer, uint32_t n_embd);
uint64_t ds4_tp_slab_out_offset(const ds4_tp *tp, uint32_t layer, uint32_t gate);
uint64_t ds4_tp_slab_in_offset(const ds4_tp *tp, uint32_t layer, uint32_t gate);
uint64_t ds4_tp_slab_batch_out_offset(const ds4_tp *tp, uint32_t layer);
uint64_t ds4_tp_slab_batch_in_offset(const ds4_tp *tp, uint32_t layer);
uint64_t ds4_tp_slab_gpu_flags_offset(const ds4_tp *tp);
int ds4_tp_attach_slab(ds4_tp *tp, void *base, char *err, size_t errlen);
/* Exchange one gate: send out[layer][gate] to the peer's in[layer][gate]
* and wait until the peer's partial for `seq` has fully landed locally.
* 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 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. */
typedef struct {
uint64_t session_id;
int32_t token;
uint32_t reserved;
} ds4_tp_batch_item;
int ds4_tp_send_session_create(ds4_tp *tp, uint64_t session_id, int ctx_size);
int ds4_tp_send_session_destroy(ds4_tp *tp, uint64_t session_id);
int ds4_tp_send_sync(ds4_tp *tp, uint64_t session_id,
const int *tokens, uint32_t n_tokens);
int ds4_tp_send_eval(ds4_tp *tp, uint64_t session_id,
uint64_t seq, int token);
int ds4_tp_send_rewind(ds4_tp *tp, uint64_t session_id, int pos);
int ds4_tp_send_invalidate(ds4_tp *tp, uint64_t session_id);
int ds4_tp_send_eval_batch(ds4_tp *tp, const ds4_tp_batch_item *items,
uint32_t count);
int ds4_tp_send_mixed_batch(ds4_tp *tp, uint64_t prefill_session_id,
const int *prompt, uint32_t prompt_count,
const ds4_tp_batch_item *items,
uint32_t count);
int ds4_tp_send_command_ack(ds4_tp *tp, uint64_t session_id, int status);
int ds4_tp_wait_command_ack(ds4_tp *tp, uint64_t session_id,
const char *operation, char *err, size_t errlen);
int ds4_tp_send_stop(ds4_tp *tp);
/* Worker: blocks for the next mirrored command. Frame types below; for
* DS4_TP_FRAME_SYNC the token array is returned in *tokens / *n_tokens
* (malloc'd, caller frees), for DS4_TP_FRAME_EVAL seq/token are filled. */
typedef enum {
DS4_TP_FRAME_ERROR = -1,
DS4_TP_FRAME_SYNC = 1,
DS4_TP_FRAME_EVAL = 2,
DS4_TP_FRAME_REWIND = 3,
DS4_TP_FRAME_INVALIDATE = 4,
DS4_TP_FRAME_STOP = 5,
DS4_TP_FRAME_HASH = 6,
DS4_TP_FRAME_RDMA_INFO = 7,
DS4_TP_FRAME_SYNC_ACK = 8,
DS4_TP_FRAME_RDMA_READY = 9,
DS4_TP_FRAME_LOGITS = 10,
DS4_TP_FRAME_VERIFY = 11,
DS4_TP_FRAME_VERIFY_COMMIT = 12,
DS4_TP_FRAME_SESSION_CREATE = 13,
DS4_TP_FRAME_SESSION_DESTROY = 14,
DS4_TP_FRAME_EVAL_BATCH = 15,
DS4_TP_FRAME_MIXED_BATCH = 16,
DS4_TP_FRAME_COMMAND_ACK = 17,
} ds4_tp_frame_type;
typedef struct {
ds4_tp_frame_type type;
uint64_t session_id;
uint64_t seq;
int value;
int *tokens;
uint32_t n_tokens;
ds4_tp_batch_item *items;
uint32_t n_items;
} ds4_tp_command;
int ds4_tp_recv_command(
ds4_tp *tp,
ds4_tp_command *command,
char *err,
size_t errlen);
void ds4_tp_command_free(ds4_tp_command *command);
/* Debug lockstep check: both sides send their hidden-state hash for a token
* and compare. Returns 0 on transport failure, -1 on hash mismatch. */
int ds4_tp_hash_check(ds4_tp *tp, uint64_t seq, uint64_t hash, char *err, size_t errlen);
/* Vocab-split output head: the worker ships its logits half to the leader
* after every eval (and after a sync) on the control socket. */
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 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, uint64_t session_id,
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);
/* Standalone worker mode entry. Loads nothing itself: the engine is already
* open. */
int ds4_tp_worker_run(ds4_engine *engine, const ds4_tp_options *opt);
#endif
+10 -1
View File
@@ -1,3 +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
View File
@@ -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_tp.o ds4_ssd.o ds4_layer_pack.o rax.o
QUALITY_OBJS := ../ds4.o ../ds4_metal.o ../ds4_distributed.o ../ds4_tp.o ../ds4_ssd.o ../ds4_layer_pack.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_tp.o ds4_ssd.o ds4_layer_pack.o rax.o
QUALITY_OBJS := ../ds4.o ../ds4_cuda.o ../ds4_distributed.o ../ds4_tp.o ../ds4_ssd.o ../ds4_layer_pack.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
+60 -1
View File
@@ -8,7 +8,7 @@ The important pieces are:
- `deepseek4-quantize.c`: C HF-safetensors to GGUF quantizer.
- `quants.[ch]`: the deliberately small local quantization implementation used
by the quantizer. It implements the DS4 output formats we actually ship:
`q8_0`, `q4_K`, `q2_K`, and `iq2_xxs`.
`q8_0`, `q8_K`, `q4_K`, `q2_K`, and `iq2_xxs`.
- `imatrix/`: dataset and instructions for collecting routed-MoE activation
importance with `ds4`.
- `quality-testing/`: prompts and scripts used to compare local GGUF variants
@@ -86,6 +86,17 @@ gguf-tools/deepseek4-quantize \
--imatrix gguf/DeepSeek-V4-Flash-chat-v2-routed-moe-ds4.dat
```
True Q8_K routed experts:
```sh
gguf-tools/deepseek4-quantize \
--hf ../deepseek-v4-quants/hf/DeepSeek-V4-Flash \
--template gguf/DeepSeek-V4-Flash-Q4KExperts-F16HC-F16Compressor-F16Indexer-Q8Attn-Q8Shared-Q8Out-chat-v2-imatrix.gguf \
--out gguf/DeepSeek-V4-Flash-Q8KExperts-F16HC-F16Compressor-F16Indexer-Q8Attn-Q8Shared-Q8Out-chat-v2.gguf \
--experts q8_K \
--threads 8
```
You can override tensor families:
```sh
@@ -108,6 +119,54 @@ gguf-tools/deepseek4-quantize \
`--compare-tensor` regenerates a single tensor and byte-compares it against the
template or `--compare-gguf`. `--threads N` controls routed-expert workers.
## Convert A DSpark Support Checkpoint
The DSpark Flash checkpoint is published as Hugging Face safetensors and stores
the draft module under `mtp.0`, `mtp.1`, and `mtp.2`. Before writing a support
GGUF, inspect the official index and verify that every DSpark tensor name is
understood by the converter:
```sh
gguf-tools/deepseek4-quantize \
--hf ../deepseek-v4-quants/hf/DeepSeek-V4-Flash-DSpark \
--dspark-manifest > /tmp/dspark-manifest.tsv
```
The manifest reads only `model.safetensors.index.json`; it does not require the
large shard files to be present. The final summary should report three DSpark
stages and zero unknown DSpark tensors before attempting a full conversion.
To build the support GGUF used by `ds4 --mtp`, run the DSpark support mode. This
mode writes standalone DSpark metadata plus the packed `mtp.*` tensor payloads;
it does not require a base-model GGUF template:
```sh
gguf-tools/deepseek4-quantize \
--hf ../deepseek-v4-quants/hf/DeepSeek-V4-Flash-DSpark \
--dspark-support \
--out DeepSeek-V4-Flash-DSpark-support.gguf
```
`--dspark-support --dry-run` reads safetensors shard headers to derive exact
GGUF shapes and types, but it does not read tensor payloads. The DSpark metadata
defaults match the published Flash DSpark config: block size 5, target layers
40,41,42, Markov rank 256, and noise token 128799. Override them with
`--dspark-block-size`, `--dspark-target-layers`, `--dspark-markov-rank`, and
`--dspark-noise-token-id` if converting a different checkpoint.
Before a full write, regenerate one support tensor and record its checksum:
```sh
gguf-tools/deepseek4-quantize \
--hf ../deepseek-v4-quants/hf/DeepSeek-V4-Flash-DSpark \
--dspark-support \
--compare-tensor mtp.0.main_proj.weight
```
This reads only the payloads needed for that tensor. Add `--compare-gguf
DeepSeek-V4-Flash-DSpark-support.gguf` to byte-compare against an existing
support GGUF.
## When No Imatrix Is Given
`iq2_xxs` requires an importance vector. If `--imatrix` is not provided and
File diff suppressed because it is too large. Load diff
+108 -13
View File
@@ -1,14 +1,36 @@
# 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/laguna-openrouter-100`: 100 Laguna S 2.1 continuations collected
through OpenRouter `poolside/laguna-s-2.1`. Poolside's endpoint does not
expose output-token logprobs.
- `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 +41,50 @@ 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
```
For Laguna S 2.1 through OpenRouter:
```sh
export OPENROUTER_API_KEY=...
python3 gguf-tools/quality-testing/collect_official.py \
--model poolside/laguna-s-2.1 \
--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/laguna-openrouter-100 \
--count 100 \
--max-tokens 24 \
--top-logprobs 0 \
--token-limit-field max_tokens \
--thinking omit \
--reasoning-effort none
```
The Laguna fixture supports local target-token NLL, first-token agreement, and
greedy-prefix comparison. API logprob-delta and top-N agreement fields remain
zero because the Poolside endpoint does not return logprobs. Its raw response
files are retained for provenance but omitted from `manifest.tsv`, avoiding
attempts to parse unavailable API logprobs during local scoring.
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 +105,14 @@ 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 response path is the optional fourth manifest column and is included only
when `--top-logprobs` is greater than zero.
## 2. Build The Local Scorer
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.
## 3. Build The Local Scorer
```sh
make -C gguf-tools quality-score
@@ -50,27 +120,42 @@ 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/laguna-openrouter-100/manifest.tsv` for Laguna S 2.1 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 +171,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.
+91 -11
View File
@@ -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,32 @@ 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,
}
if top_logprobs > 0:
payload["logprobs"] = True
payload["top_logprobs"] = top_logprobs
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 +176,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 +219,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 +268,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 +284,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"
@@ -232,13 +305,20 @@ def main() -> int:
prompt_path.write_text(prompt, encoding="utf-8")
cont_path.write_text(content, encoding="utf-8")
resp_path.write_text(json.dumps(response, ensure_ascii=False, indent=2), encoding="utf-8")
rows.append((case_id, prompt_path, cont_path, resp_path))
rows.append((case_id, prompt_path, cont_path,
resp_path if args.top_logprobs > 0 else None))
time.sleep(0.05)
with manifest.open("w", encoding="utf-8") as fp:
fp.write("# id\tprompt_file\tcontinuation_file\tresponse_file\n")
fp.write("# id\tprompt_file\tcontinuation_file")
if args.top_logprobs > 0:
fp.write("\tresponse_file")
fp.write("\n")
for row in rows:
fp.write("\t".join([row[0], str(row[1]), str(row[2]), str(row[3])]) + "\n")
fields = [row[0], str(row[1]), str(row[2])]
if row[3] is not None:
fields.append(str(row[3]))
fp.write("\t".join(fields) + "\n")
print(f"wrote {manifest}", file=sys.stderr)
return 0
+74 -2
View File
@@ -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]:
+26
View File
@@ -0,0 +1,26 @@
# 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.
- `laguna-openrouter-100`: 100 Laguna S 2.1 OpenRouter continuations. Poolside's
endpoint does not expose output-token logprobs, so these support
teacher-forced continuation scoring only.
- `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 when
the endpoint provides them.
- `manifest.tsv`: paths consumed by `score_official`; the optional raw-response
path is present only for fixtures with API logprobs.
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
Loaded 100 of 1319 files, more files were not shown because too many files have changed in this diff. Show more