[Issue]: ATOM: TP-DSA decode hang in the plugin forward-relay

[ibeattyo]

Problem Description

Bug report: DeepSeek-V3.2 TP-attention (DSA) decode HANGS on ATOM

• Date: 2026-08-27
• Severity: Blocker — DeepSeek-V3.2 TP-attention unusable on ATOM at 64k input sequence length (hangs at every batch size)
• Model: deepseek-ai/DeepSeek-V3.2 (DSA sparse MLA)
• Stack / image: ATOM rocm/atom-dev:sglang-v0.5.15.post1-nightly_20260803 (sglang 0.5.15.post1;
  also reproduced on …-nightly_20260817). DSV3.2 on ATOM is DSA-sparse via attention-backend=aiter
  (dsa_prefill/decode_backend=tilelang).
• Config: tp-size=8, no DP (--enable-dp-attention OFF), context-length=69632 (64k in / 1k out),
  kv-cache-dtype=fp8_e4m3, mem-fraction-static=0.90, page-size→64 (DSA auto), chunked-prefill-size=16384.
• Hardware: MI355X (gfx950), ROCm 7.2.

Summary

DeepSeek-V3.2 with TP-attention on ATOM starts, loads weights, allocates KV, prefills — then hangs the moment
decode begins, on a HIP device sync inside the ATOM plugin. The scheduler makes no further progress and is killed
by its 300 s watchdog. Deterministic at all batch sizes (incl. bs32). It is not memory, not the aiter tuning
lock, and not a too-short watchdog (all ruled out below).

Symptom / evidence

Server reaches serving (max_total_num_tokens=3.78M, avail 26.4 GB — no OOM), prefills to #running-req: 31,
then goes silent. After 300 s:

[TP5] Scheduler watchdog timeout (self.watchdog_timeout=300, self.soft=False)

py-spy of the (still-alive) scheduler main thread, on all 8 TP ranks, is blocked in a HIP memcpy+sync inside the
ATOM plugin's forward-payload relay:

c10::hip::memcpy_and_sync                (libtorch_hip.so)
at::native::_to_copy / at::native::to    (a .to() copy)
pool_indexed_future_store                (atom/plugin/sglang/eagle3_llama_bridge.py:755)
_relay_forward_payload                   (srt/managers/scheduler.py:3429)
run_batch                                (srt/managers/scheduler.py:3288)
event_loop_overlap                       (srt/managers/scheduler.py:1615)

The memcpy_and_sync never returns → a true GPU hang (not slow).

Scope — it is DSV3.2-TP-DSA-specific

On the same ATOM container / same 0.90 mem-fraction:

Case	Result
DSV3.2 DP-attention @64k (dp8/ep8, dp8/tp8)	✅ works (14/14) — DSA decode is fine under DP
DSR1 (dense MLA) TP-attention @64k	✅ works (8/8)
DSV3.2 TP-attention @64k (this bug)	❌ hangs, all batch sizes
So the fault is specific to DSV3.2 + TP-attention + DSA decode on ATOM.

What was ruled out (each tested)

• OOM — real at 0.95 (HIP out of memory, ~786 MiB free); fixed by 0.90 (26 GB free) — the hang persists at 0.90. Not memory.
• aiter file_baton lock (/tmp/aiter_configs/*.csv.lock) — pre-cleaned before/after each run; only transient waits remain; still hangs. (The lock only looks implicated when a rank hangs while holding it.)
• Watchdog too short — raised --watchdog-timeout 300 → 1800; still no decode progress, just delays the kill. Not slowness.

Reproduce (standalone — no framework needed)

Standard SGLang server + client on one 8×MI355X node; DeepSeek-V3.2 weights in the HF cache.

docker run --rm -it --network=host --ipc=host --shm-size=16g \
  --cap-add=SYS_PTRACE --security-opt seccomp=unconfined \
  --device=/dev/kfd --device=/dev/dri --group-add video \
  -e HF_HOME=/hf -v /path/to/hf-cache:/hf \
  rocm/atom-dev:sglang-v0.5.15.post1-nightly_20260803 bash

export SGLANG_USE_AITER=1 AITER_QUICK_REDUCE_QUANTIZATION=INT4 SGLANG_AITER_FP8_PREFILL_ATTN=0 \
       SGLANG_EXTERNAL_MODEL_PACKAGE=atom.plugin.sglang.models
# server (DSA-sparse TP-attention, 64k):
python3 -m sglang.launch_server --model-path deepseek-ai/DeepSeek-V3.2 --trust-remote-code \
  --tp-size 8 --attention-backend aiter \
  --context-length 69632 --kv-cache-dtype fp8_e4m3 --mem-fraction-static 0.90 \
  --page-size 1 --disable-radix-cache --chunked-prefill-size 16384 --max-running-requests 64
# client (once ready) — any batch (e.g. 32) => decode hang -> 300 s scheduler watchdog kill:
python3 -m sglang.bench_serving --backend sglang --dataset-name random \
  --random-input-len 65536 --random-output-len 1024 --random-range-ratio 1 \
  --num-prompts 32 --max-concurrency 32

Key knobs: attention-backend=aiter + tp-size=8 (no --enable-dp-attention) + 64k context. --cuda-graph-bs
left at default also reproduces.

Impact

DeepSeek-V3.2 cannot be served with TP-attention on ATOM at 64k at all. DP-attention is the only usable DSV3.2
attention layout on ATOM (works fine).

Requested fix

Investigate atom/plugin/sglang/eagle3_llama_bridge.py:755 pool_indexed_future_store — its HIP memcpy_and_sync
in the TP forward-relay never returns for DSV3.2 TP-DSA decode (the DP path through the same plugin is fine).

Workaround

Use DP-attention for DSV3.2 on ATOM at 64k.

Evidence

Signatures (watchdog line + py-spy stack) quoted inline; report is self-contained. Original logs archived on the
origin host: control_v32_tp_mf90_atom_sglang/8tp_64k_1k_bs32_mf90_*.stderr.log; working-DP contrast in
{control,profile}_v32_mf95_atom_sglang/dp8_*_mf90*.

Operating System

22.04.5 LTS (Jammy Jellyfish)

CPU

AMD EPYC 9575F 64-Core Processor

GPU

AMD Instinct MI355X

ROCm Version

7.2.0

ROCm Component

No response

Steps to Reproduce

No response

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

rocminfo --support output

Paste output here

Additional Information

No response