Fix deterministic DSV4 top-k fallback on ROCm

[skyguan92]

Summary

• Fall back to the existing DSV4 Top-K v1 JIT implementation when the
  packaged ROCm AOT operator is not registered for the active target.
• Make that JIT path compile with HIP and fit 64-KiB-LDS devices.
• Make ROCm selected-set output deterministic across boundary ties, signed
  zeroes and coarse-bin overflow while preserving CUDA behavior.
• Add a registered AMD test for forced fallback, page mapping, ragged rows,
  runtime Top-K sizes, overflow, long rows and repeat determinism.

Why this is needed

On an unmodified current SGLang checkout, a native gfx1100 installation whose
wheel does not contain the DSV4 AOT operator fails twice: the public Python
path raises AttributeError, and direct JIT v1 compilation calls the
CUDA-only cudaFuncSetAttribute and
cudaFuncAttributeMaxDynamicSharedMemorySize APIs.

The launchable ROCm path also needs deterministic handling at the exact
selection boundary. Atomic arrival order must not choose among equal scores,
signed -0.0/+0.0 must compare equally, and a coarse FP16 bin larger than the
bounded LDS candidate buffer must not be silently clipped.

Contract

DeepSeek-V4 uses index_topk=512. SGLang's current API supports runtime
topk in (0, 1024] and returns an unordered selected set; it does not promise
score-sorted output. This change canonicalizes the selected set by logical
index on ROCm, which provides a deterministic valid representation without
adding a score-order contract. CUDA code paths are unchanged.

Implementation

• Use the HIP function-attribute API and a 48-KiB dynamic-LDS budget on ROCm.
• Resolve exact-score ties by lower logical index.
• Normalize signed zeroes before building ROCm radix keys.
• Rescan an overflowing coarse bin with full FP32 keys instead of clipping it.
• Use deterministic short-row compaction and fixed long-row ordering.
• Check AOT operator capability before dispatching; use JIT v1 when absent.

Physical AMD validation

Exact source pair:

base:  ad47dde65c0f5e4812e72c0eb8b3e98acbba8e1a
head:  5a69fd3e5f4ae687cb969ba8930bac9099b3e740
tree:  c1652c0ac80ded83e86e71a1195872ba6b883346
diff:  3 files, +586/-1

The tested parent is the upstream main fetched immediately before this
publication gate. It is 32 commits ahead of the previous qualified parent;
none of its 247 changed paths intersects the three candidate paths. Relative
to the previous candidate 3076f2cd3, the range-diff is equal, stable patch ID
8307236dde8e is unchanged, and all three candidate-file SHA-256 values are
exact.

Immediately before push, upstream added one documentation-only commit that
renames the Qwen3.8 DSpark recipe. It changes two cookbook paths and has zero
intersection with this patch. The final publication source is therefore
d21eefc94ff8e95ea70ba54ddde83b35ff26d340..bd52a2dbf40c1aa287834416fa065dd7f187774f,
tree 4dc2bdb2a380685e28c94f40bd7adb12a2bc272e. Its range-diff and stable
patch ID are equal to the exact physical pair above; the physical timings are
not relabelled as a run of the documentation-only replay.

Environment:

GPU: Radeon PRO W7900 / native gfx1100 (independent GPU 0 and GPU 7 runs)
HSA_OVERRIDE_GFX_VERSION: unset in accepted GPU processes
HIP_VISIBLE_DEVICES: unset in accepted GPU processes
CUDA_VISIBLE_DEVICES: unset in accepted GPU processes
ROCR_VISIBLE_DEVICES: 0 or 7
ROCm: 7.2.53211-e1a6bc5663
PyTorch: 2.9.1+gitff65f5b

Unmodified parent blocker:

exact red parent: ad47dde65c0f5e4812e72c0eb8b3e98acbba8e1a
AOT operator available: false
public path: AttributeError
direct JIT: HIP compile error on the two CUDA-only function-attribute symbols

Candidate:

GPU 0: 21 passed, 5 warnings in 15.66s; exit 0; OOMKilled=false
GPU 7: 21 passed, 5 warnings in 15.69s; exit 0; OOMKilled=false

The test matrix covers forced AOT fallback, page sizes 1/64/256, signed
zeroes, exact boundary ties, dense-bin overflow, ragged rows, the
262144-token upper tested length, runtime
topk={1,7,257,513,777,1024}, and repeated output identity. The repository
registry checker passes; the file resolves to AMD
jit-kernel-unit-test-amd with a 60-second estimate.

The accepted parent and both candidate runs explicitly remove HSA/HIP/CUDA
visibility variables and select one physical card only through
ROCR_VISIBLE_DEVICES. A registry-only harness first missed the image's venv
on PATH and exited before running the checker; it is excluded. The rerun used
the explicit interpreter, and both the repository checker and suite validation
passed.

Performance guard

This is a correctness and portability fix, not a speedup claim. On the same
source bytes in the predecessor qualification, the complete candidate versus
a launchable control was +1.217% p50 for page output and +1.153% for page
plus raw output. A matched ModelScope DeepSeek-V4-Flash-0731 TP8 A-B-A run
completed 27/27 S3072 requests; the full 21-layer Indexer bucket changed by
+0.162186% and +0.179978% in the two candidate brackets, inside the
predeclared 1% guard. That model gate used the same disclosed cached-BF16
lower shim on both sides and is predecessor/source-equivalent evidence, not an
exact-current native-FP8 throughput claim.

Related work

• [BugFix][DSA] Harden top-k v1/v2 kernels against negative padded seq_lens #31123 handles negative padded seq_lens on pre-rewrite history. Its
  contract is independent; a semantic v1 compatibility port passed beside
  this patch, but that scope is intentionally not included here.
• Fix(sgl-kernel/rocm): enable RDNA3.5 (gfx1151) and cap TopK LDS #28518 targets gfx1151 AOT/LDS enablement and does not provide this JIT-v1
  fallback, HIP attribute or deterministic selected-set behavior.
• [kernel] Add opt-in deterministic CuTe-DSL DSA top-k backend (--dsa-topk-backend cutedsl) #31115 is an opt-in deterministic CuTe-DSL DSA Top-K backend for CUDA
  SM90/SM100. It does not repair the existing gfx1100 DSV4 Top-K v1 path,
  missing-AOT fallback, HIP attributes, 64-KiB-LDS behavior or AMD registered
  coverage.

Test plan

• git diff --check
• Python byte compilation
• registered-test repository checker
• registry collection and suite validation
• physical gfx1100 parent red gate
• physical gfx1100 candidate on GPU 0: 21/21
• physical gfx1100 candidate on GPU 7: 21/21
• predecessor operator latency guard
• predecessor formal-checkpoint Indexer A-B-A guard

---

CI States

Latest PR Test (Base): ❌ Run #31612796955
Latest PR Test (Extra): ❌ Run #31612796402

[skyguan92]

Lint is green, including pre-commit and the sgl-kernel clang-format check. The remaining red/skip matrix is the expected label gate: the logs report Missing required label 'run-ci', so no AMD execution job has run yet.

Could an authorized maintainer trigger the standard CI when appropriate? The new test is registered in jit-kernel-unit-test-amd (60 s). For pre-CI hardware evidence, the exact current parent reproduced the missing-AOT/HIP compile blockers on native gfx1100, while the candidate passed 21/21 independently on W7900 GPU 0 and GPU 7.