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

[Raiden-Makoto]

Summary

Replacement for #37563, which GitHub closed when its deleted support base was merged into main. The original author and review history remain credited there.

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

Combined main-based integration (RM/glm53-day0-main-integration tested at 9d0d62b744; current integration head 4e2f84ff3c; TP8/TP4 validation predates the above-cap guard, whose below-cap path is unchanged, pinned AITER d9e5ef7ce0, fresh caches): TP8 GSM8K 97.50% in 87.0s with 0.00% errors; TP4 GSM8K 96.97% in 320.2s with 0.00% errors. Decode CUDA graphs were enabled in both runs.

Recovery base: main @ 30e7a3072d.
Recovery validation on MI355X: scoped pre-commit passed; 4 fresh-cache HIP JIT tests passed.

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 #34768730151
Latest PR Test (Extra): ❌ Run #34768729972
Latest PR Test (AMD ROCm 10): ❌ Run #34768729836

[Raiden-Makoto]

Closing this replacement because @Jacob0226 will reopen and maintain the original PR, #37563. Please continue review and CI there.

[Jacob0226]

Thanks @Raiden-Makoto. Reopening as #39339 rather than #37563: #37563's base branch xinyuan/glm-5.3-flash-support has been deleted, so GitHub will not let it reopen. #39339 is off main and its changed files are byte-identical to this PR's head, including your registered-test placement.