zenlm/zen5-engine
1
0
Fork 0
mirror of https://github.com/zenlm/zen5-engine.git synced 2026-07-26 22:09:00 +00:00
Code Issues Packages Projects Releases Wiki Activity

DS4 initial release

Browse Source
This commit is contained in:
antirez
2026-05-07 17:16:28 +02:00
commit d997b56c15
49 changed files with 59189 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.gitignore
+12
View File
@@ -0,0 +1,12 @@
/ds4
/ds4-server
/ds4_native
/ds4_server_test
/ds4_test
/ds4flash.gguf
/TODO.md
/gguf/
*.o
*.dSYM/
/misc/
.*.swp
AGENT.md
+53
View File
@@ -0,0 +1,53 @@
# Agent Notes
`ds4.c` is a DeepSeek V4 Flash specific inference engine. It is not a generic
GGUF runner. The goal is a small, readable, high-performance C codebase with
Objective-C only where Metal requires it and Metal kernels under `metal/`.
## Goals
- Keep the production path as whole-model Metal graph inference.
- Keep model loading mmap-backed; do not eagerly copy the full GGUF.
- Keep the CPU backend CPU-only and use it only as reference/debug code.
- Preserve correctness before speed. Do not keep a faster path with unexplained
attention, KV cache, or logits drift.
- Make long local agent sessions practical through live KV reuse and disk KV
checkpoints.
## Quality Rules
- Comment important inference code where the model mechanics, cache lifetime,
memory policy, or API orchestration are not obvious from the local code.
- Prefer comments beside the implementation over separate design documents.
- Keep comments instructive and compact: explain why a shape, ordering, cache
boundary, or memory choice exists.
- Keep public APIs narrow. CLI/server code should not know tensor internals.
- Do not add permanent semantic variants behind flags. Diagnostic switches are
fine when they validate the one release path.
- Do not introduce C++.
## Safety
- Avoid large CPU inference runs on macOS; the CPU path has previously exposed
kernel VM failures with very large mappings.
- Do not run multiple huge model processes concurrently. The instance lock is
intentional.
- Prefer short Metal smoke tests for build verification.
## Layout
- `ds4.c`: model loading, tokenizer, CPU reference code, Metal graph scheduling,
sessions, disk-cache payload serialization.
- `ds4_cli.c`: command line, linenoise REPL, interactive transcript handling.
- `ds4_server.c`: OpenAI/Anthropic compatible HTTP API, worker queue, streaming,
tool-call mapping, disk KV cache policy.
- `ds4_metal.m`: Objective-C Metal runtime and kernel wrappers.
- `metal/*.metal`: compute kernels.
- `tests/`: unit and live integration tests.
- `misc/`: ignored notes, experiments, and old planning material.
## Testing
Use `make` for build validation. Use `make test` for unit/regression tests when a
model and Metal are available. Use live server tests only when intentionally
testing the API surface.
LICENSE
+22
View File
@@ -0,0 +1,22 @@
MIT License
Copyright (c) 2026 The ds4.c authors
Copyright (c) 2023-2026 The ggml authors
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in all
copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
SOFTWARE.
Makefile
+76
View File
@@ -0,0 +1,76 @@
CC ?= cc
CFLAGS ?= -O3 -ffast-math -mcpu=native -Wall -Wextra -std=c99
OBJCFLAGS ?= -O3 -ffast-math -mcpu=native -Wall -Wextra -fobjc-arc
LDLIBS ?= -lm -pthread
UNAME_S := $(shell uname -s)
NATIVE_LDLIBS := $(LDLIBS)
METAL_SRCS := $(wildcard metal/*.metal)
ifeq ($(UNAME_S),Darwin)
METAL_LDLIBS := $(LDLIBS) -framework Foundation -framework Metal
CORE_OBJS = ds4.o ds4_metal.o
NATIVE_CORE_OBJS = ds4_native.o
else
CFLAGS += -DDS4_NO_METAL
CORE_OBJS = ds4.o
NATIVE_CORE_OBJS = ds4_native.o
METAL_LDLIBS := $(LDLIBS)
endif
.PHONY: all clean test
all: ds4 ds4-server
ifeq ($(UNAME_S),Darwin)
ds4: ds4_cli.o linenoise.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_cli.o linenoise.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4-server: ds4_server.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_server.o $(CORE_OBJS) $(METAL_LDLIBS)
ds4_native: ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS) $(NATIVE_LDLIBS)
else
ds4: ds4_cli.o linenoise.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ $^ $(LDLIBS)
ds4-server: ds4_server.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ $^ $(LDLIBS)
ds4_native: ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_cli_native.o linenoise.o $(NATIVE_CORE_OBJS) $(LDLIBS)
endif
ds4.o: ds4.c ds4.h ds4_metal.h
$(CC) $(CFLAGS) -c -o $@ ds4.c
ds4_cli.o: ds4_cli.c ds4.h linenoise.h
$(CC) $(CFLAGS) -c -o $@ ds4_cli.c
ds4_server.o: ds4_server.c ds4.h
$(CC) $(CFLAGS) -c -o $@ ds4_server.c
ds4_test.o: tests/ds4_test.c ds4_server.c ds4.h
$(CC) $(CFLAGS) -Wno-unused-function -c -o $@ tests/ds4_test.c
linenoise.o: linenoise.c linenoise.h
$(CC) $(CFLAGS) -c -o $@ linenoise.c
ds4_native.o: ds4.c ds4.h ds4_metal.h
$(CC) $(CFLAGS) -DDS4_NO_METAL -c -o $@ ds4.c
ds4_cli_native.o: ds4_cli.c ds4.h linenoise.h
$(CC) $(CFLAGS) -DDS4_NO_METAL -c -o $@ ds4_cli.c
ds4_metal.o: ds4_metal.m ds4_metal.h $(METAL_SRCS)
$(CC) $(OBJCFLAGS) -c -o $@ ds4_metal.m
ds4_test: ds4_test.o $(CORE_OBJS)
$(CC) $(CFLAGS) -o $@ ds4_test.o $(CORE_OBJS) $(METAL_LDLIBS)
test: ds4_test
./ds4_test
clean:
rm -f ds4 ds4-server ds4_native ds4_server_test ds4_test *.o
README.md
+497
View File
@@ -0,0 +1,497 @@
# ds4.c
`ds4.c` is a small native inference engine for DeepSeek V4 Flash. It is
intentionally narrow: not a generic GGUF runner, not a wrapper around another
runtime, and not a framework. The main path is a DeepSeek V4 Flash-specific
Metal graph executor with DS4-specific loading, prompt rendering, KV state, and
server API glue.
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.
Now, back at thsi project. Why we believe DeepSeek v4 Flash to be a pretty special
model deserving a stand alone engine? Because after comparing it with powerful smaller
dense models, we can report that:
1. DeepSeek v4 Flash is faster because of less active parameters.
2. In thinking mode, if you avoid *max thinking*, it produces a thinking section that is a lot shorter than other models, even 1/5 of other models in many cases, and crucially, the thinking section length is **proportional to the problem complexity**. This makes DeepSeek v4 Flash usable with thinking enabled when other models are practically impossible to use in the same conditions.
3. The model features a context window of **1 million tokens**.
4. Being so large, it knows more things if you go sampling at the edge of knowledge. For instance asking about Italian show or political questions soon uncovers that 284B parameters are a lot more than 27B or 35B parameters.
5. It writes much better English and Italian. It *feels* a quasi-frontier model.
6. The KV cache is incredibly compress, allowing long context inference on local computers and **on disk KV cache persistence**.
7. It works well with 2-bit quantization, if quantized in a special way (read later). This allows to run it in MacBooks with 128GB of RAM.
8. We expect DeepSeek to release **updated versions of v4 Flash** in the future, even better than the current one.
That said, a few important things about this project:
* 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 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 It is actually a first class disk citizen**.
* 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. This inference engine 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 just alpha quality code, so probably we are not still there.
* This is **Metal-only**, may implement CUDA support in the future? Perhaps, but nothing more. The CPU path is only for correctness check, but **warning: current macOS versions have a bug in the virutal memory implementation that will crash the kernel** if you try to run the CPU code. Remeber? Software sucks. I was not possible to fix the CPU inference to avoid crashing, since each time there is to restart the computer, which is not funny. Help us, if you have the guts.
## Acknowledgements to llama.cpp and GGML
`ds4.c` does not link against GGML, but it **exists thanks to the path opened by the
llama.cpp project and the kernels, quantization formats, GGUF ecosystem, and hard-won
engineering knowledge developed there**.
We are thankful and indebted to [`llama.cpp`](https://github.com/ggml-org/llama.cpp)
and its contributors. Their implementation, kernels, tests, and design choices were
an essential reference while building this DeepSeek V4 Flash-specific inference path.
Some source-level pieces are retained or adapted here under the MIT license: GGUF
quant layouts and tables, CPU quant/dot logic, and certain Metal kernels. For this
reason, and because we are genuinely grateful, we keep the GGML authors copyright
notice in our `LICENSE` file.
## Model Weights
This implementation only works with the DeepSeek V4 Flash GGUFs published for
this project. It is not a general GGUF loader, and arbitrary DeepSeek/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.
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,
projections, routing) are left untouched to guarantee quality.
Download one main model:
```sh
./download_model.sh q2 # 128 GB RAM machines
./download_model.sh q4 # >= 256 GB RAM machines
```
The script downloads from `https://huggingface.co/antirez/deepseek-v4-gguf`,
stores files under `./gguf/`, resumes partial downloads with `curl -C -`, and
updates `./ds4flash.gguf` to point at the selected q2/q4 model. Authentication
is optional for public downloads, but `--token TOKEN`, `HF_TOKEN`, or the local
Hugging Face token cache are used when present.
`./download_model.sh mtp` fetches the optional speculative decoding support
GGUF. It can be used with both q2 and q4, 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.
Then build:
```sh
make
```
`./ds4flash.gguf` is the default model path used by both binaries. Pass `-m` to
select another supported GGUF from `./gguf/`. Run `./ds4 --help` and
`./ds4-server --help` for the full flag list.
## Speed
These are single-run Metal CLI numbers with the q2 GGUF, `--ctx 32768`,
`--nothink`, greedy decoding, and `-n 256`. The short prompt is a normal small
Italian story prompt. The long prompt is 11709 tokens and exercises chunked
prefill plus long-context decode.
| Machine | Prompt | Prefill | Generation |
| --- | ---: | ---: | ---: |
| MacBook Pro M3 Max, 128 GB | short | 58.52 t/s | 26.68 t/s |
| MacBook Pro M3 Max, 128 GB | 11709 tokens | 250.11 t/s | 21.47 t/s |
| Mac Studio M3 Ultra, 512 GB | short | 84.43 t/s | 36.86 t/s |
| Mac Studio M3 Ultra, 512 GB | 11709 tokens | 468.03 t/s | 27.39 t/s |
## CLI
One-shot prompt:
```sh
./ds4 -p "Explain Redis streams in one paragraph."
```
No `-p` starts the interactive prompt:
```sh
./ds4
ds4>
```
The interactive CLI is a real multi-turn DS4 chat. It keeps the rendered chat
transcript and the live Metal KV checkpoint, so each turn extends the previous
conversation. Useful commands are `/help`, `/think`, `/think-max`, `/nothink`,
`/ctx N`, `/read FILE`, and `/quit`. Ctrl+C interrupts the current generation
and returns to `ds4>`.
The CLI defaults to thinking mode. Use `/nothink` or `--nothink` for direct
answers. `--mtp MTP.gguf --mtp-draft 2` enables the optional MTP speculative
path; it is useful only for greedy decoding, currently uses a confidence gate
(`--mtp-margin`) to avoid slow partial accepts, and should be treated as an
experimental slight-speedup path.
## Server
Start a local OpenAI/Anthropic-compatible server:
```sh
./ds4-server --ctx 100000 --kv-disk-dir /tmp/ds4-kv --kv-disk-space-mb 8192
```
The server is Metal-only. It keeps one mutable graph/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 Metal worker. The current server does not batch multiple
independent requests together; concurrent requests wait their turn on the single
live graph/session.
Supported endpoints:
- `GET /v1/models`
- `GET /v1/models/deepseek-v4-flash`
- `POST /v1/chat/completions`
- `POST /v1/completions`
- `POST /v1/messages`
`/v1/chat/completions` accepts the usual OpenAI-style `messages`,
`max_tokens`/`max_completion_tokens`, `temperature`, `top_p`, `top_k`, `min_p`,
`seed`, `stream`, `stream_options.include_usage`, `tools`, and `tool_choice`.
Tool schemas are rendered into DeepSeek's DSML tool format, and generated DSML
tool calls are mapped back to OpenAI tool calls.
`/v1/messages` is the Anthropic-compatible endpoint used by Claude Code style
clients. It accepts `system`, `messages`, `tools`, `tool_choice`, `max_tokens`,
`temperature`, `top_p`, `top_k`, `stream`, `stop_sequences`, and thinking
controls. Tool uses are returned as Anthropic `tool_use` blocks.
Both APIs support SSE streaming. In thinking mode, reasoning is streamed in the
native API shape instead of being mixed into final text.
Minimal OpenAI example:
```sh
curl http://127.0.0.1:8000/v1/chat/completions \
-H 'Content-Type: application/json' \
-d '{
"model":"deepseek-v4-flash",
"messages":[{"role":"user","content":"List three Redis design principles."}],
"stream":true
}'
```
### Agent Client Usage
`ds4-server` can be used by local coding agents that speak OpenAI-compatible
chat completions. Start the server first, and set the client context limit no
higher than the `--ctx` value you started the server with:
```sh
./ds4-server --ctx 100000 --kv-disk-dir /tmp/ds4-kv --kv-disk-space-mb 8192
```
You can use larger context and larger cache if you wish. Full context of
1M tokens is going to use more or less 26GB of memory (compressed indexer
alone will be like 22GB), so configure a context which makes sense in
your system. With 128GB of RAM you would run the 2-bit quants, which are
already 81GB, 26GB are going to be likely too much, so a context window
of 100~300k tokens is wiser.
The `384000` output limit below avoids token caps since the model is able
to generate very long replies otherwise (up to 384k tokens). The server
still stops when the configured context window is full.
For **opencode**, add a provider and agent entry to
`~/.config/opencode/opencode.json`:
```json
{
"$schema": "https://opencode.ai/config.json",
"provider": {
"ds4": {
"name": "ds4.c (local)",
"npm": "@ai-sdk/openai-compatible",
"options": {
"baseURL": "http://127.0.0.1:8000/v1",
"apiKey": "dsv4-local"
},
"models": {
"deepseek-v4-flash": {
"name": "DeepSeek V4 Flash (ds4.c local)",
"limit": {
"context": 100000,
"output": 384000
}
}
}
}
},
"agent": {
"ds4": {
"description": "DeepSeek V4 Flash served by local ds4-server",
"model": "ds4/deepseek-v4-flash",
"temperature": 0
}
}
}
```
For **Pi**, add a provider to `~/.pi/agent/models.json`:
```json
{
"providers": {
"ds4": {
"name": "ds4.c local",
"baseUrl": "http://127.0.0.1:8000/v1",
"api": "openai-completions",
"apiKey": "dsv4-local",
"compat": {
"supportsStore": false,
"supportsDeveloperRole": false,
"supportsReasoningEffort": true,
"supportsUsageInStreaming": true,
"maxTokensField": "max_tokens",
"supportsStrictMode": false,
"thinkingFormat": "deepseek",
"requiresReasoningContentOnAssistantMessages": true
},
"models": [
{
"id": "deepseek-v4-flash",
"name": "DeepSeek V4 Flash (ds4.c local)",
"reasoning": true,
"thinkingLevelMap": {
"off": null,
"minimal": "low",
"low": "low",
"medium": "medium",
"high": "high",
"xhigh": "xhigh"
},
"input": ["text"],
"contextWindow": 100000,
"maxTokens": 384000,
"cost": {
"input": 0,
"output": 0,
"cacheRead": 0,
"cacheWrite": 0
}
}
]
}
}
}
```
Optionally make it the default Pi model in `~/.pi/agent/settings.json`:
```json
{
"defaultProvider": "ds4",
"defaultModel": "deepseek-v4-flash"
}
```
For **Claude Code**, use the Anthropic-compatible endpoint. A wrapper like this
matches the local `~/bin/claude-ds4` setup:
```sh
#!/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_MODEL="deepseek-v4-flash"
export ANTHROPIC_CUSTOM_MODEL_OPTION="deepseek-v4-flash"
export ANTHROPIC_CUSTOM_MODEL_OPTION_NAME="DeepSeek V4 Flash local ds4"
export ANTHROPIC_CUSTOM_MODEL_OPTION_DESCRIPTION="ds4.c local GGUF"
export ANTHROPIC_DEFAULT_SONNET_MODEL="deepseek-v4-flash"
export ANTHROPIC_DEFAULT_HAIKU_MODEL="deepseek-v4-flash"
export ANTHROPIC_DEFAULT_OPUS_MODEL="deepseek-v4-flash"
export CLAUDE_CODE_SUBAGENT_MODEL="deepseek-v4-flash"
export CLAUDE_CODE_DISABLE_NONESSENTIAL_TRAFFIC=1
export CLAUDE_CODE_DISABLE_NONSTREAMING_FALLBACK=1
export CLAUDE_STREAM_IDLE_TIMEOUT_MS=600000
exec "$HOME/.local/bin/claude" "$@"
```
Claude Code may send a large initial prompt, often around 25k tokens, before it
starts doing useful work. Keep `--kv-disk-dir` enabled: after the first expensive
prefill, the disk KV cache lets later continuations or restarted sessions reuse
the saved prefix instead of processing the whole prompt again.
## Thinking Modes
DeepSeek V4 Flash has distinct non-thinking, thinking, and Think Max modes.
The server defaults to thinking mode. `reasoning_effort=max` requests Think
Max, but it is only applied when the context size is large enough for the model
card recommendation; smaller contexts fall back to normal thinking. OpenAI
`reasoning_effort=xhigh` still maps to normal thinking, not Think Max.
For direct replies, use `thinking: {"type":"disabled"}`, `think:false`, or a
non-thinking model alias such as `deepseek-chat`.
## Disk KV Cache
Chat/completion APIs are stateless: agent clients usually resend the whole
conversation every request. `ds4-server` handles this by comparing the rendered
token stream with cached token prefixes. The live in-memory checkpoint covers
the current session; the disk KV cache makes useful prefixes survive session
switches and server restarts.
For RAM reasons there is currently only one live KV cache in memory. When a new
unrelated session replaces it, the old checkpoint can only be resumed without
re-processing if it was written to the disk KV cache. In other words, memory
cache handles the active session; disk cache is the resume mechanism for
different sessions.
Enable it with:
```sh
./ds4-server --kv-disk-dir /tmp/ds4-kv --kv-disk-space-mb 8192
```
The cache key is the SHA1 of exact token IDs, not raw text. Each token ID is
hashed as a little-endian 32-bit integer, and files are named `<sha1>.kv`.
The file is intentionally written with ordinary `read`/`write` I/O, not
`mmap`, so restoring cache entries does not add more VM mappings to a process
that already maps the model.
On disk, a cache file is:
```text
KVC fixed header, 48 bytes
u32 rendered_text_bytes
rendered_text_bytes of UTF-8-ish token text
DS4 session payload, payload_bytes from the KVC header
```
The fixed header is little-endian:
```text
0 u8[3] magic = "KVC"
3 u8 version = 1
4 u8 routed expert quant bits, currently 2 or 4
5 u8 save reason: 0 unknown, 1 cold, 2 continued, 3 evict, 4 shutdown
6 u8[2] reserved
8 u32 cached token count
12 u32 hit count
16 u32 context size the snapshot was written for
20 u8[4] reserved
24 u64 creation Unix time
32 u64 last-used Unix time
40 u64 DS4 session payload byte count
```
The rendered text is the tokenizer-decoded text for the cached token prefix.
It is stored only for observability, so humans can inspect a cache directory
without decoding token IDs. It is not used as the key and it is not trusted
when loading; after load, the stored checkpoint tokens must still match the
incoming request prefix.
The DS4 session payload starts with thirteen little-endian `u32` fields:
```text
0 magic = "DSV4"
1 payload version = 1
2 saved context size
3 prefill chunk size
4 raw KV ring capacity
5 raw sliding-window length
6 compressed KV capacity
7 checkpoint token count
8 layer count
9 raw/head KV dimension
10 indexer head dimension
11 vocabulary size
12 live raw rows serialized below
```
Then it stores:
- `u32[token_count]` checkpoint token IDs.
- `float32[vocab_size]` logits for the next token after that checkpoint.
- `u32[layer_count]` compressed attention row counts.
- `u32[layer_count]` ratio-4 indexer row counts.
- For every layer: the live raw sliding-window KV rows, written in logical
position order rather than physical ring order.
- For compressed layers: live compressed KV rows and compressor frontier
tensors.
- For ratio-4 compressed layers: live indexer compressed rows and indexer
frontier tensors.
The logits are raw IEEE-754 `float32` values from the host `ds4_session`
buffer. They are saved immediately after the checkpoint tokens so a loaded
snapshot can sample or continue from the exact next-token distribution without
running one extra decode step. MTP draft logits/state are not persisted; after
loading a disk checkpoint the draft state is invalidated and rebuilt by normal
generation.
The tensor payload is DS4-specific KV/session state, not a generic inference
graph dump. It is expected to be portable only across compatible `ds4.c`
builds for this model layout.
The cache stores checkpoints at four moments:
- `cold`: after a long first prompt reaches a stable prefix, before generation.
- `continued`: when prefill or generation advances the live conversation by the configured interval.
- `evict`: before an unrelated request replaces the live in-memory session.
- `shutdown`: when the server exits cleanly.
Cold saves intentionally trim a small token suffix and align down to a prefill
chunk boundary. This avoids common BPE boundary retokenization misses when a
future request appends text to the same prompt. The defaults are conservative:
store prefixes of at least 512 tokens, cold-save prompts up to 30000 tokens,
trim 32 tail tokens, and align to 2048-token chunks. The important knobs are:
- `--kv-cache-min-tokens`
- `--kv-cache-cold-max-tokens`
- `--kv-cache-continued-interval-tokens`
- `--kv-cache-boundary-trim-tokens`
- `--kv-cache-boundary-align-tokens`
By default, checkpoints may be reused across the 2-bit and 4-bit routed-expert
variants if the token prefix matches. Use `--kv-cache-reject-different-quant`
when you want strict same-quant reuse only.
The cache directory is disposable. If behavior looks suspicious, stop the
server and remove it. You can investigate what is cached with hexdump as
the kv cache files include the verbatim prompt cached.
## Backends
The default backend is Metal:
```sh
./ds4 -p "Hello" --metal
```
There is also a CPU reference/debug path:
```sh
./ds4 -p "Hello" --cpu
```
Do not treat the CPU path as the production target. The server is Metal-only,
and the optimized implementation lives in the Metal graph path. This may
change in the future.
## Test Vectors
`tests/test-vectors` contains short and long-context continuation vectors
captured from the official DeepSeek V4 Flash API. The requests use
`deepseek-v4-flash`, greedy decoding, thinking disabled, and the maximum
`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.
All project tests are driven by the C runner:
```sh
make test # ./ds4_test --all
./ds4_test --logprob-vectors
./ds4_test --server
```
download_model.sh
Executable
+134
View File
@@ -0,0 +1,134 @@
#!/bin/sh
set -e
REPO="antirez/deepseek-v4-gguf"
Q2_FILE="DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2.gguf"
Q4_FILE="DeepSeek-V4-Flash-Q4KExperts-F16HC-F16Compressor-F16Indexer-Q8Attn-Q8Shared-Q8Out-chat-v2.gguf"
MTP_FILE="DeepSeek-V4-Flash-MTP-Q4K-Q8_0-F32.gguf"
ROOT=$(CDPATH= cd -- "$(dirname -- "$0")" && pwd)
OUT_DIR="$ROOT/gguf"
TOKEN=${HF_TOKEN:-}
usage() {
cat <<EOF
DeepSeek V4 Flash GGUF downloader
Usage:
./download_model.sh q2 [--token TOKEN]
./download_model.sh q4 [--token TOKEN]
./download_model.sh mtp [--token TOKEN]
Targets:
q2 2-bit routed experts, about 81 GB on disk.
Main model for 128 GB RAM machines.
q4 4-bit routed experts, about 153 GB on disk.
Main model for machines with 256 GB RAM or more.
mtp Optional speculative decoding component, about 3.5 GB on disk.
It is useful with both q2 and q4, but must be enabled explicitly
with --mtp when running ds4 or ds4-server.
Options:
--token TOKEN Hugging Face token. Otherwise HF_TOKEN or the local HF token
cache is used if present.
After q2/q4 downloads the script updates:
./ds4flash.gguf -> gguf/<selected model>
Then the default commands work:
./ds4 -p "Hello"
./ds4-server --ctx 100000
After downloading mtp, enable it explicitly, for example:
./ds4 --mtp gguf/$MTP_FILE --mtp-draft 2
EOF
}
if [ $# -eq 0 ]; then
usage
exit 1
fi
MODEL=$1
shift
case "$MODEL" in
q2) MODEL_FILE=$Q2_FILE ;;
q4) MODEL_FILE=$Q4_FILE ;;
mtp) MODEL_FILE=$MTP_FILE ;;
-h|--help|help)
usage
exit 0
;;
*)
echo "Unknown model: $MODEL" >&2
echo >&2
usage >&2
exit 1
;;
esac
while [ $# -gt 0 ]; do
case "$1" in
--token)
shift
if [ $# -eq 0 ]; then
echo "Missing value after --token" >&2
exit 1
fi
TOKEN=$1
;;
*)
echo "Unknown option: $1" >&2
exit 1
;;
esac
shift
done
if [ -z "$TOKEN" ] && [ -s "$HOME/.cache/huggingface/token" ]; then
TOKEN=$(cat "$HOME/.cache/huggingface/token")
fi
download_one() {
file=$1
out="$OUT_DIR/$file"
part="$out.part"
url="https://huggingface.co/$REPO/resolve/main/$file"
mkdir -p "$OUT_DIR"
if [ -s "$out" ]; then
echo "Already downloaded: $out"
return
fi
echo "Downloading $file"
echo "from https://huggingface.co/$REPO"
if [ -n "$TOKEN" ]; then
curl -fL -C - -H "Authorization: Bearer $TOKEN" -o "$part" "$url"
else
curl -fL -C - -o "$part" "$url"
fi
mv "$part" "$out"
}
download_one "$MODEL_FILE"
if [ "$MODEL" = "mtp" ]; then
echo
echo "MTP is an optional component for both q2 and q4."
echo "Enable it explicitly, for example:"
echo " ./ds4 --mtp gguf/$MTP_FILE --mtp-draft 2"
else
cd "$ROOT"
ln -sfn "gguf/$MODEL_FILE" ds4flash.gguf
echo "Linked ./ds4flash.gguf -> gguf/$MODEL_FILE"
fi
echo
echo "Done."
ds4.c
+16640
View File
File diff suppressed because it is too large Load Diff
ds4.h
+146
View File
@@ -0,0 +1,146 @@
#ifndef DS4_H
#define DS4_H
#include <stdbool.h>
#include <stddef.h>
#include <stdint.h>
#include <stdio.h>
/* Public engine boundary.
*
* The CLI and server should treat ds4_engine as the loaded model and
* ds4_session as one mutable inference timeline. A session owns the live KV
* cache and logits; callers provide full token prefixes and let
* ds4_session_sync() reuse, extend, or rebuild the graph state. Keep this
* header narrow so HTTP/CLI code does not depend on tensor internals. */
typedef enum {
DS4_BACKEND_METAL,
DS4_BACKEND_CPU,
} ds4_backend;
typedef enum {
DS4_THINK_NONE,
DS4_THINK_HIGH,
DS4_THINK_MAX,
} ds4_think_mode;
typedef struct {
int *v;
int len;
int cap;
} ds4_tokens;
typedef struct {
int id;
float logit;
float logprob;
} ds4_token_score;
typedef struct ds4_engine ds4_engine;
typedef struct ds4_session ds4_session;
typedef void (*ds4_session_progress_fn)(void *ud, const char *event, int current, int total);
typedef struct {
const char *model_path;
const char *mtp_path;
ds4_backend backend;
int n_threads;
int mtp_draft_tokens;
float mtp_margin;
bool warm_weights;
bool quality;
} ds4_engine_options;
typedef void (*ds4_token_emit_fn)(void *ud, int token);
typedef void (*ds4_generation_done_fn)(void *ud);
typedef struct {
uint64_t total_bytes;
uint64_t raw_bytes;
uint64_t compressed_bytes;
uint64_t scratch_bytes;
uint32_t prefill_cap;
uint32_t raw_cap;
uint32_t comp_cap;
} ds4_context_memory;
int ds4_engine_open(ds4_engine **out, const ds4_engine_options *opt);
void ds4_engine_close(ds4_engine *e);
void ds4_engine_summary(ds4_engine *e);
const char *ds4_backend_name(ds4_backend backend);
bool ds4_think_mode_enabled(ds4_think_mode mode);
const char *ds4_think_mode_name(ds4_think_mode mode);
const char *ds4_think_max_prefix(void);
uint32_t ds4_think_max_min_context(void);
ds4_think_mode ds4_think_mode_for_context(ds4_think_mode mode, int ctx_size);
ds4_context_memory ds4_context_memory_estimate(ds4_backend backend, int ctx_size);
int ds4_engine_generate_argmax(ds4_engine *e, const ds4_tokens *prompt,
int n_predict, int ctx_size,
ds4_token_emit_fn emit,
ds4_generation_done_fn done,
void *emit_ud,
ds4_session_progress_fn progress,
void *progress_ud);
void ds4_engine_dump_tokens(ds4_engine *e, const ds4_tokens *tokens);
int ds4_engine_head_test(ds4_engine *e, const ds4_tokens *prompt);
int ds4_engine_first_token_test(ds4_engine *e, const ds4_tokens *prompt);
int ds4_engine_metal_graph_test(ds4_engine *e, const ds4_tokens *prompt);
int ds4_engine_metal_graph_full_test(ds4_engine *e, const ds4_tokens *prompt);
int ds4_engine_metal_graph_prompt_test(ds4_engine *e, const ds4_tokens *prompt, int ctx_size);
void ds4_tokens_push(ds4_tokens *tv, int token);
void ds4_tokens_free(ds4_tokens *tv);
void ds4_tokens_copy(ds4_tokens *dst, const ds4_tokens *src);
bool ds4_tokens_starts_with(const ds4_tokens *tokens, const ds4_tokens *prefix);
void ds4_tokenize_text(ds4_engine *e, const char *text, ds4_tokens *out);
void ds4_tokenize_rendered_chat(ds4_engine *e, const char *text, ds4_tokens *out);
void ds4_chat_begin(ds4_engine *e, ds4_tokens *tokens);
void ds4_encode_chat_prompt(
ds4_engine *e,
const char *system,
const char *prompt,
ds4_think_mode think_mode,
ds4_tokens *out);
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);
char *ds4_token_text(ds4_engine *e, int token, size_t *len);
int ds4_token_eos(ds4_engine *e);
int ds4_session_create(ds4_session **out, ds4_engine *e, int ctx_size);
void ds4_session_free(ds4_session *s);
void ds4_session_set_progress(ds4_session *s, ds4_session_progress_fn fn, void *ud);
/* Synchronize the live session to a full prompt token prefix. If the current
* checkpoint is a prefix, only the suffix is evaluated; otherwise the graph is
* refilled from scratch. */
int ds4_session_sync(ds4_session *s, const ds4_tokens *prompt, char *err, size_t errlen);
int ds4_session_common_prefix(ds4_session *s, const ds4_tokens *prompt);
int ds4_session_argmax(ds4_session *s);
int ds4_session_sample(ds4_session *s, float temperature, int top_k, float top_p, float min_p, uint64_t *rng);
int ds4_session_top_logprobs(ds4_session *s, ds4_token_score *out, int k);
int ds4_session_eval(ds4_session *s, int token, char *err, size_t errlen);
int ds4_session_eval_speculative_argmax(ds4_session *s, int first_token,
int max_tokens, int eos_token,
int *accepted, int accepted_cap,
char *err, size_t errlen);
void ds4_session_invalidate(ds4_session *s);
void ds4_session_rewind(ds4_session *s, int pos);
int ds4_session_pos(ds4_session *s);
int ds4_session_ctx(ds4_session *s);
int ds4_engine_routed_quant_bits(ds4_engine *e);
bool ds4_engine_has_mtp(ds4_engine *e);
int ds4_engine_mtp_draft_tokens(ds4_engine *e);
const ds4_tokens *ds4_session_tokens(ds4_session *s);
/* Disk KV cache payload helpers. The server owns the outer file header and
* policy; the engine owns the DS4-specific serialized graph state. */
uint64_t ds4_session_payload_bytes(ds4_session *s);
int ds4_session_save_payload(ds4_session *s, FILE *fp, char *err, size_t errlen);
int ds4_session_load_payload(ds4_session *s, FILE *fp, uint64_t payload_bytes, char *err, size_t errlen);
#endif
ds4_cli.c
+1273
View File
File diff suppressed because it is too large Load Diff
ds4_metal.h
+789
View File
@@ -0,0 +1,789 @@
#ifndef DS4_METAL_H
#define DS4_METAL_H
#include <stdbool.h>
#include <stdint.h>
/* =========================================================================
* Metal Tensor and Command Lifetime.
* =========================================================================
*
* Opaque device tensor used by the DS4-specific Metal executor.
*
* The public Metal API is tensor-resident: activations, KV state, and scratch
* buffers stay device-owned across the whole prefill/decode command sequence.
*/
typedef struct ds4_metal_tensor ds4_metal_tensor;
int ds4_metal_init(void);
void ds4_metal_cleanup(void);
ds4_metal_tensor *ds4_metal_tensor_alloc(uint64_t bytes);
ds4_metal_tensor *ds4_metal_tensor_view(const ds4_metal_tensor *base, uint64_t offset, uint64_t bytes);
void ds4_metal_tensor_free(ds4_metal_tensor *tensor);
uint64_t ds4_metal_tensor_bytes(const ds4_metal_tensor *tensor);
void *ds4_metal_tensor_contents(ds4_metal_tensor *tensor);
int ds4_metal_tensor_write(ds4_metal_tensor *tensor, uint64_t offset, const void *data, uint64_t bytes);
int ds4_metal_tensor_read(const ds4_metal_tensor *tensor, uint64_t offset, void *data, uint64_t bytes);
int ds4_metal_tensor_copy(ds4_metal_tensor *dst, uint64_t dst_offset,
const ds4_metal_tensor *src, uint64_t src_offset,
uint64_t bytes);
int ds4_metal_begin_commands(void);
int ds4_metal_flush_commands(void);
int ds4_metal_end_commands(void);
int ds4_metal_synchronize(void);
int ds4_metal_set_model_map(const void *model_map, uint64_t model_size);
int ds4_metal_set_model_map_range(const void *model_map, uint64_t model_size, uint64_t map_offset, uint64_t map_size);
void ds4_metal_set_quality(bool quality);
void ds4_metal_print_memory_report(const char *label);
/* =========================================================================
* Embeddings and Indexer Helpers.
* =========================================================================
*
* These kernels seed HC state from token embeddings and implement the ratio-4
* compressed-attention indexer that chooses visible compressed rows.
*/
int ds4_metal_embed_token_hc_tensor(
ds4_metal_tensor *out_hc,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint32_t n_vocab,
uint32_t token,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_embed_tokens_hc_tensor(
ds4_metal_tensor *out_hc,
const ds4_metal_tensor *tokens,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint32_t n_vocab,
uint32_t n_tokens,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_indexer_score_one_tensor(
ds4_metal_tensor *scores,
const ds4_metal_tensor *q,
const ds4_metal_tensor *weights,
const ds4_metal_tensor *index_comp,
uint32_t n_comp,
uint32_t n_head,
uint32_t head_dim,
float scale);
int ds4_metal_indexer_scores_prefill_tensor(
ds4_metal_tensor *scores,
const ds4_metal_tensor *q,
const ds4_metal_tensor *weights,
const ds4_metal_tensor *index_comp,
uint32_t n_comp,
uint32_t n_tokens,
uint32_t n_head,
uint32_t head_dim,
uint32_t ratio,
float scale);
int ds4_metal_indexer_scores_decode_batch_tensor(
ds4_metal_tensor *scores,
const ds4_metal_tensor *q,
const ds4_metal_tensor *weights,
const ds4_metal_tensor *index_comp,
uint32_t n_comp,
uint32_t n_tokens,
uint32_t pos0,
uint32_t n_head,
uint32_t head_dim,
uint32_t ratio,
float scale);
int ds4_metal_indexer_topk_tensor(
ds4_metal_tensor *selected,
const ds4_metal_tensor *scores,
uint32_t n_comp,
uint32_t n_tokens,
uint32_t top_k);
int ds4_metal_dsv4_topk_mask_tensor(
ds4_metal_tensor *mask,
const ds4_metal_tensor *topk,
uint32_t n_comp,
uint32_t n_tokens,
uint32_t top_k);
/* =========================================================================
* Dense Projections, Norms, RoPE, and KV Rounding.
* =========================================================================
*
* The graph uses these primitives for Q/KV projections, HC/output projections,
* attention output projections, and DS4's tail-only RoPE.
*/
int ds4_metal_matmul_q8_0_tensor(
ds4_metal_tensor *out,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint64_t in_dim,
uint64_t out_dim,
const ds4_metal_tensor *x,
uint64_t n_tok);
int ds4_metal_shared_gate_up_swiglu_q8_0_tensor(
ds4_metal_tensor *gate,
ds4_metal_tensor *up,
ds4_metal_tensor *mid,
const void *model_map,
uint64_t model_size,
uint64_t gate_offset,
uint64_t up_offset,
uint64_t in_dim,
uint64_t out_dim,
const ds4_metal_tensor *x);
int ds4_metal_matmul_f16_tensor(
ds4_metal_tensor *out,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint64_t in_dim,
uint64_t out_dim,
const ds4_metal_tensor *x,
uint64_t n_tok);
int ds4_metal_matmul_f16_pair_tensor(
ds4_metal_tensor *out_a,
ds4_metal_tensor *out_b,
const void *model_map,
uint64_t model_size,
uint64_t weight_a_offset,
uint64_t weight_b_offset,
uint64_t in_dim,
uint64_t out_dim,
const ds4_metal_tensor *x,
uint64_t n_tok);
int ds4_metal_matmul_f32_tensor(
ds4_metal_tensor *out,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint64_t in_dim,
uint64_t out_dim,
const ds4_metal_tensor *x,
uint64_t n_tok);
int ds4_metal_repeat_hc_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *row,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_rms_norm_plain_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *x,
uint32_t n,
float eps);
int ds4_metal_rms_norm_plain_rows_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *x,
uint32_t n,
uint32_t rows,
float eps);
int ds4_metal_rms_norm_weight_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *x,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint32_t n,
float eps);
int ds4_metal_rms_norm_weight_rows_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *x,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint32_t n,
uint32_t rows,
float eps);
int ds4_metal_dsv4_qkv_rms_norm_rows_tensor(
ds4_metal_tensor *q_out,
const ds4_metal_tensor *q,
const void *model_map,
uint64_t model_size,
uint64_t q_weight_offset,
uint32_t q_n,
ds4_metal_tensor *kv_out,
const ds4_metal_tensor *kv,
uint64_t kv_weight_offset,
uint32_t kv_n,
uint32_t rows,
float eps);
int ds4_metal_head_rms_norm_tensor(
ds4_metal_tensor *x,
uint32_t n_tok,
uint32_t n_head,
uint32_t head_dim,
float eps);
int ds4_metal_dsv4_fp8_kv_quantize_tensor(
ds4_metal_tensor *x,
uint32_t n_tok,
uint32_t head_dim,
uint32_t n_rot);
int ds4_metal_rope_tail_tensor(
ds4_metal_tensor *x,
uint32_t n_tok,
uint32_t n_head,
uint32_t 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);
/* Release decode fused KV finalizer: after the standalone RoPE kernel, this
* performs DS4's FP8 non-RoPE KV round trip and writes the F16-rounded raw
* attention cache row in one dispatch. */
int ds4_metal_kv_fp8_store_raw_tensor(
ds4_metal_tensor *kv,
ds4_metal_tensor *raw_cache,
uint32_t raw_cap,
uint32_t row,
uint32_t head_dim,
uint32_t n_rot);
/* Reference/raw-cache primitive kept for prefill and diagnostics. Decode uses
* ds4_metal_kv_fp8_store_raw_tensor unless a diagnostic reference path is
* explicitly selected by the graph driver. */
int ds4_metal_store_raw_kv_tensor(
ds4_metal_tensor *raw_cache,
const ds4_metal_tensor *kv,
uint32_t raw_cap,
uint32_t row,
uint32_t head_dim);
int ds4_metal_store_raw_kv_batch_tensor(
ds4_metal_tensor *raw_cache,
const ds4_metal_tensor *kv,
uint32_t raw_cap,
uint32_t pos0,
uint32_t n_tokens,
uint32_t head_dim);
/* =========================================================================
* KV Compression and Attention.
* =========================================================================
*
* Compressed layers maintain rolling score/KV state and append pooled rows at
* ratio boundaries. Attention kernels consume raw SWA rows, compressed rows,
* and optional indexer masks.
*/
int ds4_metal_compressor_update_tensor(
const ds4_metal_tensor *kv_cur,
const ds4_metal_tensor *sc_cur,
ds4_metal_tensor *state_kv,
ds4_metal_tensor *state_score,
ds4_metal_tensor *comp_cache,
const void *model_map,
uint64_t model_size,
uint64_t ape_offset,
uint32_t ape_type,
uint64_t norm_offset,
uint32_t norm_type,
uint32_t head_dim,
uint32_t ratio,
uint32_t pos,
uint32_t comp_row,
uint32_t n_rot,
uint32_t n_ctx_orig,
float freq_base,
float freq_scale,
float ext_factor,
float attn_factor,
float beta_fast,
float beta_slow,
float rms_eps);
int ds4_metal_compressor_store_batch_tensor(
const ds4_metal_tensor *kv,
const ds4_metal_tensor *sc,
ds4_metal_tensor *state_kv,
ds4_metal_tensor *state_score,
const void *model_map,
uint64_t model_size,
uint64_t ape_offset,
uint32_t ape_type,
uint32_t head_dim,
uint32_t ratio,
uint32_t pos0,
uint32_t n_tokens);
int ds4_metal_compressor_prefill_tensor(
ds4_metal_tensor *comp_cache,
ds4_metal_tensor *state_kv,
ds4_metal_tensor *state_score,
const ds4_metal_tensor *kv,
const ds4_metal_tensor *sc,
const void *model_map,
uint64_t model_size,
uint64_t ape_offset,
uint32_t ape_type,
uint64_t norm_offset,
uint32_t norm_type,
uint32_t head_dim,
uint32_t ratio,
uint32_t pos0,
uint32_t n_tokens,
uint32_t n_rot,
uint32_t n_ctx_orig,
bool quantize_fp8,
float freq_base,
float freq_scale,
float ext_factor,
float attn_factor,
float beta_fast,
float beta_slow,
float rms_eps);
int ds4_metal_compressor_prefill_ratio4_replay_tensor(
ds4_metal_tensor *comp_cache,
ds4_metal_tensor *state_kv,
ds4_metal_tensor *state_score,
const ds4_metal_tensor *kv,
const ds4_metal_tensor *sc,
const void *model_map,
uint64_t model_size,
uint64_t ape_offset,
uint32_t ape_type,
uint64_t norm_offset,
uint32_t norm_type,
uint32_t head_dim,
uint32_t pos0,
uint32_t n_tokens,
uint32_t n_rot,
uint32_t n_ctx_orig,
bool quantize_fp8,
float freq_base,
float freq_scale,
float ext_factor,
float attn_factor,
float beta_fast,
float beta_slow,
float rms_eps);
int ds4_metal_compressor_prefill_state_ratio4_tensor(
ds4_metal_tensor *state_kv,
ds4_metal_tensor *state_score,
const ds4_metal_tensor *kv_tail,
const ds4_metal_tensor *sc_tail,
const void *model_map,
uint64_t model_size,
uint64_t ape_offset,
uint32_t ape_type,
uint32_t head_dim,
uint32_t pos0);
int ds4_metal_attention_decode_heads_tensor(
ds4_metal_tensor *heads,
const void *model_map,
uint64_t model_size,
uint64_t sinks_offset,
const ds4_metal_tensor *q,
const ds4_metal_tensor *raw_kv,
uint32_t n_raw,
uint32_t raw_cap,
uint32_t raw_start,
const ds4_metal_tensor *comp_kv,
uint32_t n_comp,
const ds4_metal_tensor *comp_mask,
uint32_t use_mask,
uint32_t n_head,
uint32_t head_dim);
int ds4_metal_attention_prefill_raw_heads_tensor(
ds4_metal_tensor *heads,
const void *model_map,
uint64_t model_size,
uint64_t sinks_offset,
const ds4_metal_tensor *q,
const ds4_metal_tensor *raw_kv,
uint32_t n_tokens,
uint32_t window,
uint32_t n_head,
uint32_t head_dim);
int ds4_metal_attention_decode_raw_batch_heads_tensor(
ds4_metal_tensor *heads,
const void *model_map,
uint64_t model_size,
uint64_t sinks_offset,
const ds4_metal_tensor *q,
const ds4_metal_tensor *raw_kv,
uint32_t n_tokens,
uint32_t pos0,
uint32_t n_raw,
uint32_t raw_cap,
uint32_t raw_start,
uint32_t window,
uint32_t n_head,
uint32_t head_dim);
int ds4_metal_attention_decode_mixed_batch_heads_tensor(
ds4_metal_tensor *heads,
const void *model_map,
uint64_t model_size,
uint64_t sinks_offset,
const ds4_metal_tensor *q,
const ds4_metal_tensor *raw_kv,
const ds4_metal_tensor *comp_kv,
const ds4_metal_tensor *comp_mask,
uint32_t use_comp_mask,
uint32_t n_tokens,
uint32_t pos0,
uint32_t n_raw,
uint32_t raw_cap,
uint32_t raw_start,
uint32_t n_comp,
uint32_t window,
uint32_t ratio,
uint32_t n_head,
uint32_t head_dim);
int ds4_metal_attention_indexed_mixed_batch_heads_tensor(
ds4_metal_tensor *heads,
const void *model_map,
uint64_t model_size,
uint64_t sinks_offset,
const ds4_metal_tensor *q,
const ds4_metal_tensor *raw_kv,
const ds4_metal_tensor *comp_kv,
const ds4_metal_tensor *topk,
uint32_t n_tokens,
uint32_t pos0,
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 n_head,
uint32_t head_dim);
int ds4_metal_attention_prefill_static_mixed_heads_tensor(
ds4_metal_tensor *heads,
const void *model_map,
uint64_t model_size,
uint64_t sinks_offset,
const ds4_metal_tensor *q,
const ds4_metal_tensor *raw_kv,
const ds4_metal_tensor *comp_kv,
uint32_t n_tokens,
uint32_t n_comp,
uint32_t window,
uint32_t ratio,
uint32_t n_head,
uint32_t head_dim);
int ds4_metal_attention_prefill_masked_mixed_heads_tensor(
ds4_metal_tensor *heads,
const void *model_map,
uint64_t model_size,
uint64_t sinks_offset,
const ds4_metal_tensor *q,
const ds4_metal_tensor *raw_kv,
const ds4_metal_tensor *comp_kv,
const ds4_metal_tensor *comp_mask,
uint32_t n_tokens,
uint32_t n_comp,
uint32_t window,
uint32_t ratio,
uint32_t n_head,
uint32_t head_dim);
int ds4_metal_attention_output_q8_batch_tensor(
ds4_metal_tensor *out,
ds4_metal_tensor *low,
ds4_metal_tensor *group_tmp,
ds4_metal_tensor *low_tmp,
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,
uint64_t out_dim,
const ds4_metal_tensor *heads,
uint32_t n_tokens);
int ds4_metal_attention_output_low_q8_tensor(
ds4_metal_tensor *low,
const void *model_map,
uint64_t model_size,
uint64_t out_a_offset,
uint64_t group_dim,
uint64_t rank,
uint32_t n_groups,
const ds4_metal_tensor *heads);
/* =========================================================================
* Router, Shared Expert, and Routed MoE.
* =========================================================================
*
* These kernels implement the FFN body: router probabilities/top-k or hash
* routing, shared SwiGLU, and the IQ2_XXS/Q2_K/Q4_K routed experts.
*/
int ds4_metal_swiglu_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *gate,
const ds4_metal_tensor *up,
uint32_t n,
float clamp,
float weight);
int ds4_metal_add_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *a,
const ds4_metal_tensor *b,
uint32_t n);
int ds4_metal_router_select_tensor(
ds4_metal_tensor *selected,
ds4_metal_tensor *weights,
ds4_metal_tensor *probs,
const void *model_map,
uint64_t model_size,
uint64_t bias_offset,
uint64_t hash_offset,
uint32_t hash_rows,
uint32_t token,
uint32_t n_expert_groups,
uint32_t n_group_used,
bool has_bias,
bool hash_mode,
const ds4_metal_tensor *logits);
int ds4_metal_router_select_batch_tensor(
ds4_metal_tensor *selected,
ds4_metal_tensor *weights,
ds4_metal_tensor *probs,
const void *model_map,
uint64_t model_size,
uint64_t bias_offset,
uint64_t hash_offset,
uint32_t hash_rows,
uint32_t n_expert_groups,
uint32_t n_group_used,
bool has_bias,
bool hash_mode,
const ds4_metal_tensor *logits,
const ds4_metal_tensor *tokens,
uint32_t n_tokens);
int ds4_metal_routed_moe_one_tensor(
ds4_metal_tensor *out,
ds4_metal_tensor *gate,
ds4_metal_tensor *up,
ds4_metal_tensor *mid,
ds4_metal_tensor *experts,
const void *model_map,
uint64_t model_size,
uint64_t gate_offset,
uint64_t up_offset,
uint64_t down_offset,
uint32_t gate_type,
uint32_t down_type,
uint64_t gate_expert_bytes,
uint64_t gate_row_bytes,
uint64_t down_expert_bytes,
uint64_t down_row_bytes,
uint32_t expert_in_dim,
uint32_t expert_mid_dim,
uint32_t out_dim,
const ds4_metal_tensor *selected,
const ds4_metal_tensor *weights,
uint32_t n_expert,
float clamp,
const ds4_metal_tensor *x);
int ds4_metal_routed_moe_batch_tensor(
ds4_metal_tensor *out,
ds4_metal_tensor *gate,
ds4_metal_tensor *up,
ds4_metal_tensor *mid,
ds4_metal_tensor *experts,
const void *model_map,
uint64_t model_size,
uint64_t gate_offset,
uint64_t up_offset,
uint64_t down_offset,
uint32_t gate_type,
uint32_t down_type,
uint64_t gate_expert_bytes,
uint64_t gate_row_bytes,
uint64_t down_expert_bytes,
uint64_t down_row_bytes,
uint32_t expert_in_dim,
uint32_t expert_mid_dim,
uint32_t out_dim,
const ds4_metal_tensor *selected,
const ds4_metal_tensor *weights,
uint32_t n_expert,
float clamp,
const ds4_metal_tensor *x,
uint32_t n_tokens);
/* =========================================================================
* Hyper-Connection Kernels.
* =========================================================================
*
* HC kernels reduce four residual streams before a sublayer and expand the
* sublayer output back into four streams afterward.
*/
int ds4_metal_hc_split_sinkhorn_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *mix,
const void *model_map,
uint64_t model_size,
uint64_t scale_offset,
uint64_t base_offset,
uint32_t n_hc,
uint32_t sinkhorn_iters,
float eps);
int ds4_metal_hc_weighted_sum_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *residual_hc,
const ds4_metal_tensor *weights,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_hc_weighted_sum_split_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *residual_hc,
const ds4_metal_tensor *split,
uint32_t n_embd,
uint32_t n_hc);
/* Release decode fused HC pre-sublayer operation: split the HC mixer and
* immediately reduce four HC streams into the active 4096-wide sublayer row. */
int ds4_metal_hc_split_weighted_sum_tensor(
ds4_metal_tensor *out,
ds4_metal_tensor *split,
const ds4_metal_tensor *mix,
const ds4_metal_tensor *residual_hc,
const void *model_map,
uint64_t model_size,
uint64_t scale_offset,
uint64_t base_offset,
uint32_t n_embd,
uint32_t n_hc,
uint32_t sinkhorn_iters,
float eps);
int ds4_metal_hc_split_weighted_sum_norm_tensor(
ds4_metal_tensor *out,
ds4_metal_tensor *norm_out,
ds4_metal_tensor *split,
const ds4_metal_tensor *mix,
const ds4_metal_tensor *residual_hc,
const void *model_map,
uint64_t model_size,
uint64_t scale_offset,
uint64_t base_offset,
uint64_t norm_weight_offset,
uint32_t n_embd,
uint32_t n_hc,
uint32_t sinkhorn_iters,
float eps,
float norm_eps);
int ds4_metal_output_hc_weights_tensor(
ds4_metal_tensor *out,
const ds4_metal_tensor *pre,
const void *model_map,
uint64_t model_size,
uint64_t scale_offset,
uint64_t base_offset,
uint32_t n_hc,
float eps);
int ds4_metal_hc_expand_tensor(
ds4_metal_tensor *out_hc,
const ds4_metal_tensor *block_out,
const ds4_metal_tensor *residual_hc,
const ds4_metal_tensor *post,
const ds4_metal_tensor *comb,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_hc_expand_split_tensor(
ds4_metal_tensor *out_hc,
const ds4_metal_tensor *block_out,
const ds4_metal_tensor *residual_hc,
const ds4_metal_tensor *split,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_hc_expand_add_split_tensor(
ds4_metal_tensor *out_hc,
const ds4_metal_tensor *block_out,
const ds4_metal_tensor *block_add,
const ds4_metal_tensor *residual_hc,
const ds4_metal_tensor *split,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_shared_down_hc_expand_q8_0_tensor(
ds4_metal_tensor *out_hc,
ds4_metal_tensor *shared_out,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint64_t in_dim,
uint64_t out_dim,
const ds4_metal_tensor *shared_mid,
const ds4_metal_tensor *routed_out,
const ds4_metal_tensor *residual_hc,
const ds4_metal_tensor *split,
uint32_t n_embd,
uint32_t n_hc);
int ds4_metal_matmul_q8_0_hc_expand_tensor(
ds4_metal_tensor *out_hc,
ds4_metal_tensor *block_out,
const void *model_map,
uint64_t model_size,
uint64_t weight_offset,
uint64_t in_dim,
uint64_t out_dim,
const ds4_metal_tensor *x,
const ds4_metal_tensor *residual_hc,
const ds4_metal_tensor *split,
uint32_t n_embd,
uint32_t n_hc);
#endif
ds4_metal.m
+14491
View File
File diff suppressed because it is too large Load Diff
ds4_server.c
+6885
View File
File diff suppressed because it is too large Load Diff
linenoise.c
+2380
View File
File diff suppressed because it is too large Load Diff
linenoise.h
+120
View File
@@ -0,0 +1,120 @@
/* linenoise.h -- VERSION 1.0
*
* Guerrilla line editing library against the idea that a line editing lib
* needs to be 20,000 lines of C code.
*
* See linenoise.c for more information.
*
* ------------------------------------------------------------------------
*
* Copyright (c) 2010-2023, Salvatore Sanfilippo <antirez at gmail dot com>
* Copyright (c) 2010-2013, Pieter Noordhuis <pcnoordhuis at gmail dot com>
*
* All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are
* met:
*
* * Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
*
* * Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
* A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
* HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
* SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
* LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
* DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
* THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*/
#ifndef __LINENOISE_H
#define __LINENOISE_H
#ifdef __cplusplus
extern "C" {
#endif
#include <stddef.h> /* For size_t. */
extern char *linenoiseEditMore;
#define LINENOISE_MAX_FOLDS 16
/* The linenoiseState structure represents the state during line editing.
* We pass this state to functions implementing specific editing
* functionalities. */
struct linenoiseState {
int in_completion; /* The user pressed TAB and we are now in completion
* mode, so input is handled by completeLine(). */
size_t completion_idx; /* Index of next completion to propose. */
int ifd; /* Terminal stdin file descriptor. */
int ofd; /* Terminal stdout file descriptor. */
char *buf; /* Edited line buffer. */
size_t buflen; /* Edited line buffer size. */
size_t buflen_max; /* Max buffer size, or 0 if fixed. */
const char *prompt; /* Prompt to display. */
size_t plen; /* Prompt length. */
size_t pos; /* Current cursor position. */
size_t oldpos; /* Previous refresh cursor position. */
size_t len; /* Current edited line length. */
size_t cols; /* Number of columns in terminal. */
size_t oldrows; /* Rows used by last refrehsed line (multiline mode) */
int oldrpos; /* Cursor row from last refresh (for multiline clearing). */
int history_index; /* The history index we are currently editing. */
int fold_count; /* Number of folded ranges. */
size_t fold_start[LINENOISE_MAX_FOLDS]; /* Folded range start offsets. */
size_t fold_end[LINENOISE_MAX_FOLDS]; /* Folded range end offsets. */
};
typedef struct linenoiseCompletions {
size_t len;
char **cvec;
} linenoiseCompletions;
/* Non blocking API. */
int linenoiseEditStart(struct linenoiseState *l, int stdin_fd, int stdout_fd, char *buf, size_t buflen, const char *prompt);
char *linenoiseEditFeed(struct linenoiseState *l);
void linenoiseEditStop(struct linenoiseState *l);
void linenoiseHide(struct linenoiseState *l);
void linenoiseShow(struct linenoiseState *l);
/* Blocking API. */
char *linenoise(const char *prompt);
void linenoiseFree(void *ptr);
/* Completion API. */
typedef void(linenoiseCompletionCallback)(const char *, linenoiseCompletions *);
typedef char*(linenoiseHintsCallback)(const char *, int *color, int *bold);
typedef void(linenoiseFreeHintsCallback)(void *);
void linenoiseSetCompletionCallback(linenoiseCompletionCallback *);
void linenoiseSetHintsCallback(linenoiseHintsCallback *);
void linenoiseSetFreeHintsCallback(linenoiseFreeHintsCallback *);
void linenoiseAddCompletion(linenoiseCompletions *, const char *);
/* History API. */
int linenoiseHistoryAdd(const char *line);
int linenoiseHistorySetMaxLen(int len);
int linenoiseHistorySave(const char *filename);
int linenoiseHistoryLoad(const char *filename);
/* Other utilities. */
void linenoiseClearScreen(void);
void linenoiseSetMultiLine(int ml);
void linenoisePrintKeyCodes(void);
void linenoiseMaskModeEnable(void);
void linenoiseMaskModeDisable(void);
#ifdef __cplusplus
}
#endif
#endif /* __LINENOISE_H */
metal/argsort.metal
+266
View File
@@ -0,0 +1,266 @@
struct ds4_metal_args_argsort {
int32_t ne00;
int32_t ne01;
int32_t ne02;
int32_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne0;
int32_t ne1;
int32_t ne2;
int32_t ne3;
int32_t top_k;
};
struct ds4_metal_args_argsort_merge {
int64_t ne00;
int64_t ne01;
int64_t ne02;
int64_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne0;
int32_t ne1;
int32_t ne2;
int32_t ne3;
int32_t top_k;
int32_t len;
};
typedef void (argsort_t)(
constant ds4_metal_args_argsort & args,
device const char * src0,
device int32_t * dst,
threadgroup int32_t * shmem_i32 [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]);
// Sort one float row into an index row. DS4 only exports the descending
// instance because router and indexer selection both need top-k order.
template<ds4_sort_order order>
kernel void kernel_argsort_f32_i32(
constant ds4_metal_args_argsort & args,
device const char * src0,
device int32_t * dst,
threadgroup int32_t * shmem_i32 [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
// bitonic sort
const int col = tpitg[0];
const int ib = tgpig[0] / args.ne01;
const int i00 = ib*ntg.x;
const int i01 = tgpig[0] % args.ne01;
const int i02 = tgpig[1];
const int i03 = tgpig[2];
device const float * src0_row = (device const float *) (src0 + args.nb01*i01 + args.nb02*i02 + args.nb03*i03);
// initialize indices
shmem_i32[col] = i00 + col;
threadgroup_barrier(mem_flags::mem_threadgroup);
for (int k = 2; k <= ntg.x; k *= 2) {
for (int j = k / 2; j > 0; j /= 2) {
int ixj = col ^ j;
if (ixj > col) {
if ((col & k) == 0) {
if (shmem_i32[col] >= args.ne00 ||
(shmem_i32[ixj] < args.ne00 && (order == DS4_SORT_ORDER_ASC ?
src0_row[shmem_i32[col]] > src0_row[shmem_i32[ixj]] :
src0_row[shmem_i32[col]] < src0_row[shmem_i32[ixj]]))
) {
SWAP(shmem_i32[col], shmem_i32[ixj]);
}
} else {
if (shmem_i32[ixj] >= args.ne00 ||
(shmem_i32[col] < args.ne00 && (order == DS4_SORT_ORDER_ASC ?
src0_row[shmem_i32[col]] < src0_row[shmem_i32[ixj]] :
src0_row[shmem_i32[col]] > src0_row[shmem_i32[ixj]]))
) {
SWAP(shmem_i32[col], shmem_i32[ixj]);
}
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
}
const int64_t i0 = ib*args.top_k;
// copy the result to dst without the padding
if (i0 + col < args.ne0 && col < args.top_k) {
dst += i0 + args.ne0*i01 + args.ne0*args.ne1*i02 + args.ne0*args.ne1*args.ne2*i03;
dst[col] = shmem_i32[col];
}
}
// Host-visible sort variant used by DS4 top-k selection.
template [[host_name("kernel_argsort_f32_i32_desc")]] kernel argsort_t kernel_argsort_f32_i32<DS4_SORT_ORDER_DESC>;
typedef void (argsort_merge_t)(
constant ds4_metal_args_argsort_merge & args,
device const char * src0,
device const int32_t * tmp,
device int32_t * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]);
// Merges sorted index runs produced by kernel_argsort_f32_i32. In the DS4 graph
// this finishes top-k over router or compressed-attention score rows.
template<ds4_sort_order order>
kernel void kernel_argsort_merge_f32_i32(
constant ds4_metal_args_argsort_merge & args,
device const char * src0,
device const int32_t * tmp,
device int32_t * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
const int im = tgpig[0] / args.ne01;
const int i01 = tgpig[0] % args.ne01;
const int i02 = tgpig[1];
const int i03 = tgpig[2];
const int start = im * (2 * args.len);
const int len0 = MIN(args.len, MAX(0, args.ne0 - (int)(start)));
const int len1 = MIN(args.len, MAX(0, args.ne0 - (int)(start + args.len)));
const int total = len0 + len1;
device const int32_t * tmp0 = tmp + start
+ i01*args.ne0
+ i02*args.ne0*args.ne01
+ i03*args.ne0*args.ne01*args.ne02;
device const int32_t * tmp1 = tmp0 + args.len;
dst += start
+ i01*args.top_k
+ i02*args.top_k*args.ne01
+ i03*args.top_k*args.ne01*args.ne02;
device const float * src0_row = (device const float *)(src0
+ args.nb01*i01
+ args.nb02*i02
+ args.nb03*i03);
if (total == 0) {
return;
}
const int chunk = (total + ntg.x - 1) / ntg.x;
const int k0 = tpitg.x * chunk;
const int k1 = MIN(MIN(k0 + chunk, total), args.top_k);
if (k0 >= args.top_k) {
return;
}
if (k0 >= total) {
return;
}
int low = k0 > len1 ? k0 - len1 : 0;
int high = MIN(k0, len0);
// binary-search partition (i, j) such that i + j = k
while (low < high) {
const int mid = (low + high) >> 1;
const int32_t idx0 = tmp0[mid];
const int32_t idx1 = tmp1[k0 - mid - 1];
const float val0 = src0_row[idx0];
const float val1 = src0_row[idx1];
bool take_left;
if (order == DS4_SORT_ORDER_ASC) {
take_left = (val0 <= val1);
} else {
take_left = (val0 >= val1);
}
if (take_left) {
low = mid + 1;
} else {
high = mid;
}
}
int i = low;
int j = k0 - i;
// keep the merge fronts into registers
int32_t idx0 = 0;
float val0 = 0.0f;
if (i < len0) {
idx0 = tmp0[i];
val0 = src0_row[idx0];
}
int32_t idx1 = 0;
float val1 = 0.0f;
if (j < len1) {
idx1 = tmp1[j];
val1 = src0_row[idx1];
}
for (int k = k0; k < k1; ++k) {
int32_t out_idx;
if (i >= len0) {
while (k < k1) {
dst[k++] = tmp1[j++];
}
break;
} else if (j >= len1) {
while (k < k1) {
dst[k++] = tmp0[i++];
}
break;
} else {
bool take_left;
if (order == DS4_SORT_ORDER_ASC) {
take_left = (val0 <= val1);
} else {
take_left = (val0 >= val1);
}
if (take_left) {
out_idx = idx0;
++i;
if (i < len0) {
idx0 = tmp0[i];
val0 = src0_row[idx0];
}
} else {
out_idx = idx1;
++j;
if (j < len1) {
idx1 = tmp1[j];
val1 = src0_row[idx1];
}
}
}
dst[k] = out_idx;
}
}
// Host-visible merge variant used by DS4 top-k selection.
template [[host_name("kernel_argsort_merge_f32_i32_desc")]] kernel argsort_merge_t kernel_argsort_merge_f32_i32<DS4_SORT_ORDER_DESC>;
metal/bin.metal
+192
View File
@@ -0,0 +1,192 @@
struct ds4_metal_args_bin {
int32_t ne00;
int32_t ne01;
int32_t ne02;
int32_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne10;
int32_t ne11;
int32_t ne12;
int32_t ne13;
uint64_t nb10;
uint64_t nb11;
uint64_t nb12;
uint64_t nb13;
int32_t ne0;
int32_t ne1;
int32_t ne2;
int32_t ne3;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
uint64_t offs;
uint64_t o1[8];
};
constant short FC_bin_op [[function_constant(FC_BIN + 0)]];
constant short FC_bin_f [[function_constant(FC_BIN + 1)]];
constant bool FC_bin_rb [[function_constant(FC_BIN + 2)]];
constant bool FC_bin_cb [[function_constant(FC_BIN + 3)]];
// Generic binary elementwise op with compile-time operation and broadcast
// modes. DS4 currently instantiates this as add, multiply, scalar multiply, and
// row division in the static graph.
template <typename T0, typename T1, typename T>
kernel void kernel_bin_fuse_impl(
constant ds4_metal_args_bin & args,
device const char * src0,
device const char * src1,
device char * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
#define FC_OP FC_bin_op
#define FC_F FC_bin_f
#define FC_RB FC_bin_rb
#define FC_CB FC_bin_cb
if (FC_RB) {
const uint i0 = tgpig.y*args.ne00 + tgpig.x;
const uint i1 = FC_CB ? tgpig.x%args.ne10 : tgpig.x;
device const T0 * src0_row = (device const T0 *) (src0);
device T * dst_row = (device T *) (dst);
if (FC_F == 1) {
device const T1 * src1_row = (device const T1 *) (src1 + args.o1[0]);
if (FC_OP == 0) {
dst_row[i0] = src0_row[i0] + src1_row[i1];
}
if (FC_OP == 1) {
dst_row[i0] = src0_row[i0] - src1_row[i1];
}
if (FC_OP == 2) {
dst_row[i0] = src0_row[i0] * src1_row[i1];
}
if (FC_OP == 3) {
dst_row[i0] = src0_row[i0] / src1_row[i1];
}
} else {
T0 res = src0_row[i0];
if (FC_OP == 0) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res += ((device const T1 *) (src1 + args.o1[j]))[i1];
}
}
if (FC_OP == 1) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res -= ((device const T1 *) (src1 + args.o1[j]))[i1];
}
}
if (FC_OP == 2) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res *= ((device const T1 *) (src1 + args.o1[j]))[i1];
}
}
if (FC_OP == 3) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res /= ((device const T1 *) (src1 + args.o1[j]))[i1];
}
}
dst_row[i0] = res;
}
} else {
const int i03 = tgpig.z;
const int i02 = tgpig.y;
const int i01 = tgpig.x;
if (i01 >= args.ne01) {
return;
}
const int i13 = i03%args.ne13;
const int i12 = i02%args.ne12;
const int i11 = i01%args.ne11;
device const T0 * src0_ptr = (device const T0 *) (src0 + i03*args.nb03 + i02*args.nb02 + i01*args.nb01 + args.offs);
device T * dst_ptr = (device T *) (dst + i03*args.nb3 + i02*args.nb2 + i01*args.nb1 + args.offs);
if (FC_F == 1) {
device const T1 * src1_ptr = (device const T1 *) (src1 + args.o1[0] + i13*args.nb13 + i12*args.nb12 + i11*args.nb11);
for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
const int i10 = FC_CB ? i0%args.ne10 : i0;
if (FC_OP == 0) {
dst_ptr[i0] = src0_ptr[i0] + src1_ptr[i10];
}
if (FC_OP == 1) {
dst_ptr[i0] = src0_ptr[i0] - src1_ptr[i10];
}
if (FC_OP == 2) {
dst_ptr[i0] = src0_ptr[i0] * src1_ptr[i10];
}
if (FC_OP == 3) {
dst_ptr[i0] = src0_ptr[i0] / src1_ptr[i10];
}
}
} else {
device const T1 * src1_ptr[8];
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
src1_ptr[j] = (device const T1 *) (src1 + args.o1[j] + i13*args.nb13 + i12*args.nb12 + i11*args.nb11);
}
for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
const int i10 = FC_CB ? i0%args.ne10 : i0;
T res = src0_ptr[i0];
if (FC_OP == 0) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res += src1_ptr[j][i10];
}
}
if (FC_OP == 1) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res -= src1_ptr[j][i10];
}
}
if (FC_OP == 2) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res *= src1_ptr[j][i10];
}
}
if (FC_OP == 3) {
FOR_UNROLL (short j = 0; j < FC_F; ++j) {
res /= src1_ptr[j][i10];
}
}
dst_ptr[i0] = res;
}
}
}
#undef FC_OP
#undef FC_F
#undef FC_RB
#undef FC_CB
}
typedef decltype(kernel_bin_fuse_impl<float, float, float>) kernel_bin_fuse_t;
// Host-visible F32 binary op; function constants specialize it per use site.
template [[host_name("kernel_bin_fuse_f32_f32_f32")]] kernel kernel_bin_fuse_t kernel_bin_fuse_impl<float, float, float>;
metal/concat.metal
+62
View File
@@ -0,0 +1,62 @@
// DS4 Metal concat kernel used by the graph.
struct ds4_metal_args_concat {
int32_t ne00;
int32_t ne01;
int32_t ne02;
int32_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne10;
int32_t ne11;
int32_t ne12;
int32_t ne13;
uint64_t nb10;
uint64_t nb11;
uint64_t nb12;
uint64_t nb13;
int32_t ne0;
int32_t ne1;
int32_t ne2;
int32_t ne3;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
int32_t dim;
};
// Concatenates two float tensors along one dimension. In DS4 this is a graph
// utility for assembling attention inputs with exactly the same tensor layout
// expected by the downstream kernels.
kernel void kernel_concat(
constant ds4_metal_args_concat & args,
device const char * src0,
device const char * src1,
device char * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
const int i3 = tgpig.z;
const int i2 = tgpig.y;
const int i1 = tgpig.x;
int o[4] = {0, 0, 0, 0};
o[args.dim] = args.dim == 0 ? args.ne00 : (args.dim == 1 ? args.ne01 : (args.dim == 2 ? args.ne02 : args.ne03));
device const float * x;
for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
if (i0 < args.ne00 && i1 < args.ne01 && i2 < args.ne02 && i3 < args.ne03) {
x = (device const float *)(src0 + (i3 )*args.nb03 + (i2 )*args.nb02 + (i1 )*args.nb01 + (i0 )*args.nb00);
} else {
x = (device const float *)(src1 + (i3 - o[3])*args.nb13 + (i2 - o[2])*args.nb12 + (i1 - o[1])*args.nb11 + (i0 - o[0])*args.nb10);
}
device float * y = (device float *)(dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + i0*args.nb0);
*y = *x;
}
}
metal/cpy.metal
+57
View File
@@ -0,0 +1,57 @@
struct ds4_metal_args_cpy {
int64_t nk0;
int64_t ne00;
int64_t ne01;
int64_t ne02;
int64_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int64_t ne0;
int64_t ne1;
int64_t ne2;
int64_t ne3;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
};
// Typed copy/conversion between graph tensors. DS4 uses this for layout
// materialization and F32/F16 conversions at graph boundaries such as KV/cache
// packing and compressor pooling.
template<typename T0, typename T1>
kernel void kernel_cpy_t_t(
constant ds4_metal_args_cpy & args,
device const char * src0,
device char * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiitg[[thread_index_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
const int i03 = tgpig[2];
const int i02 = tgpig[1];
const int i01 = ntg[1] == 1 ? tgpig[0]%args.ne01 : tgpig[0]*ntg[1] + tiitg/ntg[0];
const int iw0 = ntg[1] == 1 ? tgpig[0]/args.ne01 : 0;
const int64_t n = i03*args.ne02*args.ne01*args.ne00 + i02*args.ne01*args.ne00 + i01*args.ne00;
const int64_t i3 = n/(args.ne2*args.ne1*args.ne0);
const int64_t i2 = (n - i3*args.ne2*args.ne1*args.ne0)/(args.ne1*args.ne0);
const int64_t i1 = (n - i3*args.ne2*args.ne1*args.ne0 - i2*args.ne1*args.ne0)/args.ne0;
const int64_t i0 = (n - i3*args.ne2*args.ne1*args.ne0 - i2*args.ne1*args.ne0 - i1*args.ne0);
device T1 * dst_data = (device T1 *) (dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1 + i0*args.nb0);
for (int64_t i00 = iw0*ntg[0] + tiitg%ntg[0]; i00 < args.ne00; ) {
device const T0 * src = (device T0 *)(src0 + i03*args.nb03 + i02*args.nb02 + i01*args.nb01 + i00*args.nb00);
dst_data[i00] = (T1) src[0];
break;
}
}
typedef decltype(kernel_cpy_t_t<float, float>) kernel_cpy_t;
// Host-visible copy/conversion variants used by the DS4 graph.
template [[host_name("kernel_cpy_f32_f32")]] kernel kernel_cpy_t kernel_cpy_t_t<float, float>;
template [[host_name("kernel_cpy_f32_f16")]] kernel kernel_cpy_t kernel_cpy_t_t<float, half>;
template [[host_name("kernel_cpy_f16_f32")]] kernel kernel_cpy_t kernel_cpy_t_t<half, float>;
metal/dense.metal
+1121
View File
File diff suppressed because it is too large Load Diff
metal/dsv4_hc.metal
+861
View File
@@ -0,0 +1,861 @@
struct ds4_metal_args_dsv4_hc_split_sinkhorn {
int32_t n_hc;
int32_t sinkhorn_iters;
int64_t n_rows;
int64_t mix_hc;
uint64_t nb01;
uint64_t nb1;
float eps;
};
struct ds4_metal_args_dsv4_hc_weighted_sum {
int64_t n_embd;
int64_t n_hc;
int64_t n_tokens;
uint64_t nb_x0;
uint64_t nb_x1;
uint64_t nb_x2;
uint64_t nb_w0;
uint64_t nb_w1;
uint64_t nb0;
uint64_t nb1;
};
struct ds4_metal_args_dsv4_hc_split_weighted_sum {
int64_t n_embd;
int32_t n_hc;
int32_t sinkhorn_iters;
int64_t n_rows;
int64_t mix_hc;
uint64_t nb_mix1;
uint64_t nb_split1;
uint64_t nb_x0;
uint64_t nb_x1;
uint64_t nb_x2;
uint64_t nb0;
uint64_t nb1;
float eps;
};
struct ds4_metal_args_dsv4_hc_split_weighted_sum_norm {
int64_t n_embd;
int32_t n_hc;
int32_t sinkhorn_iters;
int64_t n_rows;
int64_t mix_hc;
uint64_t nb_mix1;
uint64_t nb_split1;
uint64_t nb_x0;
uint64_t nb_x1;
uint64_t nb_x2;
uint64_t nb0;
uint64_t nb1;
uint64_t nb_norm1;
float eps;
float norm_eps;
};
struct ds4_metal_args_dsv4_hc_expand {
int64_t n_embd;
int64_t n_hc;
int64_t n_tokens;
uint64_t nb_block0;
uint64_t nb_block1;
uint64_t nb_add0;
uint64_t nb_add1;
uint64_t nb_res0;
uint64_t nb_res1;
uint64_t nb_res2;
uint64_t nb_post0;
uint64_t nb_post1;
uint64_t nb_comb0;
uint64_t nb_comb1;
uint64_t nb_comb2;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
int32_t has_add;
};
// Splits an HC mixer row into pre weights, post gates, and the HC-to-HC
// combination matrix. The 4-channel path is specialized because DS4 Flash uses
// HC=4 in normal inference, while the scalar fallback keeps diagnostics usable.
kernel void kernel_dsv4_hc_split_sinkhorn(
constant ds4_metal_args_dsv4_hc_split_sinkhorn & args,
device const float * mixes,
device const float * scale,
device const float * base,
device float * dst,
uint tid [[thread_position_in_grid]]) {
if ((int64_t) tid >= args.n_rows) {
return;
}
constexpr int HC_MAX = 16;
const int HC = args.n_hc;
if (HC <= 0 || HC > HC_MAX) {
return;
}
device const float * mix = mixes + ((int64_t) tid)*args.mix_hc;
device float * out = dst + ((int64_t) tid)*args.mix_hc;
const float epsv = args.eps;
const float pre_scale = scale[0];
const float post_scale = scale[1];
const float comb_scale = scale[2];
if (HC == 4) {
const float4 pre_z =
*((device const float4 *) mix) * pre_scale +
*((device const float4 *) base);
*((device float4 *) out) = 1.0f / (1.0f + exp(-pre_z)) + epsv;
const float4 post_z =
*((device const float4 *) (mix + 4)) * post_scale +
*((device const float4 *) (base + 4));
*((device float4 *) (out + 4)) = 2.0f / (1.0f + exp(-post_z));
float4 r0 =
*((device const float4 *) (mix + 8)) * comb_scale +
*((device const float4 *) (base + 8));
float4 r1 =
*((device const float4 *) (mix + 12)) * comb_scale +
*((device const float4 *) (base + 12));
float4 r2 =
*((device const float4 *) (mix + 16)) * comb_scale +
*((device const float4 *) (base + 16));
float4 r3 =
*((device const float4 *) (mix + 20)) * comb_scale +
*((device const float4 *) (base + 20));
const float m0 = max(max(r0.x, r0.y), max(r0.z, r0.w));
const float m1 = max(max(r1.x, r1.y), max(r1.z, r1.w));
const float m2 = max(max(r2.x, r2.y), max(r2.z, r2.w));
const float m3 = max(max(r3.x, r3.y), max(r3.z, r3.w));
r0 = exp(r0 - m0);
r1 = exp(r1 - m1);
r2 = exp(r2 - m2);
r3 = exp(r3 - m3);
r0 = r0 * (1.0f / (r0.x + r0.y + r0.z + r0.w)) + epsv;
r1 = r1 * (1.0f / (r1.x + r1.y + r1.z + r1.w)) + epsv;
r2 = r2 * (1.0f / (r2.x + r2.y + r2.z + r2.w)) + epsv;
r3 = r3 * (1.0f / (r3.x + r3.y + r3.z + r3.w)) + epsv;
float4 col_inv = 1.0f / (r0 + r1 + r2 + r3 + epsv);
r0 *= col_inv;
r1 *= col_inv;
r2 *= col_inv;
r3 *= col_inv;
for (int iter = 1; iter < args.sinkhorn_iters; ++iter) {
r0 *= 1.0f / (r0.x + r0.y + r0.z + r0.w + epsv);
r1 *= 1.0f / (r1.x + r1.y + r1.z + r1.w + epsv);
r2 *= 1.0f / (r2.x + r2.y + r2.z + r2.w + epsv);
r3 *= 1.0f / (r3.x + r3.y + r3.z + r3.w + epsv);
col_inv = 1.0f / (r0 + r1 + r2 + r3 + epsv);
r0 *= col_inv;
r1 *= col_inv;
r2 *= col_inv;
r3 *= col_inv;
}
*((device float4 *) (out + 8)) = r0;
*((device float4 *) (out + 12)) = r1;
*((device float4 *) (out + 16)) = r2;
*((device float4 *) (out + 20)) = r3;
return;
}
for (int i = 0; i < HC; ++i) {
const float z = mix[i] * pre_scale + base[i];
out[i] = 1.0f / (1.0f + exp(-z)) + epsv;
}
for (int i = 0; i < HC; ++i) {
const int off = HC + i;
const float z = mix[off] * post_scale + base[off];
out[off] = 2.0f / (1.0f + exp(-z));
}
float c[HC_MAX*HC_MAX];
for (int dst_hc = 0; dst_hc < HC; ++dst_hc) {
float row_max = -INFINITY;
for (int src_hc = 0; src_hc < HC; ++src_hc) {
const int idx = src_hc + dst_hc*HC;
const int off = 2*HC + idx;
const float v = mix[off] * comb_scale + base[off];
c[idx] = v;
row_max = max(row_max, v);
}
float row_sum = 0.0f;
for (int src_hc = 0; src_hc < HC; ++src_hc) {
const int idx = src_hc + dst_hc*HC;
const float v = exp(c[idx] - row_max);
c[idx] = v;
row_sum += v;
}
const float inv_sum = 1.0f / row_sum;
for (int src_hc = 0; src_hc < HC; ++src_hc) {
const int idx = src_hc + dst_hc*HC;
c[idx] = c[idx] * inv_sum + epsv;
}
}
for (int src_hc = 0; src_hc < HC; ++src_hc) {
float sum = 0.0f;
for (int dst_hc = 0; dst_hc < HC; ++dst_hc) {
sum += c[src_hc + dst_hc*HC];
}
const float inv_denom = 1.0f / (sum + epsv);
for (int dst_hc = 0; dst_hc < HC; ++dst_hc) {
c[src_hc + dst_hc*HC] *= inv_denom;
}
}
for (int iter = 1; iter < args.sinkhorn_iters; ++iter) {
for (int dst_hc = 0; dst_hc < HC; ++dst_hc) {
float sum = 0.0f;
for (int src_hc = 0; src_hc < HC; ++src_hc) {
sum += c[src_hc + dst_hc*HC];
}
const float inv_denom = 1.0f / (sum + epsv);
for (int src_hc = 0; src_hc < HC; ++src_hc) {
c[src_hc + dst_hc*HC] *= inv_denom;
}
}
for (int src_hc = 0; src_hc < HC; ++src_hc) {
float sum = 0.0f;
for (int dst_hc = 0; dst_hc < HC; ++dst_hc) {
sum += c[src_hc + dst_hc*HC];
}
const float inv_denom = 1.0f / (sum + epsv);
for (int dst_hc = 0; dst_hc < HC; ++dst_hc) {
c[src_hc + dst_hc*HC] *= inv_denom;
}
}
}
for (int i = 0; i < HC*HC; ++i) {
out[2*HC + i] = c[i];
}
}
// Decode-side fusion of HC split and pre-weighted HC reduction. One threadgroup
// handles one token row: lane 0 computes the HC=4 mixer split once, stores the
// post/comb data for the following HC expand, and all lanes reuse the pre
// weights from threadgroup memory to produce the embedding row.
kernel void kernel_dsv4_hc_split_weighted_sum(
constant ds4_metal_args_dsv4_hc_split_weighted_sum & args,
device const char * mixes,
device const float * scale,
device const float * base,
device const char * x,
device char * split,
device char * dst,
threadgroup float * pre_shmem [[threadgroup(0)]],
uint row [[threadgroup_position_in_grid]],
uint tid [[thread_position_in_threadgroup]],
uint ntg [[threads_per_threadgroup]]) {
if ((int64_t) row >= args.n_rows || args.n_hc != 4) {
return;
}
device const float * mix = (device const float *) (mixes + (uint64_t)row*args.nb_mix1);
device float * out = (device float *) (split + (uint64_t)row*args.nb_split1);
if (tid == 0) {
const float epsv = args.eps;
const float pre_scale = scale[0];
const float post_scale = scale[1];
const float comb_scale = scale[2];
const float4 pre_z =
*((device const float4 *) mix) * pre_scale +
*((device const float4 *) base);
const float4 pre = 1.0f / (1.0f + exp(-pre_z)) + epsv;
*((device float4 *) out) = pre;
pre_shmem[0] = pre.x;
pre_shmem[1] = pre.y;
pre_shmem[2] = pre.z;
pre_shmem[3] = pre.w;
const float4 post_z =
*((device const float4 *) (mix + 4)) * post_scale +
*((device const float4 *) (base + 4));
*((device float4 *) (out + 4)) = 2.0f / (1.0f + exp(-post_z));
float4 r0 =
*((device const float4 *) (mix + 8)) * comb_scale +
*((device const float4 *) (base + 8));
float4 r1 =
*((device const float4 *) (mix + 12)) * comb_scale +
*((device const float4 *) (base + 12));
float4 r2 =
*((device const float4 *) (mix + 16)) * comb_scale +
*((device const float4 *) (base + 16));
float4 r3 =
*((device const float4 *) (mix + 20)) * comb_scale +
*((device const float4 *) (base + 20));
const float m0 = max(max(r0.x, r0.y), max(r0.z, r0.w));
const float m1 = max(max(r1.x, r1.y), max(r1.z, r1.w));
const float m2 = max(max(r2.x, r2.y), max(r2.z, r2.w));
const float m3 = max(max(r3.x, r3.y), max(r3.z, r3.w));
r0 = exp(r0 - m0);
r1 = exp(r1 - m1);
r2 = exp(r2 - m2);
r3 = exp(r3 - m3);
r0 = r0 * (1.0f / (r0.x + r0.y + r0.z + r0.w)) + epsv;
r1 = r1 * (1.0f / (r1.x + r1.y + r1.z + r1.w)) + epsv;
r2 = r2 * (1.0f / (r2.x + r2.y + r2.z + r2.w)) + epsv;
r3 = r3 * (1.0f / (r3.x + r3.y + r3.z + r3.w)) + epsv;
float4 col_inv = 1.0f / (r0 + r1 + r2 + r3 + epsv);
r0 *= col_inv;
r1 *= col_inv;
r2 *= col_inv;
r3 *= col_inv;
for (int iter = 1; iter < args.sinkhorn_iters; ++iter) {
r0 *= 1.0f / (r0.x + r0.y + r0.z + r0.w + epsv);
r1 *= 1.0f / (r1.x + r1.y + r1.z + r1.w + epsv);
r2 *= 1.0f / (r2.x + r2.y + r2.z + r2.w + epsv);
r3 *= 1.0f / (r3.x + r3.y + r3.z + r3.w + epsv);
col_inv = 1.0f / (r0 + r1 + r2 + r3 + epsv);
r0 *= col_inv;
r1 *= col_inv;
r2 *= col_inv;
r3 *= col_inv;
}
*((device float4 *) (out + 8)) = r0;
*((device float4 *) (out + 12)) = r1;
*((device float4 *) (out + 16)) = r2;
*((device float4 *) (out + 20)) = r3;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
for (int64_t d = tid; d < args.n_embd; d += ntg) {
float acc = 0.0f;
acc += *((device const float *) (x + d*args.nb_x0 + 0*args.nb_x1 + (uint64_t)row*args.nb_x2)) * pre_shmem[0];
acc += *((device const float *) (x + d*args.nb_x0 + 1*args.nb_x1 + (uint64_t)row*args.nb_x2)) * pre_shmem[1];
acc += *((device const float *) (x + d*args.nb_x0 + 2*args.nb_x1 + (uint64_t)row*args.nb_x2)) * pre_shmem[2];
acc += *((device const float *) (x + d*args.nb_x0 + 3*args.nb_x1 + (uint64_t)row*args.nb_x2)) * pre_shmem[3];
*((device float *) (dst + d*args.nb0 + (uint64_t)row*args.nb1)) = acc;
}
}
// Decode HC-pre plus the following RMSNorm. DS4 always uses HC=4 and a
// 4096-wide sublayer row. The normal release path computes HC coefficients,
// collapses four residual streams into that row, then immediately launches a
// weighted RMSNorm over the row. This kernel keeps the HC split math identical
// to kernel_dsv4_hc_split_weighted_sum, stores the HC-pre row for diagnostics,
// and reuses the just-collapsed values from threadgroup memory for the RMSNorm
// reduction. The reduction mirrors kernel_rms_norm_mul_f32_4's 1024-thread
// float4 shape for a 4096-wide row.
kernel void kernel_dsv4_hc_split_weighted_sum_norm4(
constant ds4_metal_args_dsv4_hc_split_weighted_sum_norm & args,
device const char * mixes,
device const float * scale,
device const float * base,
device const char * x,
device char * split,
device char * dst,
device const char * norm_weight,
device char * norm_dst,
threadgroup float * shared [[threadgroup(0)]],
uint row [[threadgroup_position_in_grid]],
ushort tid [[thread_position_in_threadgroup]],
ushort sgitg [[simdgroup_index_in_threadgroup]],
ushort tiisg [[thread_index_in_simdgroup]],
ushort ntg [[threads_per_threadgroup]]) {
if ((int64_t)row >= args.n_rows || args.n_hc != 4 || args.n_embd != 4096) {
return;
}
threadgroup float4 *row_shmem = (threadgroup float4 *)shared;
threadgroup float *pre_shmem = shared + 4096;
threadgroup float *sum_shmem = pre_shmem + 4;
device const float *mix = (device const float *)(mixes + (uint64_t)row * args.nb_mix1);
device float *out = (device float *)(split + (uint64_t)row * args.nb_split1);
if (sgitg == 0) {
sum_shmem[tiisg] = 0.0f;
}
if (tid == 0) {
const float epsv = args.eps;
const float pre_scale = scale[0];
const float post_scale = scale[1];
const float comb_scale = scale[2];
const float4 pre_z =
*((device const float4 *)mix) * pre_scale +
*((device const float4 *)base);
const float4 pre = 1.0f / (1.0f + exp(-pre_z)) + epsv;
*((device float4 *)out) = pre;
pre_shmem[0] = pre.x;
pre_shmem[1] = pre.y;
pre_shmem[2] = pre.z;
pre_shmem[3] = pre.w;
const float4 post_z =
*((device const float4 *)(mix + 4)) * post_scale +
*((device const float4 *)(base + 4));
*((device float4 *)(out + 4)) = 2.0f / (1.0f + exp(-post_z));
float4 r0 =
*((device const float4 *)(mix + 8)) * comb_scale +
*((device const float4 *)(base + 8));
float4 r1 =
*((device const float4 *)(mix + 12)) * comb_scale +
*((device const float4 *)(base + 12));
float4 r2 =
*((device const float4 *)(mix + 16)) * comb_scale +
*((device const float4 *)(base + 16));
float4 r3 =
*((device const float4 *)(mix + 20)) * comb_scale +
*((device const float4 *)(base + 20));
const float m0 = max(max(r0.x, r0.y), max(r0.z, r0.w));
const float m1 = max(max(r1.x, r1.y), max(r1.z, r1.w));
const float m2 = max(max(r2.x, r2.y), max(r2.z, r2.w));
const float m3 = max(max(r3.x, r3.y), max(r3.z, r3.w));
r0 = exp(r0 - m0);
r1 = exp(r1 - m1);
r2 = exp(r2 - m2);
r3 = exp(r3 - m3);
r0 = r0 * (1.0f / (r0.x + r0.y + r0.z + r0.w)) + epsv;
r1 = r1 * (1.0f / (r1.x + r1.y + r1.z + r1.w)) + epsv;
r2 = r2 * (1.0f / (r2.x + r2.y + r2.z + r2.w)) + epsv;
r3 = r3 * (1.0f / (r3.x + r3.y + r3.z + r3.w)) + epsv;
float4 col_inv = 1.0f / (r0 + r1 + r2 + r3 + epsv);
r0 *= col_inv;
r1 *= col_inv;
r2 *= col_inv;
r3 *= col_inv;
for (int iter = 1; iter < args.sinkhorn_iters; ++iter) {
r0 *= 1.0f / (r0.x + r0.y + r0.z + r0.w + epsv);
r1 *= 1.0f / (r1.x + r1.y + r1.z + r1.w + epsv);
r2 *= 1.0f / (r2.x + r2.y + r2.z + r2.w + epsv);
r3 *= 1.0f / (r3.x + r3.y + r3.z + r3.w + epsv);
col_inv = 1.0f / (r0 + r1 + r2 + r3 + epsv);
r0 *= col_inv;
r1 *= col_inv;
r2 *= col_inv;
r3 *= col_inv;
}
*((device float4 *)(out + 8)) = r0;
*((device float4 *)(out + 12)) = r1;
*((device float4 *)(out + 16)) = r2;
*((device float4 *)(out + 20)) = r3;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
float sumf = 0.0f;
const uint n4 = 1024u;
for (uint i = tid; i < n4; i += ntg) {
device const float4 *x0 = (device const float4 *)(x + 0 * args.nb_x1 + (uint64_t)row * args.nb_x2);
device const float4 *x1 = (device const float4 *)(x + 1 * args.nb_x1 + (uint64_t)row * args.nb_x2);
device const float4 *x2 = (device const float4 *)(x + 2 * args.nb_x1 + (uint64_t)row * args.nb_x2);
device const float4 *x3 = (device const float4 *)(x + 3 * args.nb_x1 + (uint64_t)row * args.nb_x2);
const float4 v = x0[i] * pre_shmem[0] +
x1[i] * pre_shmem[1] +
x2[i] * pre_shmem[2] +
x3[i] * pre_shmem[3];
row_shmem[i] = v;
sumf += dot(v, v);
}
sumf = simd_sum(sumf);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
sum_shmem[sgitg] = sumf;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = sum_shmem[tiisg];
sumf = simd_sum(sumf);
const float norm_scale = rsqrt(sumf / 4096.0f + args.norm_eps);
device float4 *dst4 = (device float4 *)(dst + (uint64_t)row * args.nb1);
device const float4 *w4 = (device const float4 *)norm_weight;
device float4 *norm4 = (device float4 *)(norm_dst + (uint64_t)row * args.nb_norm1);
for (uint i = tid; i < n4; i += ntg) {
const float4 v = row_shmem[i];
dst4[i] = v;
norm4[i] = (v * norm_scale) * w4[i];
}
}
// Expands an embedding-sized block back into HC channels after attention/FFN.
// The post gate scales the current block, while the Sinkhorn combination matrix
// mixes residual HC channels from the previous state.
kernel void kernel_dsv4_hc_expand(
constant ds4_metal_args_dsv4_hc_expand & args,
device const char * block_out,
device const char * residual,
device const char * post,
device const char * comb,
device const char * block_add,
device char * dst,
uint gid [[thread_position_in_grid]]) {
const int64_t n_elem = args.n_embd * args.n_hc * args.n_tokens;
if ((int64_t) gid >= n_elem) {
return;
}
const int64_t d = ((int64_t) gid) % args.n_embd;
const int64_t tmp = ((int64_t) gid) / args.n_embd;
const int64_t dst_hc = tmp % args.n_hc;
const int64_t t = tmp / args.n_hc;
float block_v = *((device const float *) (block_out + d*args.nb_block0 + t*args.nb_block1));
if (args.has_add) {
block_v += *((device const float *) (block_add + d*args.nb_add0 + t*args.nb_add1));
}
const float post_v = *((device const float *) (post + dst_hc*args.nb_post0 + t*args.nb_post1));
float acc = block_v * post_v;
for (int64_t src_hc = 0; src_hc < args.n_hc; ++src_hc) {
const float comb_v = *((device const float *) (comb + dst_hc*args.nb_comb0 + src_hc*args.nb_comb1 + t*args.nb_comb2));
const float res_v = *((device const float *) (residual + d*args.nb_res0 + src_hc*args.nb_res1 + t*args.nb_res2));
acc += comb_v * res_v;
}
*((device float *) (dst + d*args.nb0 + dst_hc*args.nb1 + t*args.nb2)) = acc;
}
// HC=4 specialization of the post/expand step. One thread computes all four
// destination HC streams for one token/dimension, reusing the same block output
// and residual HC values while preserving the per-stream accumulation order.
kernel void kernel_dsv4_hc_expand4(
constant ds4_metal_args_dsv4_hc_expand & args,
device const char * block_out,
device const char * residual,
device const char * post,
device const char * comb,
device const char * block_add,
device char * dst,
uint gid [[thread_position_in_grid]]) {
if (args.n_hc != 4) {
return;
}
const int64_t n_elem = args.n_embd * args.n_tokens;
if ((int64_t) gid >= n_elem) {
return;
}
const int64_t d = ((int64_t) gid) % args.n_embd;
const int64_t t = ((int64_t) gid) / args.n_embd;
float block_v = *((device const float *) (block_out + d*args.nb_block0 + t*args.nb_block1));
if (args.has_add) {
block_v += *((device const float *) (block_add + d*args.nb_add0 + t*args.nb_add1));
}
const float r0 = *((device const float *) (residual + d*args.nb_res0 + 0*args.nb_res1 + t*args.nb_res2));
const float r1 = *((device const float *) (residual + d*args.nb_res0 + 1*args.nb_res1 + t*args.nb_res2));
const float r2 = *((device const float *) (residual + d*args.nb_res0 + 2*args.nb_res1 + t*args.nb_res2));
const float r3 = *((device const float *) (residual + d*args.nb_res0 + 3*args.nb_res1 + t*args.nb_res2));
for (int64_t dst_hc = 0; dst_hc < 4; ++dst_hc) {
float acc = block_v * *((device const float *) (post + dst_hc*args.nb_post0 + t*args.nb_post1));
acc += *((device const float *) (comb + dst_hc*args.nb_comb0 + 0*args.nb_comb1 + t*args.nb_comb2)) * r0;
acc += *((device const float *) (comb + dst_hc*args.nb_comb0 + 1*args.nb_comb1 + t*args.nb_comb2)) * r1;
acc += *((device const float *) (comb + dst_hc*args.nb_comb0 + 2*args.nb_comb1 + t*args.nb_comb2)) * r2;
acc += *((device const float *) (comb + dst_hc*args.nb_comb0 + 3*args.nb_comb1 + t*args.nb_comb2)) * r3;
*((device float *) (dst + d*args.nb0 + dst_hc*args.nb1 + t*args.nb2)) = acc;
}
}
// Decode-time FFN tail fusion:
//
// shared_out = shared_mid @ Wshared_down
// after_ffn_hc = HCPost(routed_out + shared_out, residual_hc, split)
//
// The Q8_0 dot reduction is intentionally copied from the normal matvec shape
// so the shared expert result is bit-identical. The only specialization is
// that DS4 decode has one token and HC=4, so the thread that finishes each
// shared-down output row can immediately expand it into the four HC streams.
kernel void kernel_dsv4_shared_down_hc_expand4_q8_0(
constant ds4_metal_args_mul_mv & mv,
constant ds4_metal_args_dsv4_hc_expand & hc,
device const char * weight,
device const char * shared_mid,
device char * shared_out,
device const char * routed_out,
device const char * residual,
device const char * post,
device const char * comb,
device char * dst,
threadgroup char * shmem [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
if (hc.n_hc != 4 || hc.n_tokens != 1) {
return;
}
const short NSG = FC_mul_mv_nsg;
constexpr short NW = N_SIMDWIDTH;
constexpr short NQ = 8;
constexpr short NR0 = N_R0_Q8_0;
const int nb = mv.ne00 / QK8_0;
const int row0 = tgpig.x * NR0;
const short ix = tiisg / (NW / NQ);
const short il = tiisg % (NW / NQ);
const int ib0 = sgitg * NQ + ix;
device const float *y = (device const float *)(shared_mid);
device const float *yb = y + ib0 * QK8_0 + il * NQ;
device const block_q8_0 *ax[NR0];
FOR_UNROLL(short row = 0; row < NR0; ++row) {
const uint64_t off0 = (uint64_t)(row0 + row) * mv.nb01;
ax[row] = (device const block_q8_0 *)(weight + off0);
}
float sumf[NR0] = { 0.0f };
float yl[NQ];
for (int ib = ib0; ib < nb; ib += NSG * NQ) {
FOR_UNROLL(short i = 0; i < NQ; ++i) {
yl[i] = yb[i];
}
FOR_UNROLL(short row = 0; row < NR0; ++row) {
device const int8_t *qs = ax[row][ib].qs + il * NQ;
float sumq = 0.0f;
FOR_UNROLL(short i = 0; i < NQ; ++i) {
sumq += qs[i] * yl[i];
}
sumf[row] += sumq * ax[row][ib].d;
}
yb += NSG * NQ * QK8_0;
}
threadgroup float *shmem_f32[NR0];
FOR_UNROLL(short row = 0; row < NR0; ++row) {
shmem_f32[row] = (threadgroup float *)shmem + NW * row;
if (sgitg == 0) {
shmem_f32[row][tiisg] = 0.0f;
}
sumf[row] = simd_sum(sumf[row]);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
FOR_UNROLL(short row = 0; row < NR0; ++row) {
if (tiisg == 0) {
shmem_f32[row][sgitg] = sumf[row];
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
FOR_UNROLL(short row = 0; row < NR0; ++row) {
const int d = row0 + row;
if (d >= mv.ne01) {
continue;
}
const float shared_v = simd_sum(shmem_f32[row][tiisg]);
if (tiisg == 0 && sgitg == 0) {
*((device float *)(shared_out + (uint64_t)d * sizeof(float))) = shared_v;
float block_v = *((device const float *)(routed_out + (uint64_t)d * hc.nb_block0));
block_v += shared_v;
const float r0 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 0 * hc.nb_res1));
const float r1 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 1 * hc.nb_res1));
const float r2 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 2 * hc.nb_res1));
const float r3 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 3 * hc.nb_res1));
for (int64_t dst_hc = 0; dst_hc < 4; ++dst_hc) {
float acc = block_v * *((device const float *)(post + dst_hc * hc.nb_post0));
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 0 * hc.nb_comb1)) * r0;
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 1 * hc.nb_comb1)) * r1;
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 2 * hc.nb_comb1)) * r2;
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 3 * hc.nb_comb1)) * r3;
*((device float *)(dst + (uint64_t)d * hc.nb0 + dst_hc * hc.nb1)) = acc;
}
}
}
}
// Decode-time attention output tail fusion:
//
// attn_out = attn_low @ Wob
// after_attn_hc = HCPost(attn_out, residual_hc, split)
//
// This is the no-add sibling of the shared-down/FFN fusion above. It preserves
// the exact Q8_0 matvec reduction, stores `attn_out` for diagnostics, and then
// writes the four HC streams for the same embedding dimension.
kernel void kernel_dsv4_q8_hc_expand4_q8_0(
constant ds4_metal_args_mul_mv & mv,
constant ds4_metal_args_dsv4_hc_expand & hc,
device const char * weight,
device const char * input,
device char * block_out,
device const char * residual,
device const char * post,
device const char * comb,
device char * dst,
threadgroup char * shmem [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
if (hc.n_hc != 4 || hc.n_tokens != 1) {
return;
}
const short NSG = FC_mul_mv_nsg;
constexpr short NW = N_SIMDWIDTH;
constexpr short NQ = 8;
constexpr short NR0 = N_R0_Q8_0;
const int nb = mv.ne00 / QK8_0;
const int row0 = tgpig.x * NR0;
const short ix = tiisg / (NW / NQ);
const short il = tiisg % (NW / NQ);
const int ib0 = sgitg * NQ + ix;
device const float *y = (device const float *)(input);
device const float *yb = y + ib0 * QK8_0 + il * NQ;
device const block_q8_0 *ax[NR0];
FOR_UNROLL(short row = 0; row < NR0; ++row) {
const uint64_t off0 = (uint64_t)(row0 + row) * mv.nb01;
ax[row] = (device const block_q8_0 *)(weight + off0);
}
float sumf[NR0] = { 0.0f };
float yl[NQ];
for (int ib = ib0; ib < nb; ib += NSG * NQ) {
FOR_UNROLL(short i = 0; i < NQ; ++i) {
yl[i] = yb[i];
}
FOR_UNROLL(short row = 0; row < NR0; ++row) {
device const int8_t *qs = ax[row][ib].qs + il * NQ;
float sumq = 0.0f;
FOR_UNROLL(short i = 0; i < NQ; ++i) {
sumq += qs[i] * yl[i];
}
sumf[row] += sumq * ax[row][ib].d;
}
yb += NSG * NQ * QK8_0;
}
threadgroup float *shmem_f32[NR0];
FOR_UNROLL(short row = 0; row < NR0; ++row) {
shmem_f32[row] = (threadgroup float *)shmem + NW * row;
if (sgitg == 0) {
shmem_f32[row][tiisg] = 0.0f;
}
sumf[row] = simd_sum(sumf[row]);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
FOR_UNROLL(short row = 0; row < NR0; ++row) {
if (tiisg == 0) {
shmem_f32[row][sgitg] = sumf[row];
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
FOR_UNROLL(short row = 0; row < NR0; ++row) {
const int d = row0 + row;
if (d >= mv.ne01) {
continue;
}
const float block_v = simd_sum(shmem_f32[row][tiisg]);
if (tiisg == 0 && sgitg == 0) {
*((device float *)(block_out + (uint64_t)d * sizeof(float))) = block_v;
const float r0 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 0 * hc.nb_res1));
const float r1 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 1 * hc.nb_res1));
const float r2 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 2 * hc.nb_res1));
const float r3 = *((device const float *)(residual + (uint64_t)d * hc.nb_res0 + 3 * hc.nb_res1));
for (int64_t dst_hc = 0; dst_hc < 4; ++dst_hc) {
float acc = block_v * *((device const float *)(post + dst_hc * hc.nb_post0));
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 0 * hc.nb_comb1)) * r0;
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 1 * hc.nb_comb1)) * r1;
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 2 * hc.nb_comb1)) * r2;
acc += *((device const float *)(comb + dst_hc * hc.nb_comb0 + 3 * hc.nb_comb1)) * r3;
*((device float *)(dst + (uint64_t)d * hc.nb0 + dst_hc * hc.nb1)) = acc;
}
}
}
}
// Reduces HC channels to a normal embedding row with the learned pre weights.
// This is the input adapter before the attention block and before the FFN block.
kernel void kernel_dsv4_hc_weighted_sum(
constant ds4_metal_args_dsv4_hc_weighted_sum & args,
device const char * x,
device const char * weights,
device char * dst,
uint gid [[thread_position_in_grid]]) {
const int64_t n_elem = args.n_embd * args.n_tokens;
if ((int64_t) gid >= n_elem) {
return;
}
const int64_t d = ((int64_t) gid) % args.n_embd;
const int64_t t = ((int64_t) gid) / args.n_embd;
float acc = 0.0f;
for (int64_t h = 0; h < args.n_hc; ++h) {
const float xv = *((device const float *) (x + d*args.nb_x0 + h*args.nb_x1 + t*args.nb_x2));
const float wv = *((device const float *) (weights + h*args.nb_w0 + t*args.nb_w1));
acc += xv * wv;
}
*((device float *) (dst + d*args.nb0 + t*args.nb1)) = acc;
}
metal/dsv4_kv.metal
+227
View File
@@ -0,0 +1,227 @@
constant float dsv4_e4m3fn_exp_scale[16] = {
0.0f, 0.015625f, 0.03125f, 0.0625f,
0.125f, 0.25f, 0.5f, 1.0f,
2.0f, 4.0f, 8.0f, 16.0f,
32.0f, 64.0f, 128.0f, 256.0f,
};
struct ds4_metal_args_dsv4_fp8_kv_quantize {
int64_t ne00;
int64_t ne01;
int64_t ne02;
int64_t ne03;
ulong nb00;
ulong nb01;
ulong nb02;
ulong nb03;
ulong nb0;
ulong nb1;
ulong nb2;
ulong nb3;
int n_rot;
};
struct ds4_metal_args_dsv4_kv_fp8_store {
int32_t head_dim;
int32_t n_rot;
int32_t raw_row;
};
struct ds4_metal_args_dsv4_ratio4_shift {
uint32_t width;
};
struct ds4_metal_args_dsv4_compressor_store_one {
uint32_t width;
uint32_t ratio;
uint32_t pos;
uint32_t ape_type;
};
static inline float dsv4_e4m3fn_value(int i) {
const int exp = (i >> 3) & 0x0f;
const int mant = i & 0x07;
return exp == 0
? float(mant) * 0.001953125f
: (1.0f + float(mant) * 0.125f) * dsv4_e4m3fn_exp_scale[exp];
}
static inline float dsv4_e4m3fn_dequant(float x) {
const float sign = x < 0.0f ? -1.0f : 1.0f;
const float ax = min(abs(x), 448.0f);
int lo = 0;
int hi = 126;
while (lo < hi) {
const int mid = (lo + hi + 1) >> 1;
if (dsv4_e4m3fn_value(mid) <= ax) {
lo = mid;
} else {
hi = mid - 1;
}
}
int best = lo;
if (best < 126) {
const float best_diff = abs(ax - dsv4_e4m3fn_value(best));
const float next_diff = abs(ax - dsv4_e4m3fn_value(best + 1));
if (next_diff < best_diff || (next_diff == best_diff && ((best + 1) & 1) == 0 && (best & 1) != 0)) {
best = best + 1;
}
}
return sign * dsv4_e4m3fn_value(best);
}
// Quantizes the non-RoPE part of a KV row through E4M3FN and writes the
// dequantized value back as float. DS4 uses this to match the FP8 KV-cache
// semantics while keeping the Metal graph's cache buffers float-addressable.
kernel void kernel_dsv4_fp8_kv_quantize_f32(
constant ds4_metal_args_dsv4_fp8_kv_quantize & args,
device const char * src0,
device char * dst,
threadgroup float * scratch [[threadgroup(0)]],
uint row [[threadgroup_position_in_grid]],
uint tid [[thread_position_in_threadgroup]]) {
const int64_t n_rows = args.ne01 * args.ne02 * args.ne03;
if ((int64_t) row >= n_rows) {
return;
}
const int64_t i1 = row % args.ne01;
const int64_t i2 = (row / args.ne01) % args.ne02;
const int64_t i3 = row / (args.ne01 * args.ne02);
device const char * src_base = src0 + i1*args.nb01 + i2*args.nb02 + i3*args.nb03;
device char * dst_base = dst + i1*args.nb1 + i2*args.nb2 + i3*args.nb3;
const int64_t n_nope = args.ne00 - args.n_rot;
for (int64_t off = 0; off < n_nope; off += 64) {
float v = 0.0f;
if (tid < 64) {
v = *((device const float *) (src_base + (off + tid)*args.nb00));
scratch[tid] = abs(v);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
for (uint stride = 32; stride > 0; stride >>= 1) {
if (tid < stride) {
scratch[tid] = max(scratch[tid], scratch[tid + stride]);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
const float amax = max(scratch[0], 1.0e-4f);
const float scale = exp2(ceil(log2(amax / 448.0f)));
if (tid < 64) {
const float q = dsv4_e4m3fn_dequant(clamp(v / scale, -448.0f, 448.0f)) * scale;
*((device float *) (dst_base + (off + tid)*args.nb0)) = q;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
for (int64_t i = n_nope + tid; i < args.ne00; i += 64) {
*((device float *) (dst_base + i*args.nb0)) = *((device const float *) (src_base + i*args.nb00));
}
}
// Decode-side KV finalizer after RoPE. The normal RoPE kernel intentionally
// remains separate because tiny trigonometric codegen changes can flip later
// sampled tokens. This kernel only fuses the FP8 round-trip for the non-RoPE
// prefix with the F16-rounded raw-cache row used by FlashAttention.
kernel void kernel_dsv4_kv_fp8_store_f32(
constant ds4_metal_args_dsv4_kv_fp8_store & args,
device float * kv,
device float * raw_cache,
threadgroup float * scratch [[threadgroup(0)]],
uint tid [[thread_position_in_threadgroup]]) {
const int head_dim = args.head_dim;
const int n_rot = args.n_rot;
const int n_nope = head_dim - n_rot;
if (head_dim <= 0 || n_rot < 0 || n_nope < 0 || tid >= 64) {
return;
}
device float * raw = raw_cache + (int64_t)args.raw_row * head_dim;
for (int off = 0; off < n_nope; off += 64) {
float v = 0.0f;
if (off + (int)tid < n_nope) {
v = kv[off + tid];
scratch[tid] = abs(v);
} else {
scratch[tid] = 0.0f;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
for (uint stride = 32; stride > 0; stride >>= 1) {
if (tid < stride) {
scratch[tid] = max(scratch[tid], scratch[tid + stride]);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
const float amax = max(scratch[0], 1.0e-4f);
const float fp8_scale = exp2(ceil(log2(amax / 448.0f)));
if (off + (int)tid < n_nope) {
const float q = dsv4_e4m3fn_dequant(clamp(v / fp8_scale, -448.0f, 448.0f)) * fp8_scale;
kv[off + tid] = q;
raw[off + tid] = (float)((half)q);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
}
for (int i = n_nope + tid; i < head_dim; i += 64) {
raw[i] = (float)((half)kv[i]);
}
}
// Ratio-4 compression keeps two 4-row halves of recurrent state. After an
// emitted compressed row, the second half becomes the next window's previous
// half. The old encoder expressed this as four generic copies; this DS4-specific
// kernel performs the KV and score copies together.
kernel void kernel_dsv4_ratio4_shift_f32(
constant ds4_metal_args_dsv4_ratio4_shift & args,
device float * state_kv,
device float * state_score,
uint gid [[thread_position_in_grid]]) {
const uint n = 4u * args.width;
if (gid >= n) return;
state_kv[gid] = state_kv[n + gid];
state_score[gid] = state_score[n + gid];
}
// One-token compressor frontier update. Decode appends exactly one projected KV
// row and one score row into a small recurrent state. The generic batch helper
// expresses this as APE copy, score add, and two set_rows operations; this
// kernel writes both state tensors directly while preserving the same
// score + APE arithmetic.
kernel void kernel_dsv4_compressor_store_one(
constant ds4_metal_args_dsv4_compressor_store_one & args,
device const float * kv,
device const float * score,
device const char * ape,
device float * state_kv,
device float * state_score,
uint gid [[thread_position_in_grid]]) {
if (gid >= args.width || args.width == 0 || args.ratio == 0) {
return;
}
const uint pos_mod = args.pos % args.ratio;
const uint dst_row = args.ratio == 4u ? args.ratio + pos_mod : pos_mod;
const uint dst = dst_row * args.width + gid;
const uint ape_i = pos_mod * args.width + gid;
float ape_v;
if (args.ape_type == 1u) {
ape_v = (float)(((device const half *)ape)[ape_i]);
} else {
ape_v = ((device const float *)ape)[ape_i];
}
state_kv[dst] = kv[gid];
state_score[dst] = score[gid] + ape_v;
}
metal/dsv4_misc.metal
+1043
View File
File diff suppressed because it is too large Load Diff
metal/dsv4_rope.metal
+155
View File
@@ -0,0 +1,155 @@
struct ds4_metal_args_dsv4_rope_tail {
int64_t ne00;
int64_t ne01;
int64_t ne02;
int64_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
int32_t n_dims;
int32_t mode;
int32_t n_ctx_orig;
int32_t inverse;
float freq_base;
float freq_scale;
float ext_factor;
float attn_factor;
float beta_fast;
float beta_slow;
bool src2;
};
static float rope_yarn_ramp(const float low, const float high, const int i0) {
const float y = (i0 / 2 - low) / max(0.001f, high - low);
return 1.0f - min(1.0f, max(0.0f, y));
}
// YaRN algorithm based on LlamaYaRNScaledRotaryEmbedding.py from https://github.com/jquesnelle/yarn
// MIT licensed. Copyright (c) 2023 Jeffrey Quesnelle and Bowen Peng.
static void rope_yarn(
float theta_extrap, float freq_scale, float corr_dims[2], int i0, float ext_factor, float mscale,
thread float * cos_theta, thread float * sin_theta) {
// Get n-d rotational scaling corrected for extrapolation
float theta_interp = freq_scale * theta_extrap;
float theta = theta_interp;
if (ext_factor != 0.0f) {
float ramp_mix = rope_yarn_ramp(corr_dims[0], corr_dims[1], i0) * ext_factor;
theta = theta_interp * (1 - ramp_mix) + theta_extrap * ramp_mix;
// Get n-d magnitude scaling corrected for interpolation
mscale *= 1.0f + 0.1f * log(1.0f / freq_scale);
}
*cos_theta = cos(theta) * mscale;
*sin_theta = sin(theta) * mscale;
}
// Apparently solving `n_rot = 2pi * x * base^((2 * max_pos_emb) / n_dims)` for x, we get
// `corr_fac(n_rot) = n_dims * log(max_pos_emb / (n_rot * 2pi)) / (2 * log(base))`
static float rope_yarn_corr_factor(int n_dims, int n_ctx_orig, float n_rot, float base) {
return n_dims * log(n_ctx_orig / (n_rot * 2 * M_PI_F)) / (2 * log(base));
}
static void rope_yarn_corr_dims(
int n_dims, int n_ctx_orig, float freq_base, float beta_fast, float beta_slow, float dims[2]
) {
// start and end correction dims
dims[0] = max(0.0f, floor(rope_yarn_corr_factor(n_dims, n_ctx_orig, beta_fast, freq_base)));
dims[1] = min(n_dims - 1.0f, ceil(rope_yarn_corr_factor(n_dims, n_ctx_orig, beta_slow, freq_base)));
}
// Applies DeepSeek V4's partial RoPE: the no-position prefix is copied and only
// the rotated tail is transformed. This is used for Q/K after their projections
// and before writing/reading the attention KV state.
kernel void kernel_dsv4_rope_tail_f32(
constant ds4_metal_args_dsv4_rope_tail & args,
device const char * src0,
device const char * src1,
device const char * src2,
device char * dst,
uint tid [[thread_index_in_threadgroup]],
ushort3 ntg [[threads_per_threadgroup]],
uint3 tgpig [[threadgroup_position_in_grid]]) {
const int i1 = tgpig[0];
const int i2 = tgpig[1];
const int i3 = tgpig[2];
const int n_nope = args.ne00 - args.n_dims;
if (n_nope < 0) {
return;
}
device const int32_t * pos = (device const int32_t *) src1;
float corr_dims[2];
rope_yarn_corr_dims(args.n_dims, args.n_ctx_orig, args.freq_base, args.beta_fast, args.beta_slow, corr_dims);
const float theta_base = (float) pos[i2];
const float inv_ndims = -1.f/args.n_dims;
const bool is_neox = args.mode == 2;
for (int i0 = tid; i0 < args.ne00; i0 += ntg.x) {
device const char * src_base = src0 + i3*args.nb03 + i2*args.nb02 + i1*args.nb01;
device char * dst_base = dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1;
if (i0 < n_nope) {
*((device float *) (dst_base + i0*args.nb0)) = *((device const float *) (src_base + i0*args.nb00));
continue;
}
const int r = i0 - n_nope;
if (is_neox) {
const int n_half = args.n_dims/2;
if (r >= n_half) {
continue;
}
const int ic = r;
const int rel_i0 = 2*ic;
const float theta = theta_base * pow(args.freq_base, inv_ndims*rel_i0);
const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
float cos_theta;
float sin_theta;
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, rel_i0, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
if (args.inverse) {
sin_theta = -sin_theta;
}
const int j0 = n_nope + ic;
const int j1 = n_nope + ic + n_half;
const float x0 = *((device const float *) (src_base + j0*args.nb00));
const float x1 = *((device const float *) (src_base + j1*args.nb00));
*((device float *) (dst_base + j0*args.nb0)) = x0*cos_theta - x1*sin_theta;
*((device float *) (dst_base + j1*args.nb0)) = x0*sin_theta + x1*cos_theta;
} else {
if ((r & 1) != 0) {
continue;
}
const int ic = r/2;
const float theta = theta_base * pow(args.freq_base, inv_ndims*r);
const float freq_factor = args.src2 ? ((device const float *) src2)[ic] : 1.0f;
float cos_theta;
float sin_theta;
rope_yarn(theta/freq_factor, args.freq_scale, corr_dims, r, args.ext_factor, args.attn_factor, &cos_theta, &sin_theta);
if (args.inverse) {
sin_theta = -sin_theta;
}
const int j0 = n_nope + r;
const int j1 = j0 + 1;
const float x0 = *((device const float *) (src_base + j0*args.nb00));
const float x1 = *((device const float *) (src_base + j1*args.nb00));
*((device float *) (dst_base + j0*args.nb0)) = x0*cos_theta - x1*sin_theta;
*((device float *) (dst_base + j1*args.nb0)) = x0*sin_theta + x1*cos_theta;
}
}
}
metal/flash_attn.metal
+1426
View File
File diff suppressed because it is too large Load Diff
metal/get_rows.metal
+54
View File
@@ -0,0 +1,54 @@
// DS4 Metal get-rows kernel.
struct ds4_metal_args_get_rows {
int32_t ne00t;
int32_t ne00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne10;
uint64_t nb10;
uint64_t nb11;
uint64_t nb12;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
};
// Gathers embedding/table rows by integer ids. DS4 uses this for token
// embeddings and small indexed tables such as router/hash lookup outputs.
template<typename T0, typename T>
kernel void kernel_get_rows_f(
constant ds4_metal_args_get_rows & args,
device const void * src0,
device const void * src1,
device void * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiitg[[thread_index_in_threadgroup]],
ushort3 ntg [[threads_per_threadgroup]]) {
const int32_t iw0 = tgpig.x/args.ne10;
const int32_t i10 = tgpig.x%args.ne10;
const int32_t i11 = tgpig.y;
const int32_t i12 = tgpig.z;
const int32_t r = ((const device int32_t *) ((const device char *) src1 + i12*args.nb12 + i11*args.nb11 + i10*args.nb10))[0];
const int32_t i02 = i11;
const int32_t i03 = i12;
auto psrc = (const device T0 *) ((const device char *) src0 + i03*args.nb03 + i02*args.nb02 + r*args.nb01);
auto pdst = ( device T *) (( device char *) dst + i12*args.nb3 + i11*args.nb2 + i10*args.nb1);
for (int ind = iw0*ntg.x + tiitg; ind < args.ne00t;) {
pdst[ind] = psrc[ind];
break;
}
}
typedef decltype(kernel_get_rows_f<float, float>) get_rows_f_t;
// Host-visible gather variants for F32, F16, and I32 tables.
template [[host_name("kernel_get_rows_f32")]] kernel get_rows_f_t kernel_get_rows_f<float, float>;
template [[host_name("kernel_get_rows_f16")]] kernel get_rows_f_t kernel_get_rows_f<half, float>;
template [[host_name("kernel_get_rows_i32")]] kernel get_rows_f_t kernel_get_rows_f<int32_t, int32_t>;
metal/glu.metal
+36
View File
@@ -0,0 +1,36 @@
struct ds4_metal_args_glu {
int32_t ne00;
uint64_t nb01;
int32_t ne10;
uint64_t nb11;
int32_t ne0;
uint64_t nb1;
int32_t i00;
int32_t i10;
float alpha;
float limit;
};
// SwiGLU activation for the FFN inner state: silu(gate) * up. The DS4 graph
// uses it between the gate/up expert matmuls and the down projection.
kernel void kernel_swiglu_f32(
constant ds4_metal_args_glu & args,
device const char * src0,
device const char * src1,
device char * dst,
uint tgpig[[threadgroup_position_in_grid]],
uint tpitg[[thread_position_in_threadgroup]],
uint ntg[[threads_per_threadgroup]]) {
device const float * src0_row = (device const float *) ((device const char *) src0 + tgpig*args.nb01) + args.i00;
device const float * src1_row = (device const float *) ((device const char *) src1 + tgpig*args.nb11) + args.i10;
device float * dst_row = (device float *) ((device char *) dst + tgpig*args.nb1);
for (int i0 = tpitg; i0 < args.ne0; i0 += ntg) {
const float x0 = src0_row[i0];
const float x1 = src1_row[i0];
const float silu = x0 / (1.0f + exp(-x0));
dst_row[i0] = silu*x1;
}
}
metal/moe.metal
+1737
View File
File diff suppressed because it is too large Load Diff
metal/norm.metal
+153
View File
@@ -0,0 +1,153 @@
struct ds4_metal_args_norm {
int32_t ne00;
int32_t ne00_t;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
float eps;
int32_t nef1[3];
int32_t nef2[3];
int32_t nef3[3];
uint64_t nbf1[3];
uint64_t nbf2[3];
uint64_t nbf3[3];
};
// RMSNorm over one activation row, optionally fusing the learned weight
// multiply. DS4 calls this before attention, before the FFN, and for plain
// diagnostics that need normalized but unweighted rows.
template <typename T, short F>
kernel void kernel_rms_norm_fuse_impl(
constant ds4_metal_args_norm & args,
device const char * src0,
device const char * src1_0,
device const char * src1_1,
device char * dst,
threadgroup float * shmem_f32 [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
if (sgitg == 0) {
shmem_f32[tiisg] = 0.0f;
}
const int i01 = tgpig.x;
const int i02 = tgpig.y;
const int i03 = tgpig.z;
device const T * x = (device const T *) (src0 + i03*args.nbf3[0] + i02*args.nbf2[0] + i01*args.nbf1[0]);
device const T * f0 = (device const T *) (src1_0 + (i03%args.nef3[1])*args.nbf3[1] + (i02%args.nef2[1])*args.nbf2[1] + (i01%args.nef1[1])*args.nbf1[1]);
device const T * f1 = (device const T *) (src1_1 + (i03%args.nef3[2])*args.nbf3[2] + (i02%args.nef2[2])*args.nbf2[2] + (i01%args.nef1[2])*args.nbf1[2]);
float sumf = 0.0f;
// parallel sum
for (int i00 = tpitg.x; i00 < args.ne00_t; i00 += ntg.x) {
sumf += dot(x[i00], x[i00]);
}
sumf = simd_sum(sumf);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
shmem_f32[sgitg] = sumf;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_f32[tiisg];
sumf = simd_sum(sumf);
const float mean = sumf/args.ne00;
const float scale = 1.0f/sqrt(mean + args.eps);
device T * y = (device T *) (dst + i03*args.nb3 + i02*args.nb2 + i01*args.nb1);
for (int i00 = tpitg.x; i00 < args.ne00_t; i00 += ntg.x) {
if (F == 1) {
y[i00] = (x[i00]*scale);
}
if (F == 2) {
y[i00] = (x[i00]*scale)*f0[i00];
}
if (F == 3) {
y[i00] = (x[i00]*scale)*f0[i00] + f1[i00];
}
}
}
typedef decltype(kernel_rms_norm_fuse_impl<float4, 1>) kernel_rms_norm_fuse_t;
// Host-visible RMSNorm variants: plain norm and norm multiplied by weight.
template [[host_name("kernel_rms_norm_f32_4")]] kernel kernel_rms_norm_fuse_t kernel_rms_norm_fuse_impl<float4, 1>;
template [[host_name("kernel_rms_norm_mul_f32_4")]] kernel kernel_rms_norm_fuse_t kernel_rms_norm_fuse_impl<float4, 2>;
struct ds4_metal_args_qkv_rms_norm {
int32_t q_n;
int32_t q_n4;
int32_t kv_n;
int32_t kv_n4;
uint64_t q_row_stride;
uint64_t kv_row_stride;
float eps;
};
// Normalizes DS4's q-lora row and KV row in one dispatch. The two reductions
// deliberately mirror kernel_rms_norm_mul_f32_4: Q uses the full 256-thread
// row shape for 1024 floats, while KV only has work in the first 128 lanes for
// its 512 floats. This keeps the q/kv normalization math aligned with the
// standalone kernels while removing one tiny launch from the attention setup.
kernel void kernel_dsv4_qkv_rms_norm_f32_4(
constant ds4_metal_args_qkv_rms_norm & args,
device const float4 * q_src,
device const float4 * q_weight,
device float4 * q_dst,
device const float4 * kv_src,
device const float4 * kv_weight,
device float4 * kv_dst,
threadgroup float * shmem_f32 [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
if (sgitg == 0) {
shmem_f32[tiisg] = 0.0f;
}
const uint row = tgpig.x;
const bool kv_task = tgpig.y != 0;
const int n = kv_task ? args.kv_n : args.q_n;
const int n4 = kv_task ? args.kv_n4 : args.q_n4;
const uint64_t row_stride4 = (kv_task ? args.kv_row_stride : args.q_row_stride) / sizeof(float4);
device const float4 * x = kv_task ? kv_src + row * row_stride4 : q_src + row * row_stride4;
device const float4 * w = kv_task ? kv_weight : q_weight;
device float4 * y = kv_task ? kv_dst + row * row_stride4 : q_dst + row * row_stride4;
float sumf = 0.0f;
for (int i = tpitg.x; i < n4; i += ntg.x) {
const float4 v = x[i];
sumf += dot(v, v);
}
sumf = simd_sum(sumf);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
shmem_f32[sgitg] = sumf;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_f32[tiisg];
sumf = simd_sum(sumf);
const float scale = rsqrt(sumf / float(n) + args.eps);
for (int i = tpitg.x; i < n4; i += ntg.x) {
y[i] = (x[i] * scale) * w[i];
}
}
metal/repeat.metal
+52
View File
@@ -0,0 +1,52 @@
// DS4 Metal repeat kernel used for HC embedding expansion.
struct ds4_metal_args_repeat {
int32_t ne00;
int32_t ne01;
int32_t ne02;
int32_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne0;
int32_t ne1;
int32_t ne2;
int32_t ne3;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
};
// Repeats a source row into the HC channel dimension. DS4 uses this when the
// token embedding has to become an HC activation block before layer 0.
template<typename T>
kernel void kernel_repeat(
constant ds4_metal_args_repeat & args,
device const char * src0,
device char * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
const int i3 = tgpig.z;
const int i2 = tgpig.y;
const int i1 = tgpig.x;
const int i03 = i3%args.ne03;
const int i02 = i2%args.ne02;
const int i01 = i1%args.ne01;
device const char * src0_ptr = src0 + i03*args.nb03 + i02*args.nb02 + i01*args.nb01;
device char * dst_ptr = dst + i3*args.nb3 + i2*args.nb2 + i1*args.nb1;
for (int i0 = tpitg.x; i0 < args.ne0; i0 += ntg.x) {
const int i00 = i0%args.ne00;
*((device T *)(dst_ptr + i0*args.nb0)) = *((device T *)(src0_ptr + i00*args.nb00));
}
}
typedef decltype(kernel_repeat<float>) kernel_repeat_t;
// Host-visible F32 repeat used for HC expansion of embeddings.
template [[host_name("kernel_repeat_f32")]] kernel kernel_repeat_t kernel_repeat<float>;
metal/set_rows.metal
+55
View File
@@ -0,0 +1,55 @@
// DS4 Metal set-rows kernel used for KV writes.
struct ds4_metal_args_set_rows {
int32_t nk0;
int32_t ne01;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne11;
int32_t ne12;
uint64_t nb10;
uint64_t nb11;
uint64_t nb12;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
};
// Scatters rows into the KV cache by token position. DS4 uses this after Q/K/V
// preparation so decode and later prefill chunks can attend to previous tokens.
template<typename T, typename TI>
kernel void kernel_set_rows_f(
constant ds4_metal_args_set_rows & args,
device const void * src0,
device const void * src1,
device float * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
uint tiitg[[thread_index_in_threadgroup]],
uint3 tptg [[threads_per_threadgroup]]) {
const int32_t i03 = tgpig.z;
const int32_t i02 = tgpig.y;
const int32_t i12 = i03%args.ne12;
const int32_t i11 = i02%args.ne11;
const int32_t i01 = tgpig.x*tptg.y + tiitg/tptg.x;
if (i01 >= args.ne01) {
return;
}
const int32_t i10 = i01;
const TI i1 = ((const device TI *) ((const device char *) src1 + i10*args.nb10 + i11*args.nb11 + i12*args.nb12))[0];
device T * dst_row = ( device T *) (( device char *) dst + i1*args.nb1 + i02*args.nb2 + i03*args.nb3);
const device float * src_row = (const device float *) ((const device char *) src0 + i01*args.nb01 + i02*args.nb02 + i03*args.nb03);
for (int ind = tiitg%tptg.x; ind < args.nk0; ind += tptg.x) {
dst_row[ind] = (T) src_row[ind];
}
}
typedef decltype(kernel_set_rows_f<float, int64_t>) set_rows_f_t;
// Host-visible F32/I32 scatter variant used by KV-cache writes.
template [[host_name("kernel_set_rows_f32_i32")]] kernel set_rows_f_t kernel_set_rows_f<float, int32_t>;
metal/softmax.metal
+241
View File
@@ -0,0 +1,241 @@
// DS4 Metal softmax kernel used by the compressor pooling compatibility path.
// The single-compressed-row path is intentionally left as soft_max -> mul ->
// sum_rows instead of using the fused dsv4_softmax_pool kernel.
struct ds4_metal_args_soft_max {
int32_t ne00;
int32_t ne01;
int32_t ne02;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne11;
int32_t ne12;
int32_t ne13;
uint64_t nb11;
uint64_t nb12;
uint64_t nb13;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
float scale;
float max_bias;
float m0;
float m1;
int32_t n_head_log2;
};
// Row softmax for score matrices. DS4 uses it in the literal one-compressor-row
// path where preserving the original graph operation boundary avoids drift.
template<typename T>
kernel void kernel_soft_max(
constant ds4_metal_args_soft_max & args,
device const char * src0,
device const char * src1,
device const char * src2,
device char * dst,
threadgroup float * buf [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
uint3 tpitg[[thread_position_in_threadgroup]],
uint sgitg[[simdgroup_index_in_threadgroup]],
uint tiisg[[thread_index_in_simdgroup]],
uint3 tptg[[threads_per_threadgroup]]) {
const int32_t i03 = tgpig.z;
const int32_t i02 = tgpig.y;
const int32_t i01 = tgpig.x;
const int32_t i13 = i03%args.ne13;
const int32_t i12 = i02%args.ne12;
const int32_t i11 = i01;
device const float * psrc0 = (device const float *) (src0 + i01*args.nb01 + i02*args.nb02 + i03*args.nb03);
device const T * pmask = src1 != src0 ? (device const T * ) (src1 + i11*args.nb11 + i12*args.nb12 + i13*args.nb13) : nullptr;
device const float * psrc2 = src2 != src0 ? (device const float *) (src2) : nullptr;
device float * pdst = (device float *) (dst + i01*args.nb1 + i02*args.nb2 + i03*args.nb3);
float slope = 1.0f;
if (args.max_bias > 0.0f) {
const int32_t h = i02;
const float base = h < args.n_head_log2 ? args.m0 : args.m1;
const int exp = h < args.n_head_log2 ? h + 1 : 2*(h - args.n_head_log2) + 1;
slope = pow(base, exp);
}
float lmax = psrc2 ? psrc2[i02] : -INFINITY;
for (int i00 = tpitg.x; i00 < args.ne00; i00 += tptg.x) {
lmax = MAX(lmax, psrc0[i00]*args.scale + (pmask ? slope*pmask[i00] : 0.0f));
}
float max_val = simd_max(lmax);
if (tptg.x > N_SIMDWIDTH) {
if (sgitg == 0) {
buf[tiisg] = -INFINITY;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
buf[sgitg] = max_val;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
max_val = buf[tiisg];
max_val = simd_max(max_val);
}
float lsum = 0.0f;
for (int i00 = tpitg.x; i00 < args.ne00; i00 += tptg.x) {
const float exp_psrc0 = exp((psrc0[i00]*args.scale + (pmask ? slope*pmask[i00] : 0.0f)) - max_val);
lsum += exp_psrc0;
pdst[i00] = exp_psrc0;
}
threadgroup_barrier(mem_flags::mem_none);
float sum = simd_sum(lsum);
if (tptg.x > N_SIMDWIDTH) {
if (sgitg == 0) {
buf[tiisg] = 0.0f;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
buf[sgitg] = sum;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
sum = buf[tiisg];
sum = simd_sum(sum);
}
if (psrc2) {
sum += exp(psrc2[i02] - max_val);
}
const float inv_sum = 1.0f/sum;
for (int i00 = tpitg.x; i00 < args.ne00; i00 += tptg.x) {
pdst[i00] *= inv_sum;
}
}
// Vectorized float4 row softmax for contiguous score rows whose length is a
// multiple of four; used by the same DS4 compressor/indexer graph path.
template<typename T>
kernel void kernel_soft_max_4(
constant ds4_metal_args_soft_max & args,
device const char * src0,
device const char * src1,
device const char * src2,
device char * dst,
threadgroup float * buf [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
uint3 tpitg[[thread_position_in_threadgroup]],
uint sgitg[[simdgroup_index_in_threadgroup]],
uint tiisg[[thread_index_in_simdgroup]],
uint3 tptg[[threads_per_threadgroup]]) {
const int32_t i03 = tgpig.z;
const int32_t i02 = tgpig.y;
const int32_t i01 = tgpig.x;
const int32_t i13 = i03%args.ne13;
const int32_t i12 = i02%args.ne12;
const int32_t i11 = i01;
device const float4 * psrc4 = (device const float4 *) (src0 + i01*args.nb01 + i02*args.nb02 + i03*args.nb03);
device const T * pmask = src1 != src0 ? (device const T * ) (src1 + i11*args.nb11 + i12*args.nb12 + i13*args.nb13) : nullptr;
device const float * psrc2 = src2 != src0 ? (device const float * ) (src2) : nullptr;
device float4 * pdst4 = (device float4 *) (dst + i01*args.nb1 + i02*args.nb2 + i03*args.nb3);
float slope = 1.0f;
if (args.max_bias > 0.0f) {
const int32_t h = i02;
const float base = h < args.n_head_log2 ? args.m0 : args.m1;
const int exp = h < args.n_head_log2 ? h + 1 : 2*(h - args.n_head_log2) + 1;
slope = pow(base, exp);
}
float4 lmax4 = psrc2 ? psrc2[i02] : -INFINITY;
for (int i00 = tpitg.x; i00 < args.ne00/4; i00 += tptg.x) {
lmax4 = fmax(lmax4, psrc4[i00]*args.scale + (float4)((pmask ? slope*pmask[i00] : 0.0f)));
}
const float lmax = MAX(MAX(lmax4[0], lmax4[1]), MAX(lmax4[2], lmax4[3]));
float max_val = simd_max(lmax);
if (tptg.x > N_SIMDWIDTH) {
if (sgitg == 0) {
buf[tiisg] = -INFINITY;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
buf[sgitg] = max_val;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
max_val = buf[tiisg];
max_val = simd_max(max_val);
}
float4 lsum4 = 0.0f;
for (int i00 = tpitg.x; i00 < args.ne00/4; i00 += tptg.x) {
const float4 exp_psrc4 = exp((psrc4[i00]*args.scale + (float4)((pmask ? slope*pmask[i00] : 0.0f))) - max_val);
lsum4 += exp_psrc4;
pdst4[i00] = exp_psrc4;
}
const float lsum = lsum4[0] + lsum4[1] + lsum4[2] + lsum4[3];
threadgroup_barrier(mem_flags::mem_none);
float sum = simd_sum(lsum);
if (tptg.x > N_SIMDWIDTH) {
if (sgitg == 0) {
buf[tiisg] = 0.0f;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
buf[sgitg] = sum;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
sum = buf[tiisg];
sum = simd_sum(sum);
}
if (psrc2) {
sum += exp(psrc2[i02] - max_val);
}
const float inv_sum = 1.0f/sum;
for (int i00 = tpitg.x; i00 < args.ne00/4; i00 += tptg.x) {
pdst4[i00] *= inv_sum;
}
}
typedef decltype(kernel_soft_max<float>) kernel_soft_max_t;
typedef decltype(kernel_soft_max_4<float4>) kernel_soft_max_4_t;
// Host-visible F32 softmax variants used by compressor pooling.
template [[host_name("kernel_soft_max_f32")]] kernel kernel_soft_max_t kernel_soft_max<float>;
template [[host_name("kernel_soft_max_f32_4")]] kernel kernel_soft_max_4_t kernel_soft_max_4<float4>;
metal/sum_rows.metal
+102
View File
@@ -0,0 +1,102 @@
// DS4 Metal row-sum kernel.
#define FC_SUM_ROWS 1400
#define OP_SUM_ROWS_NUM_SUM_ROWS 10
#define OP_SUM_ROWS_NUM_MEAN 11
struct ds4_metal_args_sum_rows {
int64_t ne00;
int64_t ne01;
int64_t ne02;
int64_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int64_t ne0;
int64_t ne1;
int64_t ne2;
int64_t ne3;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
};
static inline float sum(float x) {
return x;
}
static inline float sum(float4 x) {
return x[0] + x[1] + x[2] + x[3];
}
constant short FC_sum_rows_op [[function_constant(FC_SUM_ROWS + 0)]];
// Reduces each row to a sum or mean. DS4 mainly uses the sum form to preserve
// the compressor-pooling graph boundary in the single-compressor-row case.
template <typename T0, typename T>
kernel void kernel_sum_rows_impl(
constant ds4_metal_args_sum_rows & args,
device const char * src0,
device char * dst,
threadgroup char * shmem [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
#define FC_OP FC_sum_rows_op
const int i3 = tgpig.z;
const int i2 = tgpig.y;
const int i1 = tgpig.x;
threadgroup T0 * shmem_t = (threadgroup T0 *) shmem;
if (sgitg == 0) {
shmem_t[tiisg] = 0.0f;
}
device const T0 * src_row = (device const T0 *) (src0 + i1*args.nb01 + i2*args.nb02 + i3*args.nb03);
device T * dst_row = (device T *) (dst + i1*args.nb1 + i2*args.nb2 + i3*args.nb3);
T0 sumf = T0(0.0f);
for (int64_t i0 = tpitg.x; i0 < args.ne00; i0 += ntg.x) {
sumf += src_row[i0];
}
sumf = simd_sum(sumf);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (tiisg == 0) {
shmem_t[sgitg] = sumf;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
sumf = shmem_t[tiisg];
sumf = simd_sum(sumf);
if (tpitg.x == 0) {
if (FC_OP == OP_SUM_ROWS_NUM_MEAN) {
if (is_same<float4, T0>::value) {
dst_row[0] = sum(sumf) / (4*args.ne00);
} else {
dst_row[0] = sum(sumf) / args.ne00;
}
} else {
dst_row[0] = sum(sumf);
}
}
#undef FC_OP
}
typedef decltype(kernel_sum_rows_impl<float, float>) kernel_sum_rows_t;
// Host-visible F32 row reduction used by compressor pooling.
template [[host_name("kernel_sum_rows_f32_f32")]] kernel kernel_sum_rows_t kernel_sum_rows_impl<float, float>;
metal/unary.metal
+312
View File
@@ -0,0 +1,312 @@
#define FC_UNARY 1200
#define OP_UNARY_NUM_SCALE 10
#define OP_UNARY_NUM_FILL 11
#define OP_UNARY_NUM_CLAMP 12
#define OP_UNARY_NUM_SQR 13
#define OP_UNARY_NUM_SQRT 14
#define OP_UNARY_NUM_SIN 15
#define OP_UNARY_NUM_COS 16
#define OP_UNARY_NUM_LOG 17
#define OP_UNARY_NUM_LEAKY_RELU 18
#define OP_UNARY_NUM_TANH 100
#define OP_UNARY_NUM_RELU 101
#define OP_UNARY_NUM_SIGMOID 102
#define OP_UNARY_NUM_GELU 103
#define OP_UNARY_NUM_GELU_ERF 104
#define OP_UNARY_NUM_GELU_QUICK 105
#define OP_UNARY_NUM_SILU 106
#define OP_UNARY_NUM_ELU 107
#define OP_UNARY_NUM_NEG 108
#define OP_UNARY_NUM_ABS 109
#define OP_UNARY_NUM_SGN 110
#define OP_UNARY_NUM_STEP 111
#define OP_UNARY_NUM_HARDSWISH 112
#define OP_UNARY_NUM_HARDSIGMOID 113
#define OP_UNARY_NUM_EXP 114
#define OP_UNARY_NUM_SOFTPLUS 115
#define OP_UNARY_NUM_EXPM1 116
#define OP_UNARY_NUM_FLOOR 117
#define OP_UNARY_NUM_CEIL 118
#define OP_UNARY_NUM_ROUND 119
#define OP_UNARY_NUM_TRUNC 120
#define OP_UNARY_NUM_XIELU 121
struct ds4_metal_args_unary {
int32_t ne00;
int32_t ne01;
int32_t ne02;
int32_t ne03;
uint64_t nb00;
uint64_t nb01;
uint64_t nb02;
uint64_t nb03;
int32_t ne0;
int32_t ne1;
int32_t ne2;
int32_t ne3;
uint64_t nb0;
uint64_t nb1;
uint64_t nb2;
uint64_t nb3;
float slope;
float scale;
float bias;
float val;
float min;
float max;
};
constant float GELU_COEF_A = 0.044715f;
constant float GELU_QUICK_COEF = -1.702f;
constant float SQRT_2_OVER_PI = 0.79788456080286535587989211986876f;
constant float SQRT_2_INV = 0.70710678118654752440084436210484f;
// based on Abramowitz and Stegun formula 7.1.26 or similar Hastings' approximation
// ref: https://www.johndcook.com/blog/python_erf/
constant float p_erf = 0.3275911f;
constant float a1_erf = 0.254829592f;
constant float a2_erf = -0.284496736f;
constant float a3_erf = 1.421413741f;
constant float a4_erf = -1.453152027f;
constant float a5_erf = 1.061405429f;
template<typename T>
inline T erf_approx(T x) {
T sign_x = sign(x);
x = fabs(x);
T t = 1.0f / (1.0f + p_erf * x);
T y = 1.0f - (((((a5_erf * t + a4_erf) * t) + a3_erf) * t + a2_erf) * t + a1_erf) * t * exp(-x * x);
return sign_x * y;
}
template<typename T> T elu_approx(T x);
template<> inline float elu_approx<float>(float x) {
return (x > 0.f) ? x : (exp(x) - 1);
}
template<> inline float4 elu_approx<float4>(float4 x) {
float4 res;
res[0] = (x[0] > 0.0f) ? x[0] : (exp(x[0]) - 1.0f);
res[1] = (x[1] > 0.0f) ? x[1] : (exp(x[1]) - 1.0f);
res[2] = (x[2] > 0.0f) ? x[2] : (exp(x[2]) - 1.0f);
res[3] = (x[3] > 0.0f) ? x[3] : (exp(x[3]) - 1.0f);
return res;
}
constant short FC_unary_op [[function_constant(FC_UNARY + 0)]];
constant bool FC_unary_cnt[[function_constant(FC_UNARY + 1)]];
// Generic unary elementwise op selected by function constant. DS4 only uses a
// small subset in inference, mainly sigmoid, SiLU, softplus, sqrt, clamp,
// scale, and fill.
template <typename T0, typename T, typename TC>
kernel void kernel_unary_impl(
constant ds4_metal_args_unary & args,
device const char * src0,
device char * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort3 tpitg[[thread_position_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
#define FC_OP FC_unary_op
#define FC_CNT FC_unary_cnt
device const T0 * src0_ptr;
device T * dst_ptr;
int i0;
if (FC_CNT) {
i0 = tgpig.x;
src0_ptr = (device const T0 *) (src0);
dst_ptr = (device T *) (dst);
} else {
const int i03 = tgpig.z;
const int i02 = tgpig.y;
const int k0 = tgpig.x/args.ne01;
const int i01 = tgpig.x - k0*args.ne01;
i0 = k0*ntg.x + tpitg.x;
src0_ptr = (device const T0 *) (src0 + i03*args.nb03 + i02*args.nb02 + i01*args.nb01);
dst_ptr = (device T *) (dst + i03*args.nb3 + i02*args.nb2 + i01*args.nb1 );
}
{
if (!FC_CNT) {
if (i0 >= args.ne0) {
return;
}
}
const TC x = (TC) src0_ptr[i0];
if (FC_OP == OP_UNARY_NUM_SCALE) {
dst_ptr[i0] = (T) (args.scale * x + args.bias);
}
if (FC_OP == OP_UNARY_NUM_FILL) {
dst_ptr[i0] = (T) args.val;
}
if (FC_OP == OP_UNARY_NUM_CLAMP) {
dst_ptr[i0] = (T) clamp(x, args.min, args.max);
}
if (FC_OP == OP_UNARY_NUM_SQR) {
dst_ptr[i0] = (T) (x * x);
}
if (FC_OP == OP_UNARY_NUM_SQRT) {
dst_ptr[i0] = (T) sqrt(x);
}
if (FC_OP == OP_UNARY_NUM_SIN) {
dst_ptr[i0] = (T) sin(x);
}
if (FC_OP == OP_UNARY_NUM_COS) {
dst_ptr[i0] = (T) cos(x);
}
if (FC_OP == OP_UNARY_NUM_LOG) {
dst_ptr[i0] = (T) log(x);
}
if (FC_OP == OP_UNARY_NUM_LEAKY_RELU) {
dst_ptr[i0] = (T) (TC(x > 0)*x + TC(x <= 0)*(x * args.slope));
}
if (FC_OP == OP_UNARY_NUM_TANH) {
dst_ptr[i0] = (T) precise::tanh(x);
}
if (FC_OP == OP_UNARY_NUM_RELU) {
dst_ptr[i0] = (T) fmax(0, x);
}
if (FC_OP == OP_UNARY_NUM_SIGMOID) {
dst_ptr[i0] = (T) (1 / (1 + exp(-x)));
}
if (FC_OP == OP_UNARY_NUM_GELU) {
dst_ptr[i0] = (T) (0.5*x*(1 + precise::tanh(SQRT_2_OVER_PI*x*(1 + GELU_COEF_A*x*x))));
}
if (FC_OP == OP_UNARY_NUM_GELU_ERF) {
dst_ptr[i0] = (T) (0.5*x*(1 + erf_approx(SQRT_2_INV*x)));
}
if (FC_OP == OP_UNARY_NUM_GELU_QUICK) {
dst_ptr[i0] = (T) (x * (1/(1 + exp(GELU_QUICK_COEF*x))));
}
if (FC_OP == OP_UNARY_NUM_SILU) {
dst_ptr[i0] = (T) (x / (1 + exp(-x)));
}
if (FC_OP == OP_UNARY_NUM_ELU) {
dst_ptr[i0] = (T) elu_approx(x);
}
if (FC_OP == OP_UNARY_NUM_NEG) {
dst_ptr[i0] = (T) -x;
}
if (FC_OP == OP_UNARY_NUM_ABS) {
dst_ptr[i0] = (T) fabs(x);
}
if (FC_OP == OP_UNARY_NUM_SGN) {
dst_ptr[i0] = T(x > 0) - T(x < 0);
}
if (FC_OP == OP_UNARY_NUM_STEP) {
dst_ptr[i0] = T(x > 0);
}
if (FC_OP == OP_UNARY_NUM_HARDSWISH) {
dst_ptr[i0] = (T) (x * fmax(0, fmin(1, x/6 + 0.5)));
}
if (FC_OP == OP_UNARY_NUM_HARDSIGMOID) {
dst_ptr[i0] = (T) fmax(0, fmin(1, x/6 + 0.5));
}
if (FC_OP == OP_UNARY_NUM_EXP) {
dst_ptr[i0] = (T) exp(x);
}
if (FC_OP == OP_UNARY_NUM_SOFTPLUS) {
dst_ptr[i0] = (T) select(log(1 + exp(x)), x, x > 20);
}
if (FC_OP == OP_UNARY_NUM_EXPM1) {
// Metal target profiles used here do not all expose expm1(); this
// generic unary branch is not used by the DS4 inference graph.
dst_ptr[i0] = (T) (exp(x) - 1);
}
if (FC_OP == OP_UNARY_NUM_FLOOR) {
dst_ptr[i0] = (T) floor(x);
}
if (FC_OP == OP_UNARY_NUM_CEIL) {
dst_ptr[i0] = (T) ceil(x);
}
if (FC_OP == OP_UNARY_NUM_ROUND) {
dst_ptr[i0] = (T) round(x);
}
if (FC_OP == OP_UNARY_NUM_TRUNC) {
dst_ptr[i0] = (T) trunc(x);
}
if (FC_OP == OP_UNARY_NUM_XIELU) {
const TC xi = x;
const TC gate = TC(xi > TC(0.0f));
const TC clamped = fmin(xi, TC(args.val));
const TC y_pos = TC(args.scale) * xi * xi + TC(args.bias) * xi;
const TC y_neg = (exp(clamped) - TC(1.0f) - xi) * TC(args.slope) + TC(args.bias) * xi;
dst_ptr[i0] = (T) (gate * y_pos + (TC(1.0f) - gate) * y_neg);
}
}
#undef FC_OP
#undef FC_CNT
}
typedef decltype(kernel_unary_impl<float, float, float>) kernel_unary_t;
// Decode router probability transform. The generic path applies softplus and
// sqrt as two elementwise kernels; DS4 decode always transforms one 256-wide
// expert-logit row, so this vectorized kernel does both in one pass.
kernel void kernel_dsv4_softplus_sqrt_f32_4(
constant ds4_metal_args_unary & args,
device const char *src,
device char *dst,
uint3 tgpig [[threadgroup_position_in_grid]],
ushort3 tpitg [[thread_position_in_threadgroup]],
ushort3 ntg [[threads_per_threadgroup]]) {
const int k0 = tgpig.x/args.ne01;
const int i01 = tgpig.x - k0*args.ne01;
const int i0 = k0*ntg.x + tpitg.x;
if (i0 >= args.ne0) return;
device const float4 *s = (device const float4 *)(src + i01*args.nb01);
device float4 *d = (device float4 *)(dst + i01*args.nb1);
const float4 x = s[i0];
const float4 sp = select(log(1.0f + exp(x)), x, x > 20.0f);
d[i0] = sqrt(sp);
}
// Host-visible unary variants. Function constants select the actual DS4 op.
template [[host_name("kernel_unary_f32_f32")]] kernel kernel_unary_t kernel_unary_impl<float, float, float>;
template [[host_name("kernel_unary_f32_f32_4")]] kernel kernel_unary_t kernel_unary_impl<float4, float4, float4>;
template [[host_name("kernel_unary_f16_f16")]] kernel kernel_unary_t kernel_unary_impl<half, half, float>;
tests/ds4_test.c
+665
View File
@@ -0,0 +1,665 @@
#define DS4_SERVER_TEST
#define DS4_SERVER_TEST_NO_MAIN
#include "../ds4_server.c"
#ifndef DS4_NO_METAL
#include "../ds4_metal.h"
#include <math.h>
static ds4_engine *test_engine_fast;
static ds4_engine *test_engine_quality;
static const char *test_model_path(void) {
const char *model_path = getenv("DS4_TEST_MODEL");
return (model_path && model_path[0]) ? model_path : "ds4flash.gguf";
}
static ds4_engine *test_get_engine(bool quality) {
ds4_engine **slot = quality ? &test_engine_quality : &test_engine_fast;
if (*slot) return *slot;
ds4_engine_options opt = {
.model_path = test_model_path(),
.backend = DS4_BACKEND_METAL,
.quality = quality,
};
TEST_ASSERT(ds4_engine_open(slot, &opt) == 0);
return *slot;
}
static void test_close_engines(void) {
ds4_engine_close(test_engine_fast);
ds4_engine_close(test_engine_quality);
test_engine_fast = NULL;
test_engine_quality = NULL;
}
static void test_close_engine(bool quality) {
ds4_engine **slot = quality ? &test_engine_quality : &test_engine_fast;
ds4_engine_close(*slot);
*slot = NULL;
}
static uint64_t test_round_up_u64(uint64_t n, uint64_t align) {
return (n + align - 1) & ~(align - 1);
}
static uint16_t test_float_to_f16(float f) {
union {
float f;
uint32_t u;
} v = { .f = f };
uint32_t sign = (v.u >> 16) & 0x8000u;
int32_t exp = (int32_t)((v.u >> 23) & 0xffu) - 127 + 15;
uint32_t mant = v.u & 0x7fffffu;
if (exp <= 0) {
if (exp < -10) return (uint16_t)sign;
mant |= 0x800000u;
uint32_t shift = (uint32_t)(14 - exp);
uint32_t half_mant = mant >> shift;
if ((mant >> (shift - 1)) & 1u) half_mant++;
return (uint16_t)(sign | half_mant);
}
if (exp >= 31) return (uint16_t)(sign | 0x7c00u);
uint32_t half = sign | ((uint32_t)exp << 10) | (mant >> 13);
if (mant & 0x1000u) half++;
return (uint16_t)half;
}
static void test_metal_f16_matvec_fast_nr0_4(void) {
/*
* This is the short regression for the long-context repetition failure.
* Decode uses one-token F16 matvecs for several DS4 projections; the fast
* nr0=4 variant must be numerically equivalent to the plain kernel.
*/
const uint32_t in_dim = 4096;
const uint32_t out_dim = 512;
const uint64_t weight_bytes = (uint64_t)in_dim * out_dim * sizeof(uint16_t);
const uint64_t weight_alloc = test_round_up_u64(weight_bytes, (uint64_t)getpagesize());
void *weights_raw = NULL;
TEST_ASSERT(posix_memalign(&weights_raw, (size_t)getpagesize(), (size_t)weight_alloc) == 0);
if (!weights_raw) return;
uint16_t *weights = weights_raw;
memset(weights, 0, (size_t)weight_alloc);
for (uint32_t o = 0; o < out_dim; o++) {
for (uint32_t i = 0; i < in_dim; i++) {
float w = (float)((int)((o * 3u + i * 5u) % 23u) - 11) / 64.0f;
weights[(uint64_t)o * in_dim + i] = test_float_to_f16(w);
}
}
ds4_metal_tensor *x = ds4_metal_tensor_alloc((uint64_t)in_dim * sizeof(float));
ds4_metal_tensor *out = ds4_metal_tensor_alloc((uint64_t)out_dim * sizeof(float));
TEST_ASSERT(x != NULL);
TEST_ASSERT(out != NULL);
if (!x || !out) {
ds4_metal_tensor_free(x);
ds4_metal_tensor_free(out);
free(weights_raw);
return;
}
float *x_host = malloc((size_t)in_dim * sizeof(float));
float *out_host = malloc((size_t)out_dim * sizeof(float));
TEST_ASSERT(x_host != NULL);
TEST_ASSERT(out_host != NULL);
if (!x_host || !out_host) {
free(x_host);
free(out_host);
ds4_metal_tensor_free(x);
ds4_metal_tensor_free(out);
free(weights_raw);
return;
}
for (uint32_t i = 0; i < in_dim; i++) {
x_host[i] = (float)((int)(i % 31u) - 15) / 32.0f;
}
TEST_ASSERT(ds4_metal_tensor_write(x, 0, x_host, (uint64_t)in_dim * sizeof(float)) != 0);
TEST_ASSERT(ds4_metal_set_model_map(weights_raw, weight_alloc) != 0);
ds4_metal_set_quality(false);
TEST_ASSERT(ds4_metal_matmul_f16_tensor(out, weights_raw, weight_alloc, 0,
in_dim, out_dim, x, 1) != 0);
TEST_ASSERT(ds4_metal_tensor_read(out, 0, out_host, (uint64_t)out_dim * sizeof(float)) != 0);
float max_abs = 0.0f;
for (uint32_t o = 0; o < out_dim; o++) {
float ref = 0.0f;
for (uint32_t i = 0; i < in_dim; i++) {
float w = (float)((int)((o * 3u + i * 5u) % 23u) - 11) / 64.0f;
ref += w * x_host[i];
}
float err = fabsf(out_host[o] - ref);
if (err > max_abs) max_abs = err;
}
TEST_ASSERT(max_abs < 0.02f);
free(x_host);
free(out_host);
ds4_metal_tensor_free(x);
ds4_metal_tensor_free(out);
free(weights_raw);
}
static char *test_read_file(const char *path) {
FILE *fp = fopen(path, "rb");
if (!fp) return NULL;
if (fseek(fp, 0, SEEK_END) != 0) {
fclose(fp);
return NULL;
}
long len = ftell(fp);
if (len < 0) {
fclose(fp);
return NULL;
}
rewind(fp);
char *s = malloc((size_t)len + 1);
if (!s) {
fclose(fp);
return NULL;
}
size_t nread = fread(s, 1, (size_t)len, fp);
fclose(fp);
if (nread != (size_t)len) {
free(s);
return NULL;
}
s[len] = '\0';
return s;
}
static int test_count_substr(const char *s, const char *needle) {
int count = 0;
size_t n = strlen(needle);
const char *p = s;
while ((p = strstr(p, needle)) != NULL) {
count++;
p += n;
}
return count;
}
static int test_hex_digit(char c) {
if (c >= '0' && c <= '9') return c - '0';
if (c >= 'a' && c <= 'f') return 10 + c - 'a';
if (c >= 'A' && c <= 'F') return 10 + c - 'A';
return -1;
}
static bool test_hex_to_bytes(const char *hex, unsigned char *out, int cap, int *len) {
int n = 0;
while (*hex && !isspace((unsigned char)*hex)) {
int hi = test_hex_digit(hex[0]);
int lo = test_hex_digit(hex[1]);
if (hi < 0 || lo < 0 || n >= cap) return false;
out[n++] = (unsigned char)((hi << 4) | lo);
hex += 2;
}
*len = n;
return true;
}
static bool test_token_bytes_equal(ds4_engine *engine, int token,
const unsigned char *want, int want_len) {
size_t got_len = 0;
char *got = ds4_token_text(engine, token, &got_len);
bool eq = got && got_len == (size_t)want_len &&
memcmp(got, want, (size_t)want_len) == 0;
free(got);
return eq;
}
static void test_long_prefill_progress(void *ud, const char *event, int current, int total) {
(void)ud;
if (strcmp(event, "prefill_chunk")) return;
if (current == 0 || current == total || current % 8192 == 0) {
fprintf(stderr, "ds4-test: long-context prefill %d/%d\n", current, total);
}
}
static void test_long_security_continuation(void) {
const char *prompt_path = getenv("DS4_TEST_LONG_PROMPT");
if (!prompt_path || !prompt_path[0]) {
prompt_path = "tests/long_context_security_prompt.txt";
}
char *prompt_text = test_read_file(prompt_path);
TEST_ASSERT(prompt_text != NULL);
if (!prompt_text) return;
ds4_engine *engine = test_get_engine(false);
if (!engine) {
free(prompt_text);
return;
}
ds4_tokens prompt = {0};
ds4_tokenize_rendered_chat(engine, prompt_text, &prompt);
TEST_ASSERT(prompt.len > 30000);
ds4_session *session = NULL;
TEST_ASSERT(ds4_session_create(&session, engine, 100000) == 0);
if (!session) {
ds4_tokens_free(&prompt);
free(prompt_text);
return;
}
char err[160];
ds4_session_set_progress(session, test_long_prefill_progress, NULL);
TEST_ASSERT(ds4_session_sync(session, &prompt, err, sizeof(err)) == 0);
ds4_session_set_progress(session, NULL, NULL);
buf out = {0};
uint64_t rng = 12345;
int generated = 0;
bool decode_ok = true;
for (; generated < 700; generated++) {
int token = ds4_session_sample(session, 0.8f, 40, 0.95f, 0.05f, &rng);
if (token == ds4_token_eos(engine)) break;
size_t piece_len = 0;
char *piece = ds4_token_text(engine, token, &piece_len);
buf_append(&out, piece, piece_len);
free(piece);
if (ds4_session_eval(session, token, err, sizeof(err)) != 0) {
decode_ok = false;
break;
}
}
const char *text = out.ptr ? out.ptr : "";
TEST_ASSERT(decode_ok);
TEST_ASSERT(generated > 0);
TEST_ASSERT(strstr(text, "</think>") != NULL);
TEST_ASSERT(test_count_substr(text, "</think>") == 1);
TEST_ASSERT(test_count_substr(text, "The most critical security issue") == 1);
TEST_ASSERT(strstr(text, "arbitrary file") != NULL);
buf_free(&out);
ds4_session_free(session);
ds4_tokens_free(&prompt);
free(prompt_text);
}
#define TEST_VEC_MAX_STEPS 16
#define TEST_VEC_MAX_TOP 32
#define TEST_VEC_MAX_TOKEN_BYTES 128
typedef struct {
unsigned char bytes[TEST_VEC_MAX_TOKEN_BYTES];
int len;
float logprob;
} test_vec_top;
typedef struct {
unsigned char selected[TEST_VEC_MAX_TOKEN_BYTES];
int selected_len;
int ntop;
test_vec_top top[TEST_VEC_MAX_TOP];
} test_vec_step;
typedef struct {
char id[96];
char prompt_path[512];
int ctx;
int nsteps;
test_vec_step steps[TEST_VEC_MAX_STEPS];
} test_vec_case;
static char *test_trim_line(char *line) {
while (*line && isspace((unsigned char)*line)) line++;
size_t n = strlen(line);
while (n && isspace((unsigned char)line[n - 1])) line[--n] = '\0';
return line;
}
static bool test_read_vector_case(FILE *fp, test_vec_case *vc) {
char line[2048];
memset(vc, 0, sizeof(*vc));
while (fgets(line, sizeof(line), fp)) {
char *p = test_trim_line(line);
if (!p[0] || p[0] == '#') continue;
if (sscanf(p, "case %95s %d %d %511s",
vc->id, &vc->ctx, &vc->nsteps, vc->prompt_path) == 4) {
TEST_ASSERT(vc->nsteps > 0 && vc->nsteps <= TEST_VEC_MAX_STEPS);
return true;
}
TEST_ASSERT(!"unexpected line before vector case");
}
return false;
}
static bool test_fill_vector_case(FILE *fp, test_vec_case *vc) {
char line[2048];
int step_index = -1;
int top_index = 0;
while (fgets(line, sizeof(line), fp)) {
char *p = test_trim_line(line);
if (!p[0] || p[0] == '#') continue;
if (!strcmp(p, "end")) return true;
if (!strncmp(p, "step ", 5)) {
char hex[TEST_VEC_MAX_TOKEN_BYTES * 2 + 2];
int ntop = 0;
if (sscanf(p, "step %d %257s %d", &step_index, hex, &ntop) != 3) {
TEST_ASSERT(!"bad vector step line");
return false;
}
TEST_ASSERT(step_index >= 0 && step_index < vc->nsteps);
TEST_ASSERT(ntop >= 0 && ntop <= TEST_VEC_MAX_TOP);
vc->steps[step_index].ntop = ntop;
TEST_ASSERT(test_hex_to_bytes(hex,
vc->steps[step_index].selected,
TEST_VEC_MAX_TOKEN_BYTES,
&vc->steps[step_index].selected_len));
top_index = 0;
continue;
}
if (!strncmp(p, "top ", 4)) {
char hex[TEST_VEC_MAX_TOKEN_BYTES * 2 + 2];
float lp = 0.0f;
TEST_ASSERT(step_index >= 0 && step_index < vc->nsteps);
TEST_ASSERT(top_index < vc->steps[step_index].ntop);
if (sscanf(p, "top %257s %f", hex, &lp) != 2) {
TEST_ASSERT(!"bad vector top line");
return false;
}
test_vec_top *top = &vc->steps[step_index].top[top_index++];
top->logprob = lp;
TEST_ASSERT(test_hex_to_bytes(hex, top->bytes,
TEST_VEC_MAX_TOKEN_BYTES, &top->len));
continue;
}
TEST_ASSERT(!"unexpected vector line");
return false;
}
TEST_ASSERT(!"unterminated vector case");
return false;
}
static void test_logprob_vector_case(ds4_engine *engine, const test_vec_case *vc) {
char *prompt_text = test_read_file(vc->prompt_path);
TEST_ASSERT(prompt_text != NULL);
if (!prompt_text) return;
ds4_tokens prompt = {0};
ds4_encode_chat_prompt(engine, "", prompt_text, DS4_THINK_NONE, &prompt);
free(prompt_text);
ds4_session *session = NULL;
TEST_ASSERT(ds4_session_create(&session, engine, vc->ctx) == 0);
if (!session) {
ds4_tokens_free(&prompt);
return;
}
char err[160];
TEST_ASSERT(ds4_session_sync(session, &prompt, err, sizeof(err)) == 0);
ds4_token_score scores[20];
for (int i = 0; i < vc->nsteps; i++) {
const test_vec_step *step = &vc->steps[i];
int nscore = ds4_session_top_logprobs(session, scores, 20);
int token = ds4_session_argmax(session);
if (!test_token_bytes_equal(engine, token, step->selected, step->selected_len)) {
fprintf(stderr, "ds4-test: vector %s step %d selected token mismatch\n",
vc->id, i);
TEST_ASSERT(false);
}
for (int t = 0; t < step->ntop; t++) {
bool found = false;
float local_lp = 0.0f;
for (int j = 0; j < nscore; j++) {
if (scores[j].id < 0) continue;
if (test_token_bytes_equal(engine, scores[j].id,
step->top[t].bytes,
step->top[t].len)) {
found = true;
local_lp = scores[j].logprob;
break;
}
}
if (!found) {
fprintf(stderr, "ds4-test: vector %s step %d official top token missing locally\n",
vc->id, i);
TEST_ASSERT(false);
} else if (fabsf(local_lp - step->top[t].logprob) > 4.0f) {
fprintf(stderr,
"ds4-test: vector %s step %d logprob delta too high: local=%g official=%g\n",
vc->id, i, local_lp, step->top[t].logprob);
TEST_ASSERT(false);
}
}
if (i + 1 < vc->nsteps) {
TEST_ASSERT(ds4_session_eval(session, token, err, sizeof(err)) == 0);
}
}
ds4_session_free(session);
ds4_tokens_free(&prompt);
}
static void test_official_logprob_vectors(void) {
const char *path = getenv("DS4_TEST_VECTOR_FILE");
if (!path || !path[0]) path = "tests/test-vectors/official.vec";
FILE *fp = fopen(path, "rb");
TEST_ASSERT(fp != NULL);
if (!fp) return;
ds4_engine *engine = test_get_engine(false);
if (!engine) {
fclose(fp);
return;
}
test_vec_case vc;
while (test_read_vector_case(fp, &vc)) {
if (!test_fill_vector_case(fp, &vc)) break;
fprintf(stderr, "ds4-test: vector %s\n", vc.id);
test_logprob_vector_case(engine, &vc);
}
fclose(fp);
}
static const char *test_tool_call_request_json(void) {
return
"{"
"\"model\":\"deepseek-v4-flash\","
"\"messages\":[{\"role\":\"user\",\"content\":\"List the files in the current directory. Use the provided tool; do not answer in prose.\"}],"
"\"tools\":[{\"type\":\"function\",\"function\":{"
"\"name\":\"list_files\","
"\"description\":\"List files in a directory.\","
"\"parameters\":{\"type\":\"object\",\"properties\":{"
"\"path\":{\"type\":\"string\",\"description\":\"Directory path to list.\"}"
"},\"required\":[\"path\"]}"
"}}],"
"\"tool_choice\":\"auto\","
"\"think\":false,"
"\"temperature\":0,"
"\"max_tokens\":256,"
"\"stream\":false"
"}";
}
static void test_tool_call_quality_one(bool quality) {
ds4_engine *engine = test_get_engine(quality);
if (!engine) return;
request r;
char err[160];
TEST_ASSERT(parse_chat_request(engine, test_tool_call_request_json(),
512, 32768, &r, err, sizeof(err)));
ds4_session *session = NULL;
TEST_ASSERT(ds4_session_create(&session, engine, 32768) == 0);
if (!session) {
request_free(&r);
return;
}
TEST_ASSERT(ds4_session_sync(session, &r.prompt, err, sizeof(err)) == 0);
buf text = {0};
uint64_t rng = 123;
bool decode_ok = true;
for (int i = 0; i < r.max_tokens; i++) {
int token = ds4_session_sample(session, r.temperature, r.top_k,
r.top_p, r.min_p, &rng);
size_t piece_len = 0;
char *piece = ds4_token_text(engine, token, &piece_len);
buf_append(&text, piece, piece_len);
free(piece);
if (tool_calls_finished(text.ptr ? text.ptr : "")) break;
if (ds4_session_eval(session, token, err, sizeof(err)) != 0) {
decode_ok = false;
break;
}
}
char *content = NULL;
char *reasoning = NULL;
tool_calls calls = {0};
bool parsed = parse_generated_message(text.ptr ? text.ptr : "",
&content, &reasoning, &calls);
TEST_ASSERT(decode_ok);
TEST_ASSERT(parsed);
TEST_ASSERT(calls.len > 0);
TEST_ASSERT(calls.len > 0 && !strcmp(calls.v[0].name, "list_files"));
free(content);
free(reasoning);
tool_calls_free(&calls);
buf_free(&text);
ds4_session_free(session);
request_free(&r);
}
static void test_tool_call_quality(void) {
fprintf(stderr, "ds4-test: tool-call quality fast path\n");
test_tool_call_quality_one(false);
test_close_engine(false);
fprintf(stderr, "ds4-test: tool-call quality exact path\n");
test_tool_call_quality_one(true);
test_close_engine(true);
}
#endif
static void test_server_unit_group(void) {
ds4_server_unit_tests_run();
}
typedef void (*test_fn)(void);
typedef struct {
const char *flag;
const char *name;
const char *desc;
test_fn fn;
} ds4_test_entry;
static const ds4_test_entry test_entries[] = {
#ifndef DS4_NO_METAL
{"--long-context", "long-context", "long Metal continuation regression", test_long_security_continuation},
{"--tool-call-quality", "tool-call-quality", "model emits valid DSML tool calls", test_tool_call_quality},
{"--logprob-vectors", "logprob-vectors", "official API top-logprob vector comparison", test_official_logprob_vectors},
{"--metal-kernels", "metal-kernels", "isolated Metal kernel numeric regressions", test_metal_f16_matvec_fast_nr0_4},
#endif
{"--server", "server", "server parser/rendering/cache unit tests", test_server_unit_group},
};
static void test_print_help(const char *prog) {
printf("Usage: %s [--all | TEST...]\n\n", prog);
puts("Tests:");
puts(" --all");
puts(" Run every test. This is the default, ordered from slower to faster.");
for (size_t i = 0; i < sizeof(test_entries) / sizeof(test_entries[0]); i++) {
printf(" %-20s %s\n", test_entries[i].flag, test_entries[i].desc);
}
puts(" --list");
puts(" Print test names only.");
puts(" -h, --help");
puts(" Show this help.");
puts("\nEnvironment:");
puts(" DS4_TEST_MODEL=FILE Model path. Default: ds4flash.gguf");
puts(" DS4_TEST_LONG_PROMPT=FILE Rendered long-context regression prompt.");
puts(" DS4_TEST_VECTOR_FILE=FILE Simple official-vector fixture.");
}
static const ds4_test_entry *test_find_entry(const char *arg) {
for (size_t i = 0; i < sizeof(test_entries) / sizeof(test_entries[0]); i++) {
if (!strcmp(arg, test_entries[i].flag)) return &test_entries[i];
}
return NULL;
}
static void test_run_entry(const ds4_test_entry *entry) {
int before = test_failures;
fprintf(stderr, "%s:\n", entry->name);
entry->fn();
fprintf(stderr, "%s: %s%s%s\n",
entry->name,
test_failures == before ? "\033[32m" : "\033[31m",
test_failures == before ? "OK" : "ERR",
"\033[0m");
}
int main(int argc, char **argv) {
bool run_all = argc == 1;
bool selected[sizeof(test_entries) / sizeof(test_entries[0])] = {0};
for (int i = 1; i < argc; i++) {
if (!strcmp(argv[i], "--all")) {
run_all = true;
} else if (!strcmp(argv[i], "--list")) {
for (size_t j = 0; j < sizeof(test_entries) / sizeof(test_entries[0]); j++) {
puts(test_entries[j].flag);
}
return 0;
} else if (!strcmp(argv[i], "-h") || !strcmp(argv[i], "--help")) {
test_print_help(argv[0]);
return 0;
} else {
const ds4_test_entry *entry = test_find_entry(argv[i]);
if (!entry) {
fprintf(stderr, "ds4-test: unknown test switch: %s\n", argv[i]);
test_print_help(argv[0]);
return 2;
}
selected[(size_t)(entry - test_entries)] = true;
}
}
if (run_all) {
for (size_t i = 0; i < sizeof(test_entries) / sizeof(test_entries[0]); i++) {
test_run_entry(&test_entries[i]);
}
} else {
for (size_t i = 0; i < sizeof(test_entries) / sizeof(test_entries[0]); i++) {
if (selected[i]) test_run_entry(&test_entries[i]);
}
}
#ifndef DS4_NO_METAL
test_close_engines();
#endif
if (test_failures) {
fprintf(stderr, "ds4 tests: %d failure(s)\n", test_failures);
return 1;
}
puts("ds4 tests: ok");
return 0;
}
tests/long_context_security_prompt.txt
+1762
View File
File diff suppressed because one or more lines are too long
tests/test-vectors/README.md
+40
View File
@@ -0,0 +1,40 @@
# DeepSeek V4 Flash Test Vectors
These vectors were captured from the official DeepSeek V4 Flash API using
`deepseek-v4-flash`, greedy decoding, thinking disabled, and
`top_logprobs=20`. The hosted API does not expose full logits, so these files
store the best logprob slice the API provides.
Files:
- `prompts/*.txt`: exact user prompts.
- `official/*.official.json`: official API continuations and top-logprobs.
- `official.vec`: compact C-test fixture generated from the official JSON.
Regenerate official vectors:
```sh
DEEPSEEK_API_KEY=... ./tests/test-vectors/fetch_official_vectors.py
```
Running the fetcher without `--only` also regenerates `official.vec`.
The C runner consumes `official.vec` directly:
```sh
./ds4_test --logprob-vectors
```
`official.vec` is intentionally trivial to parse from C: each case points to a
prompt file and each expected token is hex-encoded by bytes. The official JSON
files remain in the tree so the compact fixture can be audited against the raw
API response.
To inspect a local top-logprob dump manually:
```sh
./ds4 --metal --nothink -sys "" --temp 0 -n 4 --ctx 16384 \
--prompt-file tests/test-vectors/prompts/long_code_audit.txt \
--dump-logprobs /tmp/long_code_audit.ds4.json \
--logprobs-top-k 20
```
tests/test-vectors/fetch_official_vectors.py
Executable
+277
View File
@@ -0,0 +1,277 @@
#!/usr/bin/env python3
"""Fetch small DeepSeek V4 Flash logprob vectors from the official API.
The API exposes top-logprobs, not full logits. These vectors are therefore
golden continuation slices: useful for catching tokenizer/template/attention
regressions, but not a replacement for a full internal logit dump.
"""
from __future__ import annotations
import argparse
import json
import os
import sys
import time
import urllib.error
import urllib.request
from pathlib import Path
MODEL = "deepseek-v4-flash"
ENDPOINT = "https://api.deepseek.com/chat/completions"
TOP_LOGPROBS = 20
MAX_TOKENS = 4
CTX_BY_ID = {
"short_italian_fact": 16384,
"short_code_completion": 4096,
"short_reasoning_plain": 4096,
"long_memory_archive": 16384,
"long_code_audit": 16384,
}
def long_memory_prompt() -> str:
block = (
"Record {i:03d}: the archive entry says that component alpha keeps a "
"compressed index, component beta keeps raw observations, and component "
"gamma reports anomalies only after the checksum phrase appears. "
"Do not summarize yet; retain the exact final question.\n"
)
body = "".join(block.format(i=i) for i in range(72))
return (
"You are checking a long technical archive. Read the repeated records "
"and answer only the final question with one short sentence.\n\n"
+ body
+ "\nFinal question: which component reports anomalies after the checksum phrase appears?"
)
def long_code_prompt() -> str:
stanza = (
"Function f_{i} validates a queue entry, calls normalize_path(), then "
"appends a compact audit line. The invariant is that strlen() must not "
"be recomputed when a trusted length returned by snprintf() is already "
"available. Security note {i}: reject negative sizes before casting.\n"
)
body = "".join(stanza.format(i=i) for i in range(68))
return (
"Review this generated C-code audit log. After the log, complete the "
"sentence with the most likely next words.\n\n"
+ body
+ "\nCompletion target: The most important code quality issue is"
)
PROMPTS = [
{
"id": "short_italian_fact",
"kind": "short",
"prompt": "Rispondi in italiano con una frase: chi era Ada Lovelace?",
},
{
"id": "short_code_completion",
"kind": "short",
"prompt": "Complete the C statement with the next exact token only:\nreturn snprintf(buf, sizeof(buf), \"%d\", value",
},
{
"id": "short_reasoning_plain",
"kind": "short",
"prompt": "Answer with only the number: 2048 divided by 128 is",
},
{
"id": "long_memory_archive",
"kind": "long",
"prompt": long_memory_prompt(),
},
{
"id": "long_code_audit",
"kind": "long",
"prompt": long_code_prompt(),
},
]
def token_bytes(token: str, value) -> list[int]:
if isinstance(value, list):
return [int(x) for x in value]
return list(token.encode("utf-8"))
def request_vector(api_key: str, prompt: str) -> dict:
payload = {
"model": MODEL,
"messages": [{"role": "user", "content": prompt}],
"temperature": 0,
"max_tokens": MAX_TOKENS,
"logprobs": True,
"top_logprobs": TOP_LOGPROBS,
"thinking": {"type": "disabled"},
"stream": False,
}
req = urllib.request.Request(
ENDPOINT,
data=json.dumps(payload).encode("utf-8"),
headers={
"Authorization": f"Bearer {api_key}",
"Content-Type": "application/json",
},
method="POST",
)
try:
with urllib.request.urlopen(req, timeout=120) as fp:
return json.loads(fp.read().decode("utf-8"))
except urllib.error.HTTPError as e:
body = e.read().decode("utf-8", "replace")
raise RuntimeError(f"DeepSeek API HTTP {e.code}: {body}") from e
def normalize_record(prompt_spec: dict, response: dict) -> dict:
choice = response["choices"][0]
logprob_items = choice.get("logprobs", {}).get("content", []) or []
steps = []
for step, item in enumerate(logprob_items):
top = []
for alt in item.get("top_logprobs", []) or []:
tok = alt.get("token", "")
top.append(
{
"token": {
"text": tok,
"bytes": token_bytes(tok, alt.get("bytes")),
},
"logprob": alt.get("logprob"),
}
)
tok = item.get("token", "")
steps.append(
{
"step": step,
"token": {
"text": tok,
"bytes": token_bytes(tok, item.get("bytes")),
},
"logprob": item.get("logprob"),
"top_logprobs": top,
}
)
return {
"schema": "ds4-official-logprobs-v1",
"source": "deepseek-official-api",
"model": MODEL,
"endpoint": ENDPOINT,
"created_at": time.strftime("%Y-%m-%dT%H:%M:%SZ", time.gmtime()),
"id": prompt_spec["id"],
"kind": prompt_spec["kind"],
"prompt": prompt_spec["prompt"],
"request": {
"model": MODEL,
"temperature": 0,
"max_tokens": MAX_TOKENS,
"logprobs": True,
"top_logprobs": TOP_LOGPROBS,
"thinking": {"type": "disabled"},
"messages": [{"role": "user", "content": prompt_spec["prompt"]}],
},
"usage": response.get("usage"),
"finish_reason": choice.get("finish_reason"),
"message": choice.get("message", {}),
"logits_available": False,
"steps": steps,
}
def hex_bytes(values: list[int]) -> str:
return "".join(f"{int(x):02x}" for x in values)
def write_compact_fixture(root: Path, manifest: dict) -> None:
lines = [
"# ds4-official-logprob-vectors-v1",
"# case <id> <ctx> <steps> <prompt-file>",
"# step <index> <selected-hex> <top-count>",
"# top <token-hex> <official-logprob>",
"",
]
for prompt in manifest["prompts"]:
vector_id = prompt["id"]
record = json.loads((root / prompt["official_file"]).read_text(encoding="utf-8"))
steps = record["steps"]
prompt_file = root / prompt["prompt_file"]
lines.append(f"case {vector_id} {CTX_BY_ID[vector_id]} {len(steps)} {prompt_file}")
for i, step in enumerate(steps):
top = []
for alt in step.get("top_logprobs", []):
lp = float(alt.get("logprob", -9999))
if lp <= -1000:
continue
token_hex = hex_bytes(alt["token"]["bytes"])
if token_hex:
top.append((token_hex, lp))
lines.append(f"step {i} {hex_bytes(step['token']['bytes'])} {len(top)}")
for token_hex, lp in top:
lines.append(f"top {token_hex} {lp:.9g}")
lines.append("end")
lines.append("")
(root / "official.vec").write_text("\n".join(lines), encoding="ascii")
def main() -> int:
parser = argparse.ArgumentParser(description=__doc__)
parser.add_argument("--out", default="tests/test-vectors", help="output directory")
parser.add_argument("--only", action="append", help="fetch only the named prompt id")
args = parser.parse_args()
api_key = os.environ.get("DEEPSEEK_API_KEY")
if not api_key:
print("DEEPSEEK_API_KEY is required", file=sys.stderr)
return 2
root = Path(args.out)
prompt_dir = root / "prompts"
official_dir = root / "official"
prompt_dir.mkdir(parents=True, exist_ok=True)
official_dir.mkdir(parents=True, exist_ok=True)
wanted = set(args.only or [])
manifest = {
"schema": "ds4-test-vector-manifest-v1",
"source": "deepseek-official-api",
"model": MODEL,
"endpoint": ENDPOINT,
"top_logprobs": TOP_LOGPROBS,
"max_tokens": MAX_TOKENS,
"prompts": [],
}
for spec in PROMPTS:
if wanted and spec["id"] not in wanted:
continue
prompt_path = prompt_dir / f"{spec['id']}.txt"
prompt_path.write_text(spec["prompt"], encoding="utf-8")
response = request_vector(api_key, spec["prompt"])
record = normalize_record(spec, response)
out_path = official_dir / f"{spec['id']}.official.json"
out_path.write_text(json.dumps(record, ensure_ascii=False, indent=2) + "\n", encoding="utf-8")
manifest["prompts"].append(
{
"id": spec["id"],
"kind": spec["kind"],
"prompt_file": str(prompt_path.relative_to(root)),
"official_file": str(out_path.relative_to(root)),
"prompt_chars": len(spec["prompt"]),
"steps": len(record["steps"]),
}
)
print(f"wrote {out_path}")
(root / "manifest.json").write_text(json.dumps(manifest, indent=2) + "\n", encoding="utf-8")
if not wanted:
write_compact_fixture(root, manifest)
return 0
if __name__ == "__main__":
raise SystemExit(main())
tests/test-vectors/manifest.json
+50
View File
@@ -0,0 +1,50 @@
{
"schema": "ds4-test-vector-manifest-v1",
"source": "deepseek-official-api",
"model": "deepseek-v4-flash",
"endpoint": "https://api.deepseek.com/chat/completions",
"top_logprobs": 20,
"max_tokens": 4,
"prompts": [
{
"id": "short_italian_fact",
"kind": "short",
"prompt_file": "prompts/short_italian_fact.txt",
"official_file": "official/short_italian_fact.official.json",
"prompt_chars": 57,
"steps": 4
},
{
"id": "short_code_completion",
"kind": "short",
"prompt_file": "prompts/short_code_completion.txt",
"official_file": "official/short_code_completion.official.json",
"prompt_chars": 102,
"steps": 4
},
{
"id": "short_reasoning_plain",
"kind": "short",
"prompt_file": "prompts/short_reasoning_plain.txt",
"official_file": "official/short_reasoning_plain.official.json",
"prompt_chars": 51,
"steps": 1
},
{
"id": "long_memory_archive",
"kind": "long",
"prompt_file": "prompts/long_memory_archive.txt",
"official_file": "official/long_memory_archive.official.json",
"prompt_chars": 18503,
"steps": 4
},
{
"id": "long_code_audit",
"kind": "long",
"prompt_file": "prompts/long_code_audit.txt",
"official_file": "official/long_code_audit.official.json",
"prompt_chars": 18851,
"steps": 4
}
]
}
tests/test-vectors/official.vec
+53
View File
@@ -0,0 +1,53 @@
# ds4-official-logprob-vectors-v1
# case <id> <ctx> <steps> <prompt-file>
# step <index> <selected-hex> <top-count>
# top <token-hex> <official-logprob>
case short_italian_fact 16384 4 tests/test-vectors/prompts/short_italian_fact.txt
step 0 416461 1
top 416461 0
step 1 204c6f76 1
top 204c6f76 0
step 2 656c 1
top 656c 0
step 3 616365 1
top 616365 0
end
case short_code_completion 4096 4 tests/test-vectors/prompts/short_code_completion.txt
step 0 606060 1
top 606060 0
step 1 63 1
top 63 0
step 2 0a 1
top 0a 0
step 3 72657475726e 1
top 72657475726e 0
end
case short_reasoning_plain 4096 1 tests/test-vectors/prompts/short_reasoning_plain.txt
step 0 3136 1
top 3136 0
end
case long_memory_archive 16384 4 tests/test-vectors/prompts/long_memory_archive.txt
step 0 436f6d706f6e656e74 1
top 436f6d706f6e656e74 0
step 1 2067616d6d61 1
top 2067616d6d61 0
step 2 207265706f727473 1
top 207265706f727473 0
step 3 20616e6f6d616c696573 1
top 20616e6f6d616c696573 0
end
case long_code_audit 16384 4 tests/test-vectors/prompts/long_code_audit.txt
step 0 546865 1
top 546865 0
step 1 206d6f7374 1
top 206d6f7374 0
step 2 20696d706f7274616e74 1
top 20696d706f7274616e74 0
step 3 20636f6465 1
top 20636f6465 0
end
tests/test-vectors/official/long_code_audit.official.json
+1138
View File
File diff suppressed because one or more lines are too long
tests/test-vectors/official/long_memory_archive.official.json
+1075
View File
File diff suppressed because one or more lines are too long
tests/test-vectors/official/short_code_completion.official.json
+1038
View File
File diff suppressed because it is too large Load Diff
tests/test-vectors/official/short_italian_fact.official.json
+982
View File
@@ -0,0 +1,982 @@
{
"schema": "ds4-official-logprobs-v1",
"source": "deepseek-official-api",
"model": "deepseek-v4-flash",
"endpoint": "https://api.deepseek.com/chat/completions",
"created_at": "2026-05-06T22:09:33Z",
"id": "short_italian_fact",
"kind": "short",
"prompt": "Rispondi in italiano con una frase: chi era Ada Lovelace?",
"request": {
"model": "deepseek-v4-flash",
"temperature": 0,
"max_tokens": 4,
"logprobs": true,
"top_logprobs": 20,
"thinking": {
"type": "disabled"
},
"messages": [
{
"role": "user",
"content": "Rispondi in italiano con una frase: chi era Ada Lovelace?"
}
]
},
"usage": {
"prompt_tokens": 21,
"completion_tokens": 4,
"total_tokens": 25,
"prompt_tokens_details": {
"cached_tokens": 0
},
"prompt_cache_hit_tokens": 0,
"prompt_cache_miss_tokens": 21
},
"finish_reason": "length",
"message": {
"role": "assistant",
"content": "Ada Lovelace"
},
"logits_available": false,
"steps": [
{
"step": 0,
"token": {
"text": "Ada",
"bytes": [
65,
100,
97
]
},
"logprob": 0.0,
"top_logprobs": [
{
"token": {
"text": "Ada",
"bytes": [
65,
100,
97
]
},
"logprob": 0.0
},
{
"token": {
"text": "August",
"bytes": [
65,
117,
103,
117,
115,
116
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Lady",
"bytes": [
76,
97,
100,
121
]
},
"logprob": -9999.0
},
{
"token": {
"text": "**",
"bytes": [
42,
42
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Una",
"bytes": [
85,
110,
97
]
},
"logprob": -9999.0
},
{
"token": {
"text": "La",
"bytes": [
76,
97
]
},
"logprob": -9999.0
},
{
"token": {
"text": " Ada",
"bytes": [
32,
65,
100,
97
]
},
"logprob": -9999.0
},
{
"token": {
"text": "E",
"bytes": [
69
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Consider",
"bytes": [
67,
111,
110,
115,
105,
100,
101,
114
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Ad",
"bytes": [
65,
100
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Chi",
"bytes": [
67,
104,
105
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Mat",
"bytes": [
77,
97,
116
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Fig",
"bytes": [
70,
105,
103
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Cert",
"bytes": [
67,
101,
114,
116
]
},
"logprob": -9999.0
},
{
"token": {
"text": "L",
"bytes": [
76
]
},
"logprob": -9999.0
},
{
"token": {
"text": "A",
"bytes": [
65
]
},
"logprob": -9999.0
},
{
"token": {
"text": "\"",
"bytes": [
34
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Il",
"bytes": [
73,
108
]
},
"logprob": -9999.0
},
{
"token": {
"text": "È",
"bytes": [
195,
136
]
},
"logprob": -9999.0
},
{
"token": {
"text": "C",
"bytes": [
67
]
},
"logprob": -9999.0
}
]
},
{
"step": 1,
"token": {
"text": " Lov",
"bytes": [
32,
76,
111,
118
]
},
"logprob": 0.0,
"top_logprobs": [
{
"token": {
"text": " Lov",
"bytes": [
32,
76,
111,
118
]
},
"logprob": 0.0
},
{
"token": {
"text": "<｜｜end▁of▁sentence｜｜>",
"bytes": [
60,
239,
189,
156,
239,
189,
156,
101,
110,
100,
226,
150,
129,
111,
102,
226,
150,
129,
115,
101,
110,
116,
101,
110,
99,
101,
239,
189,
156,
239,
189,
156,
62
]
},
"logprob": -9999.0
},
{
"token": {
"text": "!",
"bytes": [
33
]
},
"logprob": -9999.0
},
{
"token": {
"text": "\"",
"bytes": [
34
]
},
"logprob": -9999.0
},
{
"token": {
"text": "#",
"bytes": [
35
]
},
"logprob": -9999.0
},
{
"token": {
"text": "%",
"bytes": [
37
]
},
"logprob": -9999.0
},
{
"token": {
"text": "&",
"bytes": [
38
]
},
"logprob": -9999.0
},
{
"token": {
"text": "'",
"bytes": [
39
]
},
"logprob": -9999.0
},
{
"token": {
"text": "(",
"bytes": [
40
]
},
"logprob": -9999.0
},
{
"token": {
"text": ")",
"bytes": [
41
]
},
"logprob": -9999.0
},
{
"token": {
"text": "*",
"bytes": [
42
]
},
"logprob": -9999.0
},
{
"token": {
"text": "+",
"bytes": [
43
]
},
"logprob": -9999.0
},
{
"token": {
"text": ",",
"bytes": [
44
]
},
"logprob": -9999.0
},
{
"token": {
"text": "-",
"bytes": [
45
]
},
"logprob": -9999.0
},
{
"token": {
"text": ".",
"bytes": [
46
]
},
"logprob": -9999.0
},
{
"token": {
"text": "/",
"bytes": [
47
]
},
"logprob": -9999.0
},
{
"token": {
"text": "1",
"bytes": [
49
]
},
"logprob": -9999.0
},
{
"token": {
"text": "2",
"bytes": [
50
]
},
"logprob": -9999.0
},
{
"token": {
"text": "3",
"bytes": [
51
]
},
"logprob": -9999.0
},
{
"token": {
"text": " &",
"bytes": [
32,
38
]
},
"logprob": -9999.0
}
]
},
{
"step": 2,
"token": {
"text": "el",
"bytes": [
101,
108
]
},
"logprob": 0.0,
"top_logprobs": [
{
"token": {
"text": "el",
"bytes": [
101,
108
]
},
"logprob": 0.0
},
{
"token": {
"text": "<｜｜end▁of▁sentence｜｜>",
"bytes": [
60,
239,
189,
156,
239,
189,
156,
101,
110,
100,
226,
150,
129,
111,
102,
226,
150,
129,
115,
101,
110,
116,
101,
110,
99,
101,
239,
189,
156,
239,
189,
156,
62
]
},
"logprob": -9999.0
},
{
"token": {
"text": "!",
"bytes": [
33
]
},
"logprob": -9999.0
},
{
"token": {
"text": "ateg",
"bytes": [
97,
116,
101,
103
]
},
"logprob": -9999.0
},
{
"token": {
"text": "\"",
"bytes": [
34
]
},
"logprob": -9999.0
},
{
"token": {
"text": "#",
"bytes": [
35
]
},
"logprob": -9999.0
},
{
"token": {
"text": "%",
"bytes": [
37
]
},
"logprob": -9999.0
},
{
"token": {
"text": "&",
"bytes": [
38
]
},
"logprob": -9999.0
},
{
"token": {
"text": "'",
"bytes": [
39
]
},
"logprob": -9999.0
},
{
"token": {
"text": "(",
"bytes": [
40
]
},
"logprob": -9999.0
},
{
"token": {
"text": "í",
"bytes": [
195,
173
]
},
"logprob": -9999.0
},
{
"token": {
"text": "che",
"bytes": [
99,
104,
101
]
},
"logprob": -9999.0
},
{
"token": {
"text": "ari",
"bytes": [
97,
114,
105
]
},
"logprob": -9999.0
},
{
"token": {
"text": "*",
"bytes": [
42
]
},
"logprob": -9999.0
},
{
"token": {
"text": ",",
"bytes": [
44
]
},
"logprob": -9999.0
},
{
"token": {
"text": "-",
"bytes": [
45
]
},
"logprob": -9999.0
},
{
"token": {
"text": ".",
"bytes": [
46
]
},
"logprob": -9999.0
},
{
"token": {
"text": "/",
"bytes": [
47
]
},
"logprob": -9999.0
},
{
"token": {
"text": "0",
"bytes": [
48
]
},
"logprob": -9999.0
},
{
"token": {
"text": "1",
"bytes": [
49
]
},
"logprob": -9999.0
}
]
},
{
"step": 3,
"token": {
"text": "ace",
"bytes": [
97,
99,
101
]
},
"logprob": 0.0,
"top_logprobs": [
{
"token": {
"text": "ace",
"bytes": [
97,
99,
101
]
},
"logprob": 0.0
},
{
"token": {
"text": "<｜｜end▁of▁sentence｜｜>",
"bytes": [
60,
239,
189,
156,
239,
189,
156,
101,
110,
100,
226,
150,
129,
111,
102,
226,
150,
129,
115,
101,
110,
116,
101,
110,
99,
101,
239,
189,
156,
239,
189,
156,
62
]
},
"logprob": -9999.0
},
{
"token": {
"text": "ateg",
"bytes": [
97,
116,
101,
103
]
},
"logprob": -9999.0
},
{
"token": {
"text": "'",
"bytes": [
39
]
},
"logprob": -9999.0
},
{
"token": {
"text": "(",
"bytes": [
40
]
},
"logprob": -9999.0
},
{
"token": {
"text": "che",
"bytes": [
99,
104,
101
]
},
"logprob": -9999.0
},
{
"token": {
"text": "*",
"bytes": [
42
]
},
"logprob": -9999.0
},
{
"token": {
"text": ",",
"bytes": [
44
]
},
"logprob": -9999.0
},
{
"token": {
"text": "-",
"bytes": [
45
]
},
"logprob": -9999.0
},
{
"token": {
"text": ".",
"bytes": [
46
]
},
"logprob": -9999.0
},
{
"token": {
"text": "/",
"bytes": [
47
]
},
"logprob": -9999.0
},
{
"token": {
"text": "ike",
"bytes": [
105,
107,
101
]
},
"logprob": -9999.0
},
{
"token": {
"text": "ole",
"bytes": [
111,
108,
101
]
},
"logprob": -9999.0
},
{
"token": {
"text": "ances",
"bytes": [
97,
110,
99,
101,
115
]
},
"logprob": -9999.0
},
{
"token": {
"text": "2",
"bytes": [
50
]
},
"logprob": -9999.0
},
{
"token": {
"text": "4",
"bytes": [
52
]
},
"logprob": -9999.0
},
{
"token": {
"text": "8",
"bytes": [
56
]
},
"logprob": -9999.0
},
{
"token": {
"text": "9",
"bytes": [
57
]
},
"logprob": -9999.0
},
{
"token": {
"text": "<",
"bytes": [
60
]
},
"logprob": -9999.0
},
{
"token": {
"text": "way",
"bytes": [
119,
97,
121
]
},
"logprob": -9999.0
}
]
}
]
}
tests/test-vectors/official/short_reasoning_plain.official.json
+287
View File
@@ -0,0 +1,287 @@
{
"schema": "ds4-official-logprobs-v1",
"source": "deepseek-official-api",
"model": "deepseek-v4-flash",
"endpoint": "https://api.deepseek.com/chat/completions",
"created_at": "2026-05-06T22:09:34Z",
"id": "short_reasoning_plain",
"kind": "short",
"prompt": "Answer with only the number: 2048 divided by 128 is",
"request": {
"model": "deepseek-v4-flash",
"temperature": 0,
"max_tokens": 4,
"logprobs": true,
"top_logprobs": 20,
"thinking": {
"type": "disabled"
},
"messages": [
{
"role": "user",
"content": "Answer with only the number: 2048 divided by 128 is"
}
]
},
"usage": {
"prompt_tokens": 18,
"completion_tokens": 1,
"total_tokens": 19,
"prompt_tokens_details": {
"cached_tokens": 0
},
"prompt_cache_hit_tokens": 0,
"prompt_cache_miss_tokens": 18
},
"finish_reason": "stop",
"message": {
"role": "assistant",
"content": "16"
},
"logits_available": false,
"steps": [
{
"step": 0,
"token": {
"text": "16",
"bytes": [
49,
54
]
},
"logprob": 0.0,
"top_logprobs": [
{
"token": {
"text": "16",
"bytes": [
49,
54
]
},
"logprob": 0.0
},
{
"token": {
"text": "204",
"bytes": [
50,
48,
52
]
},
"logprob": -9999.0
},
{
"token": {
"text": "<｜｜end▁of▁sentence｜｜>",
"bytes": [
60,
239,
189,
156,
239,
189,
156,
101,
110,
100,
226,
150,
129,
111,
102,
226,
150,
129,
115,
101,
110,
116,
101,
110,
99,
101,
239,
189,
156,
239,
189,
156,
62
]
},
"logprob": -9999.0
},
{
"token": {
"text": "To",
"bytes": [
84,
111
]
},
"logprob": -9999.0
},
{
"token": {
"text": "15",
"bytes": [
49,
53
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Let",
"bytes": [
76,
101,
116
]
},
"logprob": -9999.0
},
{
"token": {
"text": "The",
"bytes": [
84,
104,
101
]
},
"logprob": -9999.0
},
{
"token": {
"text": "14",
"bytes": [
49,
52
]
},
"logprob": -9999.0
},
{
"token": {
"text": "We",
"bytes": [
87,
101
]
},
"logprob": -9999.0
},
{
"token": {
"text": "1",
"bytes": [
49
]
},
"logprob": -9999.0
},
{
"token": {
"text": "\n",
"bytes": [
10
]
},
"logprob": -9999.0
},
{
"token": {
"text": "64",
"bytes": [
54,
52
]
},
"logprob": -9999.0
},
{
"token": {
"text": "8",
"bytes": [
56
]
},
"logprob": -9999.0
},
{
"token": {
"text": "4",
"bytes": [
52
]
},
"logprob": -9999.0
},
{
"token": {
"text": "Since",
"bytes": [
83,
105,
110,
99,
101
]
},
"logprob": -9999.0
},
{
"token": {
"text": "0",
"bytes": [
48
]
},
"logprob": -9999.0
},
{
"token": {
"text": "I",
"bytes": [
73
]
},
"logprob": -9999.0
},
{
"token": {
"text": "18",
"bytes": [
49,
56
]
},
"logprob": -9999.0
},
{
"token": {
"text": "10",
"bytes": [
49,
48
]
},
"logprob": -9999.0
},
{
"token": {
"text": "6",
"bytes": [
54
]
},
"logprob": -9999.0
}
]
}
]
}
tests/test-vectors/prompts/long_code_audit.txt
+72
View File
@@ -0,0 +1,72 @@
Review this generated C-code audit log. After the log, complete the sentence with the most likely next words.
Function f_0 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 0: reject negative sizes before casting.
Function f_1 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 1: reject negative sizes before casting.
Function f_2 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 2: reject negative sizes before casting.
Function f_3 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 3: reject negative sizes before casting.
Function f_4 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 4: reject negative sizes before casting.
Function f_5 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 5: reject negative sizes before casting.
Function f_6 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 6: reject negative sizes before casting.
Function f_7 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 7: reject negative sizes before casting.
Function f_8 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 8: reject negative sizes before casting.
Function f_9 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 9: reject negative sizes before casting.
Function f_10 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 10: reject negative sizes before casting.
Function f_11 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 11: reject negative sizes before casting.
Function f_12 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 12: reject negative sizes before casting.
Function f_13 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 13: reject negative sizes before casting.
Function f_14 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 14: reject negative sizes before casting.
Function f_15 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 15: reject negative sizes before casting.
Function f_16 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 16: reject negative sizes before casting.
Function f_17 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 17: reject negative sizes before casting.
Function f_18 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 18: reject negative sizes before casting.
Function f_19 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 19: reject negative sizes before casting.
Function f_20 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 20: reject negative sizes before casting.
Function f_21 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 21: reject negative sizes before casting.
Function f_22 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 22: reject negative sizes before casting.
Function f_23 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 23: reject negative sizes before casting.
Function f_24 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 24: reject negative sizes before casting.
Function f_25 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 25: reject negative sizes before casting.
Function f_26 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 26: reject negative sizes before casting.
Function f_27 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 27: reject negative sizes before casting.
Function f_28 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 28: reject negative sizes before casting.
Function f_29 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 29: reject negative sizes before casting.
Function f_30 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 30: reject negative sizes before casting.
Function f_31 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 31: reject negative sizes before casting.
Function f_32 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 32: reject negative sizes before casting.
Function f_33 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 33: reject negative sizes before casting.
Function f_34 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 34: reject negative sizes before casting.
Function f_35 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 35: reject negative sizes before casting.
Function f_36 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 36: reject negative sizes before casting.
Function f_37 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 37: reject negative sizes before casting.
Function f_38 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 38: reject negative sizes before casting.
Function f_39 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 39: reject negative sizes before casting.
Function f_40 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 40: reject negative sizes before casting.
Function f_41 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 41: reject negative sizes before casting.
Function f_42 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 42: reject negative sizes before casting.
Function f_43 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 43: reject negative sizes before casting.
Function f_44 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 44: reject negative sizes before casting.
Function f_45 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 45: reject negative sizes before casting.
Function f_46 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 46: reject negative sizes before casting.
Function f_47 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 47: reject negative sizes before casting.
Function f_48 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 48: reject negative sizes before casting.
Function f_49 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 49: reject negative sizes before casting.
Function f_50 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 50: reject negative sizes before casting.
Function f_51 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 51: reject negative sizes before casting.
Function f_52 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 52: reject negative sizes before casting.
Function f_53 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 53: reject negative sizes before casting.
Function f_54 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 54: reject negative sizes before casting.
Function f_55 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 55: reject negative sizes before casting.
Function f_56 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 56: reject negative sizes before casting.
Function f_57 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 57: reject negative sizes before casting.
Function f_58 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 58: reject negative sizes before casting.
Function f_59 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 59: reject negative sizes before casting.
Function f_60 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 60: reject negative sizes before casting.
Function f_61 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 61: reject negative sizes before casting.
Function f_62 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 62: reject negative sizes before casting.
Function f_63 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 63: reject negative sizes before casting.
Function f_64 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 64: reject negative sizes before casting.
Function f_65 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 65: reject negative sizes before casting.
Function f_66 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 66: reject negative sizes before casting.
Function f_67 validates a queue entry, calls normalize_path(), then appends a compact audit line. The invariant is that strlen() must not be recomputed when a trusted length returned by snprintf() is already available. Security note 67: reject negative sizes before casting.
Completion target: The most important code quality issue is
tests/test-vectors/prompts/long_memory_archive.txt
+76
View File
@@ -0,0 +1,76 @@
You are checking a long technical archive. Read the repeated records and answer only the final question with one short sentence.
Record 000: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 001: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 002: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 003: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 004: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 005: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 006: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 007: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 008: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 009: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 010: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 011: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 012: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 013: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 014: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 015: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 016: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 017: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 018: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 019: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 020: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 021: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 022: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 023: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 024: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 025: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 026: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 027: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 028: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 029: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 030: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 031: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 032: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 033: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 034: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 035: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 036: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 037: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 038: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 039: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 040: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 041: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 042: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 043: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 044: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 045: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 046: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 047: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 048: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 049: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 050: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 051: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 052: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 053: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 054: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 055: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 056: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 057: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 058: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 059: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 060: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 061: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 062: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 063: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 064: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 065: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 066: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 067: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 068: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 069: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 070: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Record 071: the archive entry says that component alpha keeps a compressed index, component beta keeps raw observations, and component gamma reports anomalies only after the checksum phrase appears. Do not summarize yet; retain the exact final question.
Final question: which component reports anomalies after the checksum phrase appears?
tests/test-vectors/prompts/short_code_completion.txt
+2
View File
@@ -0,0 +1,2 @@
Complete the C statement with the next exact token only:
return snprintf(buf, sizeof(buf), "%d", value
tests/test-vectors/prompts/short_italian_fact.txt
+1
View File
@@ -0,0 +1 @@
Rispondi in italiano con una frase: chi era Ada Lovelace?
tests/test-vectors/prompts/short_reasoning_plain.txt
+1
View File
@@ -0,0 +1 @@
Answer with only the number: 2048 divided by 128 is