refactor(rocke): hoist the shared dispatch selector and kernel-id builder - #22
Open
AviralGoelAMD wants to merge 32 commits into
Open
AviralGoelAMD wants to merge 32 commits into
AviralGoelAMD wants to merge 32 commits into
Conversation
…lder ISSUE ID : AICK-1513 Every operator family reimplemented two identical pieces of dispatch logic: the explicit algorithm/spec_id pin selector, and the KernelId builder that gives one pick its stable identity. Two copies today, byte-identical, and a third arriving with every new family -- which is how pin semantics and cache identity drift apart between families that are supposed to share them. Both now live in rocke.dispatch.core as selector_matches() and make_kernel_id(op=...); attention and kda call them and keep their private names as thin aliases, so no call site outside these files changes. Identity is preserved, not assumed. make_kernel_id reads candidate.family where the private copies read each family's _FAMILY constant, and KernelId.family feeds the cache key -- so the two agree only while every candidate is registered with its own family constant. test_core_helpers.py asserts exactly that over the real registries, which is what makes the migration provably key-identical and keeps it so for families added later. The rest of the new test file points straight at the helpers: pin matching including the auto/case/whitespace rules and both rejection reasons, and kernel-id determinism, spec sensitivity, and field provenance. The family suites only exercised these indirectly, so a helper regression used to surface as a confusing failure downstream. Also drops a typing import in kda/common.py left unused by the change. Verification: library/tests/dispatch 334 passed, 216 subtests. The migration is mutation-verified -- forcing selector_matches to reject turns 46 kda dispatch tests red, proving the family really routes through the shared helper rather than a surviving private copy.
AviralGoelAMD
changed the base branch from
develop
to
users/avirgoel/rocke/gdn-quadperm
September 15, 2026 04:16
AviralGoelAMD
changed the base branch from
users/avirgoel/rocke/gdn-quadperm
to
develop
September 15, 2026 04:17
There was a problem hiding this comment.
🟢 Approval recommended
The refactor preserves existing dispatch identities and selector behavior with focused regression coverage.
Pull request overview
Hoists duplicated dispatch selector and kernel-ID logic into the shared dispatch core without changing kernel behavior.
Changes:
- Adds shared
selector_matchesandmake_kernel_idhelpers. - Routes attention and KDA dispatch through the shared helpers.
- Adds direct helper and registry-invariant tests.
File summaries
| File | Description |
|---|---|
rocke/dispatch/core.py |
Adds shared dispatch helpers. |
tests/dispatch/test_core_helpers.py |
Tests helper behavior and family consistency. |
dispatch/kda/common.py |
Reuses the shared selector. |
dispatch/kda/__init__.py |
Reuses the shared kernel-ID builder. |
dispatch/attention/common.py |
Reuses the shared selector. |
dispatch/attention/__init__.py |
Reuses the shared kernel-ID builder. |
Review details
- Files reviewed: 6/6 changed files
- Comments generated: 0
- Review effort level: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
ISSUE ID : AICK-1513 Three findings from @arai713 on ROCm#12090. 1. selector_matches read the pin fields through getattr with an auto default, so a request type that never declared algorithm/spec_id went from raising AttributeError -- what every family's private copy did -- to matching every candidate. A pin silently ignored is the worst failure this function has, and the base OperatorRequest does not declare either field, so the quiet path was reachable by any new family. Both are read directly again, with a test that a request missing them raises. 2. Only algorithm's case/whitespace normalization was tested, though the helper normalizes spec_id on its own line. spec_id is the likelier victim -- a short hand-typed tag like b4 in a config or env override -- and the line could have been deleted with the suite still green. Now asserted too. 3. The family-constant test could not fail: candidates() returns what register() already filtered, and register() rejects a family mismatch. A test that cannot fail hides the day someone removes the guard it was standing in for. Replaced with a test of that guard -- registering a foreign-family candidate must raise -- which is the property the hoist's identity-preservation actually rests on. All three are mutation-verified: restoring the getattr defaults, dropping spec_id normalization, or removing register()'s family check each turns the matching test red. library/tests/dispatch 335 passed.
…h (ALMIOPEN-2509) (ROCm#12084) Two defects kept the hipDNN test suite from working in an **installed** tree: `ctest` found no tests at all, and one test fixture was addressed by an absolute path into the build machine's source checkout. Both were invisible in a normal build-tree run and only surfaced when the kernel ingestor was enabled in a TheRock packaged build, where the tests run from an artifact that carries no checkout. This PR fixes both. It touches test wiring and test sources only — no product code changes.
Updated all pinned TheRock workflow and source refs to `dd17fd5` due to submodule bump Co-authored-by: therockbot <therockbot@amd.com>
ROCm#12070) 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 (ROCm#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 ROCm#11807 has merged, and this PR is now based directly on `develop`.
JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-4714 ## Summary Fix CDNA5 Layer 2 exclusive after/before overlap detection: claim the full unclamped after demand (wmmaWindowsNeeded + latencyWmmaBudget) instead of the promote-clamped issue window. Previously, after barriers whose base threshold exceeded issuedCount could miss overlap with earlier before-barriers, so they stayed pinned at the region end. Clean up Layer 2 structs/helpers without changing scheduling policy beyond the claim-window fix; update unit tests accordingly. ## Motivation Layer 2 decides whether exclusive after/before barrier groups fight for the same WMMA windows. The after claim window used overlapWmmaWindow, which is capped to adjustedAfterThreshold / issuedCount for promote timing. When base after demand was large (issue windows + drain latency) but the region had few WMMAs, the claimed interval shrank and failed to overlap a nearby before group — so no reconcile ran and the after barrier stayed at the end. ## Technical Details Overlap claim (behavior change): after exclusive groups now claim wmmaWindowsNeeded + latencyWmmaBudget (unclamped). Promote threshold still clamps to issuedCount. Proportional split demand: after still uses issue-only wmmaWindowsNeeded; before uses its claim window. Cleanup: remove dead overlapWmmaWindow; rename summary fields (claimWindow / splitNeeded / pendingThreshold); split after/before group builders; fold pending thresholds and descendants into the group; write setGroupThreshold once after optional tensor-load spacing. Tests: rebuild the “separate windows” fixture so lastOverlap > wmmaWindowsNeeded still yields non-overlap under the new claim; fold end-to-end merge checks into that test; drop duplicated/dummy smoke tests. ## Test plan unit_tests --gtest_filter='CDNA5ReadyQueueTest.*:DAGSchedulerPassTest.*:InFlightQueue*' DAGSchedulerPassTest.Layer2* hipblaslt / tensilelite kernel spot-check (e.g. mxf4 barrier placement) as needed ## Test Result <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
…OCm#12065) ## Motivation - 4 StreamK gfx1250 tests carry the `ffm_fail` marker (expected-fail under FFM emulation) but now PASS under emulation. With `xfail_strict = True`, an XPASS is a hard failure, so the FFM tensilelite leg fails on every PR: - `sk_dynamic_sgemm_quick` - `sk_dynamic_hgemm_quick` - `sk_tdm_general_batched` - `sk_hybrid_quick` ## Technical Details - Remove stale markers on above StreamK gfx1250 tests - FFM-only, HW unaffected. ## Test Plan - FFM tensilelite gfx1250 leg should clear the 4 XPASS(strict) failures. - The other `ffm_fail` StreamK tests (`tdm_multicast`) still correctly XFAIL and are untouched. ## Test Result - [ ] Verify FFM test passes: TBD ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low — test-marker only; no kernel/codegen change. Re-adding the marker is trivial if a test later regresses on FFM.
## Motivation Wrong a_type was generated for TF32 with config yaml option. Fixed to return the correct a_type. ## Test Plan Added small unit test to verify the correctness of the output. ## Test Result All tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. JIRA ID : AIHPBLAS-4759
…#11928) ISSUE ID: AICK-1763 ## Motivation Forward conv and wgrad both have a fast path for 1×1 convolutions. dgrad never got one. For 1×1 with stride 1 and no padding, the address math is trivial, the offsets are simple multiply-adds. But dgrad was still running the general-case code: a runtime integer divide plus a bounds check that can never fail, and both of them inside the K-loop, so the cost repeats on every iteration. 1×1 shapes are common in real backward-data workloads, so this was worth fixing. ## Technical Details Four commits, branched off users/aledudek/transpose_load_bwd_weight_data_gfx1250 @ f02fad7. 1. The fast path. When the shape is 1×1, stride 1, unpadded and ungrouped, use the simple offset formula directly instead of the general one. Two early-returns in the dY and W descriptors, mirrored in the C++ engine in the same commit. The 1×1 kernel goes from 1295 to 991 instructions, and most of the integer divides and bounds selects disappear. 2. A vector_size_b bug. That field is meant as an upper limit, not a hard request, wgrad treats it that way. The K-outer path clamped it properly, the M-outer path didn't, so a too-wide value passed validation and then crashed the builder. Both engines now clamp. This changes no emitted kernel. 3. Two benchmark flags. --lds-k-outer {auto,on,off} so the two LDS layouts can be compared (the automatic choice is the same for every config of a given shape, so a normal sweep has nothing to compare against), and --csv-top N because the CSV writer was hardcoded to 5 rows and ignored --top. Both keep the old behaviour by default. 4. Docs and a driver. Both LDS case studies end their trace command with "single-config driver", that driver didn't exist, so it's added here. Also a short write-up of two other optimisations that were built, measured, and dropped because they didn't help. ## Test Plan - tools/check_byte_identity.py at llvm20 and llvm22 - tests/instances/differential/run_diff.py --only conv_implicit_gemm_dgrad - tests/instances/test_conv_dgrad_correctness.py - Numeric check of the fast path against dgrad_reference on three 1×1 shapes plus a 3×3 control - hsaco comparison of 3×3 stride-1 and stride-2 with and without the change - Emission-neutrality of commit 2 by hashing lowered IR across the full sweep grid, before and after ## Test Result - Byte-identity GREEN, 69 configs, at both LLVM flavors. dgrad parity config 4 is pointwise, so the gate covers the new path. - Differential GREEN (69 families; dgrad 14 configs). - dgrad correctness: 22 passed. - 3×3 stride-1 and stride-2 hsacos byte-identical before/after. only pointwise emission moves. - Fast path is numerically exact: relative error against dgrad_reference identical to the generic path on all four cases. - Commit 2: zero IR-hash changes across 4564 previously-building configs, zero regressions, 428 previously-unbuildable configs now build. - check_golden reports no drift. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
…m#12079) ## Motivation Removing a leftover from an experiment to run our test suites on a non-default stream. ## Technical Details Removes the creation of a stream rather than using the default one. ## Test Plan A test is in place to make sure we won't regress. ## Test Result unit tests are passing. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
…Kernels (ROCm#11037) ## Motivation Add the ability to switch between 32/64-bit integer arithmetic for the remaining kernels (Stockham, Bluestein, Real/Complex copy kernels, Twiddles, and Chirp). JIRA ID: AIFFT-462 ## Technical Details As previously implemented for the transpose kernel, the actual problem size is checked to decide the tightest integral type for the kernels involved in the computations. A genuine effort has been made to determine if the integral types can be made 32-bit regardless of the FFT size, or if they need to be widened to 64-bit for large enough FFTs: - Integral types in the Stockham kernels were always 64-bit, but this is not really required if the FFT size is not large enough. - The real/complex copy kernels were hardcoded to 32-bit and they now can dynamically change depending on the FFT size. - Single kernel Bluestein kernels were inheriting the arithmetic from the base Stockham kernels, and now the integral type is also dynamic. Same for fused and non-fused Bluestein kernels. - Chirp and twiddle kernels were also fixed at 64-bit integral types and now are also dynamic, although they are not performance critical since they are on the plan path (not the hot FFT execution path). This PR is IMO a step forward if, in the future, we decide to completely eliminate 64-bit integral types from the Stockham kernels, i.e., by passing an already offset pointer to memory in such a way that the kernel only needs 32-bit indexing to reach the work dataset. This is a larger change since it requires a complete redesign on how the kernels actually compute their indices work set. There should be no observable performance changes for the FFTs we test. The Stockham kernels now have less register pressure on reasonably sized FFTs due to the switch to 32-bit arithmetic. The changes here, however, are not enough to affect occupancy - but, if callbacks or additional user computations are fused in the FFT kernels then there is more headroom to still maintain the same occupancy. Other minor changes: - The previous type "index_type" has been renamed to a more suitable "integer_type" for the generated kernels, and IndexType now maps to KIntType. - Fixed an issue with the chirp generator that would produce incorrect results with sizes larger than 32-bits. There is a pre-existing issue with large twiddle tables where, for example, a very large 1D, unbatched transform with length > 2^32 `./rocfft-bench --length 8589934592 -t 0 --precision single` fail to run and throw an exception because the twiddle table look up calculation (TWLstep1 to TWLStep4) does not handle a length larger than 32 bit: ``` if(node.large1D > (size_t)256 * 256 * 256 * 256) throw std::runtime_error("large1D twiddle size too large error"); ``` I will add a separate issue and fix for this one. ## Test Plan Existing tests should exercise the changes in this PR. ## Test Result The tests should pass without any issues. No performance regressions should be observed. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
JIRA ID : ALMIOPEN-2622 ## Motivation MIOpen is currently subscribed to the shared `.github/workflows/clang-tidy.yml` workflow. The clang-tidy workflow cannot be made a required check for MIOpen at this time, which prevents reliable signal from the check and blocks retirement of Jenkins CI for MIOpen. Retiring Jenkins CI requires clang-tidy to be a required check for MIOpen, and the only path to make the check required is to run it as part of the component CI that is already gated for MIOpen. This PR moves the MIOpen clang-tidy steps out of the shared workflow and into `.github/workflows/component-ci-miopen.yml` so the check can be enforced via component CI. The change is reversible once the shared clang-tidy workflow can be made required. ## Technical Details Removed MIOpen from the shared clang-tidy workflow: - `.github/workflows/clang-tidy.yml`: removed LIBS entry for miopen with path_filters and checkout_paths, removed `Install deps (miopen)` step, removed `Run clang-tidy (miopen)` step. Added dedicated clang-tidy job to MIOpen component CI: - `.github/workflows/component-ci-miopen.yml`: added `miopen-clang-tidy` job with `runs-on: azure-linux-scale-rocm`. - Steps replicate the former clang-tidy workflow: checkout rocm-libraries sparse `projects/miopen`/`shared/ctest`, checkout TheRock `df3d451a3c054e14705ddf94e58498e1208df8d5`, install base dependencies, install ROCm via TheRock artifacts, install MIOpen deps (`cmake<4`, cget, rbuild, clang-tidy-23) and prepare deps via rbuild, then run `cmake -S projects/miopen -B build-miopen -G Ninja -DMIOPEN_BACKEND=HIP -DBUILD_DEV=ON -DMIOPEN_USE_MLIR=OFF -DCLANG_TIDY_EXE=/usr/bin/clang-tidy-23 …` followed by `ninja -C build-miopen -k 0 analyze`. Trade-off: clang-tidy now runs in component CI which is already required, at the cost of duplicating setup steps previously shared. Revert plan exists if shared clang-tidy becomes required. ## Risk Assessment **Risk Level: 🟢 Low** ### Impacted Components - `.github/workflows/clang-tidy.yml` – MIOpen removed from matrix - `.github/workflows/component-ci-miopen.yml` – new `miopen-clang-tidy` job added - MIOpen CI gating – clang-tidy now runs as part of component CI ### Potential Side Effects - Clang-tidy failures will now block component CI instead of the shared clang-tidy workflow. - Duplicate setup steps increase job runtime slightly. - No code changes; only workflow configuration. ### Mitigation Steps - YAML validated for both workflows. - Change is reversible; can be reverted once shared clang-tidy can be made required. - Existing component CI gating provides required-check enforcement. ## Test Plan - Validate YAML syntax for both modified workflows using `python3 -c "import yaml; yaml.safe_load(open(...))"`. - Verify `git diff develop...HEAD --stat` shows only the two workflow files changed. - Confirm `grep miopen .github/workflows/clang-tidy.yml` returns no matches. - Confirm `miopen-clang-tidy` job exists in `.github/workflows/component-ci-miopen.yml` with expected steps. ## Test Result - YAML parsing succeeded for both `.github/workflows/clang-tidy.yml` and `.github/workflows/component-ci-miopen.yml`. - Diff shows 50 deletions in clang-tidy.yml and 54 insertions in component-ci-miopen.yml. - MIOpen references removed from clang-tidy workflow; new job present in component CI. - No tests were added; change is CI configuration only. Follow-up work: monitor first component CI runs for clang-tidy failures.
…ROCm#9627) ## Motivation Closes the warp-layer part of ROCm#7934. The FP64 `v_mfma_f64_16x16x4f64` matrix-core instruction is already exposed by the ck_tile **arch** layer for CDNA2/3/4 (`mfma_gfx9.hpp`, `enable_if_target_id_t<…, GFX90A, GFX942, GFX950>`), but there was no **warp**-layer chain to surface it to a GEMM. This PR adds that wiring so ck_tile can express a double-precision GEMM, which teams doing dense linear algebra, FEM/stencil solvers, and state-vector quantum simulation need and currently have to drop back to the legacy `DeviceGemmXdl` API for. Per @adityas-amd's guidance in the issue, this is the **16x16x4 MVP**: lowest-risk, maps to a single builtin, and reuses the well-tested FP32 16x16x4 per-lane distribution. ## Technical Details Adds the warp attribute → alias → dispatcher → test chain for FP64: | File | Change | |---|---| | `include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp` | New `WarpGemmAttributeMfmaImplF64F64F64M16N16K4` | | `include/ck_tile/ops/gemm/warp/warp_gemm.hpp` | `WarpGemmMfmaF64F64F64M16N16K4` + `…K16` aliases | | `include/ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp` | Dispatcher specializations for `fp64_t` (K=4, K=16) | | `test/ck_tile/warp_gemm/test_f64_16x16x4_mfma.cpp` | New unit test vs CPU reference GEMM | | `test/ck_tile/warp_gemm/CMakeLists.txt` | Registers the test, scoped to gfx90a/gfx942/gfx950 | **C-tile layout:** the FP64 output does **not** follow the 4×N block layout of other MFMA instructions. Per the MI300 (CDNA3) and MI350 (CDNA4) ISAs — both **§7.1.3.4 (`V_MFMA_F64_16X16X4_F64`)** — "the output rows are packed contiguously across the lanes." Accordingly the C distribution uses `kCM0PerLane=4, kCM1PerLane=1` (not the FP32 `1/4`). `blgp=0` (for f64 builtins BLGP is A/B/C negation, not lane-group select). The guard is `__gfx90a__ || __gfx942__ || __gfx950__` — deliberately **not** `__gfx9__`, since gfx908/CDNA1 has no FP64 matrix core. **Scope / follow-up:** ROCm#7934 also asked for an `example/ck_tile/03_gemm` FP64 instance and a `DeviceGemmXdl`-baseline comparison. Building those surfaced **two pre-existing gaps in shared, precision-agnostic infrastructure** that nothing had exercised with an 8-byte element type: **1.** `ck_tile::numeric_traits<double>` doesn't exist (required by `check_err.hpp`'s tolerance helpers; existing specializations use `uint32_t` masks too narrow for a 52-bit mantissa) **2.** a compile-time divide-by-zero in `gemm_universal_pipeline_ag_bg_cr_policy.hpp` when `sizeof(ADataType)=8` in the ColumnMajor + Wave64 path. Neither is caused by this change. I'd like to land this self-contained warp-layer piece first and address those in a follow-up so the example + baseline can build on a fixed base. ## Test Plan - Unit test `test_ck_tile_wg_f64_16x16x4_mfma` compares the warp-GEMM device output against `ck_tile::reference_gemm` (CPU) with random inputs. - Built and run on **gfx942 (MI300X)** hardware. - Built for **gfx950 (MI350X)** target; runtime left to CI (no MI350X hardware on my side). - **gfx90a (MI250)** runtime left to CI (no MI250 hardware on my side). - **gfx908** guard-regression: confirmed the test is **not** registered and no FP64 device code compiles. ## Test Result - **gfx942 (MI300X): PASS** — compiles cleanly under `-Werror -Weverything` and passes at runtime: ``` [ RUN ] WarpGemmF64.MFMA_16x16x4 [ OK ] WarpGemmF64.MFMA_16x16x4 (152 ms) [ PASSED ] 1 test. ``` - **gfx950:** compiles cleanly; the CDNA4 ISA's f64 16x16x4 output layout is byte-identical to CDNA3's (both §7.1.3.4) and the arch layer uses one shared definition for both, so the same result is expected. Awaiting CI runtime confirmation. - **gfx90a:** covered by the same shared arch-layer definition and the existing `DeviceGemmXdl` f64 path. Awaiting CI runtime confirmation. - **gfx908:** correctly excluded (no FP64 code compiled). ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. - [x] Added a unit test under `test/ck_tile/warp_gemm/`. - [x] Runtime-verified on gfx942; gfx908 correctly excluded. - [ ] CI green on gfx90a / gfx950 (draft → ready once confirmed). --------- Signed-off-by: Danila <55054065+danila-permogorskii@users.noreply.github.com> Signed-off-by: Danila Permogorsky <developer.permogorsky@gmail.com> Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com>
## Motivation Improvements for coverage, performance and benchmarking for direct conv. ## Technical Details - new DirectDepthwiseSpatialSpec - support for stride =2 - timeout for benchmarking - disable for wgrad and dgrad - compilation limitation for some implicit gemm pipelines (out of scope but needed for benchmarking) ## Test Plan test_direct_conv_coretcness ## Test Result Pass ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. JIRA ID: AICK-2235
ROCm#11342) ## Motivation The fused GEMM + all-to-all epilogue redistributes the leading run of D's feature dimension across the ranks of a device communicator from the GEMM's own store path, removing a separate collective pass. This lands the host side ahead of the TensileLite kernel so callers can build and validate a request, and so the family's enum and attribute ranges are reserved before the other fused-epilogue families arrive. JIRA ID : ROCM-27524 ## Technical Details - Public API: `HIPBLASLT_FUSEABLE_EPILOGUE_A2A_PREFIX` (7, leaving 0-6 to the chainable families) and its five attributes (12-16), plus `hipblasLtSdmaQueue_t`, `hipblasLtA2ACompletionMode_t`, and the shared `hipblasLtFusedEpilogueCreate` / `Add` / `SetAttribute` / `Destroy` builder. The builder documents family-neutral rules; per-family specifics stay on the enums. - `hipblasLtSetDeviceComm` registers a per-handle view of a device communicator, allocating the fine-grained flag regions and resolving peer addresses through a caller-supplied allgather, with IPC mapping for out-of-process peers. The flag-block layout is twinned with `FUSED_A2A_FLAG_BLOCK_BYTES` in the kernel's `FusedA2AKernArg.hpp`. - The descriptor attaches through `HIPBLASLT_MATMUL_DESC_FUSED_EPILOGUE` (106, after develop's `UNIFORM_SUMMATION_ORDER_EXT`) and is carried by `_rocblaslt_matmul_desc` as a non-owning pointer. - Validation is split three ways: stage composition at `Add`, required attributes at attach, and communicator, shape, and layout constraints before the heuristic and again at launch. No kernel is wired, so a well-formed request reports `HIPBLAS_STATUS_NOT_SUPPORTED`. ## Test Plan Build `hipblaslt-test` and run `--gtest_filter='FusedA2A*'` on gfx950, covering descriptor lifecycle, builder composition rules, attribute range and width checks, communicator registration, attach-time completeness, and each dispatch-time rejection. ## Test Result 40 tests across 6 suites pass on gfx950. Nothing dispatches a GEMM: a valid request reports `NOT_SUPPORTED`, confirming the validation chain runs to completion. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Ding, Yi <yi.ding@amd.com> Co-authored-by: Alex Vasile <48962821+Alex-Vasile@users.noreply.github.com>
## Motivation Reduce runtime of callback tests. ## Technical Details Remove unconditional emission of callback cases from partial-pass tests and instead move a few lengths into the callback suite. It's not terribly interesting to test every single partial-pass case with callbacks so it's not really worth the time spent. Move callback test cases to only test batch-2, as it's difficult to imagine batch-1 or batch > 2 making any difference to them working. Cache the JIT callbacks compiled by the tests in memory since we really should have fewer than 10 alive in any test run. There's no reason to recompile them for every test case. This reduces the run time of the callback tests by about a half on its own on my workstation. Increase default callback_prob to 0.2 as the above changes have reduced redundant cases. ## Test Plan Existing test cases run. ## Test Result Tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
…docs/sphinx (ROCm#11569) Bumps [tornado](https://github.com/tornadoweb/tornado) from 6.5.7 to 6.5.8. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/tornadoweb/tornado/blob/master/docs/releases.rst">tornado's changelog</a>.</em></p> <blockquote> <h1>Release notes</h1> <p>.. toctree:: :maxdepth: 2</p> <p>releases/v6.5.8 releases/v6.5.7 releases/v6.5.6 releases/v6.5.5 releases/v6.5.4 releases/v6.5.3 releases/v6.5.2 releases/v6.5.1 releases/v6.5.0 releases/v6.4.2 releases/v6.4.1 releases/v6.4.0 releases/v6.3.3 releases/v6.3.2 releases/v6.3.1 releases/v6.3.0 releases/v6.2.0 releases/v6.1.0 releases/v6.0.4 releases/v6.0.3 releases/v6.0.2 releases/v6.0.1 releases/v6.0.0 releases/v5.1.1 releases/v5.1.0 releases/v5.0.2 releases/v5.0.1 releases/v5.0.0 releases/v4.5.3 releases/v4.5.2 releases/v4.5.1 releases/v4.5.0 releases/v4.4.3 releases/v4.4.2 releases/v4.4.1 releases/v4.4.0 releases/v4.3.0 releases/v4.2.1 releases/v4.2.0 releases/v4.1.0 releases/v4.0.2 releases/v4.0.1 releases/v4.0.0 releases/v3.2.2</p> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/tornadoweb/tornado/commit/a55abe3e3bf59994f29b2f7084c46341f0d4f6a7"><code>a55abe3</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3704">#3704</a> from bdarnell/security-6.5.8</li> <li><a href="https://github.com/tornadoweb/tornado/commit/fc794885f0ccf9c33f3a66d890abcc237dd50b3c"><code>fc79488</code></a> docs: add additional credit to release notes</li> <li><a href="https://github.com/tornadoweb/tornado/commit/7b017630d3139ca0d1ebdf6ac3b3ffe7725a7129"><code>7b01763</code></a> Fix test_strip_headers_on_redirect's URL-embedded-credentials cases</li> <li><a href="https://github.com/tornadoweb/tornado/commit/d72fff8d7b9b8f6aa68505847e5483d600e3184c"><code>d72fff8</code></a> release notes and version bump for 6.5.8</li> <li><a href="https://github.com/tornadoweb/tornado/commit/b168818f8aae39808b981878fb358cbe02a6238e"><code>b168818</code></a> auth: Formally deprecated OpenIDMixin</li> <li><a href="https://github.com/tornadoweb/tornado/commit/da284767eae8e1f0484f123b8c3225f6465b09c7"><code>da28476</code></a> web: Also check for semicolons in deprecated mixed-case cookie args</li> <li><a href="https://github.com/tornadoweb/tornado/commit/8d6363ed7b69d5f0da806efe34d256627a2191de"><code>8d6363e</code></a> httputil: Enforce a new limit on the number of arguments in a request</li> <li><a href="https://github.com/tornadoweb/tornado/commit/de85b3f87446e323e881bbaa3d5a74f4b76e5f05"><code>de85b3f</code></a> httputil: Apply multipart max_parts limit earlier</li> <li>See full diff in <a href="https://github.com/tornadoweb/tornado/compare/v6.5.7...v6.5.8">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) You can disable automated security fix PRs for this repo from the [Security Alerts page](https://github.com/ROCm/rocm-libraries/network/alerts). </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…/docs/sphinx (ROCm#11570) Bumps [tornado](https://github.com/tornadoweb/tornado) from 6.5.7 to 6.5.8. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/tornadoweb/tornado/blob/master/docs/releases.rst">tornado's changelog</a>.</em></p> <blockquote> <h1>Release notes</h1> <p>.. toctree:: :maxdepth: 2</p> <p>releases/v6.5.8 releases/v6.5.7 releases/v6.5.6 releases/v6.5.5 releases/v6.5.4 releases/v6.5.3 releases/v6.5.2 releases/v6.5.1 releases/v6.5.0 releases/v6.4.2 releases/v6.4.1 releases/v6.4.0 releases/v6.3.3 releases/v6.3.2 releases/v6.3.1 releases/v6.3.0 releases/v6.2.0 releases/v6.1.0 releases/v6.0.4 releases/v6.0.3 releases/v6.0.2 releases/v6.0.1 releases/v6.0.0 releases/v5.1.1 releases/v5.1.0 releases/v5.0.2 releases/v5.0.1 releases/v5.0.0 releases/v4.5.3 releases/v4.5.2 releases/v4.5.1 releases/v4.5.0 releases/v4.4.3 releases/v4.4.2 releases/v4.4.1 releases/v4.4.0 releases/v4.3.0 releases/v4.2.1 releases/v4.2.0 releases/v4.1.0 releases/v4.0.2 releases/v4.0.1 releases/v4.0.0 releases/v3.2.2</p> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/tornadoweb/tornado/commit/a55abe3e3bf59994f29b2f7084c46341f0d4f6a7"><code>a55abe3</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3704">#3704</a> from bdarnell/security-6.5.8</li> <li><a href="https://github.com/tornadoweb/tornado/commit/fc794885f0ccf9c33f3a66d890abcc237dd50b3c"><code>fc79488</code></a> docs: add additional credit to release notes</li> <li><a href="https://github.com/tornadoweb/tornado/commit/7b017630d3139ca0d1ebdf6ac3b3ffe7725a7129"><code>7b01763</code></a> Fix test_strip_headers_on_redirect's URL-embedded-credentials cases</li> <li><a href="https://github.com/tornadoweb/tornado/commit/d72fff8d7b9b8f6aa68505847e5483d600e3184c"><code>d72fff8</code></a> release notes and version bump for 6.5.8</li> <li><a href="https://github.com/tornadoweb/tornado/commit/b168818f8aae39808b981878fb358cbe02a6238e"><code>b168818</code></a> auth: Formally deprecated OpenIDMixin</li> <li><a href="https://github.com/tornadoweb/tornado/commit/da284767eae8e1f0484f123b8c3225f6465b09c7"><code>da28476</code></a> web: Also check for semicolons in deprecated mixed-case cookie args</li> <li><a href="https://github.com/tornadoweb/tornado/commit/8d6363ed7b69d5f0da806efe34d256627a2191de"><code>8d6363e</code></a> httputil: Enforce a new limit on the number of arguments in a request</li> <li><a href="https://github.com/tornadoweb/tornado/commit/de85b3f87446e323e881bbaa3d5a74f4b76e5f05"><code>de85b3f</code></a> httputil: Apply multipart max_parts limit earlier</li> <li>See full diff in <a href="https://github.com/tornadoweb/tornado/compare/v6.5.7...v6.5.8">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) You can disable automated security fix PRs for this repo from the [Security Alerts page](https://github.com/ROCm/rocm-libraries/network/alerts). </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
…docs/sphinx (ROCm#11572) Bumps [tornado](https://github.com/tornadoweb/tornado) from 6.5.7 to 6.5.8. <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/tornadoweb/tornado/blob/master/docs/releases.rst">tornado's changelog</a>.</em></p> <blockquote> <h1>Release notes</h1> <p>.. toctree:: :maxdepth: 2</p> <p>releases/v6.5.8 releases/v6.5.7 releases/v6.5.6 releases/v6.5.5 releases/v6.5.4 releases/v6.5.3 releases/v6.5.2 releases/v6.5.1 releases/v6.5.0 releases/v6.4.2 releases/v6.4.1 releases/v6.4.0 releases/v6.3.3 releases/v6.3.2 releases/v6.3.1 releases/v6.3.0 releases/v6.2.0 releases/v6.1.0 releases/v6.0.4 releases/v6.0.3 releases/v6.0.2 releases/v6.0.1 releases/v6.0.0 releases/v5.1.1 releases/v5.1.0 releases/v5.0.2 releases/v5.0.1 releases/v5.0.0 releases/v4.5.3 releases/v4.5.2 releases/v4.5.1 releases/v4.5.0 releases/v4.4.3 releases/v4.4.2 releases/v4.4.1 releases/v4.4.0 releases/v4.3.0 releases/v4.2.1 releases/v4.2.0 releases/v4.1.0 releases/v4.0.2 releases/v4.0.1 releases/v4.0.0 releases/v3.2.2</p> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/tornadoweb/tornado/commit/a55abe3e3bf59994f29b2f7084c46341f0d4f6a7"><code>a55abe3</code></a> Merge pull request <a href="https://redirect.github.com/tornadoweb/tornado/issues/3704">#3704</a> from bdarnell/security-6.5.8</li> <li><a href="https://github.com/tornadoweb/tornado/commit/fc794885f0ccf9c33f3a66d890abcc237dd50b3c"><code>fc79488</code></a> docs: add additional credit to release notes</li> <li><a href="https://github.com/tornadoweb/tornado/commit/7b017630d3139ca0d1ebdf6ac3b3ffe7725a7129"><code>7b01763</code></a> Fix test_strip_headers_on_redirect's URL-embedded-credentials cases</li> <li><a href="https://github.com/tornadoweb/tornado/commit/d72fff8d7b9b8f6aa68505847e5483d600e3184c"><code>d72fff8</code></a> release notes and version bump for 6.5.8</li> <li><a href="https://github.com/tornadoweb/tornado/commit/b168818f8aae39808b981878fb358cbe02a6238e"><code>b168818</code></a> auth: Formally deprecated OpenIDMixin</li> <li><a href="https://github.com/tornadoweb/tornado/commit/da284767eae8e1f0484f123b8c3225f6465b09c7"><code>da28476</code></a> web: Also check for semicolons in deprecated mixed-case cookie args</li> <li><a href="https://github.com/tornadoweb/tornado/commit/8d6363ed7b69d5f0da806efe34d256627a2191de"><code>8d6363e</code></a> httputil: Enforce a new limit on the number of arguments in a request</li> <li><a href="https://github.com/tornadoweb/tornado/commit/de85b3f87446e323e881bbaa3d5a74f4b76e5f05"><code>de85b3f</code></a> httputil: Apply multipart max_parts limit earlier</li> <li>See full diff in <a href="https://github.com/tornadoweb/tornado/compare/v6.5.7...v6.5.8">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) You can disable automated security fix PRs for this repo from the [Security Alerts page](https://github.com/ROCm/rocm-libraries/network/alerts). </details> Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
## Motivation We want to be able to test AMD vs. NVIDIA's (rocSOLVER vs cuBLAS) implementation of `getrf_batched` using a benchmark framework unified with the rest of hipSOLVER. JIRA ID: AISOLVE-61 ## Technical Details We add `hipsolver_getrfXBatched` to hipSOLVER's API, implement a backend for AMD and NVIDIA that reconciles the differences in input parameters, and add it to our testing/benchmarking client. We make new function headers for `getrf_batched` in the client's internal API and the testing files as opposed to integrating the batched case into a function with the same function header. This is to account for the differences of `getrf_batched` between cuSOLVER and cuBLAS, but breaks the pattern we typically follow in hipSOLVER. ### NOTE: This change is made alongside changes to add `geqrf_batched`, `gels_batched`, `getrs_batched`, and `getri_batched`. You can find PRs for each of these routines, each of them reading extremely similarly. ## Test Plan Run new tests on AMD and NVIDIA hardware. Try invoking `getrf_batched` from the benchmark client. ## Test Result Tests pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Troy Alderson <58866654+tfalders@users.noreply.github.com> Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>
AviralGoelAMD
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Sep 18, 2026
…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#22.
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
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.
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.
attentionandkdaeach 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.pygainsselector_matches()andmake_kernel_id(op=...). Both families now call them and keep their old private names as aliases: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 readsfamily=candidate.family.Those are different expressions.
familyis part ofselection_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 ofcompile_key, which isarch: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:
Tests
library/tests/dispatch/test_core_helpers.pyis 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.selector_matchesto 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.blackclean. Onetypingimport inkda/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 (ROCm#12070) → this → GDN decode → GDN prefill.
It does not depend on ROCm#12070 — that one touches crosslane lowering in
platform/, this one touches dispatch. Both sit ondevelop. ROCm#11795's version of this hoist coveredkdaonly;attentionis 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: #22.