feat(BACKEND-ROCM): kKdaGatedDeltaRule ROCm kernel — the per-K-channel-decay recurrence - #3120
Merged
Merged
Conversation
…hannel-decay recurrence The second of #2942's four unregistered ROCm ops. The CUDA donor (cuda_gdn.cu:3139 KdaScanKernel) is byte-for-byte the GDN scan kernel except the state decay is per-K-channel: GDN uses one scalar expf(g[t*hv_n+hv]); KDA stages a dk-vector expf(g[(t*hv_n+hv)*dk+i]) in shared memory. The ROCm GDN kernel (rocm_gdn_scan.hip, 181 lines) is a clean hand-translation of the CUDA GDN kernel; the KDA port is the same hand-translation, one step further. The op lands unreached: glm5_next_kda.cpp:322 refuses any non-CPU queue, so the model's only call site (:404) cannot run on ROCm. Lifting that refusal is W9c-2's job, and it depends on W9b (keep-quant residency). This is a staged slice: the commit body and PR body name what is unreached, the owning row, and the tracking issue. Issue: #2942, which STAYS OPEN. It tracks four unregistered ops; this lands the second. No closing keyword appears in this body, deliberately. Row: BACKEND-ROCM. Spec: .agents/specs/rocm-kda-gated-delta-rule.md. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:regolo-glm5.2 [maki]
…l-decay recurrence Hand-translates the CUDA KdaScanKernel (cuda_gdn.cu:3139) into HIP as src/vt/rocm/rocm_kda_scan.hip, mirroring the existing ROCm GDN kernel (rocm_gdn_scan.hip). The KDA scan is byte-for-byte GdnScanK except the state decay is per-K-channel: GDN uses one scalar expf(g[t*hv_n+hv]); KDA stages a dk-vector expf(g[(t*hv_n+hv)*dk+i]) in shared memory and applies s_row[ki] *= d_sh[ki]. Three differences from GdnScanK: shared memory is 3*dk (adds a per-K decay vector), decay is per-channel not scalar, and state math indexes d_sh[ki] instead of a scalar. State is float only (no f16/bf16 state dispatch, unlike GDN); the launcher passes nullptr for state_idx. Registers the op in rocm_ops.hip and adds the .hip TU to both VLLM_CPP_HIP source lists in CMakeLists.txt. The cross-device test case in test_backend_cross_device.cpp checks NMSE on output and state vs the CPU oracle, OpRegistered unconditionally on a ROCm build, and that GetReferenceTierHits does not increase. This is a staged slice: the op is unreached by the model's forward on its default configuration. glm5_next_kda.cpp:322 refuses any non-CPU queue, so the only call site (:404) cannot run on ROCm. Lifting that refusal is W9c-2's job (spec glm5-next-flash.md), depending on W9b (keep-quant residency). The cross-device test exercises the kernel directly through vt::KdaGatedDeltaRule, the production entry point for the op. CPU-only build and test pass (39 cases, 44 assertions, 0 failures). The ROCm device gate is PENDING: no HIP toolchain or AMD device in this session. The .hip TU is PENDING compilation for the same reason. Issue: #2942, which STAYS OPEN. It tracks four unregistered ops; this lands the second. No closing keyword appears in this body, deliberately. Row: BACKEND-ROCM. Spec: .agents/specs/rocm-kda-gated-delta-rule.md. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus [maki]
The spec's ## Now said "implementation pending." The implementation landed in 8d0186d, so that line is stale. This commit updates ## Now to reflect the committed state: CPU-only build and test pass, the ROCm .hip TU and the gfx1151 test case are pending a HIP toolchain and AMD device. Issue: #2942. Row: BACKEND-ROCM. Spec: .agents/specs/rocm-kda-gated-delta-rule.md. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus [maki]
…nd spec The KDA kernel header and spec cited the CPU oracle at cpu_ops.cpp:4046, but KdaGatedDeltaRuleKernel is at :2241 (recurrence math in KdaHeadTokenStep :2208). Line 4046 is QkvSplitKernel. Found by fresh review. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:regolo-glm5.2 [maki]
mudler
force-pushed
the
row/BACKEND-ROCM-KDA
branch
2 times, most recently
from
September 10, 2026 14:48
8b1e7ff to
c86332c
Compare
KDA dispatches Tin/Tout in {float, __hip_bfloat16} only and TState is
always float, so the __half overloads carried over from the GDN template
are dead code. With -Werror they trip -Wunused-function on gfx1151.
Drop them and keep the __hip_bfloat16 overloads that the bf16 dispatch
paths need.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:regolo-glm5.2 [maki]
mudler
force-pushed
the
row/BACKEND-ROCM-KDA
branch
from
September 10, 2026 15:08
ae8242c to
70cd2bc
Compare
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
Ports
kKdaGatedDeltaRuleto AMD ROCm/HIP. This is the second of fourops required for GLM-5.3-Flash on ROCm (issue #2942).
The kernel is a hand-translation of the CUDA
KdaScanKernel(
cuda_gdn.cu:3139-3235) to HIP, readable side by side against thetemplate
rocm_gdn_scan.hip(GdnScanK). Three differences from GDN:3*dknot2*dk(extrad_shbuffer forper-K-channel decay).
g[(t*hv_n+hv)*dk+i]instead ofg[t*hv_n+hv].State dtype is float only (no f16/bf16 state dispatch unlike GDN). The
launcher passes
nullptrforstate_idx.Staged slice
The op lands unreached by the model's forward.
glm5_next_kda.cpp:322refuses any non-CPU queue, so the only call site (
:404) cannot run onROCm. Lifting that refusal is W9c-2's job, which depends on W9b
(keep-quant residency). The spec lists this wiring under
## Owed.Issue #2942 stays open — this lands op 2 of 4.
Gate status
test_backend_cross_deviceKDA case, 39/39assertions green. The test reaches the op through the production ABI
(
vt::KdaGatedDeltaRule->GetOpregistry), not hand-construction.check-commit-style.py).check-commit-trailers.py --filled).hipcctoolchain on this host. The.hipTU is not compiled and the GPU kernel has not run. The CPU-onlytest compares CPU-vs-CPU (trivially exact); the meaningful
ROCm-vs-CPU-oracle verification is owed.
Files
src/vt/rocm/rocm_kda_scan.hip(NEW, 155 lines) — KDA scan kernel,launcher, entry point
src/vt/rocm/rocm_ops.hip(+10 lines) —RegisterOpforkKdaGatedDeltaRuleonkROCMCMakeLists.txt(+2 lines) — addedrocm_kda_scan.hipto both HIPsource lists
tests/vt/test_backend_cross_device.cpp(+94 lines) — cross-devicetest case mirroring the GDN test
.agents/specs/rocm-kda-gated-delta-rule.md(NEW) — specFresh review
A fresh reviewer (separate agent) reviewed the implementation statically
and performed mutation analysis. Verdict: NEEDS_REPAIR -> repaired (one
wrong CPU oracle citation
cpu_ops.cpp:4046->:2241). The kernel,registration, CMake, test, and commit trailers were confirmed correct.
All five mutations (decay indexing, shared memory, state math, deleted
registration, deleted test call) would be caught by the test or build on
a ROCm build.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:regolo-glm5.2 [maki]