---

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:

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:

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.

5. gfx1151 F16 Embedding Dispatch & Wave32 Warp Shuffles (src/rocm/)

• F16 Token Embedding Launch: Restored direct F16 embedding kernels in src/rocm/ds4_rocm_embedding_launch.cuh, eliminating a bug where F16 float tables were falsely decoded as Q4_K packed nibbles on full models.
• Wave32 Shuffle Intrinsics: Replaced 64-bit masked CUDA-emulated shuffles with native HIP intrinsics (__shfl, __shfl_down, __shfl_xor) across attention, router, and MoE kernels for exact reductions and broadcasts.
• 128 GB APU Memory Safeguards: Removed artificial ulimit -v virtual-address caps in run-deepseek-v4.sh and calibrated physical cache defaults (32GB), ensuring full-model execution stays safely within 55–68 GB of physical unified RAM.

---

📦 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

---

🔧 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=32 ./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=32 ./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

Measured July 2026 on Ryzen AI MAX+ 395, ROCm 7.2.4, Radeon high (2.9 GHz), context 8,192, temp 0.

---

❓ Troubleshooting

---

📚 Documentation

🌐 Also on Hugging Face

• Code (this repo, mirrored): julianmb/ds4fa
• Model weights: julianmb/DeepSeek-V4-Flash-0731-IQ2XXS-STRIX

🤝 Acknowledgements

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

Local AI should be the default, not a privilege.