Skip to content

[AMD] [GLM-5.3-Flash Day 0] Build the fused DSA k-pool top-k JIT kernel on HIP - #37563

Closed
Jacob0226 wants to merge 2 commits into
sgl-project:xinyuan/glm-5.3-flash-supportfrom
Jacob0226:jacob/glm53-rocm-kpool-hip
Closed

Jacob0226 wants to merge 2 commits into
sgl-project:xinyuan/glm-5.3-flash-supportfrom
Jacob0226:jacob/glm53-rocm-kpool-hip

Conversation

@Jacob0226

@Jacob0226 Jacob0226 commented Sep 2, 2026

Copy link
Copy Markdown
Contributor

Summary

jit/csrc/dsa/kpool_topk_transform.cuh includes <cuda_fp16.h> directly and reaches ::cudaFuncSetAttribute through no alias, so hipcc cannot compile it. This is a hard stop on ROCm rather than a slow path: GLM-5.3-Flash pools index_topk=2048 over index_kpool=4, giving group_topk=512, and kpool_fp8_index routes every pooled budget in (128, 160, 192, 224, 256, 512) to this JIT kernel — the fast_topk_v2 fallback covers group_topk=2048 alone, so the indexer cannot run at all.

Symbol Fix
__half, __float2half_rn, __half_as_ushort give the <cuda_fp16.h> include a USE_ROCM branch that reaches for <hip/hip_fp16.h>, matching jit/csrc/elementwise/fused_metadata_copy.cuh in the same tree. The CUDA line is untouched
cudaFuncSetAttribute alias it in sgl_kernel/utils.cuh, beside the shim's existing cudaGetLastError and cudaDeviceGetAttribute entries, rather than guarding one call site

Scope

One include wrapped in USE_ROCM branches, two aliases added to the shim's existing USE_ROCM list. Every CUDA line is unchanged context in the diff and resolves exactly as before. No kernel body, launch configuration, tile size or numeric path moves; no environment variable, no architecture gate.

fast_kpool_topk_transform_fused has no direct test on any platform. CUDA reaches it end to end through test_glm53_flash_b200.py, which serves GLM-5.3-Flash and therefore JIT-compiles it every commit; AMD has no equivalent, so this adds focused unit coverage on the mi35x per-PR suite.

Test plan

rocm/sgl-dev:v0.5.18-rocm720-mi35x-20260901, the image family pr-test-amd-rocm720.yml resolves to, on xinyuan/glm-5.3-flash-support @ 515e865189 with no other PRs applied. Both builds run the same new test file; the shim aliases are the only variable. JIT caches were cleared between builds, and the build hashes differ, so neither result is a stale artifact.

Build Result
Baseline 4 errors — every case dies in the JIT build, group_topk 128 and 512 alike, on 'cuda_fp16.h' file not found and no member named 'cudaFuncSetAttribute' in the global namespace
This PR 4 passed on gfx950

CI States

Latest PR Test (Base): ❌ Run #33619800731
Latest PR Test (Extra): ❌ Run #33628108935
Latest PR Test (AMD ROCm 7.2): ❌ Run #33619800720

jit/csrc/dsa/kpool_topk_transform.cuh includes <cuda_fp16.h> directly and
reaches ::cudaFuncSetAttribute through no alias, so hipcc cannot compile
it. On gfx950 the JIT build fails with "fatal error: 'cuda_fp16.h' file
not found", and after that with "no member named 'cudaFuncSetAttribute'
in the global namespace".

This is a hard stop rather than a slow path. GLM-5.3-Flash pools
index_topk=2048 over index_kpool=4, so group_topk is 512, and
kpool_fp8_index routes every pooled budget in (128, 160, 192, 224, 256,
512) to this JIT kernel -- the fast_topk_v2 fallback covers group_topk
2048 alone. The indexer cannot run at all on ROCm.

Give the fp16 include a USE_ROCM arm, matching
jit/csrc/elementwise/fused_metadata_copy.cuh in the same tree, and add
cudaFuncSetAttribute to the USE_ROCM alias list in sgl_kernel/utils.cuh
beside the existing cudaGetLastError and cudaDeviceGetAttribute entries.
Both CUDA lines stay untouched.

Add a focused parity test. The kernel has no direct coverage on any
platform; CUDA reaches it end to end through
test/registered/models_e2e/test_glm53_flash_b200.py, which serves
GLM-5.3-Flash and therefore JIT-compiles it every commit, while AMD has
no equivalent. Registered on the mi35x per-PR suite only.

Co-authored-by: Cursor <cursoragent@cursor.com>
@Jacob0226

Copy link
Copy Markdown
Contributor Author

Recovered as #39339, off main. This one cannot be reopened — its base branch xinyuan/glm-5.3-flash-support no longer exists. The replacement #38542 is closed; please continue review and CI on #39339.

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

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants