[AMD] Enable HiSparse on ROCm - #26639
Conversation
|
Warning You have reached your daily quota limit. Please wait up to 24 hours and I will start processing your requests again! |
|
/tag-and-rerun-ci |
|
GSM8k also shows 95% accuracy on MI300X for both attention backends. @HaiShaw |
|
cc @Jacob0226 @Lzy17 |
|
hi @clintg6 , you mentioned "Both TileLang and AITER keep ROCm HiSparse FP8 KV cache in the raw MLA layout: nope 512 fp8 + rope 64 fp8 = 576", so this is for supporting DPSK-V3.2, currently NOT for V4, right? |
|
This PR enables HiSparse on ROCm for DSA models (DeepSeek V3.2, GLM-5). DeepSeek V4 HiSparse support for ROCm would be a follow-up PR once the DSv4 ROCm integration lands on main. @HaiShaw @amd-danli103 |
|
@amd-bot ci-status |
CI Status for PR #26639PR: [AMD] Enable HiSparse on ROCm AMD: 2 failures (1 likely related) | Others: 1 failure (0 related) (The AMD CI Failures
Other CI Failures
Details🔴 🟢 The two AMD VLM failures and the H20 disaggregation failure are in code paths completely untouched by this PR and look like pre-existing flakes / infra issues. Safe to ignore from this PR's perspective, but worth a retry.Generated by amd-bot using Claude Code CLI |
hi @clintg6 glad to know that. Currently most of the V4 optimization jobs have been merged into main as far as I know, except for several ongoing ones. May I know if there will be a schedule/plan for your follow-up PR to support V4? I would like to avoid duplicate work if your next PR is already underway, thank you! |
|
@amd-bot ci-status |
CI Status for PR #26639 — [AMD] Enable HiSparse on ROCmMerge verdict: Caution This PR's main value (ROCm HiSparse: wavefront64 swap-in kernel in
Changed files (13): AMD: startup_failure (0 jobs ran, infra-wide) · Others: ~18 failures (0 clearly PR-caused; 1 needs author check) AMD CI
Other CI Failures
( Details / what to do before merge
Generated by amd-bot using Claude Code CLI |
hi Hai @HaiShaw, currently this PR supports the original separate packed KV memory layout, not the unified memory layout, where the compressed region lives in the unified pool ( I'm tracking this feature since we have a high-priority use case for customer-T. Now I already persuade them to shift to use the unified KV attention, so I'm actually developing the HiSparse feature with unified KV on top of this PR. |
Add a temporary guard for DSv4 HiSparse on unified-KV path.
|
@amd-danli103 @HaiShaw Added a temporary guard in validate_hisparse: HiSparse + unified_kv_triton on DSv4 now errors at startup instead of crashing with a cryptic AssertionError. Scoped to just that combo and doesn't affect NV HiSparse path. @amd-danli103 please remove it when your unified-KV HiSparse support lands. |
Thank you @clintg6 , the guard makes sence. Thanks for letting me know that! |
|
@amd-bot ci-status |
CI Status for PR #26639Merge verdict: ❌ Not ready to merge. PR CI is incomplete (fast-fail cascades skipped downstream NVIDIA jobs; one AMD job Caution This PR's headline functionality is effectively untested by PR CI. The two new end-to-end accuracy tests ( Changed files (13): Executed CI failure attribution: AMD: 3 executed failures (0 related) + 1 queued · Others: ~10 failures across NVIDIA/NPU/XPU/CPU (1 related — CPU). Fast-fail cascade jobs collapsed into their root cause; skipped downstream jobs counted as completeness gaps, not failures. AMD Executed Failures
Other Executed Failures
Fast-fail cascade jobs (not independent failures): base-b small(7), 2gpu-large(2)/(3) — all skipped via Details / what to do before merge
Generated by amd-bot using Claude Code CLI |
|
1test/registered/unit/server_args/test_server_args.py` was fixed (but CI [build-test (xeon-gnr, base-b-test-cpu) is stuck). |
Co-authored-by: clintg6 <7388379+clintg6@users.noreply.github.com> Co-authored-by: HAI <hixiao@gmail.com>
Enable HiSparse for DSA models on ROCm
Summary
This change enables HiSparse for DeepSeek-style DSA models on ROCm, validated on MI355X, while keeping CUDA behavior unchanged. ROCm uses TileLang as the default HiSparse DSA backend, supports FP8 and BF16 KV cache, and can also run with AITER when the user explicitly selects the AITER DSA backend.
The main work is in three areas: making the shared HiSparse swap-in kernel correct for AMD wavefront64, adding the ROCm-specific HiSparse allocator/coordinator paths, and wiring ROCm DSA backend policy so TileLang and AITER use the KV-cache layout their kernels expect.
Why This Is Needed
HiSparse keeps a hot working set of KV cache on GPU and stores the rest in host memory. The existing implementation was built around CUDA assumptions: warp32 behavior, CUDA inline PTX copies, and allocator flows that did not cover ROCm’s
page_size == 1path.On ROCm these assumptions can break swap-in correctness and leak temporary device mappings during decode. For DSA models, ROCm also needs a backend choice that matches the KV-cache layout. TileLang is currently the fastest validated option, so it remains the default. AITER is supported as an explicit opt-in backend for comparison and debugging.
Backend Behavior
With
--enable-hisparseon ROCm, the default DSA backend istilelangfor both prefill and decode.If the user explicitly selects:
AITER is accepted on ROCm. With BF16 KV cache, it uses the normal AITER sparse path. With FP8 KV cache, it automatically uses AITER’s BF16-Q/FP8-KV persistent sparse MLA metadata path. No extra environment flag is required.
Both TileLang and AITER keep ROCm HiSparse FP8 KV cache in the raw MLA layout:
This avoids the CUDA scaled FP8 layout, which is 656 bytes per token and is incompatible with the ROCm sparse MLA views used here.
CUDA Impact
This PR is intended to be behaviorally unchanged on CUDA.
The wavefront64 kernel fixes are guarded with
USE_ROCM, and the CUDA branch keeps the original assumptions: warp32 masks, CUDA inline PTX copy, and the prior LRU writeback loop. The coordinator cleanup for stale temporary HiSparse mappings is also ROCm-gated because CUDA’s swap-in path consumes onlytop_k_device_locs, where those stale mapping entries are harmless.CUDA HiSparse backend selection remains dtype-specific:
Implementation
The swap-in JIT kernel now has ROCm-safe wavefront handling in
python/sglang/jit_kernel/csrc/hisparse.cuh. The code uses a platform-dependentWARP_SIZE, ballot mask type, full-warp mask, and popcount helper. ROCm gets a HIP byte-copy fallback instead of CUDA inline PTX. The inclusive-scan window and large LRU writeback path are fixed underUSE_ROCMso wavefront64 lanes do not write outside the current iteration’s intended shared-memory window.The HiSparse allocator now supports ROCm’s
page_size == 1path.HiSparseTokenToKVPoolAllocator.alloc()allocates both logical and HiSparse device indices and rolls back the logical allocation if the device allocation fails.alloc_extend()no longer assertspage_size > 1, allowing ROCm decode/extend flows to use the normal allocation path.The HiSparse coordinator now frees stale temporary device mappings during ROCm decode remap before replacing them with the reserved device-buffer slot. Without this cleanup, temporary page-size-1 mappings can leak and later corrupt swap-in lookups.
The model-runner KV-cache sizing keeps ROCm TileLang and AITER in the raw 576-wide MLA layout for FP8 KV cache. This is required for both backends and prevents the AITER sparse MLA view failure caused by the CUDA scaled layout.
The AITER DSA path now handles FP8 KV cache without user environment flags. For explicit AITER + FP8 KV, both decode and extend build AITER persistent sparse MLA metadata and pass BF16 Q with FP8 KV into
mla_decode_fwd.Tests And Validation
The packaged changes add focused unit and regression coverage:
page_size == 1allocator path and decode-remap cleanup.Local checks run on the packaged files:
All pass.
Runtime validation on MI355X showed:
TileLang is therefore kept as the default ROCm HiSparse backend. AITER remains useful for explicit comparison and debugging, and it no longer requires a separate environment flag for FP8 KV cache.
Recommended ROCm Starting Point
For MI300X-class GPUs, the validated GLM-5 HiSparse configuration starts from:
These are test/serving recommendations, not forced runtime defaults.
Notes
TileLang remains the recommended ROCm HiSparse backend because it is faster in current testing. AITER is accepted when explicitly selected and uses BF16-Q/FP8-KV persistent sparse MLA automatically for FP8 KV cache.
CI States
Latest PR Test (Base): ❌ Run #27799270152
Latest PR Test (Extra): ❌ Run #27799269960