license	mit
tags	deepseek moe amd strix-halo rocm gguf rocmfpx
pipeline_tag	text-generation

⚡ 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.

./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

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

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

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

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/
├── deepseek-v4-flash-0731/          # native route target
└── draft/                           # speculative-draft models

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:

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:

DS4_ROCM_STREAM_MODEL_CACHE_GB=48 ./ds4-server -m ds4flash.gguf -c 8192 \
  --port 8000 --ssd-streaming --ssd-streaming-cache-experts 32GB

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	ROCm install, GRUB params, TTM priority, hardware notes
FORK_NOTES.md	Audit of what was retained/rejected from upstream

🤝 Acknowledgements

DeepSeek V4 Flash · antirez/ds4 · Lucebox / ROCmFPX · llama.cpp / GGML · tekosML

Local AI should be the default, not a privilege.