⚡ DeepSeek V4 Flash on AMD Strix Halo
Up to 32 tok/s decode — a tuned gfx1151 fork of antirez/ds4 with ROCmFPX tooling, SSD expert streaming, and native DeepSeek-V4-Flash-0731 support
284B MoE · 43 routed layers · 256 experts/layer · 128 GB unified memory · 100% local
--- ## 🚀 The Short Version Run a **284B-parameter** model at **up to 32 tok/s** on a single AMD APU — no cloud, no discrete GPU, no per-token bill. This fork turns a Strix Halo mini-PC (Ryzen AI MAX+ 395 / Radeon 8060S) into a local DeepSeek V4 Flash inference machine. That number is not a fantasy. It's the same hardware configuration that took the **LocalMaxxing** DeepSeek V4 Flash record: | Engine / Quant | Decode (tok/s) | | :--- | ---: | | DwarfStar · Q2_K | 15.60 | | HipFire · MQ2 + MTP | 18.99 | | **This stack · ROCmFPX + DSpark** | **32.00** | That's **2.05× faster** than the previous unified-memory leader and **68.5% ahead** of the runner-up — measured on the exact silicon this repo targets. --- ## 🧠 What Makes ROCmFPX So Good ### The format **ROCmFPX is not one quantization format — it's a family of block formats designed for AMD ROCm/HIP silicon.** Each block holds **32 weights** as packed low-bit codes plus one or two tiny scales, and the GPU kernels are written for exactly that byte layout: | Variant | Block size | Bits / weight | Typical use | | :--- | ---: | ---: | --- | | **ROCmFP2** | 10 B / 32 | **2.50** | Routed expert gate & up matrices (the biggest tensors) | | **ROCmFP3** | — | **3.50** | Expert down projections | | **ROCmFP4** | — | **4.25** | Dense / sensitive projections | A Strix-specific mixed-precision recipe combines them: the enormous routed-expert gate/up matrices at ROCmFP2, expert down at ROCmFP3, and dense projections at ROCmFP4+. With an importance matrix during quantization, the full DeepSeek V4 Flash target lands at **~2.88 bits per parameter** in a **102.3 GB** file — just under 95.3 GiB, so the whole model fits in Strix Halo's 128 GB unified pool with room to spare. ### Why it's fast The format is inseparable from the kernel that eats it: 1. **Register-resident codebooks.** Kernels expand the tiny packed codebooks **in GPU registers** using AMD's byte-permute instruction (`v_perm_b32`) instead of doing a separate gather from memory. No indirection, no extra loads. 2. **Integer dot products.** Packed blocks feed `dp4a`-style integer dot products directly — the fastest path on RDNA3/3.5 — instead of dequantizing to floats first. 3. **Designed as one path.** The file layout and the HIP kernel are specified together, so decode is a straight memory-traffic-bound stream of weights through fixed-purpose hardware. At batch one, every generated token streams the active experts across all 43 layers, so **decode is memory-traffic-limited** — and ROCmFPX is built to maximize useful bytes per load. ### The engine side ROCmFPX only handles the weight traffic. The full 32 tok/s profile also uses: - **DSpark draft + fused q=4 verification** — a small 3-layer drafter proposes up to 3 tokens; the 284B target verifies 4 positions in one fused HIP graph pass (26.4% faster than autoregressive). - **Weight reuse across verification columns** — each packed dense weight is decoded once and applied to all 4 verify columns (+2.1–2.3%). - **Indexed sparse prefill** — ~250 tok/s on 8K prompts via DeepSeek's learned indexer. --- ## 🛠️ What This Fork Improves Over Upstream `antirez/ds4` The upstream repo brought the initial DeepSeek V4 ROCm backend. This fork makes it actually run *well* on Strix Halo: ### 1. Fixed the SSD Expert Streaming Slab Allocator (`src/ds4.c`) Mixed-precision GGUFs have layers with different per-expert sizes (e.g. 0731's Layer 26 uses `IQ2_S` gate/up at 82 B/256 vs `IQ2_XXS` at 66 B/256). Upstream pinned the streaming cache to the *first* layer's size, which bounced Layer 26 into pageable mapped views — an MMU fault on ROCm. Now the slab is sized to the **maximum** across all 43 layers, so **0/43 layers fall off the fast path**. ### 2. New GPU kernels (`src/rocm/`) - **`Q4_K` token embedding** kernels (`embed_token_hc_q4_k_kernel`, `embed_tokens_hc_q4_k_kernel`) - **`Q4_K` dense matmul** kernels (`matmul_q4_k_f32_sharedx_warp_rows_w32_kernel`, `matmul_q4_k_f32_batch_warp8_kernel`) - **`BF16` dense matmul** kernel (`matmul_bf16_ordered_chunks_kernel`) ### 3. `MXFP4` → `Q2_K` requantization tool (`gguf-tools/requant_down_q2k.c`) In-place GGUF converter with bit-exact `MXFP4` dequantization and correct element interleaving — converts `IQ3_XXS`/`MXFP4` down experts to custom `Q2_K` in under 2 minutes. ### 4. ROCm 7.2.x diagnostics & TTM auto-sizing `make rocm-diag`, `make rocm-doctor`, `make rocm-smoke`, `make rocm-bench-quick`, and `DS4_ROCM_TTM_AUTORAISE=1` — detect bad configs and fix them automatically. --- ## 📦 Two Ways to Run DeepSeek-V4-Flash-0731 ### Route A — High-throughput ROCmFPX (~32 tok/s) The ROCmFPX/ROCmFP2 target (`Q2_0_ROCMFPX`, ~98–102 GB) plus the DSpark drafter. This is the LocalMaxxing record path. ```sh ./download_model.sh rocmfpx-strix # 102.3 GB ROCmFP2-STRIX target ./download_model.sh dspark-drafter # 11.3 GB DSpark draft ./run-deepseek-v4.sh # 32 tok/s high-throughput server ``` ### Route B — Native ds4fa quant recipe (~13 tok/s, zero kernel gaps) Uses standard GGML quants that the ds4fa engine supports natively: - **Model**: `tekosML/DeepSeek-V4-Flash-0731-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-imatrix.gguf` (**86.72 GB**) - **HF model page**: [tekosML/DeepSeek-V4-Flash-0731-GGUF-GX10](https://huggingface.co/tekosML/DeepSeek-V4-Flash-0731-GGUF-GX10) | Tensor group | Type | Status | | :--- | :--- | :--- | | Attention projections | `Q8_0` | ✅ | | Shared experts | `Q8_0` | ✅ | | Output language head | `Q8_0` | ✅ | | Token embedding | `F16` | ✅ | | Routed gate/up experts | `IQ2_XXS` | ✅ | | Routed down experts | `Q2_K` | ✅ | --- ## 🔧 Install & Run (Self-Hosted) ### 1. Clone & one-shot setup ```sh git clone https://github.com/julianmb/ds4fa.git ds4-strix-halo cd ds4-strix-halo bash misc/strix-halo-setup.sh # GRUB gttsize/pages_limit, udev, tuned profile — reboot after ``` ### 2. Install the toolchain ```sh sudo apt-get update sudo apt-get install -y hipcc rocminfo rocm-smi libamdhip64-dev \ libhipblas-dev libhipblaslt-dev librocblas-dev \ librocwmma-dev libhipcub-dev aria2 git clone --depth 1 --branch rocm-7.2.3 https://github.com/ROCm/rocWMMA.git /tmp/rocWMMA sudo cp -a /tmp/rocWMMA/library/include/rocwmma /usr/local/include/ ``` ### 3. Build for gfx1151 ```sh make strix-halo -j"$(nproc)" make rocm-doctor # verify TTM/GTT limit; warns + suggests the exact amd-ttm fix ``` ### 4. Download the 0731 model (86.72 GB, ~110 MB/s with 16 connections) Models live under `gguf/` in organized subdirectories (see [gguf/README.md](gguf/README.md)): ``` gguf/ ├── deepseek-v4-flash-0731/ # native route target └── draft/ # speculative-draft models ``` ```sh aria2c -x 16 -s 16 -k 1M -j 16 -c --file-allocation=none \ -d gguf/deepseek-v4-flash-0731 -o DeepSeek-V4-Flash-0731-IQ2XXS-STRIX.gguf \ "https://huggingface.co/tekosML/DeepSeek-V4-Flash-0731-GGUF-GX10/resolve/main/DeepSeek-V4-Flash-0731-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-imatrix.gguf" ln -sf gguf/deepseek-v4-flash-0731/DeepSeek-V4-Flash-0731-IQ2XXS-STRIX.gguf ds4flash.gguf ``` ### 5. Chat **Interactive CLI:** ```sh DS4_ROCM_STREAM_MODEL_CACHE_GB=48 ./ds4 -m ds4flash.gguf -c 512 \ --ssd-streaming --ssd-streaming-cache-experts 32GB \ -p "What is the capital of France?" --think --tokens 60 ``` **OpenAI-compatible server:** ```sh DS4_ROCM_STREAM_MODEL_CACHE_GB=48 ./ds4-server -m ds4flash.gguf -c 8192 \ --port 8000 --ssd-streaming --ssd-streaming-cache-experts 32GB ``` ```bash curl -X POST http://127.0.0.1:8000/v1/chat/completions \ -H "Content-Type: application/json" \ -d '{ "model": "ds4flash", "messages": [{"role": "user", "content": "What is the capital of France?"}], "temperature": 0.6, "max_tokens": 512 }' ``` --- ## 📊 Performance Reference | Stage | Measured | | :--- | ---: | | Decode (ROCmFPX + DSpark q=4, top-k 4) | **32.0 tok/s** | | Decode (ROCmFPX autoregressive) | 25.3 tok/s | | Sparse prefill (indexed, 8K) | **~250 tok/s** | | Exact prefill (short prompts) | 22.5–23 tok/s | *Measured July 2026 on Ryzen AI MAX+ 395, ROCm 7.2.4, Radeon `high` (2.9 GHz), context 8,192, temp 0.* --- ## ❓ Troubleshooting | Problem | Fix | | :--- | :--- | | `raw KV batch store failed` | Set `DS4_ROCM_STREAM_MODEL_CACHE_GB=48` and pass `--ssd-streaming` | | `pageable-memory access disabled` | `sudo amd-ttm --set-pages 8126464` or `DS4_ROCM_TTM_AUTORAISE=1` | | Garbage output | Use `--think` (or `temperature: 0.6`) so DeepSeek reasoning format is respected | | `missing gfx1151` | Build with `make strix-halo`; ensure ROCm 7.2.x | | rocWMMA header errors | Install the matching `rocwmma` tree (step 2) | --- ## 📚 Documentation | Document | Description | | :--- | :--- | | [STRIXHALO.md](STRIXHALO.md) | ROCm install, GRUB params, TTM priority, hardware notes | | [FORK_NOTES.md](FORK_NOTES.md) | Audit of what was retained/rejected from upstream | ## 🤝 Acknowledgements DeepSeek V4 Flash · [antirez/ds4](https://github.com/antirez/ds4) · [Lucebox / ROCmFPX](https://github.com/Luce-Org/lucebox) · [llama.cpp / GGML](https://github.com/ggml-org/llama.cpp) · [tekosML](https://huggingface.co/tekosML) *Local AI should be the default, not a privilege.*