feat(rocke): quad_perm crosslane intrinsic + kernarg-packer precompile by AviralGoelAMD · Pull Request #12070 · ROCm/rocm-libraries

AviralGoelAMD · 2026-09-14T20:49:36Z

ISSUE ID : AICK-1513

What this is

PR 1 of a 4-PR stack that splits the GDN gfx950 work (originally one 51-file PR) into reviewable layers. This is the foundation layer: it adds a new crosslane hardware primitive and a launch-path speedup, with no GDN code and no dependency on the other three PRs.

Stack (bottom → top): quad_perm (this) → dispatch.core hoist → GDN decode → GDN prefill.

What changed

1. quad_perm — an intra-quad DPP permute intrinsic, added to both engines.
quad_perm(data, [p0,p1,p2,p3]) lets lane 4q+i read data from lane 4q+perm[i] on the VALU (via v_mov_b32_dpp / llvm.amdgcn.update.dpp.i32), with no LDS crossbar and no lgkmcnt wait. Also adds warp_shuffle_xor_quad, a fast path for the two xor-masks that stay inside a quad: masks 1 and 2 lower to quad_perm, and any other mask is rejected — callers needing a wider mask use warp_shuffle_xor, which goes through ds_swizzle.

Python engine: core/ir.py (builder + validation), core/lower_hip.py, core/lower_llvm.py.
C++ engine (live): core/ir/ir_flow.cpp, core/ir/core_types.cpp (opcode + purity), core/lower_hip/lower_hip_mma.cpp, core/lower_llvm/crosslane.cpp, include/rocke/ir.h.
Coverage: Python tests/test_rocke.py, C++ tests/core/future_intrinsic_lowering.cpp, and cross-engine parity tests/instances/parity/target_intrinsics_emit.{c,py} (identical IR from both lowerers).

2. Kernarg-packer precompile — a launch-path speedup for every kernel.
runtime/packing.py gains compile_packer(signature); runtime/launcher.py compiles the packer once at module-load and calls it per launch instead of re-deriving the layout each time. Byte-identical to pack_args.

What it removes is kernarg layout reconstruction — the offset/alignment walk, the per-argument type dispatch, and the format-string assembly. It is not a saving on format compilation: CPython's struct module already caches recently used format strings, so re-packing the same format is a cache lookup, not a recompile.

Review round 1 — what changed since the first push

Four items were raised by @yraparti and @tenpercent. All four are addressed.

1. Lowerers accepted out-of-range ctrl (bffb5fd911f). All four quad_perm lowering sites masked the control word with 0xFF instead of validating it. ctrl packs four two-bit lane selectors (p0 | p1<<2 | p2<<4 | p3<<6), so 0..255 is the whole legal range — and masking silently rewrote malformed IR into a different, valid permute: 256 became 0 ([0,0,0,0], a lane-0 broadcast) and -1 became 255 ([3,3,3,3]). A wrong reduction then computed wrong numbers instead of failing.

The builders already validate selectors, but IR reaching a lowerer by another route (deserialized, rewritten by a pass, hand-built) skips them. Now rejected in the Python lowerers (ValueError) and the C++ ones (ROCKE_ERR_VALUE), with the mask dropped so the check cannot be bypassed. Tests on both sides are mutation-verified: with the mask restored, the Python subtests and the four C++ assertions fail.

2. Wave-size semantics undocumented (cfeec0a1315, 8cea213942f). quad_perm makes no wave-size assumption, and both engines now say so: the control word applies within every four-lane group, and four divides both 32 and 64, so a lane never addresses outside its own quad. Wave size changes only the number of quads (8 in wave32, 16 in wave64).

Deliberately scoped — wave-size-independent is not architecture-independent. The op still needs DPP-capable hardware; the useful point is that base-DPP quad_perm is available on CDNA where dpp_xor's RDNA-only row_xmask is not. Also recorded: the op has no lane targeting (the control is broadcast to every quad, row/bank masks fixed at 15, 15), so selecting a subset of quads is the caller's job.

A first draft of this docstring cited dpp_xor as an op whose partner lane can leave the wave; 8cea213942f corrects that. dpp_xor caps xor_mask at 1..15 and its partner stays inside a 16-lane row, which also divides both wave sizes. The accurate counterexample is warp_shuffle_xor at lane_xor = 32.

3. "The dominant Python cost" was unbacked (786f6c827b0). Correct: a packing-only microbenchmark cannot establish a share of launcher overhead. Both sites now describe the mechanism instead of a magnitude -- the packer precomputes the fixed argument layout once per launcher, so a launch does not rebuild the offset table, re-dispatch on argument types, or re-assemble the format string. struct already caches recently used format strings, so the saving is that surrounding work rather than the format compile itself.

The profiler written to measure the share now lives in its own PR (#12168, draft), and performance claims are deferred until its gate is fixed: its chunk-size-sensitivity check is a ratio, so a chunk size large enough to saturate both comparison windows accepts a back-pressured run.

4. PR description misdescribed warp_shuffle_xor_quad. It said the helper "leaves wider masks on ds_swizzle", implying a dispatcher. It is a specialist: masks other than 1 and 2 raise. Corrected in the What-changed section above.

Scope note (disclosed)

Item 2 is functionally independent of item 1. It is kept here rather than split for one concrete reason: its byte-identity test (test_compile_packer_matches_pack_args_byte_for_byte) lives in tests/test_rocke.py, the same file that holds the quad_perm tests — splitting would put two PRs on one test file. Reviewers should weigh the packer's global blast radius (it changes the launch path for every kernel), not only its decode benefit.

Why it's safe

No kernel IR changes — this only adds an opcode and a runtime fast path.
quad_perm is validated at the builder (selectors ∈ 0..3, i32-only) and now at every lowerer (ctrl ∈ 0..255), and is marked pure in both engines.
llvm.amdgcn.update.dpp.i32 is an already-shipped overload (used by mov_dpp8/ds_swizzle), not a new intrinsic surface.
The packer is asserted byte-identical to the existing pack_args.

Why it's first in the stack

GDN decode's in-quad xor-butterfly reduction lowers to quad_perm (3 call sites), so this must land before the decode PR. It touches only platform/, so it reviews without any GDN context.

Verification (run, not planned)

Both engines built and exercised locally on this diff:

Python, with the C++ extension importable: pytest tests/test_rocke.py -k "quad_perm or warp_shuffle_xor_quad or compile_packer" → 8 passed, 0 skipped (the cross-engine assertions no longer skip). Full file: 291 passed, 14 skipped, 43 subtests.
C++ engine: rocke_future_intrinsic_lowering → 31 case(s) OK.
Cross-engine byte-identity: tools/check_byte_identity.py --only target_intrinsics → GREEN, configs=8, bad=0 — Python and C++ emit identical .ll.
test_rocke_ci_static.py 5 passed.
Formatters: black clean; clang-format applied (local binary is v20; repo pre-commit pins v18.1.4, so CI may adjust whitespace).

GPU numerical execution of quad_perm itself and the full all-family byte-identity gate were not run here.

Stacking / base

#11807 has merged, and this PR is now based directly on develop.

Specifies the Gated DeltaNet operator, its relationship to KDA, and the two kernel pathways: the single-token decode kernel and the chunkwise prefill forward that ships as a mode of the KDA chunkwise kernel. Covers the gated delta rule, why a scalar forget gate makes GDN a special case of KDA's per-channel gate, the parallel decomposition and dataflow of each kernel, why decode cannot be prefill at chunk size 1, and the asymmetry that makes the reverse reuse a kernel change rather than a configuration.

Correct claims a code-vs-doc audit flagged: scope the decode state-pool validation to the driver/prepare() path (dispatch selects specs, not tensors); fix the tile-selection trend to match the non-monotone banded table; correct the KDA-on-GDN reuse attribution (raw-prep and the value_splits knob are pre-existing KDA; the value_splits selection table is GDN-only); reconcile the state-scan EV x DK notation with the DV x DK state; correct the triangular-solve serial-work scaling to linear in solve_block; and tighten the softplus-overflow, fusion-flag, l2norm-flag, idle-wave-Kt, and tensor-contract wordings.

… the KDA algorithm Section 5 had silently adopted kda/ALGORITHM.md's transposed state convention while sections 0 and 4 kept the DV x DK orientation that both kernels actually allocate ([pool, HV, DV, DK] for decode, [BH, DV, DK] for the prefill scan). The scan equations therefore did not type-check, and the state update had Diag(dec) on the left, where a DK x DK diagonal cannot multiply an EV x DK state. Keep the DV x DK orientation everywhere, drop the stray transposes on S, and move Diag(dec) to the right. State the S_kda = S^T mapping once so the two docs can be read together. Also replace the near-verbatim restatement of the six-tile factorization and the midpoint-factored decay with a link to kda/ALGORITHM.md, keeping only what GDN changes: the scalar gate broadcasts across DK, so Gamma is channel-constant and dec is a DK vector with equal entries. Renumber the rest of section 5 and its cross-references and the contents block. Addresses review comments r3975876924, r3975879718, r3975887752.

…-domain mapping Section 1 named no target model, so every constant in the doc -- the tile bands, the value_splits bands, the (Hv, Hk) shapes in the validation section -- had no provenance. Add 1.4 with the published config.json geometry for the two deployments this family serves, following the builders/mla/DESIGN.md precedent: a per-model row plus a named out-of-scope entry with the config keys that exclude it. The table pins three things the rest of the doc assumed silently: DK == DV on every supported row, kv_group = 2 as the only shipping GDN grouping (kv_group = 4 in section 7 is a stress point), and where each model lands in the tuned tables -- at Hv = 32 the prefill value_splits bands put batch 1-2 on 8 splits and batch 5+ unsplit, with batches 2 and 4 exactly on band edges. Also correct the log-domain mapping added in the previous commit: kda/ALGORITHM.md reconstructs the decay ratio with exp(Gc - Gref), so its Gc is the natural log of Gamma, not log2. The base-2 exponent domain there is the hardware exp2 detail (the sum is pre-scaled by log2(e)), not the symbol's definition. Addresses review comment r3975904504.

Section 1.2 gives the gate as a decay multiplier in (0, 1); section 2.1 gives KDA's bound as (lower_bound, 0) without noting that the domain changed to log space in between, so the two ranges read as a contradiction. Name the domain and give the bridge (decay = exp(gate)). Also correct 'unbounded' to 'unbounded below': GDN's gate is -exp(A_log)*softplus(.), and softplus is positive, so the value is always <= 0 -- which is what makes section 8's decay-range limit a one-sided contract. The raw per-token inputs g and a are called out as distinct from the gate value so the sentence carries no second meaning for g. Addresses review comment r3975918258.

Every arch-scoped family directory under builders/gfx950 documents itself with an ALGORITHM.md + README.md pair (attention, kda; kda adds CASE_STUDY.md). The only DESIGN.md precedent, builders/mla, is an arch-less design spike with no kernel code, so GDN_DESIGN.md took its name from the one directory whose shape does not match this one. Rename to ALGORITHM.md and state the hierarchy in the header: this file is the family entry point, README.md is host-side tooling. The GDN decode walkthrough that the kernel PR adds under the same name is folded into this file there, so the directory ends with the standard pair rather than three overlapping documents. Addresses review comment r3975928486.

The blocked-solve entry said the scalar substitution shrinks as the square of the block size. That is the per-block cost; shrinking the block also multiplies the block count, so the total is ~C*b/2 and is linear in solve_block. Halving the block halves the scalar work rather than quartering it, which is how a reader sizes the knob. Markdown only, so this PR stays docs-only. Addresses review comment r3975891138.

Four review fixes, all in gdn/ALGORITHM.md: - §7 now names the prefill head shapes it validates -- (4,4), (8,4) and (32,8) -- so §1.4's pointer at the kv_group=4 stress point resolves. Shapes match the parametrize list in test_kda_gdn_gfx950_numeric.py. - §5.5 uses EV for DV/value_splits, the name already defined in §0 and used in §5.3. ev_slice was prose-only and is gone. - §4.6 prints the batch-banded tile table instead of describing it. The non-monotone num_warps claim (4 -> 2 -> 1 -> 8) is now checkable from the rows. Values are _TUNED_TILES in dispatch/gdn/gfx950.py. The band edges are also marked as interpolated between the four measured batch anchors, since the table now prints them as concrete numbers. - §5.2 cites the code site enforcing the wave-0-only solve (the tid < 64 scf_if and _emit_idle_during_solve), and notes that §8 follow-up 4 would have to change that constraint.

…ker precompile Adds the DPP quad_perm intrinsic across both engines (C++ ir_flow/lowering and the Python IR/lowerers) with parity and golden coverage, plus warp_shuffle_xor_quad. Also precompiles the hot-path kernarg packer in the launcher (byte-identical to pack_args) -- a launch-cost win for all kernels, notably tiny decode launches. No kernel IR changes; foundation for GDN decode.

tenpercent

Reviewed the PR diff at 188fbddf80d570aecd11149ce6f0a89c6f9c386f against stacked base 5bebfc552291515ff7d51b31bcec290edf24c612, including the existing discussion. No blocking defect found; these are non-blocking suggestions.

Validate raw control values in the lowerers (existing thread). Both builders validate selectors, and every value in 0..255 is a legal quad-permute control, including repeated selectors for broadcasts. The remaining concern is malformed IR bypassing the builders: I reproduced ctrl=256 becoming 0 and ctrl=-1 becoming 255. Reject controls outside 0..255 in the Python and C++ HIP/LLVM lowerers, with matching rejection tests. This is defensive hardening; valid builder-generated IR is unaffected.
Document wave-size semantics (existing thread). State that the mapping applies independently within each four-lane group in both wave32 and wave64 on supported DPP targets. Avoid an unrestricted architecture-independent claim. Generated HIP and LLVM probes compiled and linked for gfx1151 and gfx950, and disassembly contained the intended quad-permute instructions; this is code-generation evidence, not GPU numerical validation.
Qualify the packing-cost claim (existing thread). Please substantiate or soften “the dominant Python cost” in launcher.py and packing.py. A packing-only microbenchmark does not establish the fraction of complete launcher overhead. Python's struct module already caches recently used formats, so the concrete explanation should emphasize hoisting layout reconstruction, type dispatch, and format-string construction. For example: “Precompute the fixed argument layout once per launcher to avoid rebuilding it on each call.”
Correct the PR description's wider-mask behavior. warp_shuffle_xor_quad rejects masks other than 1 and 2; it does not automatically route wider masks to ds_swizzle. Callers must explicitly use warp_shuffle_xor for those masks.

Validation: focused new Python tests had 5 passes and 2 skips because the Python C++ extension was unavailable; all 10 packing ABI tests passed; 1,365 additional signature combinations packed identically. A fresh C++ build and rocke_future_intrinsic_lowering passed. All eight target-intrinsics configurations were byte-identical between Python and C++ for LLVM flavors 20, 22, and 23. HIP/LLVM compile-and-link probes used ROCm 7.1. GPU numerical execution and the full all-family gate were not run.

— codex / astra / medium

@yraparti

ISSUE ID : AICK-1513 The four quad_perm lowering sites masked the DPP control word with 0xFF instead of validating it. ctrl packs four two-bit lane selectors (p0 | p1<<2 | p2<<4 | p3<<6), so 0..255 is the whole legal range and masking silently rewrote malformed IR into a different, valid permute: 256 became 0 ([0,0,0,0], a lane-0 broadcast) and -1 became 255 ([3,3,3,3]). A wrong reduction then computed wrong numbers instead of failing. The builders already validate selectors, but IR reaching a lowerer by another route (deserialized, rewritten by a pass, hand-built) skips them; the lowerer is the last gate before text reaches the compiler. Reject in the Python LLVM/HIP lowerers (ValueError) and the C++ ones (ROCKE_ERR_VALUE), and drop the mask so the check cannot be bypassed. Tests added on both sides are mutation-verified: with the mask restored the Python subtests and the four C++ assertions fail. Raised in review by @yraparti and confirmed by @tenpercent.

@yraparti

…ines ISSUE ID : AICK-1513 quad_perm's docstring said what lane 4q+i reads without saying whether the op assumes a wavefront size. It does not: the control word applies within every four-lane group, and four divides both 32 and 64, so a lane never addresses outside its own quad. Wave size changes only the number of quads (8 in wave32, 16 in wave64), not the permutation any quad performs -- unlike dpp_xor or ds_swizzle_xor, whose partner lane can leave the wave for a large enough mask. Scoped deliberately: wave-size-independent is not architecture- independent. The op still needs DPP-capable hardware; the point is that base-DPP quad_perm is available on CDNA where dpp_xor's RDNA-only row_xmask is not. Also records that the op has no lane targeting -- the control word is broadcast to every quad, with row/bank masks fixed at 15, 15 -- so a caller wanting a subset of quads must select afterwards. Raised in review by @yraparti; scope wording per @tenpercent.

ISSUE ID : AICK-1513 dpp_xor was cited as an op whose partner lane can leave the wave. It cannot: it validates xor_mask to 1..15 and its partner stays inside a 16-lane DPP row, and 16 divides both 32 and 64 -- so dpp_xor is as wave-safe as quad_perm and the contrast contradicted its own contract. The accurate counterexample is warp_shuffle_xor at lane_xor = 32, which addresses a lane that exists in wave64 and does not exist in wave32. The separate, correct dpp_xor mention (row_xmask is RDNA-only, base-DPP quad_perm is not) is unchanged.

@yraparti

ISSUE ID : AICK-1513 The launcher comment and the compile_packer docstring called per-call layout derivation "the dominant Python cost" on tiny decode launches. A packing-only microbenchmark cannot establish a share of total launcher overhead, so the claim was not backed by what was measured. Both sites now state the mechanism instead: the fixed argument layout is precomputed once per launcher, so a launch does not rebuild the offset table, re-dispatch on argument types, or re-assemble the format string. Also records that struct already caches recently used formats -- the saving is the layout reconstruction around that compile, not the compile itself, which the original wording implied. Byte-identity with pack_args is unchanged and still asserted by test_compile_packer_matches_pack_args_byte_for_byte. Raised in review by @yraparti and @tenpercent.

AviralGoelAMD · 2026-09-15T03:42:51Z

Reviewed the PR diff at 188fbddf80d570aecd11149ce6f0a89c6f9c386f against stacked base 5bebfc552291515ff7d51b31bcec290edf24c612, including the existing discussion. No blocking defect found; these are non-blocking suggestions.

Validate raw control values in the lowerers (existing thread). Both builders validate selectors, and every value in 0..255 is a legal quad-permute control, including repeated selectors for broadcasts. The remaining concern is malformed IR bypassing the builders: I reproduced ctrl=256 becoming 0 and ctrl=-1 becoming 255. Reject controls outside 0..255 in the Python and C++ HIP/LLVM lowerers, with matching rejection tests. This is defensive hardening; valid builder-generated IR is unaffected.

Document wave-size semantics (existing thread). State that the mapping applies independently within each four-lane group in both wave32 and wave64 on supported DPP targets. Avoid an unrestricted architecture-independent claim. Generated HIP and LLVM probes compiled and linked for gfx1151 and gfx950, and disassembly contained the intended quad-permute instructions; this is code-generation evidence, not GPU numerical validation.

Qualify the packing-cost claim (existing thread). Please substantiate or soften “the dominant Python cost” in launcher.py and packing.py. A packing-only microbenchmark does not establish the fraction of complete launcher overhead. Python's struct module already caches recently used formats, so the concrete explanation should emphasize hoisting layout reconstruction, type dispatch, and format-string construction. For example: “Precompute the fixed argument layout once per launcher to avoid rebuilding it on each call.”

Correct the PR description's wider-mask behavior. warp_shuffle_xor_quad rejects masks other than 1 and 2; it does not automatically route wider masks to ds_swizzle. Callers must explicitly use warp_shuffle_xor for those masks.

Validation: focused new Python tests had 5 passes and 2 skips because the Python C++ extension was unavailable; all 10 packing ABI tests passed; 1,365 additional signature combinations packed identically. A fresh C++ build and rocke_future_intrinsic_lowering passed. All eight target-intrinsics configurations were byte-identical between Python and C++ for LLVM flavors 20, 22, and 23. HIP/LLVM compile-and-link probes used ROCm 7.1. GPU numerical execution and the full all-family gate were not run.

— codex / astra / medium

addressed - added a check
addressed - edited the comments as suggested
addressed - data in comments and reproducible script added.
addressed - corrected the PR description

therock-pr-bot · 2026-09-15T04:02:50Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

…obust ISSUE ID : AICK-1513 A silicon re-run surfaced two measurement defects, both in checks added by the previous commit. 1. Both gate statistics were a single-sample ratio of two one-shot timings, and the tolerance sits inside a shared node's jitter envelope -- so one scheduler stall landing in the numerator and not the denominator fires the gate. Repeating the statistic 15 times on one node showed it firing on a denominator that had passed in the real run, and not firing on the one that had failed. Since the gate INVALIDATES a run, a false positive discards a good measurement. Both checks now take the minimum over repetitions, for the same reason micro() already does: a minimum times the operation rather than the operation plus whatever else the machine did. Back-pressure is systematic and survives the minimum; a stall does not. in_chunk_stepup also takes the drain callback, so its repetitions each start from an empty queue. 2. The remainder-residual check compared against the run's own noise floor alone, which made it stricter the better the measurement got: on a very quiet run a residual of 0.1% of the remainder read as a 4x violation. The tolerance is now max(noise_floor, 0.5% of remainder), and the verdict is computed in the helper as remainder_within_tol rather than left for a reader to apply. Passing no noise floor reports the residual and judges nothing, so a caller cannot read a verdict that was never computed. Both are mutation-verified: restoring the single-sample statistic makes the stall test fail, and restoring the floor-only tolerance fails the quiet-run test.

…pose ISSUE ID : AICK-1513 801 -> 669 lines. The script answers one question -- what fraction of the Python launch path is kernarg packing -- and three pieces were not needed to answer it. Dropped the cProfile view. It charges per Python frame while the single largest real cost in this path is one ctypes FFI call it barely charges for, so it only ever produced an upper bound that then had to be explained away. The wall-clock A/B and the component isolation answer the question without it, and they agree with each other. Dropped the signature-substitution model. substitute_share carried a swap/pack_target path for re-expressing the share against a signature that was never launched, and the shipped script never used it -- its only call passed no swap. What remains is packing_share: each arm's share from its own measured total, plus the remainder-residual consistency check. Same numbers, no unused mode. Dropped sig_shape and stepup_check. sig_shape existed for the substitution; stepup_check compared chunk means, which cannot see a queue that saturates identically inside every drained chunk -- the blind spot in_chunk_stepup and chunk_size_sensitivity were added to cover. Keeping a check that cannot see the failure it names is worse than not having it. The gate itself is untouched: segment_costs, in_chunk_stepup, chunk_size_sensitivity, chunked and micro are byte-identical, so the silicon run validating them still applies. Tests follow the API: 7 pass, full perf suite 157.

tenpercent · 2026-09-16T02:15:48Z

Would you consider moving the launch-overhead profiler and its tests into a separate PR? The intrinsic and packer changes look ready from this review, while the profiler still has a measurement-validity issue: with a sufficiently large --chunk-size, both comparison windows can be saturated, their cost ratio approaches 1, and the gate accepts the run. A separate PR would let us fix and test that without holding up the runtime changes.

I'd keep quad_perm, the precompiled packer, and their correctness/parity tests here. The description can explain the packer's mechanism and link the profiler follow-up for measurement evidence, with performance claims deferred until the gate is resolved.

The GDN algorithm document would also fit better in the GDN implementation PR: it currently describes behavior as shipped and references files/APIs that are absent from this tree.

ISSUE ID : AICK-1513 Requested in review. The profiler has an open measurement-validity issue -- its chunk-size-sensitivity check is a ratio, so a chunk size large enough to saturate BOTH comparison windows returns a ratio near 1.0 and the gate accepts a fully back-pressured run -- and fixing that should not hold up the intrinsic and the packer, which are independent of it. The quad_perm intrinsic, the precompiled kernarg packer, and their correctness and cross-engine parity tests stay here. The packer is described by its mechanism: it precomputes the fixed argument layout once per launcher, so a launch does not rebuild the offset table, re-dispatch on argument types, or re-assemble the format string. Byte-identity with pack_args is asserted by test_compile_packer_matches_pack_args_byte_for_byte, which is unaffected by this move.

AviralGoelAMD · 2026-09-16T02:26:10Z

Would you consider moving the launch-overhead profiler and its tests into a separate PR? The intrinsic and packer changes look ready from this review, while the profiler still has a measurement-validity issue: with a sufficiently large --chunk-size, both comparison windows can be saturated, their cost ratio approaches 1, and the gate accepts the run. A separate PR would let us fix and test that without holding up the runtime changes.

I'd keep quad_perm, the precompiled packer, and their correctness/parity tests here. The description can explain the packer's mechanism and link the profiler follow-up for measurement evidence, with performance claims deferred until the gate is resolved.

The GDN algorithm document would also fit better in the GDN implementation PR: it currently describes behavior as shipped and references files/APIs that are absent from this tree.

Yes, I extracted the profiler script into its own PR. (I may take a while to actually make it perfect). We can now just focus on the core changes of the PR. The docs are not part of the PR, its just some GitHub glitch that they show up here.

tenpercent

thanks!

Copilot

🟡 Changes recommended

Unresolved lowering validation issues and missing GPU numeric coverage remain, along with documentation corrections.

Get a fresh assessment by requesting another Copilot review.

Pull request overview

Adds the quad_perm crosslane intrinsic across Python and C++ engines, plus reusable kernarg-packer compilation for launcher performance.

Changes:

Adds IR builders, lowerers, validation, purity metadata, and parity tests for quad_perm.
Adds compile_packer() and integrates it into KernelLauncher.
Updates GDN and KDA documentation.

File summaries

File	Summary
`dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py`	Intrinsic and packer tests
`dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/target_intrinsics_emit.py`	Python parity fixture
`dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/parity/target_intrinsics_emit.c`	C parity fixture
`dnn-providers/hip-kernel-provider/rocke/platform/tests/core/future_intrinsic_lowering.cpp`	C++ lowering tests
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/packing.py`	Precompiled packer implementation
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/launcher.py`	Launcher integration
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py`	Python LLVM lowering
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py`	Python HIP lowering
`dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py`	Python IR builders and validation
`dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/ir.h`	C++ API declarations
`dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/crosslane.cpp`	C++ LLVM lowering
`dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp`	C++ HIP lowering
`dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/ir_flow.cpp`	C++ IR builder
`dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/ir/core_types.cpp`	Opcode and purity metadata
`dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/kda/CASE_STUDY.md`	KDA documentation
`dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/gdn/ALGORITHM.md`	GDN algorithm documentation

Review details

Suppressed comments (9)

dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/gdn/ALGORITHM.md:205

This present-tense implementation claim is false in this foundation PR: library/kernels/gfx950/kda_chunkwise.py has no gate_kind field, and the tree has no GDN implementation or dispatcher. Since the PR explicitly adds no GDN code, defer this implementation-specific text to the GDN code PR or rewrite it as planned design rather than advertising an API and cache-key contract that does not exist yet.

This is why GDN prefill is implemented as `gate_kind="gdn"` inside `kda_chunkwise.py` rather than
as a second engine. The direction matters and is not symmetric — see §6.

dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/gdn/ALGORITHM.md:396

The cross-reference points to ../kda/ALGORITHM.md, but that file does not exist in this tree; the KDA directory currently contains CASE_STUDY.md instead. Point this section at an existing document or add the referenced design document so the core prefill explanation is navigable.

GDN prefill is the KDA chunkwise kernel in `gate_kind="gdn"` mode, so the factorization is KDA's,
unchanged: see `../kda/ALGORITHM.md` for the six state-independent tiles (`A`, `GK`, `GQ`, `Aqk`,
`Kt`, `dec`), the chunk-parallel / state-serial split, and the midpoint-factored decay that keeps

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mma.cpp:293

The C++ HIP lowerer hard-codes __builtin_amdgcn_update_dpp for i32 but does not validate the operand type. Raw IR can bypass the builder and pass f32 data, producing uncompilable HIP instead of ROCKE_ERR_VALUE; add the same i32 check used by the builder contract.

    const rocke_value_t* data = op->operands[0];
    const rocke_value_t* r = h_res(op);

dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/crosslane.cpp:330

The C++ LLVM lowerer hard-codes the i32 intrinsic but does not check the operand type. A raw ROCKE_OP_TILE_QUAD_PERM can bypass rocke_b_quad_perm and reach this handler with f32 data, yielding an invalid i32 call rather than ROCKE_ERR_VALUE; validate the operand before emitting.

    const rocke_value_t* data = op->operands[0];

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:1940

This HIP lowerer always emits the i32 __builtin_amdgcn_update_dpp overload but does not validate data.type; malformed IR can reach this path without the builder check and produce invalid HIP source. Reject non-i32 data here as well, matching the intrinsic contract.

        (data,) = op.operands

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:1943

int(...) silently converts malformed immediate attributes such as 1.9 or '1' into a valid control word. This HIP lowerer must reject non-integer ctrl values before range checking, otherwise raw IR can silently select a different lane permutation.

        ctrl = int(op.attrs["ctrl"])
        if not 0 <= ctrl <= 255:
            raise ValueError(f"tile.quad_perm: ctrl must be in 0..255, got {ctrl}")

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py:3800

This lowerer always emits the ...update.dpp.i32 overload but does not validate data.type; malformed IR can therefore produce an i32 intrinsic call with an f32 or other operand instead of failing. The builder check is bypassable for deserialized or hand-built IR, so reject non-i32 data at this lowering boundary too.

        (data,) = op.operands

dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py:3804

int(...) silently converts malformed immediate attributes such as 1.9 or '1' into a valid control word. Since this lowerer is the validation boundary for raw/deserialized IR, check that ctrl is an actual integer before the range check; otherwise malformed IR selects a different lane permutation without an error.

        ctrl = int(op.attrs["ctrl"])
        if not 0 <= ctrl <= 255:
            raise ValueError(f"tile.quad_perm: ctrl must be in 0..255, got {ctrl}")

dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:3272

These tests only compare emitted strings. Because the Python/C++ parity fixture asserts the same control word, both lowerers could share an incorrect hardware lane mapping and still pass; the PR also reports no GPU execution of quad_perm. Add a GPU-gated numeric test that feeds distinct per-lane values through the intrinsic and verifies all four outputs in each quad.

    def test_quad_perm_encodes_lane_selection(self):
        # [1,0,3,2] -> 1 | (0 << 2) | (3 << 4) | (2 << 6) == 177.
        ll = self._lower("qperm", lambda b: b.quad_perm(b.const_i32(1), [1, 0, 3, 2]))
        self.assertIn("declare i32 @llvm.amdgcn.update.dpp.i32(", ll)
        self.assertIn(
            "call i32 @llvm.amdgcn.update.dpp.i32("
            "i32 1, i32 1, i32 177, i32 15, i32 15, i1 true)",
            ll,
        )

Files reviewed: 16/16 changed files
Comments generated: 2
Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

assistant-librarian · 2026-09-17T19:13:53Z

TheRock Submodule Bump Activity

Newest first

2026-09-17 — Included in TheRock via ROCm/TheRock#8289.

…lder (#12090) ISSUE ID : AICK-1513 Two small functions were copy-pasted into every operator family. This moves them into one shared place. No behaviour changes, no kernel code. ## The two functions **The pin check.** A caller can force a specific kernel instead of letting rocKE route (`algorithm="chunk_scan"`). Something has to compare that request against each candidate and say yes/no with a reason. `"auto"`, the default, matches anything. **The id builder.** Once a kernel is picked it needs an identity — operator, family, candidate, algorithm, arch, ABI, plus hashes of the request and the spec. Logs, tuning records and benchmark rows all refer to a pick by that id. `attention` and `kda` each had their own copy of both, character for character. GDN decode and prefill (PRs 3 and 4 of this stack) would have made copies three and four. ## What changed `rocke/dispatch/core.py` gains `selector_matches()` and `make_kernel_id(op=...)`. Both families now call them and keep their old private names as aliases: ```python _selector_matches = selector_matches # both common.py files def _kernel_id(req, candidate, spec): # both __init__.py files return make_kernel_id(req, candidate, spec, op="attention") ``` No call site outside these four files changes. ## The one thing worth checking closely The old builders set `family=_FAMILY` (each family's own constant). The shared one reads `family=candidate.family`. Those are different expressions. `family` is part of `selection_key`, the identity tuning records and benchmark rows index by — so if they ever disagreed for some candidate, that candidate's identity would shift silently and old records would stop matching. (It is *not* part of `compile_key`, which is `arch:abi_version:spec_hash`, so nothing recompiles and no wrong binary can be dispatched.) They agree today: every candidate in both registries is registered with its own family constant. The new test asserts that over the real registries, so it also holds for families added later: ```python for candidates, family in ((attention_candidates(), ATTENTION_FAMILY), (kda_candidates(), KDA_FAMILY)): assert candidates, "registry is empty -- the check would pass vacuously" for candidate in candidates: assert candidate.family == family ``` ## Tests `library/tests/dispatch/test_core_helpers.py` is new and tests the helpers directly — the family suites only reached them indirectly, so a bug in one used to surface as a confusing failure somewhere downstream. It covers pin matching (`auto`, exact pins, both rejection reasons, case/whitespace tolerance), id determinism and spec sensitivity, and the invariant above. ## Verification - `library/tests/dispatch` → 334 passed, 216 subtests. - Mutation check: forcing `selector_matches` to reject everything turns 46 kda dispatch tests red, with the injected reason appearing in the real "no candidate supports request" error. That proves the families route through the shared function rather than a surviving private copy. Reverted, green again. - `black` clean. One `typing` import in `kda/common.py`, left unused by the change, removed. ## Context PR 2 of a 4-PR stack splitting the GDN gfx950 work (originally one 51-file PR): quad_perm (#12070) → **this** → GDN decode → GDN prefill. It does not depend on #12070 — that one touches crosslane lowering in `platform/`, this one touches dispatch. Both sit on `develop`. #11795's version of this hoist covered `kda` only; `attention` is included here because it held an identical copy, and leaving it would have undercut the point of the change. Self-reviewed on the fork first: AviralGoelAMD/rocm-libraries#22.

ISSUE ID : AICK-2229 Adds **GDN (Gated DeltaNet) chunkwise prefill** for gfx950 — the many-tokens-at-once path that runs when a sequence first arrives — plus dispatch wiring, tests and docs. 19 files, +2214. ## The one design decision **GDN prefill is not a new kernel.** It is the shared KDA chunkwise pair run in a new gate mode. The chunkwise algorithm is identical for both families — chunk factorization, triangular solve, state scan. Only the decay gate differs: | | KDA | GDN | |---|---|---| | gate | `lower_bound * sigmoid(exp(A_log) * (g + dt_bias))` | `-exp(A_log) * softplus(a + dt_bias)` | | shape | per channel | scalar per `(token, head)`, broadcast across `DK` | | range | bounded in `(0,1)` | **unbounded** | Forking the emitter would have duplicated three pieces of real algebra to vary one expression. Instead `KdaChunkPrepSpec` gains `gate_kind`, defaulting to `"kda"`. ## Because this touches a shipped family, the default path must not move | flavour | existing cases | added | **existing SHAs moved** | |---|---:|---|---:| | llvm20 | 18 | 2 GDN cases | **0** | | llvm22 | 18 | 2 GDN cases | **0** | | llvm23 | 18 | 2 GDN cases | **0** | Every KDA spec emits byte-identical IR. The golden fixture gains the two new GDN cases and nothing else. ## Two launches, and `auto` is refused A GDN prefill is `chunk_prep` then `chunk_scan`. There is no fused single-kernel path, so `dispatch_gdn_prefill` **rejects** `algorithm="auto"` rather than resolving it to one half — dispatching half a computation and returning successfully is the failure mode worth preventing. ```python prep = dispatch_gdn_prefill(GdnPrefillRequest(..., algorithm="chunk_prep")) scan = dispatch_gdn_prefill(GdnPrefillRequest(..., algorithm="chunk_scan")) ``` ## What is in the diff | area | files | |---|---| | emitter | `kernels/gfx950/kda_chunkwise.py` — `gate_kind` field, GDN gate lowering, validation | | driver | `builders/gfx950/kda/gdn_prefill.py` — host driver + fp32 reference | | dispatch | `dispatch/gdn/{__init__,prefill_common,prefill_gfx950}.py` | | tests | prefill wiring, decay guard, KDA chunkwise spec, KDA golden + fixture, on-device GDN numeric | | docs | `gdn.md` prefill section, `index.md` entry points, README inventory, `SUPPORT_MATRIX` | | benchmark | `sweep_prefill_value_splits.py` | The unbounded gate is why `test_gdn_prefill_decay_guard.py` exists: KDA's gate cannot leave `(0,1)`, GDN's can, so the supported envelope is asserted rather than assumed. ### One unrelated change, declared `kda_chunkwise.py:286-293` rewrites a comment about the triangular solve's cost: the old text said the scalar part "shrinks as the square of the block size", which is wrong -- there are `C/b` blocks, so the total scalar work is `~C*b/2`, linear in `solve_block`. Comment only, zero codegen effect, and the dispatcher pins `solve_block=8` so nothing here exercises it. It is not GDN work; it is kept rather than reverted because leaving a known-wrong comment in a shipped file costs the next reader more than the scope purity saves a reviewer. ## Stack PR 4 of 4: #12070 -> #12090 -> #12172 (decode) -> **this**, based on #12172. CI is red until the first two merge; verification below ran on a tree carrying all of them. ## Verification Gates run on this branch at head `8edf4a84007`, with the base branch as the baseline: | gate | this PR | base branch | verdict | |---|---|---|---| | installed golden IR (llvm20/22/23) | PASS | PASS | clean | | `platform/tests` | 3 failed / 865 passed | 3 failed / 865 passed on `develop` | pre-existing, identical | | `library/tests` * | 18 failed / 1276 passed | 18 failed / 1268 passed | **0 new failures, +8 passed** | | On-device gfx950 (Slurm 872, `smci355-ccs-aus-n01-09`, commit `45399628b51`) | 15 numeric + 18 wiring + 2 guard passed, 0 skipped | — | see caveats | | KDA regression fingerprint | 371 keys, **0 moved** — LLVM text + hsaco bytes + float results | — | same machine code | \* `--continue-on-collection-errors` is required. At this head two GDN suites error at collection and six decode tests fail, all of them `make_kernel_id` / `warp_shuffle_xor_quad` — symbols that arrive with #12090 and #12070. Without that flag pytest aborts the session and runs nothing. `check_byte_identity.py` is the third named gate and was NOT run: it needs the C++ engine, which is not built on this node. ### Three caveats **Criterion 3 is PARTIAL.** The decay guard works and is silicon-verified, but `launch_gdn` is the dev/test launcher -- no production launcher exists for prefill yet, and the decay lives in tensors the request does not carry, so the check cannot move into the dispatcher. Whoever adds one must call `reject_if_decay_out_of_range`. **KDA is undisturbed, measured not argued.** Both trees fingerprinted on one GPU in one session -- LLVM text, hsaco bytes, float results: **371 keys, 0 moved.** **Past its limit the gate saturates, it does not fault.** A bounded wrong answer, and the output degrades while the carried state stays clean -- so a loop validating only its state sees nothing. That is what the guard exists to stop. ### Known limitation Small-batch prefill is not yet optimized; that work is tracked as follow-up. Measured comparisons live in the internal perf-data record, not here.

AviralGoelAMD added 9 commits September 8, 2026 15:28

Merge branch 'develop' into users/avirgoel/rocke/gdn-design-doc

5bebfc5

AviralGoelAMD requested review from a team as code owners September 14, 2026 20:49

github-actions Bot added project: hip-kernel-provider rocke: platform Touches rocKE platform labels Sep 14, 2026

AviralGoelAMD force-pushed the users/avirgoel/rocke/gdn-quadperm branch from 11cc537 to 188fbdd Compare September 14, 2026 20:54

AviralGoelAMD self-assigned this Sep 14, 2026

assistant-librarian Bot added the organization: ROCm label Sep 14, 2026

yraparti reviewed Sep 14, 2026

View reviewed changes

Comment thread dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py Resolved

yraparti reviewed Sep 14, 2026

View reviewed changes

Comment thread dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/ir.py Resolved

yraparti reviewed Sep 14, 2026

View reviewed changes

Comment thread dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/launcher.py Resolved

tenpercent reviewed Sep 15, 2026

View reviewed changes

AviralGoelAMD added 4 commits September 14, 2026 23:22

Base automatically changed from users/avirgoel/rocke/gdn-design-doc to develop September 15, 2026 03:58

AviralGoelAMD requested a review from a team as a code owner September 15, 2026 03:58

AviralGoelAMD added 2 commits September 15, 2026 16:10

AviralGoelAMD requested a review from tenpercent September 16, 2026 02:04

AviralGoelAMD mentioned this pull request Sep 16, 2026

test(rocke): add a launch-path kernarg-packing A/B profiler #12168

Closed

AviralGoelAMD requested a lite review from Copilot September 16, 2026 02:50

Copilot started reviewing on behalf of AviralGoelAMD September 16, 2026 02:51

View session

tenpercent approved these changes Sep 16, 2026

View reviewed changes

Copilot AI reviewed Sep 16, 2026

View reviewed changes

Comment thread dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/gdn/ALGORITHM.md Resolved

Comment thread dnn-providers/hip-kernel-provider/rocke/library/builders/gfx950/gdn/ALGORITHM.md Resolved

AviralGoelAMD enabled auto-merge (squash) September 16, 2026 04:02

AviralGoelAMD mentioned this pull request Sep 16, 2026

feat(rocke): GDN decode kernel and dispatch for gfx950 #12172

Merged

AviralGoelAMD merged commit bc358d8 into develop Sep 16, 2026
111 of 112 checks passed

AviralGoelAMD deleted the users/avirgoel/rocke/gdn-quadperm branch September 16, 2026 05:17

AviralGoelAMD mentioned this pull request Sep 17, 2026

feat(rocke): GDN chunkwise prefill for gfx950 #12225

Merged

AviralGoelAMD mentioned this pull request Sep 17, 2026

feat(hip-kernel-provider): add gfx950 KDA decode #12280

Closed

AviralGoelAMD mentioned this pull request Sep 28, 2026

feat(hip-kernel-provider): add gfx950 KDA decode #12645

Closed

AnaghaRaoAMD mentioned this pull request Oct 2, 2026

feat(rocke): add gfx942 bf16 MLA prefill kernel and dispatch #12464

Merged

This was referenced Oct 17, 2026

feat(rocke): L2-locality-aware block orders for dense attention #12714

Merged

feat(rocke/gfx1151): swapqk prefill + paged split-K decode attention #12891

Open

perf(rocke): add opt-in gfx942 BF16 D128 K LDS double buffering and V register prefetch #13205

Merged

yraparti mentioned this pull request Nov 3, 2026

feat(rocke): add MLA decode kernel for gfx942 and gfx950 #13923

Open

8 of 9 tasks

Navigation Menu

Repository navigation

feat(rocke): quad_perm crosslane intrinsic + kernarg-packer precompile - #12070

feat(rocke): quad_perm crosslane intrinsic + kernarg-packer precompile#12070
AviralGoelAMD merged 20 commits into
developfrom
users/avirgoel/rocke/gdn-quadperm

AviralGoelAMD commented Sep 14, 2026 •

edited

Loading

Uh oh!

Uh oh!

Uh oh!

Uh oh!

tenpercent left a comment

Uh oh!

AviralGoelAMD commented Sep 15, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 15, 2026 •

edited

Loading

Uh oh!

tenpercent commented Sep 16, 2026

Uh oh!

AviralGoelAMD commented Sep 16, 2026

Uh oh!

tenpercent left a comment

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 17, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

AviralGoelAMD commented Sep 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What this is

What changed

Review round 1 — what changed since the first push

Scope note (disclosed)

Why it's safe

Why it's first in the stack

Verification (run, not planned)

Stacking / base

Uh oh!

Uh oh!

Uh oh!

Uh oh!

tenpercent left a comment

Choose a reason for hiding this comment

Uh oh!

AviralGoelAMD commented Sep 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

therock-pr-bot Bot commented Sep 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

tenpercent commented Sep 16, 2026

Uh oh!

AviralGoelAMD commented Sep 16, 2026

Uh oh!

tenpercent left a comment

Choose a reason for hiding this comment

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Suppressed comments (9)

Uh oh!

Uh oh!

Uh oh!

Uh oh!

assistant-librarian Bot commented Sep 17, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

AviralGoelAMD commented Sep 14, 2026 •

edited

Loading

AviralGoelAMD commented Sep 15, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 15, 2026 •

edited

Loading