Conversation
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 the GLM-5.3-Flash B200 serving test, which JIT-compiles it every commit, while AMD has no equivalent. Registered on the mi35x per-PR suite only. Recovers sgl-project#37563, which GitHub closed when its base branch xinyuan/glm-5.3-flash-support was merged into main, and supersedes the sgl-project#38542 replacement. The test-file placement under the registered-test taxonomy is Raiden-Makoto's from sgl-project#38542. Co-Authored-By: Raiden-Makoto <Raiden-Makoto@users.noreply.github.com>
This was referenced Sep 14, 2026
Jacob0226
marked this pull request as ready for review
September 14, 2026 07:56
Jacob0226
requested review from
DarkSharpness,
HaiShaw,
HydraQYH,
celve and
yuan-luo
as code owners
September 14, 2026 07:56
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
GLM-5.3-Flash cannot run on ROCm: the kernel that picks its sparse-attention candidates fails to build under hipcc.
Both come from the same file assuming CUDA. This is a hard stop rather than a slow path — GLM-5.3-Flash's pooled top-k has no other kernel to fall back to, so the indexer cannot run at all.
Scope
Two edits, both inside
USE_ROCMblocks that already exist: give the fp16 include a HIP arm, and add two names to the shim's alias list. Every CUDA line in the diff is unchanged context. No kernel body, launch configuration, tile size or numeric path moves; no environment variable, no architecture gate.The kernel has no direct test on any platform — CUDA only reaches it through a full GLM-5.3-Flash serving test, which JIT-compiles it every commit. This adds a focused one on the mi35x per-PR suite.
Test plan
Unit base
rocm/sgl-dev:v0.5.18-rocm720-mi35x-20260901, the image familypr-test-amd-rocm720.ymlresolves to, onxinyuan/glm-5.3-flash-support@515e865189with 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.group_topk128 and 512 alike, on'cuda_fp16.h' file not foundandno member named 'cudaFuncSetAttribute' in the global namespaceEnd to end, since no single PR here can serve the model on its own:
mainplus the eight open Day-0 PRs — #39338, #39339, #39340, #39341, #38545, #38546, #38547 and #39317 — which all merge cleanly onto it.rocm/sgl-dev:v0.5.19-rocm720-mi35x-20260909, MI355X TP4, decode CUDA graphs on, 1,319 GSM8K examples, one scoring pass each. Only the checkpoint differs between the two rows:zai-org/GLM-5.3-Flash, block FP8amd/GLM-5.3-Flash-Quark-MXFP4Accuracy: not attributable to this PR in isolation, and there is no standalone baseline — GLM-5.3-Flash does not start on ROCm without the rest of the stack.
CI States
Latest PR Test (Base): ❌ Run #34937436095
Latest PR Test (Extra): ❌ Run #34937435834
Latest PR Test (AMD ROCm 10): ➖ No AMD PR run found for this commit.