Skip to content

refactor(rocke): hoist the shared dispatch selector and kernel-id builder - #22

Open
AviralGoelAMD wants to merge 32 commits into
developfrom
users/avirgoel/rocke/gdn-dispatch-core
Open

AviralGoelAMD wants to merge 32 commits into
developfrom
users/avirgoel/rocke/gdn-dispatch-core

Conversation

@AviralGoelAMD

@AviralGoelAMD AviralGoelAMD commented Sep 15, 2026

Copy link
Copy Markdown
Owner

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:

_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:

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 (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 on develop. ROCm#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: #22.

…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.
@github-actions github-actions Bot added documentation Improvements or additions to documentation project: hip-kernel-provider github actions labels Sep 15, 2026
@AviralGoelAMD
AviralGoelAMD changed the base branch from develop to users/avirgoel/rocke/gdn-quadperm September 15, 2026 04:16
@AviralGoelAMD
AviralGoelAMD changed the base branch from users/avirgoel/rocke/gdn-quadperm to develop September 15, 2026 04:17
@AviralGoelAMD
AviralGoelAMD requested a balanced review from Copilot September 15, 2026 04:18

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟢 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_matches and make_kernel_id helpers.
  • 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.
CMiservaAMD and others added 18 commits September 16, 2026 02:49
…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 />


[![Dependabot compatibility
score](https://dependabot-badges.githubapp.com/badges/compatibility_score?dependency-name=tornado&package-manager=pip&previous-version=6.5.7&new-version=6.5.8)](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 />


[![Dependabot compatibility
score](https://dependabot-badges.githubapp.com/badges/compatibility_score?dependency-name=tornado&package-manager=pip&previous-version=6.5.7&new-version=6.5.8)](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 />


[![Dependabot compatibility
score](https://dependabot-badges.githubapp.com/badges/compatibility_score?dependency-name=tornado&package-manager=pip&previous-version=6.5.7&new-version=6.5.8)](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.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment