[AMD/MI300A] DeepSeek-V4 (deepseek_v4) serving support — tracking

[xzyaoi]

Tracking issue for serving DeepSeek-V4 (model_type: deepseek_v4) on AMD MI300A (gfx942).

Architecture (why it's new vs V3/Kimi)

• FP4/mxfp4 routed experts + FP8 (block 128×128) for the rest. Flash: 256 experts / hidden 4096; Pro: 384 experts / hidden 7168; top-6, 1 shared.
• Hybrid attention: Compressed Sparse Attention (CSA) + Heavily Compressed Attention (HCA) + a DSA indexer (top-512 Flash / top-1024 Pro), per-layer compress_ratios, sliding_window: 128, o_lora_rank.
• mHC (Manifold-Constrained Hyper-Connections: hc_mult, hc_sinkhorn_iters), sqrtsoftplus routing, MTP, YaRN→1M.

Findings (single-node probe on MI300A, 2026-06-11)

Much more AMD-ready than V3 was — the model code imports and constructs cleanly on gfx942 (no import wall; deepseek_v4.py feature-detects kernels and gates NVIDIA-only paths on cc=10 with fallbacks). Confirmed working at construction: model import, attention + DSA indexer construction, and the mxfp4 MoE kernel at ep_size=1 (selects + allocates).

Blockers

• [AMD/MI300A] mxfp4 fused-MoE: support expert parallelism (ep_size>1) on gfx942 (DeepSeek-V4) #26 — mxfp4 fused-MoE: no expert-parallelism (ep_size > 1) on gfx942. PRIMARY/confirmed. The gluon/triton mxfp4 backends are gated to ep_size <= 1; V4's 256–384 experts must shard across GPUs/nodes to fit. (At ep=1 the kernel works but the model OOMs.)
• [AMD/MI300A] DeepSeek-V4 DSA indexer + attention: validate/provide gfx942 forward path #27 — DSA indexer + V4 attention: forward unvalidated on gfx942. Constructs fine, but forward couldn't be reached (OOM behind [AMD/MI300A] mxfp4 fused-MoE: support expert parallelism (ep_size>1) on gfx942 (DeepSeek-V4) #26). Indexer uses NVIDIA deep_gemm.fp8_fp4_mqa_logits + a cuda-only mxfp4 paged-gather; AMD fallback unverified.

Not blockers (confirmed OK at construction)

Model import, attention/indexer module construction, mxfp4 MoE backend selection + weight alloc (ep=1).

Notes

• FlashInfer-AMD (amd-flashinfer, gfx942) does not help here: v0.2.5 / ROCm 7.1.1 / torch 2.8 (ABI-mismatched with our 7.2/2.11 image) and predates V4's FP4-MoE + sparse-indexer kernels.
• This mirrors the Kimi/V3 bring-up (which needed an AMD INT4-W4A16 MoE + AITER MLA): same shape, smaller remaining lift.

Weights / repro

V4-Pro staged at infra01/hf_models/.../DeepSeek-V4-Pro (806 GB); Flash not staged (MIT, ~150 GB). Probes: jobs/serve-v4pro-1node-probe{,2}.sbatch on beverin. HW: MI300A gfx942 / ROCm 7.2 / torch 2.11.

[xzyaoi]

Shipped a self-contained building block toward V4-on-MI300A: a gfx942 Triton mxfp4_paged_gather — the gather half of #27 (the Triton attention variant had no mxfp4 paged gather; only the CUDA indexer_mxfp4_paged_gather existed).

What it does: gathers the DSA indexer's top-512 (Flash) / top-1024 (Pro) selected KV positions out of a paged (block-table) mxfp4 KV cache and dequantizes to bf16. One Triton program per (seq, top-k slot); E2M1 nibbles decoded arithmetically (no LUT), E8M0 group scale (2^(b-127), 0xFF→0) broadcast; padded slots (sel_pos<0) → zero row. New xkernels.mxfp4_paged_gather op.

Validation

• Local (interpreter): ruff clean; TRITON_INTERPRET=1 pytest 106 passed / 3 skipped (6 new tests; exact at V4 head_dim=128, topk=512).
• On-device, beverin MI300A gfx942 (job 381968): GPU bf16 correctness 6/6 pass; decode-shape timing 0.028–0.042 ms across V4 indexer shapes with max|err| = 0.0000 (mxfp4 dequant is exact — only the bf16 round contributes).

Caveats / still open

• Does not provide the deep_gemm.fp8_fp4_mqa_logits indexer-logits kernel (other CUDA-only piece of [AMD/MI300A] DeepSeek-V4 DSA indexer + attention: validate/provide gfx942 forward path #27).
• Does not touch [AMD/MI300A] mxfp4 fused-MoE: support expert parallelism (ep_size>1) on gfx942 (DeepSeek-V4) #26 (mxfp4 fused-MoE expert parallelism) — still the PRIMARY blocker for fitting V4 on MI300A. A full V4 forward needs [AMD/MI300A] mxfp4 fused-MoE: support expert parallelism (ep_size>1) on gfx942 (DeepSeek-V4) #26 + the indexer-logits path on top of this gather.

PR (draft): #29 · docs: docs/issue-28-mxfp4-paged-gather.md

[xzyaoi]

Update — V4-Flash (fp8 experts) bring-up walked the full blocker chain on MI300A:

1. DSA indexer (selection) — DONE on gfx942: [AMD/MI300A] DeepSeek-V4 DSA indexer + attention: validate/provide gfx942 forward path #27/feat(attention): DeepSeek-V4 DSA indexer logits — gfx942 forward path (issue #27) #31 (dsa_indexer_logits/dsa_indexer_topk) + feat(gather): Triton mxfp4 paged KV gather for DeepSeek-V4 DSA indexer on gfx942 (issue #28/#27) #29 (mxfp4_paged_gather), wired into the runtime in feat(amd): DeepSeek-V4 gfx942 bring-up — DSA indexer + o_proj FP8 fallback tokenspeed#2, parity-validated (Jaccard=1.0).
2. o_proj is_bmm FP8 einsum — bf16 AMD fallback in tokenspeed#2.
3. MoE — not a blocker for Flash: its experts are block-fp8 (Fp8TritonBackend, no ep gate, arch=any), TP=4 works. ([AMD/MI300A] mxfp4 fused-MoE: support expert parallelism (ep_size>1) on gfx942 (DeepSeek-V4) #26/feat(moe): expert parallelism (ep_size>1) for the quantized fused-MoE GEMM (issue #26) #30 mxfp4-EP only matters for the mxfp4 V4-Pro variant.)
4. V4 sparse-MLA attention compute — the remaining gating blocker → new issue [AMD/MI300A] DeepSeek-V4 sparse-MLA attention compute on gfx942 (flash_mla_sparse_fwd / flash_mla_with_kvcache) #32 (flash_mla_sparse_fwd / flash_mla_with_kvcache are NVIDIA-Hopper+ only; need a gfx942 sparse-MLA kernel).

Correction to the earlier "indexer is the sole blocker" note (that came from a load-only probe): the indexer is the selection; the sparse-attention compute (#32) is the larger NVIDIA-only piece. With #32, V4-Flash should serve single-node on MI300A (TP=4).

[xzyaoi]

Session update (2026-06-12): V4-Flash forward bring-up — blocker chain peeled, paused mid-fp8

Walked the DeepSeek-V4-Flash serve forward on beverin (MI300A, gfx942, TP=4, eager).
The forward now advances deep into the model. Blockers peeled this session, in
execution order:

1. ✅ Backend registration — deepseek_v4 was unregistered on AMD (the cluster ran a
   stale rsync snapshot pre-dating the registration); fixed by a full code sync.
2. ✅ KV-pool OOM — --gpu-memory-utilization 0.9 → 0.5 (MI300A unified memory, 4
   ranks share ~512 GB).
3. ✅ NCCL — Failed to initialize any NET plugin → add the aws-ofi-rccl plugin to
   LD_LIBRARY_PATH (the myofi EDF forces OFI even single-node).
4. ✅ sparse-MLA compute — feat(attention): DeepSeek-V4 sparse-MLA attention compute on gfx942 (issue #32) #33 merged (main), bound in tokenspeed
   (ops/attention/flash_mla AMD branch + decode cache-split adapter; decode op
   generalization in feat(attention): sparse-MLA decode accepts block_table=None (physical indices) #35).
5. ✅ MHC prenorm GEMM — feat(mhc): DeepSeek-V4 MHC hidden-compression prenorm GEMM on gfx942 (issue #36) #37 merged, bound via new tokenspeed_kernel/ops/mhc.
6. ✅ MHC TileLang PDL — tl.pdl_sync/tl.pdl_trigger (Hopper launch-overlap) don't
   lower on the AMD TileLang backend → guarded off on AMD (_MHC_USE_PDL); the
   _mhc_pre_big_fuse_tilelang fusion now compiles + runs on all ranks.
7. 🔧 fp8 block-scale dense GEMM — the MLA projections select the portable
   torch_mm_fp8_blockscale reference; its online-quant branch was missing → added
   (now functional, slow). Performant gfx942 kernel tracked in [AMD/MI300A] fp8 block-scale dense GEMM on gfx942 — only the slow torch reference path exists #38.

Current position: the forward reaches the block-fp8 MLA projections; the fp8 fix was
synced but the validating serve was paused before it completed (user request).

Remaining checklist

• Validate the fp8-projection forward → coherent prefill (max_tokens=1).
• Validate decode end-to-end (max_tokens=48) — needs feat(attention): sparse-MLA decode accepts block_table=None (physical indices) #35 merged + pinned.
• [AMD/MI300A] fp8 block-scale dense GEMM on gfx942 — only the slow torch reference path exists #38 — performant gfx942 fp8 block-scale GEMM (torch reference is slow).
• [AMD/MI300A] performance pass for the V4 sparse-MLA (#33) + MHC (#37) kernels on gfx942 #39 — perf pass on the V4 sparse-MLA (feat(attention): DeepSeek-V4 sparse-MLA attention compute on gfx942 (issue #32) #33) + MHC (feat(mhc): DeepSeek-V4 MHC hidden-compression prenorm GEMM on gfx942 (issue #36) #37) kernels for CDNA3.
• HIP-graph decode capture for V4 (after eager correctness).
• Any further NVIDIA-only ops the (unvalidated) forward surfaces past the fp8 layer.

Wiring state (tokenspeed)

• PR Fused parallel dual RMSNorm (MLA q_a / kv_a latents) #2 (amd/deepseek-v4-indexer): sparse-MLA binding + decode adapter + skills
  committed/pushed.
• Local/synced, pending commit: MHC binding (ops/mhc), the PDL guard
  (deepseek_v4_mhc.py), the fp8 online-quant branch (ops/gemm/__init__.py), pin bump
  to xkernels main (feat(attention): DeepSeek-V4 sparse-MLA attention compute on gfx942 (issue #32) #33+feat(mhc): DeepSeek-V4 MHC hidden-compression prenorm GEMM on gfx942 (issue #36) #37).
• kernels PR feat(attention): sparse-MLA decode accepts block_table=None (physical indices) #35 (decode block_table=None physical indices): draft, pending V4
  decode validation.

The sparse-MLA and MHC kernels (the two original gating blockers, #32/#36) are done and
proven to integrate through the forward — the bring-up is now a layer-by-layer grind
through correctness-first fp8/projection paths, each re-serve costing ~18 min to reload
the 149 GB model.