Skip to content

[Feat] [AMD] DeepSeek-v4.1: Support prefill context parallel - #43465

Open
At1a8 wants to merge 21 commits into
sgl-project:mainfrom
At1a8:fangyuan/dsv4p1_pcp
Open

At1a8 wants to merge 21 commits into
sgl-project:mainfrom
At1a8:fangyuan/dsv4p1_pcp

Conversation

@At1a8

@At1a8 At1a8 commented Oct 10, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Prefill CP for DeepSeek-V4.1 already works on CUDA, but on AMD it's still blocked.
Long prompts (e.g. 100k) are slow to prefill on plain TP8, this PR turns it on for AMD HIP platform, validated on gfx950.

Modifications

  • deepseek_v4_hook.py: stop rejecting V4.1 for prefill CP on HIP, including with --enable-decoder-swa-bounded-replay (V4 is still rejected).
  • deepseek_v4_backend_hip_radix.py:
    • Add a HIP _forward_low_ratio_sources_cp instead of previous CUDA version:
      • Compressor: all-gather the prompt and let every rank compress all tokens, so each rank has the full main KV and indexer K. The full extend reuses the same fused kernel as non-CP cases.
      • Indexer: only score this rank's rows.
    • _low_ratio_index_topk: drop the assert that rejected rows_per_request and pass it down to the HIP indexer.
    • Add _first_rows: on gfx950 q_lora can be an FP8 activation wrapper that can't be sliced with [:n].
    • Bounded-replay tail under CP:
      • Reuse CUDA's _late_layer_tail_cp_layout to build the tail's own CP layout.
      • init_forward_metadata_prefill takes an optional cp_metadata for the tail: pad swa_replay_start and reindex rows with local_index.
      • enter_late_layer_tail / exit_late_layer_tail swap in and restore the CP metadata and DP buffer length, and cut candidate masks by each rank's local tail rows.
  • low_ratio_backend_hip.py: pass rows_per_request into the HIP indexer so it knows how many local rows each request has.
  • deepseek_v4.py: with HIP CP, put the tail outputs (including DSpark aux) back into the rank's full layout, because the CP runner gathers by that layout.
  • interleave.py: on HIP, skip CP when a rank would get fewer than 16 tokens. aiter's group32 MXFP8 GEMM returns NaN for wqkv_a (N=1792, K=5120) at M=5~8. Left a TODO to drop this once the kernel is fixed.
  • Remove the stale "V4.1 is rejected on HIP" case from test_platform_prefill_cp_deprecation.py.

Accuracy Tests

settings accuracy
TP8 0.902
TP8CP8 on this pr 0.907

Server Command

TP=${TP:-8}
CP=${CP:-8}
MODEL=${MODEL:-/data2/models/DeepSeek-V4.1-Flash}
MEM_FRACTION=${MEM_FRACTION:-0.85}
CHUNKED_PREFILL_SIZE=${CHUNKED_PREFILL_SIZE:-8192}
ENGRAM_HOST=${ENGRAM_HOST:-0}
CP_ARGS=""
if [ "${CP}" -gt 1 ]; then
    CP_ARGS="--enable-prefill-cp --cp-strategy interleave"
fi

echo "TP: ${TP}"
echo "CP_ARGS: ${CP_ARGS:-<无>}"
echo "MODEL: ${MODEL}  ENGRAM_HOST: ${ENGRAM_HOST}"
echo "MEM_FRACTION: ${MEM_FRACTION}"
echo "CHUNKED_PREFILL_SIZE: ${CHUNKED_PREFILL_SIZE}"

export AITER_BF16_FP8_MOE_BOUND=0
export SGLANG_ENABLE_DSV41_ENGRAM_HOST_TABLE=${ENGRAM_HOST}
export SGLANG_DSV41_ENGRAM_HOST_TABLE_LAYOUT=per_rank
export SGLANG_USE_ROCM700A=0
export TORCH_BLAS_PREFER_HIPBLASLT=1

sglang serve \
  --trust-remote-code \
  --model-path ${MODEL} \
  --tp ${TP} \
  --mem-fraction-static ${MEM_FRACTION} \
  --speculative-algorithm DSPARK \
  --speculative-dspark-block-size 5 \
  --enforce-shared-experts-fusion \
  --fp8-gemm-backend aiter \
  --cuda-graph-max-bs-decode 128 \
  --max-running-requests 128 \
  --enable-decoder-swa-bounded-replay \
  --cuda-graph-backend-prefill disabled \
  --chunked-prefill-size ${CHUNKED_PREFILL_SIZE} \
  --reasoning-parser auto \
  --tool-call-parser auto \
  --host 0.0.0.0 \
  --port 30000 \
  --json-model-override-args '{"vision_n_layers":0}' \
  ${CP_ARGS} ${EXTRA_ARGS:-}

Speed Tests and Profiling

I/O=100k/500

Concurrency Parallel Setting Total token throughput (tok/s) Mean TTFT (ms) Mean TPOT (ms)
8 TP8 68626.53 3745.77 15.22
8 TP8CP8 82251.05 (1.20x) 3133.05 (1.20x) 12.97 (1.17x)
16 TP8 73536.87 6059.17 31.19
16 TP8CP8 89493.44 (1.22x) 4670.26 (1.30x) 26.30 (1.19x)
32 TP8 76043.26 9506.36 65.31
32 TP8CP8 93275.17 (1.23x) 7639.09 (1.24x) 53.35 (1.22x)
64 TP8 77483.01 16432.47 132.68
64 TP8CP8 95447.03 (1.23x) 13245.00 (1.24x) 107.79 (1.23x)

Checklist

Review and Merge Process

  1. Ping Merge Oncalls to start the process. See the PR Merge Process.
  2. Get approvals from CODEOWNERS and other reviewers.
  3. Trigger CI tests with comments or contact authorized users to do so.
    • Common commands include /tag-and-rerun-ci, /tag-run-ci-label, /rerun-failed-ci
  4. After green CI and required approvals, ask Merge Oncalls or people with Write permission to merge the PR.

CI States

Latest PR Test (Base): ⏳ Run #41991791828
Latest PR Test (Extra): ❌ Run #41991791640
Latest PR Test (AMD ROCm 10): ⏳ Run #41991792278

Tip

Base CI only covers the core cases, and its workload may not exercise your change. Use /rerun-test to run the tests your change touches, or copy this prompt for an agent:

Find the registered tests this PR affects and run them with /rerun-test, following https://github.com/sgl-project/sglang/blob/main/.agents/skills/rerun-test/SKILL.md

@At1a8
At1a8 marked this pull request as ready for review October 10, 2026 03:49
@bingxche bingxche added the run-ci CI: run the baseline test suite on this PR label Oct 10, 2026
ChangLi811 pushed a commit to ChangLi811/sglang that referenced this pull request Oct 13, 2026
Carried unchanged from sgl-project#43465 (by @At1a8, open at the time of writing) so this
PR is self-contained; only the imports were re-resolved against current main.
Drop this commit once sgl-project#43465 is merged.
ChangLi811 added a commit to ChangLi811/sglang that referenced this pull request Oct 13, 2026
… TP, skip zero-row indexer groups

- --enable-cp-decode-attn-tp narrows the attention linears per rank. On gfx950
  the V4.1 block-fp8 (ue8m0, 32-wide) linears run from scale copies derived at
  load time (weight_scale_mx_e8m0, weight_scale_inv_mx, weight_bf16); slice
  them on the same dim as the weight, and reject the MXFP8-native route whose
  weight layout cannot be narrowed.
- With prefill CP a request group can own no rows on a rank (a one-token
  extend lands on one rank). Skip that group in the HIP low-ratio indexer
  before scoring: the fp4 prefill logits schedule still sends its CTAs to row 0
  of a freshly allocated (null) page table, and top-k cannot launch on an empty
  batch. Seen as a GPU memory fault at concurrency 64 with mixed short and long
  extends.
- sgl-project#43465's HIP small-shard floor (interleave CP needs 16 rows per rank) now
  lives in interleave_cp_min_tokens(), which also bounds the CP breakable-graph
  capture buckets: a bucket below the floor never activates CP, so its CP-local
  capture asserted at startup on ROCm.
- The HIP BCG metadata test follows sgl-project#43465's apply_cp_reindex call, which
  passes the CP layout's local_index (the stub gains it, and the expected call
  carries it).

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
@At1a8 At1a8 changed the title [AMD] Support prefill context parallel for DeepSeek V4.1 [Feat] [AMD] DeepSeek-v4.1: Support prefill context parallel Oct 14, 2026
@bingxche bingxche added the bypass-fail-fast CI: a failing job no longer aborts its siblings (lint still gates) label Oct 15, 2026
AMD-yanfeiwang added a commit to AMD-yanfeiwang/sglang that referenced this pull request Oct 15, 2026
…oject#43465 by At1a8)

Applied sgl-project#43465 (head 5c5862f, author @At1a8, open) on top of
the HIP CP+BCG port (sgl-project#43819). Conflict resolution in
deepseek_v4_backend_hip_radix.py: both PRs pad swa_replay_start and pass
cp_metadata.local_index to apply_cp_reindex; kept the single copy from sgl-project#43819,
which also keeps num_write_tokens for breakable prefill graphs.
@SAXEM1997

Copy link
Copy Markdown

CUDA: prefill CP + decoder SWA bounded replay + DSpark crashes at the first large prefill (tail rows are not scattered back before the CP gather)

We reproduced a failure on CUDA/SM90 with the combination enabled. The analysis points at the same tail-layout gap that the "Bounded-replay tail under CP" part of this PR closes on HIP.

Setup: DeepSeek-V4.1-Flash, 8×H20-3e (SM90), CUDA, sglang v0.5.22, --tp 8 --ep 4 --enable-prefill-cp --cp-strategy interleave --enable-decoder-swa-bounded-replay --speculative-algorithm DSPARK --speculative-dspark-block-size 5 --chunked-prefill-size 32768 (prefill CUDA graph disabled).

Symptom: the server boots, but the first 128K-token prefill kills the scheduler:

RuntimeError: Interleave CP gather received an unexpected local token count:
  rank=3, got=16, expected_at_least=4096, total=32768, cp_size=8
  File ".../srt/layers/cp/interleave.py", line 288, in _gather_interleaved_tensor
  File ".../srt/layers/cp/utils.py", line 246, in cp_gather_after_forward
  File ".../srt/model_executor/runner/eager_runner.py", line 405, in _execute_extend_cp
  File ".../srt/speculative/dspark_components/dspark_worker_v2.py", line 650, in _forward_prefill

Analysis

  • With replay, the late layers run over each request's SWA window (128 tokens); after interleave sharding a rank owns 16 of them.
  • In DeepseekV4Model.forward, the _scatter_tail_rows(...) call is guarded by not capture_dspark. With DSpark (capture_dspark=True), the model returns tail rows (and tail-row aux tensors) un-scattered.
  • eager_runner._execute_extend_cp then gathers with the full-batch CP metadata (per_rank_logical_token = 32768 / 8 = 4096), and interleave.py:288 raises.

Launch-time guard gap: arg_groups/deepseek_v4_hook.py only rejects --enable-decoder-swa-bounded-replay when attn_dp_enabled; a DP=1 pure-CP configuration passes the check and crashes at the first large prefill instead of being rejected at startup. The comment there already says "Attention TP/CP shards need a separate tail padding layout".

Fix we validated: in deepseek_v4.py, when is_cp_active(forward_batch) and a replay tail exists, scatter hidden_states, pre_hc_head and the DSpark aux tensors back into the rank's full (sharded) layout before the runner gathers (22 lines; non-tail rows zero-filled). With that:

  • Cold TTFT 128K / 256K / 512K: 3.37 s / 6.86 s / 15.95 s, vs 5.41 / 11.71 / 28.66 s for the same CP + DSpark config without replay.
  • Greedy outputs on a 5-prompt set are token-identical to the non-CP production config with replay on.
  • Decode replay (36 real sessions): c8 ITL p90 86 ms vs 160 ms (CP without replay), output throughput 259 vs 176 tok/s.

Is this the same gap #43465 closes on HIP, and is a CUDA-side counterpart planned? Happy to send the 22-line patch if useful.

@At1a8

At1a8 commented Oct 20, 2026

Copy link
Copy Markdown
Contributor Author

CUDA: prefill CP + decoder SWA bounded replay + DSpark crashes at the first large prefill (tail rows are not scattered back before the CP gather)

We reproduced a failure on CUDA/SM90 with the combination enabled. The analysis points at the same tail-layout gap that the "Bounded-replay tail under CP" part of this PR closes on HIP.

Setup: DeepSeek-V4.1-Flash, 8×H20-3e (SM90), CUDA, sglang v0.5.22, --tp 8 --ep 4 --enable-prefill-cp --cp-strategy interleave --enable-decoder-swa-bounded-replay --speculative-algorithm DSPARK --speculative-dspark-block-size 5 --chunked-prefill-size 32768 (prefill CUDA graph disabled).

Symptom: the server boots, but the first 128K-token prefill kills the scheduler:

RuntimeError: Interleave CP gather received an unexpected local token count:
  rank=3, got=16, expected_at_least=4096, total=32768, cp_size=8
  File ".../srt/layers/cp/interleave.py", line 288, in _gather_interleaved_tensor
  File ".../srt/layers/cp/utils.py", line 246, in cp_gather_after_forward
  File ".../srt/model_executor/runner/eager_runner.py", line 405, in _execute_extend_cp
  File ".../srt/speculative/dspark_components/dspark_worker_v2.py", line 650, in _forward_prefill

Analysis

  • With replay, the late layers run over each request's SWA window (128 tokens); after interleave sharding a rank owns 16 of them.
  • In DeepseekV4Model.forward, the _scatter_tail_rows(...) call is guarded by not capture_dspark. With DSpark (capture_dspark=True), the model returns tail rows (and tail-row aux tensors) un-scattered.
  • eager_runner._execute_extend_cp then gathers with the full-batch CP metadata (per_rank_logical_token = 32768 / 8 = 4096), and interleave.py:288 raises.

Launch-time guard gap: arg_groups/deepseek_v4_hook.py only rejects --enable-decoder-swa-bounded-replay when attn_dp_enabled; a DP=1 pure-CP configuration passes the check and crashes at the first large prefill instead of being rejected at startup. The comment there already says "Attention TP/CP shards need a separate tail padding layout".

Fix we validated: in deepseek_v4.py, when is_cp_active(forward_batch) and a replay tail exists, scatter hidden_states, pre_hc_head and the DSpark aux tensors back into the rank's full (sharded) layout before the runner gathers (22 lines; non-tail rows zero-filled). With that:

  • Cold TTFT 128K / 256K / 512K: 3.37 s / 6.86 s / 15.95 s, vs 5.41 / 11.71 / 28.66 s for the same CP + DSpark config without replay.
  • Greedy outputs on a 5-prompt set are token-identical to the non-CP production config with replay on.
  • Decode replay (36 real sessions): c8 ITL p90 86 ms vs 160 ms (CP without replay), output throughput 259 vs 176 tok/s.

Is this the same gap #43465 closes on HIP, and is a CUDA-side counterpart planned? Happy to send the 22-line patch if useful.

Hi @SAXEM1997, we dont have Nvidia device for your settings, feel free to submit a new pr for your issues.

ChangLi811 pushed a commit to ChangLi811/sglang that referenced this pull request Oct 20, 2026
Squashed from sgl-project#43465 by @At1a8 (head c55ba0c), unchanged.
It is carried here only until sgl-project#43465 merges and will be dropped on rebase.
Resolve the conflicts with sgl-project#45339 in deepseek_v4_backend_hip_radix.py: keep
main's tail entry (the source layer runs inside the tail after
write_tail_source_kv) and add the CP metadata / DP buffer switch on top of it.
The HIP CP low-ratio sources take q_lora rows only for the indexer, since
write_tail_source_kv runs the compressor alone with q_lora=None.
@bingxche bingxche added the parallel-stages CI: stages dispatch together instead of waiting on each other label Oct 21, 2026
SAXEM1997 added a commit to SAXEM1997/sglang that referenced this pull request Oct 21, 2026
…ill CP gather with DSpark

When prefill CP, decoder SWA bounded replay and DSpark run together, the late
layers return only each rank's SWA tail rows (16 of 32768), while
_eager_runner._execute_extend_cp gathers with the full-batch CP layout (4096
rows per rank). The capture_dspark guard skipped the tail scatter, so the first
large prefill aborted in interleave.py:

  RuntimeError: Interleave CP gather received an unexpected local token count:
  rank=3, got=16, expected_at_least=4096, total=32768, cp_size=8

Add _needs_cp_tail_scatter() for the CUDA/CP gate and scatter hidden_states,
pre_hc_head and the tail-layer DSpark aux tensors back into the rank's full
layout before the CP gather. Non-tail rows are zero-filled (defensive; they are
never read). Scoped to CUDA so it composes with the HIP counterpart in sgl-project#43465.

This covers eager prefill only: the breakable prefill CUDA graph with prefill
CP is guarded by sgl-project#44141.

Tested on 8x H20-3e (SM90), sglang v0.5.22: 5/5 greedy outputs match the
non-CP production baseline; cold TTFT 128K/256K/512K = 3.37/6.86/15.95 s vs
5.41/11.71/28.66 s with replay off.
SAXEM1997 added a commit to SAXEM1997/sglang that referenced this pull request Oct 21, 2026
…ill CP gather with DSpark

When prefill CP, decoder SWA bounded replay and DSpark run together, the late
layers return only each rank's SWA tail rows (16 of 32768), while
_eager_runner._execute_extend_cp gathers with the full-batch CP layout (4096
rows per rank). The capture_dspark guard skipped the tail scatter, so the first
large prefill aborted in interleave.py:

  RuntimeError: Interleave CP gather received an unexpected local token count:
  rank=3, got=16, expected_at_least=4096, total=32768, cp_size=8

Add _needs_cp_tail_scatter() for the CUDA/CP gate and scatter hidden_states,
pre_hc_head and the tail-layer DSpark aux tensors back into the rank's full
layout before the CP gather. Non-tail rows are zero-filled (defensive; they are
never read). Scoped to CUDA so it composes with the HIP counterpart in sgl-project#43465.

This covers eager prefill only: the breakable prefill CUDA graph with prefill
CP is guarded by sgl-project#44141.

Tested on 8x H20-3e (SM90), sglang v0.5.22: 5/5 greedy outputs match the
non-CP production baseline; cold TTFT 128K/256K/512K = 3.37/6.86/15.95 s vs
5.41/11.71/28.66 s with replay off.
@At1a8

At1a8 commented Oct 21, 2026

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@bingxche

Copy link
Copy Markdown
Collaborator

@At1a8

CI Status for PR #43465

Verdict: None of the 6 executed test failures was caused by this PR. Four AMD failures fail the same test case the same way on comparable main runs. The two CPU failures are in MoE code the PR doesn't touch: one comes from main breakage that is already fixed, the other from the Xeon job installing the PR branch, which is 7 commits behind main. The core PR CI suites have finished. The opt-in suites pr-test-amd-extra and pr-test-extra did not run because their pr-gate failed.

Note

The PR updates test/registered/amd/test_dsv4_hip_bcg_metadata.py::test_cp_prefill_attaches_before_reindex (D19, D20), but that test never ran. The file runs with fail-fast, and it stopped after the first test, TestAiterSparseLengthFoldPerStep, which already fails on main (the log shows Ran 1 test). So in this CI run, nothing checked the PR's HIP prefill context-parallel metadata change.

Executed failures: AMD: 4 (0 related) · Others: 2 (0 related)

AMD Executed Failures

Job Test Failure On main Related? Why
stage-b-test-large-8-gpu-mi35x-disaggregation-amd (rocm10) test/registered/amd/disaggregation/test_disaggregation_basic.py::TestDisaggregationAccuracy.setUpClass Qwen3-8B prefill server exits at startup: torch.AcceleratorError: CUDA error: out of memory in scheduler.run_event_loop, then Server process exited with code -9 Same out-of-memory crash on 4 of 4 runs that reached it: 39570716186 and 39522599360 (same image build), 39421857090 and 39348992034 🟢 Same GPU, same launch config, same failure mode. The model is Qwen3-8B, and the PR only changes DeepSeek-V4/V4.1 HIP code and a HIP-only gate in context-parallel can_apply. Against: none found.
same job test/registered/amd/disaggregation/test_disaggregation_pp.py::TestDisaggregationDecodePPAccuracy.setUpClass Same out-of-memory crash, then prefill server exit -9 (TP2/PP2) Same mode on 39522599360, 39421857090, 39348992034; not reached in 39570716186 🟢 Same launch config and mode on 3 comparable runs. The model is outside the PR's scope.
stage-b-test-1-gpu-small-amd-mi35x (rocm10, 1) test/registered/amd/test_dsv4_hip_bcg_metadata.py::TestAiterSparseLengthFoldPerStep::test_graph_bucket_refold_follows_the_new_lengths AttributeError: 'types.SimpleNamespace' object has no attribute 'swa_req_ring_size' at deepseek_v4_backend.py:4696 (get_dspark_swa_page_indices), same on both attempts Same error at the same frame, on both attempts, on 39522599360, 39421857090, 39348992034 🟢 The failing frame is in deepseek_v4_backend.py, which the PR doesn't change. The test's mock lacks swa_req_ring_size on main too. Against: the PR does edit deepseek_v4_backend_hip_radix.py, which is in the stack, but only line numbers moved (2009 → 2045).
stage-b-test-1-gpu-small-amd-mi35x (rocm10, 0) test/registered/attention/unittests/dsv4/test_dsv41_bcg_hip.py::TestHipBreakableGraphCaptureReplay::test_capture_replay_matches_eager AssertionError: 2 != 1 at line 168, same on both attempts Same 2 != 1 on both attempts on 39522599360, 39421857090 (same image build), 39348992034 🟢 Same assertion and values on every comparable main run. The PR's new paths need context parallelism to be active, and this single-GPU unit test doesn't turn it on. Against: none found.

Other Executed Failures

Job Test Failure On main Related? Why
build-test (xeon-spr, …, 4) test/registered/unit/layers/moe/test_standard_dispatcher_runner_backend.py (all 7 cases) TypeError: MoeRunnerConfig.__init__() got an unexpected keyword argument 'runner_backend' No comparable main run (cancelled or not in run) 🟢 Version mismatch inside CI, not the PR's code (details below the table).
base-a-test-cpu (7) test/registered/unit/layers/moe/test_standard_dispatcher.py::TestLayerRunnerDispatch::test_gfx1250_internal_router_requires_full_expert_weights AssertionError: ValueError not raised Earlier main runs passed only because the test didn't exist yet. It was added in #44993 🟢 Known main breakage, already fixed (details below the table).

Xeon (xeon-spr) failure:

  • The Xeon image installs sglang from the PR branch head fangyuan/dsv4p1_pcp (174fc987) into site-packages (--build-arg VER_SGLANG=$PR_HEAD_REF).
  • The tests run from the merge checkout (/sglang-checkout).
  • The branch is 7 commits behind main. Its merge base is c1e43b95.
  • MoeRunnerConfig.runner_backend was added by 4f8f06fa (#44993) after that merge base. base.py at the PR head confirms the field is missing.
  • The PR doesn't touch MoE code.

base-a-test-cpu (7) failure:

What to do before merge

  • Rebase onto current main and re-run the CPU jobs. This should clear both CPU failures: the Xeon image would pick up MoeRunnerConfig.runner_backend, and base-a-test-cpu would pick up the #45476 test fix. If either still fails after the rebase, the 🟢 verdict for it is wrong and needs another look.
  • Run the PR's updated test on its own on MI35x, because fail-fast cut it off in CI:
    python3 -m pytest test/registered/amd/test_dsv4_hip_bcg_metadata.py -k "cp_prefill or cp_reindex or cp_unified"
    It should pass. A failure there would be the first evidence against this PR.
  • For the actual feature (V4.1 prefill context parallelism on HIP), run a DeepSeek-V4.1 prefill-CP accuracy check on MI35x. As far as these logs show, CI never exercised that path.

Generated by amd-bot using Claude Code CLI

At1a8 and others added 6 commits October 23, 2026 02:57
Resolve the conflict with sgl-project#45783 in deepseek_v4_backend_hip_radix.py: main
removed the SGLANG_DSV41_TORCH_PREFILL_INDEXER check from the extend branch of
_low_ratio_index_topk and the PR removed the prefill CP assert next to it, so
drop both.
…io_sources

sgl-project#45407 slices q_lora with [:n] on the hoisted-rows path of
forward_low_ratio_sources. On gfx950 the indexer's wq_b operand is an
Mxfp8Activation or Fp8GridActivation, which refuses slicing, so a DeepSeek-V4.1
DSpark server crashes while capturing the target-verify CUDA graph. Move
_first_rows from the HIP backend into the shared backend and take the rows
through it.
Gate the gfx950 q_lora row helper behind get_platform().is_hip and keep it in the
HIP backend, so the CUDA path still slices q_lora with [:n].
SAXEM1997 pushed a commit to SAXEM1997/sglang that referenced this pull request Oct 25, 2026
Decoder SWA bounded replay keeps only each request's 128-token tail in the
late layers. With prefill CP, the runner's gather expects this rank's full
shard layout, so CUDA needs the tail rows put back even when DSpark captures
aux hidden states; the DSpark aux tensors must be scattered the same way.
HIP keeps rejecting replay + CP until sgl-project#43465 lands.

Co-Authored-By: Claude Code <noreply@anthropic.com>
@At1a8

At1a8 commented Oct 28, 2026

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@bingxche

Copy link
Copy Markdown
Collaborator

@At1a8

CI Status for PR #43465

Warning

Verdicts the evidence does not support (treat them as unknown):

Verdict: Neither executed failure comes from this PR. Both are bugs that were already on main at the PR's merge base 26a1e511, and both were fixed on main after that base. Rebasing onto current main should clear them. PR CI is not complete: the pr-test-amd-extra.yml and pr-test-extra.yml gates did not pass, so their test jobs were skipped.

Note

Each failure has a fix on main that is not in this PR's base: #46236 (5b0be031) for the ROCm 10 V4.1 k-norm-rope build, and #46441 (8db75fc6) for the MLX stub's alloc_memory_pool signature. The compare API reports both as ahead of 26a1e511.

Executed failures: AMD: 1 (0 related) · Others: 1 (0 related)

AMD Executed Failures

Job Test Failure On main Related? Why
stage-b-test-1-gpu-small-amd-mi35x (rocm10, 1) test/registered/kernels/ops/attention/dsv4/test_v41_kv_store.py::TestFusedKNormRopeFlashMLA.test_query_rope_in_the_k_launch The JIT build failed on gfx950: RuntimeError: Failed to build JIT module sgl_kernel_jit_dpsk_v4_main_k_norm_rope_flashmla_…_V41_false. The cause is a static_assert in ROCm 10 amd_warp_sync_functions.h:307 (sizeof(MaskT) == 8, evaluated as 4 == 8), triggered by __shfl_xor_sync(0xffffffffu, …) at main_norm_rope.cuh:513 (log lines 13755–13830) Same error on run 40250572724 (tested 8252e3af, before the fix). Passed on 40333555627 (316ff4dc) and 40427865857 (5884e4fe), both after the fix ⚪ unverified Comparable recurrence and independence. The PR changes no .cuh/JIT file; its diff (D1–D22) is Python only, in the DSv4 HIP backend, CP and model code. This is a compile error raised before any PR code runs. main on the same GPU and image line failed in the same way until #46236 changed exactly these __shfl_xor_sync calls to __shfl_xor on ROCm. That fix is in the 9 commits between 26a1e511 and 8252e3af…316ff4dc, not in this PR's base. Against: none. The PR's own modified test, test/registered/amd/test_dsv4_hip_bcg_metadata.py, ran in this same job and passed (Ran 19 tests, log line 14170).

Other Executed Failures

Job Test Failure On main Related? Why
stage-a-unit-test-mlx test/registered/unit/hardware_backend/mlx/test_mlx_runner_pool_contract.py::TestMlxRunnerPoolContract.test_worker_startup_uses_mlx_token_capacity TypeError: MlxModelRunnerStub.alloc_memory_pool() got an unexpected keyword argument 'draft_parallel_layout', the same in both attempts (log lines 1229–1298) No completed non-PR baseline (G1) 🟢 Independence. tp_worker.py:409-416 passes draft_parallel_layout= to the model runner, but at the tested checkout the stub's signature at python/sglang/srt/hardware_backend/mlx/model_runner_stub.py:340 is still alloc_memory_pool(self, memory_pool_config=None). The PR touches neither file, nor anything under MLX. #46441 (8db75fc6) adds exactly this keyword to the stub and landed after 26a1e511. Against: there is no main run to show the same failure directly.

The failed gate and coordinator jobs (call-gate / pr-gate, pr-test-amd-finish, pr-test-amd-extra-finish, pr-test-extra-finish, pr-test-mlx-finish) are follow-on failures of the rows above, or gates that skipped their suites. They get no verdict of their own.

What to do before merge

  • Rebase onto current main and re-run pr-test-amd.yml and pr-test-mlx.yml. The new base will contain #46236 and #46441. If both tests pass afterwards, the 🟢 verdicts are confirmed. If either still fails with the same error after the rebase, revisit the verdict.
  • If the DSv4 HIP CP change should also be covered by the extra suites, get pr-test-amd-extra.yml past its gate. Its 4 test jobs were skipped on this head.

Generated by amd-bot using Claude Code CLI

@At1a8

At1a8 commented Oct 29, 2026

Copy link
Copy Markdown
Contributor Author

@amd-bot ci-status

@bingxche

Copy link
Copy Markdown
Collaborator

@At1a8

CI Status for PR #43465

Warning

Verdicts the evidence does not support (treat them as unknown):

Verdict: None of the three executed AMD failures were caused by this PR. Two of them came from a main regression in the ROCm KV memory pool that was already in this PR's merge base. That regression has since been fixed on main by #46861 (86f72e5b). The third, test_disaggregation_pp.py, is a known flaky failure. PR CI is incomplete: the opt-in pr-test-amd-extra and pr-test-extra tiers did not run because their call-gate / pr-gate jobs failed.

Note

CI tested d9e3234, which merges the PR head into main at 15d39bd8. That base already contains #32272 (1a120754e), which stores K and V in one shared allocation (memory_pool.py:2630). It does not yet contain the ROCm fix #46861. Merging current main and re-running should clear F1 and F3.

Executed failures: AMD: 3 (0 related) · Others: 0 (0 related)

AMD Executed Failures

Job Test Failure On main Related? Why
stage-b-test-large-8-gpu-mi35x-disaggregation-amd (rocm10) test/registered/amd/disaggregation/test_disaggregation_basic.py::TestDisaggregationAccuracy.test_gsm8k AssertionError: 0.01 not greater than 0.7. Accuracy 0.010, invalid 0.400 on attempt 1; accuracy 0.005, invalid 0.455 on attempt 2. Both attempts fail with garbage output. Comparable main runs in the pack passed at 0.95–0.955 (F1.M2, F1.M5, F1.M6), but all ran on bases from before #32272. On 2026-10-28 08:20–10:24 UTC, 19 other PRs' runs of this job failed the same way (accuracy 0.00–0.015) on bases 990ee898…edd270fd. Every run on a base from 86f72e5b onward scored 0.94–0.965, e.g. 123212405384. 🟢 Independence. The test serves Qwen3-8B with TP=1, PD disaggregation and no context parallel. Most of the diff is DeepSeek-V4-only (D1–D15, D18, D19). The one shared hunk (D17, interleave.py::can_apply) only adds an early return False. It is reached only through is_cp_active() in CP-extend mode (layers/cp/utils.py:121-135). The PR does not touch memory_pool.py. The fix commit message says "same-node Mooncake PD on ROCm decodes garbage when V sits inside K's allocation", which matches this failure. Against: the pack's main records only show passes, so this rests on the other PRs' runs listed here.
stage-b-test-large-8-gpu-mi35x-disaggregation-amd (rocm10) test/registered/amd/disaggregation/test_disaggregation_pp.py::TestDisaggregationDecodePPAccuracy.test_gsm8k RuntimeError: Decode transfer failed … KVTransferError … Failed to get kvcache on both attempts F2.M6 (run 40250572724, same GPU, image line and launch config) failed both attempts with the same Decode transfer failed error. F2.M2 and F2.M5 passed. On 2026-10-28 it also failed this way in at least 15 other PRs' runs, including bases that already have #46861, e.g. 123095532409. ⚪ unverified Comparable recurrence. Same failure mode and same frequency as on main, both attempts included. Nothing here is worse than on main (no eval metrics either way). The PR's TP2/PP2 Qwen3 path does not use DeepSeek-V4 or CP.
stage-b-test-1-gpu-small-amd (rocm10, 0) none (runner lost during Run test, step 7) Runner linux-mi300-1gpu-sglang-runner-gs6qx lost after about 36 min (02:15→02:51 UTC). The log returns 404. Passed on main (F3.M2, F3.M5, F3.M6). The same shard lost its runner in 16 other PRs' runs between 2026-10-28 08:20 and 10:43 UTC: different runners, each about 36 min, no failed step, same annotation, e.g. 122707472072. Those runs had pre-#46861 bases. Once bases included 86f72e5b, the shard passed in about 12–15 min. 🟢 Independence. This is a shard that takes its runner down, and it happens across many unrelated PRs during the same main-regression window as F1. That points at the same main code, not at this PR. Neither test the PR modifies runs in this suite: test_dsv4_hip_bcg_metadata.py is in stage-b-test-1-gpu-small-amd-mi35x, and test_platform_prefill_cp_deprecation.py is in base-a-test-cpu. The PR changes no CI or image files. Against: this PR had only one attempt, and there is no log showing which test was running.

pr-test-amd-finish, pr-test-amd-extra-finish and pr-test-extra-finish failed only as roll-ups of the jobs above or of the gate. The two call-gate / pr-gate jobs exited with code 1 (opt-in gate), so every extra-* job was skipped.

What to do before merge

  • F1 and F3: merge or rebase onto current main (it includes 86f72e5b) and re-run pr-test-amd.yml. If test_disaggregation_basic.py scores about 0.95 and stage-b-test-1-gpu-small-amd (rocm10, 0) finishes in about 15 min, the verdicts hold. If basic accuracy stays near 0 on a fixed base, look at the PR again.
  • F2: re-run the disaggregation job. A pass, or the same Decode transfer failed error that other PRs get, confirms it is unrelated.
  • Extra tiers: if this DeepSeek-V4.1 CP change needs the pr-test-amd-extra suites, satisfy the opt-in gate and re-run it.

Generated by amd-bot using Claude Code CLI

ChangLi811 added a commit to ChangLi811/sgl-amd-mono that referenced this pull request Oct 29, 2026
With --enable-prefill-cp SGLang keeps all 64 attention heads on every rank
(attention TP 1) and the whole KV and index caches on every rank, but decode and
verify steps never run CP. K1 and K2 therefore run this rank's TP share of the
heads as without CP: wq_b's and wo_a's rows are views of the replicated weights,
wo_b's columns (in aiter's (16, 16) preshuffle too) and its block codes a copy.
The kernels are unchanged. The gate accepts CP only as prefill CP over the whole
TP group; decode CP or a narrower CP group is refused.

The README names sgl-project/sglang#43465 (prefill CP on HIP) among its
requirements.
SAXEM1997 added a commit to SAXEM1997/sglang that referenced this pull request Nov 1, 2026
Decoder SWA bounded replay keeps only each request's 128-token tail in the
late layers. With prefill CP, the runner's gather expects this rank's full
shard layout, so CUDA needs the tail rows put back even when DSpark captures
aux hidden states; the DSpark aux tensors must be scattered the same way.
HIP keeps rejecting replay + CP until sgl-project#43465 lands.

Co-Authored-By: Claude Code <noreply@anthropic.com>
@kkHuang-amd

Copy link
Copy Markdown
Collaborator

Reviewed the full diff. Most of this is a faithful mirror of the CUDA prefill-CP work onto HIP — the cp_metadata parameter and swa_replay_start padding in init_forward_metadata_prefill, the tail layout in _build_late_layer_tail_metadata, and the save/restore of attn_cp_metadata plus the DP buffer length in enter_late_layer_tail / exit_late_layer_tail all match what deepseek_v4_backend.py already does. Those parts look right.

Three things I think should block, one of which cannot execute at all.

1. The copied forward_low_ratio_sources is a stale fork, and it is dead code

The PR drops the inherited bindings:

-    forward_low_ratio_sources = DeepseekV4AttnBackend.forward_low_ratio_sources
-    _forward_low_ratio_sources_cp = DeepseekV4AttnBackend._forward_low_ratio_sources_cp

and replaces them with a ~50-line copy of the CUDA body, justified in the docstring as:

The CUDA backend's body, except that the hoisted q_lora rows go through _first_rows: gfx950 hands q_lora over as an fp8-grid or MXFP8 activation, which refuses plain slicing.

That reason no longer holds at this PR's own merge base (15d39bd840, 2026-10-27):

  • _first_rows is already in deepseek_v4_backend.py:333, and it was written for HIP — its docstring reads "cut field by field for HIP's quantized activations (Mxfp8Activation, Fp8GridActivation), which refuse a slice". It landed on 2026-10-23 in [Fix] [AMD] dsv4.1: Fix decode graph capture with quantized low-ratio activations #46151, "[Fix] [AMD] dsv4.1: Fix decode graph capture with quantized low-ratio activations".
  • The CUDA forward_low_ratio_sources already calls it for both operands:
            n = min(hoisted_pos.shape[0], positions.shape[0])
            if n < positions.shape[0]:
                x, positions = _first_rows(x, n), positions[:n]
                if q_lora is not None:
                    q_lora = _first_rows(q_lora, n)
            req, pos = hoisted_req[:n], hoisted_pos[:n]

So inheriting it would already have done what the fork exists to do. Meanwhile the fork has drifted in three places:

  • Wrong arity. The copy calls _cuda_backend._bcg_prefill_graph_full_topk(layer, x, q_lora, pos) with four arguments and consumes a return value, but the signature at the merge base is _prefill_graph_full_topk(layer, x, q_lora, pos, table) — five arguments, returning None (changed in [Refactor] dsv41: Write the candidate table into caller-held prefill graph buffers #46399, 2026-10-25). This is a TypeError.
  • Missing the use_fused_gate branch. The CUDA body selects _low_ratio_compress_fused when layer.compressor.use_fused_gate is set; the copy only has the _low_ratio_compress_torch arm.
  • x[:n] instead of _first_rows(x, n). The copy applies _first_rows only to q_lora, which is the opposite of what its own docstring argues for.

And the whole block is unreachable on HIP anyway: low_ratio_prefill_graph = False is a class attribute on DeepseekV4HipRadixBackend (line 894 at the merge base), so _low_ratio_in_prefill_graph() is always False here and the ~35 lines of prefill-graph handling can never run. The arity bug is therefore latent rather than live, but the lines are pure maintenance debt that will keep diverging from the CUDA body.

Suggest deleting the copy and restoring forward_low_ratio_sources = DeepseekV4AttnBackend.forward_low_ratio_sources. The only piece that genuinely needs a HIP override is _forward_low_ratio_sources_cp, for the tail is None case that uses _low_ratio_compress instead of _low_ratio_compress_torch.

2. The new _first_rows in the HIP file duplicates and weakens the existing one

if isinstance(act, Fp8GridActivation):
    return Fp8GridActivation(act.x[:n])
if isinstance(act, Mxfp8Activation):
    return Mxfp8Activation(act.q[:n], act.scale[:n])

The existing helper uses msgspec.structs.replace over __struct_fields__, so it follows these structs if a field is ever added. This version hard-codes the field names and builds positionally, so a new field would be silently dropped. Two functions with the same name and the same purpose in two modules is also worth avoiding — importing the existing one covers both cases.

3. The DSpark CP branch in deepseek_v4.py cannot run

if hip_cp_tail and capture_dspark:
    dspark_aux_hidden_states = [
        _scatter_tail_rows(tail=tail, rows=aux, num_tokens=input_ids.shape[0])
        for aux in dspark_aux_hidden_states
    ]

if capture_dspark:
    return (hidden_states, pre_hc_head), dspark_aux_hidden_states.finalize()

dspark_aux_hidden_states is annotated Optional[AuxHiddenStatePacker] (deepseek_v4.py:3977). AuxHiddenStatePacker defines __len__ but neither __iter__ nor __getitem__, so the comprehension raises TypeError: 'AuxHiddenStatePacker' object is not iterable. Even if it iterated, rebinding the name to a plain list breaks the .finalize() call two lines below.

The packer exists to write every capture into a single preallocated [tokens, K * hidden] buffer, so the scatter belongs on the finalized buffer rather than on each capture:

aux = dspark_aux_hidden_states.finalize()
if hip_cp_tail:
    aux = _scatter_tail_rows(tail=tail, rows=aux, num_tokens=input_ids.shape[0])
return (hidden_states, pre_hc_head), aux

One thing to confirm with that shape: AuxHiddenStatePacker.for_batch aliases forward_batch.aux_hidden_states_buffer, and the scatter produces a new tensor. That should be fine for prefill, but it is worth checking that no caller depends on the returned tensor being that buffer.

Other

The gfx950 NaN workaround is broader than the bug it works around.

if _is_hip and seq_len < cp_size * _HIP_MIN_TOKENS_PER_RANK:
    return False

This lives in the generic interleave CP file and is gated only on is_hip(), so it disables interleave CP for short prompts on every HIP model, while the failing kernel is DeepSeek-V4.1's wqkv_a (N=1792, K=5120). Separately, seq_len >= cp_size * 16 is an average rather than a per-rank guarantee: with multiple requests the interleave shares are uneven, so a rank can still end up below the threshold while the total clears it. Since the NaN is known at M=5..8, deriving the check from min(interleave_rows_per_request(...)) would say what is meant, and scoping it to the model (or the GEMM) would keep it off unrelated HIP runs. The TODO acknowledges the workaround, but not its blast radius.

Test coverage. 260 added lines, and the only test changes are two mechanical signature updates. Nothing exercises the HIP CP tail layout, _forward_low_ratio_sources_cp, the enter/exit save-restore of attn_cp_metadata and the DP buffer length, or the DSpark scatter — which is presumably how point 3 went unnoticed. test_dsv4_hip_bcg_metadata.py already has the _check_cp_prefill_order harness that could be extended for most of this.

The relaxed validation is only tested by subtraction. test_platform_prefill_cp_deprecation.py drops the ("support DeepSeek-V4.1", dict(model_type="deepseek_v41")) rejection case but adds no positive case asserting that V4.1 with enable_decoder_swa_bounded_replay is now accepted. Nothing pins the newly allowed combination.

@At1a8

At1a8 commented Nov 5, 2026 •

Copy link
Copy Markdown
Contributor Author

Reviewed the full diff. Most of this is a faithful mirror of the CUDA prefill-CP work onto HIP — the cp_metadata parameter and swa_replay_start padding in init_forward_metadata_prefill, the tail layout in _build_late_layer_tail_metadata, and the save/restore of attn_cp_metadata plus the DP buffer length in enter_late_layer_tail / exit_late_layer_tail all match what deepseek_v4_backend.py already does. Those parts look right.

Three things I think should block, one of which cannot execute at all.

1. The copied forward_low_ratio_sources is a stale fork, and it is dead code

The PR drops the inherited bindings:

-    forward_low_ratio_sources = DeepseekV4AttnBackend.forward_low_ratio_sources
-    _forward_low_ratio_sources_cp = DeepseekV4AttnBackend._forward_low_ratio_sources_cp

and replaces them with a ~50-line copy of the CUDA body, justified in the docstring as:

The CUDA backend's body, except that the hoisted q_lora rows go through _first_rows: gfx950 hands q_lora over as an fp8-grid or MXFP8 activation, which refuses plain slicing.

That reason no longer holds at this PR's own merge base (15d39bd840, 2026-10-27):

  • _first_rows is already in deepseek_v4_backend.py:333, and it was written for HIP — its docstring reads "cut field by field for HIP's quantized activations (Mxfp8Activation, Fp8GridActivation), which refuse a slice". It landed on 2026-10-23 in [Fix] [AMD] dsv4.1: Fix decode graph capture with quantized low-ratio activations #46151, "[Fix] [AMD] dsv4.1: Fix decode graph capture with quantized low-ratio activations".
  • The CUDA forward_low_ratio_sources already calls it for both operands:
            n = min(hoisted_pos.shape[0], positions.shape[0])
            if n < positions.shape[0]:
                x, positions = _first_rows(x, n), positions[:n]
                if q_lora is not None:
                    q_lora = _first_rows(q_lora, n)
            req, pos = hoisted_req[:n], hoisted_pos[:n]

So inheriting it would already have done what the fork exists to do. Meanwhile the fork has drifted in three places:

  • Wrong arity. The copy calls _cuda_backend._bcg_prefill_graph_full_topk(layer, x, q_lora, pos) with four arguments and consumes a return value, but the signature at the merge base is _prefill_graph_full_topk(layer, x, q_lora, pos, table) — five arguments, returning None (changed in [Refactor] dsv41: Write the candidate table into caller-held prefill graph buffers #46399, 2026-10-25). This is a TypeError.
  • Missing the use_fused_gate branch. The CUDA body selects _low_ratio_compress_fused when layer.compressor.use_fused_gate is set; the copy only has the _low_ratio_compress_torch arm.
  • x[:n] instead of _first_rows(x, n). The copy applies _first_rows only to q_lora, which is the opposite of what its own docstring argues for.

And the whole block is unreachable on HIP anyway: low_ratio_prefill_graph = False is a class attribute on DeepseekV4HipRadixBackend (line 894 at the merge base), so _low_ratio_in_prefill_graph() is always False here and the ~35 lines of prefill-graph handling can never run. The arity bug is therefore latent rather than live, but the lines are pure maintenance debt that will keep diverging from the CUDA body.

Suggest deleting the copy and restoring forward_low_ratio_sources = DeepseekV4AttnBackend.forward_low_ratio_sources. The only piece that genuinely needs a HIP override is _forward_low_ratio_sources_cp, for the tail is None case that uses _low_ratio_compress instead of _low_ratio_compress_torch.

2. The new _first_rows in the HIP file duplicates and weakens the existing one

if isinstance(act, Fp8GridActivation):
    return Fp8GridActivation(act.x[:n])
if isinstance(act, Mxfp8Activation):
    return Mxfp8Activation(act.q[:n], act.scale[:n])

The existing helper uses msgspec.structs.replace over __struct_fields__, so it follows these structs if a field is ever added. This version hard-codes the field names and builds positionally, so a new field would be silently dropped. Two functions with the same name and the same purpose in two modules is also worth avoiding — importing the existing one covers both cases.

3. The DSpark CP branch in deepseek_v4.py cannot run

if hip_cp_tail and capture_dspark:
    dspark_aux_hidden_states = [
        _scatter_tail_rows(tail=tail, rows=aux, num_tokens=input_ids.shape[0])
        for aux in dspark_aux_hidden_states
    ]

if capture_dspark:
    return (hidden_states, pre_hc_head), dspark_aux_hidden_states.finalize()

dspark_aux_hidden_states is annotated Optional[AuxHiddenStatePacker] (deepseek_v4.py:3977). AuxHiddenStatePacker defines __len__ but neither __iter__ nor __getitem__, so the comprehension raises TypeError: 'AuxHiddenStatePacker' object is not iterable. Even if it iterated, rebinding the name to a plain list breaks the .finalize() call two lines below.

The packer exists to write every capture into a single preallocated [tokens, K * hidden] buffer, so the scatter belongs on the finalized buffer rather than on each capture:

aux = dspark_aux_hidden_states.finalize()
if hip_cp_tail:
    aux = _scatter_tail_rows(tail=tail, rows=aux, num_tokens=input_ids.shape[0])
return (hidden_states, pre_hc_head), aux

One thing to confirm with that shape: AuxHiddenStatePacker.for_batch aliases forward_batch.aux_hidden_states_buffer, and the scatter produces a new tensor. That should be fine for prefill, but it is worth checking that no caller depends on the returned tensor being that buffer.

Other

The gfx950 NaN workaround is broader than the bug it works around.

if _is_hip and seq_len < cp_size * _HIP_MIN_TOKENS_PER_RANK:
    return False

This lives in the generic interleave CP file and is gated only on is_hip(), so it disables interleave CP for short prompts on every HIP model, while the failing kernel is DeepSeek-V4.1's wqkv_a (N=1792, K=5120). Separately, seq_len >= cp_size * 16 is an average rather than a per-rank guarantee: with multiple requests the interleave shares are uneven, so a rank can still end up below the threshold while the total clears it. Since the NaN is known at M=5..8, deriving the check from min(interleave_rows_per_request(...)) would say what is meant, and scoping it to the model (or the GEMM) would keep it off unrelated HIP runs. The TODO acknowledges the workaround, but not its blast radius.

Test coverage. 260 added lines, and the only test changes are two mechanical signature updates. Nothing exercises the HIP CP tail layout, _forward_low_ratio_sources_cp, the enter/exit save-restore of attn_cp_metadata and the DP buffer length, or the DSpark scatter — which is presumably how point 3 went unnoticed. test_dsv4_hip_bcg_metadata.py already has the _check_cp_prefill_order harness that could be extended for most of this.

The relaxed validation is only tested by subtraction. test_platform_prefill_cp_deprecation.py drops the ("support DeepSeek-V4.1", dict(model_type="deepseek_v41")) rejection case but adds no positive case asserting that V4.1 with enable_decoder_swa_bounded_replay is now accepted. Nothing pins the newly allowed combination.

Thanks for @kkHuang-amd 's comments — all three blocking points were right.

  1. Stale forward_low_ratio_sources fork. Agreed. Main branch landed the shared _first_rows while this PR was open, so the fork had nothing left to do. Deleted it and restored forward_low_ratio_sources = DeepseekV4AttnBackend.forward_low_ratio_sources; HIP now overrides only _forward_low_ratio_sources_cp, whose non-tail case runs the fused _low_ratio_compress.

  2. Duplicate _first_rows. Removed; the HIP backend imports the one in deepseek_v4_backend.

  3. DSpark CP branch. It was written while the captures were still a list; Main branch switched them to AuxHiddenStatePacker and the merge had no conflict. The HIP CP path now finalizes the packer and scatters the packed buffer.

  4. NaN workaround scope. Moved off the generic HIP path: interleave.py reads SGLANG_CP_MIN_TOKENS_PER_RANK (default 0, so nothing changes for other models or platforms), and the DeepSeek-V4 hook sets it to 16 only for V4.1 prefill CP on HIP, unless the user set it. On per-rank vs average: interleave deals tokens by global index (interleave_rows_per_request), so rank shares differ by at most one row and the smallest shard holds seq_len // cp_size rows; the check is now written that way.

  5. Tests. Added test/registered/amd/accuracy/mi35x/test_deepseek_v41_flash_cp_eval_mi35x.py: GSM8K (1319 questions, 8-shot completion) on MI35x at TP4 + interleave CP4 with DSpark and decoder SWA bounded replay, in the same nightly suite as the non-CP V4.1 test. Local run: 0.912 (non-CP test: 0.911); threshold 0.89 as in the non-CP test.

image

@At1a8

At1a8 commented Nov 6, 2026

Copy link
Copy Markdown
Contributor Author

/rerun-failed-ci

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

bypass-fail-fast CI: a failing job no longer aborts its siblings (lint still gates) deepseek parallel-stages CI: stages dispatch together instead of waiting on each other run-ci CI: run the baseline test suite on this PR

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants