[aarnetalman]

Motivation

Hybrid Gated DeltaNet (GDN) models (Qwen3-Next, Qwen3.5 and later) are not batch-invariant for short prompts under --enable-deterministic-inference: the same prompt gives different log-probabilities depending on what it is batched with.

The cause is one line in chunk_fwd_o, the last kernel of the GDN chunked prefill:

BT = min(chunk_size, max(16, triton.next_power_of_2(T)))

For varlen input (cu_seqlens), T is the token count of the whole batch, not of one sequence. A 9-token prompt prefilled alone gets a 16-token tile; the same prompt prefilled together with others gets a 32- or 64-token tile. The output is reduced over a different tile, so the last bits change. Every other stage of the GDN prefill (L2 norm, gate cumsum, the intra-chunk solve, the chunk states) gave bitwise-identical results in our tests.

We found it on MI350X with Qwen3.8-27B. The bitwise check below sends the same short prompt 1 to 50 times in one request and compares every response with the batch-size-1 one: token IDs were always identical, but log-probabilities matched in only 33-48 of 1,275 responses at TP2 (triton attention, AITER off; 8 runs), and 50-51 of 1,275 at TP1 with AITER on. The first token's log-probability already differs. Disabling the radix cache, CUDA graphs, or the overlap scheduler, or forcing RCCL for the all-reduce, did not change that.

The same Triton kernels run on NVIDIA in deterministic mode (FlashInfer GDN prefill is refused there), so this is not specific to ROCm.

Modifications

• chunk_fwd_o: use the full chunk size (64) as the tile for varlen input. Fixed-length batches keep the smaller tile, since there T is the sequence length.
• test_chunk_gated_delta_rule.py: new case test_short_sequence_output_independent_of_batchmates: sequences of 10-16 tokens (Qwen3.8's head layout: 16 query/key heads, 48 value heads, 128 wide), 6 seeds each, must give bitwise the same output alone and packed with a copy.
• This also applies outside deterministic mode, but only to prefill batches with at most 32 tokens in total (above that the tile was already 64), where the output changes in the last bits.

Accuracy Tests

MI350X, rocm/sgl-dev:v0.5.21-rocm724-mi35x-20261012 with this change applied.

Unit test, python3 -m pytest test/registered/attention/test_chunk_gated_delta_rule.py: before, the new case fails (1 failed, 30 passed); after, 31 passed.

Kernel level: for 1,896 random sequences (lengths 2-80, 12 seeds, Qwen3.8-27B head counts at TP1 and TP2), each sequence alone vs packed with a copy, every stage of chunk_gated_delta_rule compared bitwise. Before: only chunk_fwd_o differs, in 26 cases, all 10-16 tokens long. After: 0 cases.

End to end, Qwen3.8-27B BF16, --attention-backend triton --linear-attn-prefill-backend triton --linear-attn-decode-backend triton --enable-deterministic-inference, #43788, #34543 and #44076 applied. Bitwise check: 1,275 greedy responses (the single prompt sent 1 to 50 times per request, 100 tokens each), token IDs / log-probabilities identical to batch size 1:

	SGLANG_USE_AITER	before	after
TP1	0	1275 / 1275	1275 / 1275
TP1	1	1275 / 50-51 (2 runs)	1275 / 1275
TP2	0	1275 / 33-48 (8 runs)	1275 / 1275
TP2	1	1275 / 1275	1275 / 1275

python3 -m sglang.test.test_deterministic modes single, prefix, prefix --return-logprob and radix_cache (50 trials) return one output per group in all four configurations after the change; the bitwise result repeated on TP1 (AITER off and on) and TP2 (AITER off). These runs follow the existing deterministic end-to-end tests (TestDeterministicBase, e.g. test_qwen35_deterministic.py), which are CUDA-only and compare decoded text only; here log-probabilities and the bitwise check are compared as well. Happy to add an AMD end-to-end test if wanted.

GSM8K (all 1,319 questions, temperature 0, sgl-eval 0.1.2), Qwen3.8-27B: 95.91% in normal mode (--attention-backend aiter, without these changes), 96.06% with --enable-deterministic-inference and this change together with #43788, #34543, #44076 and #44195.

Speed Tests and Profiling

Only prefill batches with at most 32 tokens in total take a different tile; larger ones are unchanged (the tile was already 64). Time of a whole chunk_gated_delta_rule call for such batches (Qwen3.8-27B shapes, triton.testing.do_bench, median, before/after alternated twice):

Batch	before	after
1 x 9 tokens	205-209 us	211-213 us
2 x 12 tokens	202 us	209-211 us
4 x 8 tokens	202-206 us	209-213 us
1 x 64 tokens (unaffected)	204-207 us	211-216 us
4 x 128 tokens (unaffected)	207-208 us	213-216 us

The unaffected batches moved by the same 2-4%, so that is run-to-run variation; at these sizes the call is bound by kernel launches.

Checklist

• Format your code according to the Format code with pre-commit.
• Add unit tests according to the Run and add unit tests.
• Update documentation according to Write documentations.
• Provide accuracy and speed benchmark results according to Test the accuracy and Benchmark the speed.
• Follow the SGLang code style guidance.

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 #40674924130
Latest PR Test (Extra): ❌ Run #40674924228
Latest PR Test (AMD ROCm 10): ❌ Run #40674923828

[rkarhila-amd]

It's a tiny but necessary change. Again more test code than code changes.