Repository navigation
feat(hip-kernel-provider): route SDPA to unified and dense attention kernels - #13115
ThruptiRajLakshmanaGowda wants to merge 79 commits into
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Pre-commit check failed⛔ pre-commit failed Please run locally:
This repo uses |
|
🎉 All checks passed! This PR is ready for review. |
…orward rocKE had no request type that could describe a hipDNN SDPA forward node: AttentionRequest has no strides, scale source, bias, sequence-length tensors, per-tensor dtype, h_v, or general band bounds. SdpaRequest represents every field of the node. Operands are SdpaTensor (dtype, dims, strides, ragged, alignment); optional subsystems (bias, dropout, paged KV, FP8) are their own frozen dataclasses. Construction normalizes to one spelling per meaning (negative bounds -> -1, scale 1.0 unless baked, bias right-aligned to rank 4 with stride 0 on size-1 dims), so normalized() is a stable cache key and round-trips via from_normalized(). __post_init__ rejects only what the hipDNN frontend rejects; valid-but-unimplemented fields are reported by features() so candidates decline them through Capability.check instead of dropping them. from_hipdnn() builds a request from the flatbuffer JSON hipDNN's JSON helpers and golden generator write, applying getMaskType precedence for the deprecated causal flags and enforcing scale exclusivity. Every checked-in SdpaFwd golden bundle graph translates and round-trips. The case table in tests/sdpa/cases.py reports every feature at least once. Co-Authored-By: Claude <noreply@anthropic.com>
The SDPA streams need one numeric spec to test kernels against. keep() is
hipDNN's band mask (CpuFpReferenceSdpa::isMasked): a negative bound is
unbounded, and left_bound = W keeps W + 1 keys. sdpa_reference() runs the
hipDNN forward pipeline in FP32 for an SdpaRequest: scale and Q/K descale,
unscaled bias broadcast by query head, band mask, softmax with fully
masked rows giving O = 0 and LSE = -inf, and V descale.
It handles strides (flat buffers are read through the request's strides),
GQA including h_k != h_v, d_qk != d_v, ragged packed operands, and
per-batch seq_len padding. As in hipDNN's golden generator, the diagonal
offset uses each batch's effective lengths and padded rows are O = 0,
LSE = -inf. Dropout, paged KV, block masks, sinks, ALiBi, FP8 S/O scaling
and max/sum_exp raise NotImplementedError by name, as hipDNN's CPU
reference rejects them too.
platform cannot import library, so the request is read by attribute.
Tests: keep() matches a transliteration of isMasked for S <= 8, both
diagonals, bounds in {-1, 0, 1, 3}; sdpa_reference matches a scalar fp64
transliteration of the hipDNN forward loop on every case-table case, and
PyTorch's math SDPA (the golden generator's backend) per batch.
Co-Authored-By: Claude <noreply@anthropic.com>
…ntion bodies Seam SA for the SDPA streams. Bias needs each score cell's position and LSE needs each row's final max and denominator; the MFMA and WMMA forward bodies exposed neither. extra_score_transform is now called with keyword q_row and k_col, the query and key positions the mask already computes, so passing them adds no IR. A new optional epilogue_fn(b, *, m_log2, l, row, row_valid) runs per accumulator row slot after O is stored; row_valid is None because these bodies have no Q-length bound. With epilogue_fn=None the body emits nothing new. sage_attention's transform, the only in-tree one on these bodies, takes the new keywords as **_coords. Default IR is unchanged: the byte-identity gate is GREEN for all 76 families at llvm20, llvm22 and llvm23. Tests drive the real fmha_mfma kernel on gfx950, gfx942 and gfx1151: an identity transform yields identical IR, and the epilogue sees every row slot. Co-Authored-By: Claude <noreply@anthropic.com>
Two new change types in KERNEL_AUTHORING.md. A new operator that has a hipDNN graph node needs a lossless request type, a from_hipdnn translator, a reference oracle with hipDNN semantics, and a case table, with unimplemented fields reported as features rather than dropped. A new operand or output (bias, scale tensor, sequence lengths, LSE) is routed only to families that declare the feature and bind it to a kernel argument, behind a default-off flag, and tested against the reference. Co-Authored-By: Claude <noreply@anthropic.com>
Conflicts in mfma_attention.py and sage_attention.py: develop (#12710) already passes the query and key positions to extra_score_transform as positional arguments. Take develop's form for the score hook and keep only epilogue_fn from this branch. The wave32 body now has a Q-length bound, so epilogue_fn receives it as row_valid there. Co-Authored-By: Claude <noreply@anthropic.com>
AnaghaRaoAMD
left a comment
There was a problem hiding this comment.
The PR body says 11/11 passed, but tests/sdpa/test_unified_sdpa_wiring.py now defines ~30 tests. Please refresh the count from the current head.
…sdpa/st0-contract # Conflicts: # dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/__init__.py # dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/cases.py # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/mfma_attention.py # dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/numeric/references.py
…op's sdpa_reference #13575 added rocke.numeric.sdpa_reference (keep, sdpa_reference) and a kernel case table at library/tests/sdpa/cases.py. Redefining keep and sdpa_reference in references.py would shadow them for their existing callers, so: - sdpa_reference(req, tensors) becomes sdpa_reference_from_request in numeric/sdpa_request_reference.py. It decodes strides, ragged buffers, the scale source and FP8 descales, then delegates to develop's float64 oracle with an explicit scale (hipDNN's default is 1.0, not 1/sqrt(d)). - The request case table moves to tests/sdpa/request_cases.py. - The epilogue_fn tests go: AttnFwdExt.epilogue_hook replaces the hook. The score-coordinates no-IR test stays. - TestKeep moves to develop's test_sdpa_reference.py as a _band_mask vs keep agreement check over S <= 8 and bounds {-1, 0, 1, 3}. Co-Authored-By: Claude <noreply@anthropic.com>
…ently drop from_hipdnn read the node with .get() and ignored any other key, so a new, unknown or misspelt hipDNN field left the request unchanged. That breaks the lossless rule the request contract exists for. Allow-list every key hipDNN's JSON helpers write (node, inputs, outputs, attributes, tensor entries) and raise on anything else. Absent keys still read as hipDNN's null. A stats_tensor_uid without generate_stats now raises: the hipDNN frontend only creates a stats output when generate_stats is true, so that combination is contradictory input, not a no-op. A test checks the allow-lists against the IngestorGenerator SDPA graph template, so the two cannot drift. All 57 SdpaFwd golden-bundle graphs still translate. Co-Authored-By: Claude <noreply@anthropic.com>
…d canonicalize SdpaRequest arch Bottom-right band: _mask_features treated the diagonal as moved only when s_q != s_kv. With padded, ragged or paged lengths the offset is s_kv[b] - s_q[b] per batch, so a request with equal max extents reported causal_tl and a top-left-only candidate could admit it and mask the wrong keys. Such requests now report causal_br / window_br. Arch: AttentionRequest canonicalizes arch since #12275, SdpaRequest did not, so "gfx950:sramecc+" failed Capability.check's exact arch match. Use the same base_arch_from_target_id canonical form. Document, for the per-family follow-ups, that candidates must leave Capability.layouts empty and that lowering to AttentionRequest needs left_bound + 1 (hipDNN left_bound=W keeps W+1 keys). Co-Authored-By: Claude <noreply@anthropic.com>
None of library/tests/sdpa ran in CI: neither the request-contract and oracle tests nor develop's forward/backward SDPA case-table tests. Add rocke_library_sdpa_pytest, pointed at the directory like rocke_library_attention_pytest, so a new file there runs without a CMake edit. Every file in it is host-only (torch is optional), and the directory runs in about 20 s. rocke_sdpa_reference_unit_pytest is not reused: the foreach over (sdpa conv) already defines it for test_sdpa_reference_contract.py, so a second add_test with that name would be a duplicate. Co-Authored-By: Claude <noreply@anthropic.com>
…, case tables and LSE hook The E/F change types named the oracle in references.py, one shared case table, and an epilogue_fn hook. The request oracle is now an adapter over numeric/sdpa_reference.py, the request table sits beside develop's kernel coverage table, and LSE emission goes through AttnFwdExt.epilogue_hook with the _lse_store factories. Also state that from_hipdnn rejects keys hipDNN does not write. Co-Authored-By: Claude <noreply@anthropic.com>
…sdpa/st0-contract
|
Automated overlap notice (rocKE merge monitor) Hi @ThruptiRajLakshmanaGowda, this PR and #14929 each add the same bounds fix to The two copies, here at
Both clamp to the same row and page wherever the kernel reads, so dropping yours loses nothing. Your no-key-row O = 0 ( Owner: #14929, the bounds fix itself. Neither PR puts the clamps in a shared home, so the PR that lands first owns them, and #14929 is approved and first in the merge plan (cgmillette's priority). Resolution: this PR drops its two clamps, keeps its no-key-row O = 0, and re-records Suggested order: #14929 first, then this PR: the merge plan (develop Rebase, after #14929 lands (checked with
Resolved this way, the shard re-records at all four flavors, and If this PR lands first (not the plan's order), take #14929's text now in place of your clamps: its If you think this PR should own it instead, reply here and on #14929, and the reviews will follow the new decision. Automated overlap check: one owner per shared piece of code, coordinated on the PRs. Nothing else changed. |
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13115 — route SDPA to unified and dense attention kernels (ThruptiRajLakshmanaGowda, AICK-2613, AICK-2614) — 🟠 Follow-up: fixes hold; 14929 rebase pending
Utility: 🟢 Positive — routes hipDNN SDPA requests to unified and dense attention kernels. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: delta since |
…of truncating; test sdpa __all__ Co-Authored-By: Claude <noreply@anthropic.com>
|
Follow-up on the 11-04 23:10Z review, pushed as 33013ea:
|
…key-row O = 0, re-record attention_d128_swa Co-Authored-By: Claude <noreply@anthropic.com>
|
#14929 landed, so I merged develop (head below) and resolved N13 as agreed:
I haven't re-run the gfx942 / gfx950 GPU suites on the merged head yet. |
|
GPU re-run on the merged head 6daf7e6 (single process each:
|
yraparti
left a comment
There was a problem hiding this comment.
rocKE checklist re-review of the fixes since d19c06a6d342
ROCm/rocm-libraries#13115 · 6daf7e6b96b0 · model gemini-4-argon, fallback gpt-5.5
PASS · graded 238/238 · BLOCKER 0, HIGH 0, MEDIUM 0, LOW 0, NEEDS-INFO 25
The diff was truncated before grading. Items that needed a missing hunk are NEEDS-INFO.
Since the last review (d19c06a6d342)
Resolved 1 · still open 0 · new 0
- ✅ §S.4 was FAIL, now PASS — had asked: Add a CI-run test that iterates
dispatch.attention.__all__and asserts each listed name resolves from the module, e.g. viagetattr(dispatch.attention, name)or star-import coverage.
Findings
- No FAIL rows.
Needs info
- §A.1 Unresolved since previous review:
dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6017and:6372modify the gfx942 4-warp attention builders (_build_gfx942_4warp_gqa_leanandbuild_gfx942_4warp_gqa) without a C++ update in the PR diff, andknown_gaps.mdandcpp/instances/are not in the diff to verify whether the 4-warp GQA body has a C++ twin or is Python-only. - §C.9 In
dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:5862-5865,6017,6067,6159-6161,6372-6394,6576, and6749-6751, the Q-row/key clamping andno_keyzeroing fixes were applied across both the 4-warp lean and 4-warp GQA bodies. — no diff line supports this verdict - §E.9 New post-review diff content includes changes outside the SDPA/rocKE scope, such as projects/composablekernel/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp:1824 and projects/hipblaslt/clients/common/include/testing_matmul.hpp:1238. — no diff line supports this verdict
- §E.10 Not resolved since the previous review: the diff still does not include the requested merge-simulation record listing open PRs that touch the same helpers/tests/goldens, agreed land order, and affected-suite reruns on current develop plus those PRs.
- §E.14 New post-review changes include omitted deletion hunks, so the diff provided is insufficient to verify that removed upstream code, comments, or KNOWN LIMITATION notes were intentionally ported after the rebase/port.
- §F.6 Still unresolved since previous review: local
pre-commitexecution status or CI lint check output is not visible in the diff. - §F.7 Still unresolved since previous review: GitHub review-thread state and ticket IDs for any deferred review comments are not included in the provided PR context.
- §G.3 In the new commits,
attention_d128_swa.jsonis updated across all four LLVM flavors (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8) andattention_d256.jsonis updated (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8), but the tail ofattention_d256.jsonis still truncated in the diff and local non-GPU gate sweep execution results cannot be verified from the diff alone.
Section notes
- §D: Re-reviewed items 1-11 against current head 6daf7e6. No new issues introduced by recent commits.
- §E: Re-review of Structure & hygiene against the current diff.
- §G: Re-reviewed items 1-17 of Section G against head 6daf7e6.
- §H: Trace of device-side masks/bounds: (1) Q row load in
_build_gfx942_4warp_gqa_leanandbuild_gfx942_4warp_gqavisits aTOKBLK(128-row) compile-time Q block; extent comes from runtimecu_seqlens_q(cu_q_stop,qlen); indices are clamped via_clamp_to_last(..., q_last)(q_last = cu_q_stop - 1,attention_tiled_2d.py:6067,6373,6382,6388) and O stores are guarded byq_inseq < qlen(attention_tiled_2d.py:6162,6752) — agree. (2) KV page/key lookup inphys_keyvisitsBN-key tiles (BN=32or64); extent comes from runtimeseqused_k(klen); lean clampskeytok_last(attention_tiled_2d.py:6017) and 4-warp GQA clampspagetok_last_page = (klen - 1) >> BS_LOG2(attention_tiled_2d.py:6372,6393), while trailing rows of the last page are zero-initialized ink_cache/v_cache(bind_unified.py:352-358) and rows with no visible keys (context_off + q_inseq < 0) have O zeroed (attention_tiled_2d.py:6159-6160,6749-6750), though the in-loop score… - §J: The diff includes Markdown changes in
dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md. Shapes | Formula / Op | Operands | Inner dims | Result vs Declared | |---|---|---|---| | None | — | — | No tensor equations inKERNEL_AUTHORING.md| Symbols | Symbol | Defined at | Domain | |---|---|---| |SDPA_<FAMILY>_FEATURES|KERNEL_AUTHORING.md:609| Set of carried SDPA feature names per family | |UNIFIED_ALGORITHMS|KERNEL_AUTHORING.md:639(dispatch/attention/__init__.py:85) | Frozenset of unified route algorithm strings | |DENSE_ALGORITHMS|KERNEL_AUTHORING.md:641(dispatch/attention/candidate.py:51) | Frozenset of dense route algorithm strings | Novelty | Claim | Symbol / Path | Status | |---|---|---| | Routed family SDPA package |library/dispatch/sdpa/(bind_unified.py,bind_dense.py,bind_shared.py) | Added in this PR | | Exported route algorithm sets |dispatch.attention.UNIFIED_ALGORITHMS… - §K: Identity table: no knob axes, declared defaults, canonicalization rules, or knob-driven builder semantics are added or modified in this PR (
unified_rules.pyonly renames_problemtoattention_problem, and theattention_tiled_2d.pyedits fix unconditional bounds clamping in the 4-warp GQA emitter regardless of knobs);config_keyandTUNING_ID_VERSIONare unchanged.
Graded with the rocKE PR checklist, one model pass per section. FAIL rows that sit on a diff line are inline comments.
§A graded 14/14
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | Unresolved since previous review: dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6017 and :6372 modify the gfx942 4-warp attention builders (_build_gfx942_4warp_gqa_lean and build_gfx942_4warp_gqa) without a C++ update in the PR diff, and known_gaps.md and cpp/instances/ are not in the diff to verify whether the 4-warp GQA body has a C++ twin or is Python-only. |
| 2 | N/A | The PR does not modify any core lowering ops or IR op-handlers under core/lower_* or core/ir. |
| 3 | PASS | New dispatch and routing logic is confined to dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1 (plus public exports in… |
| 4 | N/A | No C++ kernel twin or C++ IR emitter code is modified in this PR. |
| 5 | PASS | Both dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8 re-bless only the… |
| 6 | N/A | No C++ twin is touched and the PR does not claim cross-engine parity from a golden update. |
| 7 | PASS | The bounds-clamping and no-key fixes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6017, :6067, :6159, :6372, and :6749 apply only to the 4warp_gqa builders, and the golden diffs in… |
| 8 | N/A | No rocKE C++ files are modified or added in this PR. |
| 9 | N/A | No shared op-handler or lowerer arch divergence fix is included in this PR. |
| 10 | N/A | No new kernel IR feature flag is introduced; the PR only adds host-side SDPA dispatch bindings and fixes bounds clamping in… |
| 11 | N/A | No new IR op is added in this PR. |
| 12 | N/A | core/backend.py and ROCKE_BACKEND handling are not touched in this PR. |
| 13 | N/A | LLVM .ll emission and debug-location handling are not modified in this PR. |
| 14 | N/A | No helper files under helpers/ or platform/cpp/helpers/ are modified in this PR. |
§B graded 6/6
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:174-189 classifies implementation-shortcut declines with NOT_YET_IMPLEMENTED: separately from ABI/kernel limits, and… |
| 2 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:211-214 probes each dense candidate with opt_in_probe, checks candidate.admits(probe), and builds the candidate spec only after admits succeeds… |
| 3 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:219-226 declines requests when no unified candidate covers them, and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:432-434 and… |
| 4 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/hipdnn.py:220-222 rejects a node missing its 'type' key instead of defaulting to 'SdpaAttributes', and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:243-248 and 260-265… |
| 5 | PASS | dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:49-57 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_numeric.py:43-51 run GPU numeric verification over every case admitted by sdpa_unified_admits… |
| 6 | N/A | The diff adds SDPA dispatch admission/binding logic and kernel bug fixes, without adding or modifying an is_valid_spec or… |
§C graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Shared predicates, decline reasons, and translation logic across unified and dense SDPA bindings are factored into dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:44-244, torch_dtype and device_tensor are consolidated in… |
| 2 | PASS | The SDPA bindings route to existing public entrypoints (bind_dense_attention_torch in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:342-344 and run_unified_attention_torch in… |
| 3 | N/A | No new kernel builder functions or kernel_name() definitions are introduced in this PR. |
| 4 | PASS | Arch-independent SDPA request/binding logic lives in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:1-326, while gfx942-specific kernel fixes stay in… |
| 5 | PASS | All touched files in the PR are directly scoped to SDPA dispatch/binding (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:1-152), SDPA tests (dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:869-871… |
| 6 | PASS | Shared SDPA request, band, decline-table, and operand abstractions are cleanly factored into dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:44-326, and single-source public attention-layer accessors (attention_band… |
| 7 | N/A | This PR is a feature routing and bug-fix change rather than a pure no-op cleanup. |
| 8 | PASS | The PR reuses existing develop helpers including opt_in_probe, empty_workspace, and resolve_stream (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:31-32), validate_lengths… |
| 9 | NEEDS-INFO | In dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:5862-5865, 6017, 6067, 6159-6161, 6372-6394, 6576, and 6749-6751, the Q-row/key clamping and no_key zeroing fixes were applied across both the 4-warp lean and 4-warp GQA bodies. — no diff line supports this verdict |
| 10 | PASS | Every feature in SDPA_FEATURES is explicitly partitioned into carried vs declined sets in both dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:84-130 and… |
§D graded 21/21
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR does not add or modify any KernelSpec knobs or environment variable reads. |
| 2 | PASS | Coverage limits and algorithm constants are imported from kernel/attention modules rather than transcribed (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:38 imports UNIFIED_ALGORITHMS and UNIFIED_HEAD_SIZES… |
| 3 | PASS | All candidate admission checks go through candidate.admits(probe) or ATTENTION_ROUTE_REGISTRY.supported(...) and never call _supports directly (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:212… |
| 4 | PASS | Selection and admission are pure functions of SdpaRequest taking req.arch directly without any device probes (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:72… |
| 5 | N/A | This PR does not introduce or modify any KernelSpec identity, spec_hash, compile_key, or kernel_name() methods. |
| 6 | N/A | This PR does not introduce or modify cohort routing, spec-builders, or kernel cache keys. |
| 7 | N/A | This PR does not add hand-pinned heuristic overrides. |
| 8 | N/A | This PR does not define or modify any kernel cohort predicates. |
| 9 | N/A | No new KernelCandidate or architecture is registered; the PR routes SdpaRequest onto existing ATTENTION_ROUTE_REGISTRY candidates. |
| 10 | PASS | Every SdpaRequest field and feature is translated into AttentionRequest and routed to the unified and dense bindings with dispatcher-level wiring, entry-point, and numeric tests (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:44… |
| 11 | PASS | Existing opt-in dense candidates remain opt-in and are queried explicitly via opt_in_probe(attn, candidate) in the SDPA dense gate (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:211). |
| 12 | PASS | Routing priority order is preserved and tested: dense SDPA asserts the first gfx950 variant in registry order is chosen at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_wiring.py:312, and unified SDPA asserts returned candidates are sorted… |
| 13 | PASS | Pinned SDPA requests are explicitly declined instead of falling back to auto at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:231, and tested for known and unknown pins at… |
| 14 | N/A | This PR does not modify compiled-binary cache keys, selection-key logging, or any candidate kernarg ABI version. |
| 15 | N/A | No new KernelCandidate or arch registration module is added; SDPA routes through existing attention candidates. |
| 16 | PASS | registered_sdpa_combos, sdpa_sweep_space, and dispatch_sdpa_all short-circuit on _request_errors(req) before delegating at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/init.py:119, with coverage in… |
| 17 | N/A | No new KernelCandidate is registered in this PR, so per-candidate CPU routing and family fixture invariants do not apply. |
| 18 | N/A | This PR does not change registry ratchets, route-only exemptions, or candidate binding requirements. |
| 19 | PASS | The unified SDPA binding accurately documents and uses run_unified_attention_torch with no tuning spec (backend="auto") at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:440, while dense SDPA launches through… |
| 20 | N/A | The selection/routing changes are confined to the Python library dispatch layer (library/dispatch/), which has no C++ twin. |
| 21 | PASS | Band and scale mismatches are rejected in the chooser/admission gates rather than hacking the emitter, at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:151 and… |
§E graded 17/17
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Under platform, the current diff adds installed-test wiring in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:867 and torch runtime helpers in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/torch_interop.py:96, with no platform… |
| 2 | PASS | Library SDPA binding code imports torch glue from rocke.runtime.torch_interop rather than importing torch directly, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:32 and… |
| 3 | PASS | Torch-dependent tests gate torch with pytest.importorskip, e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_binding.py:22, and runtime torch access stays behind helper functions such as… |
| 4 | PASS | New dispatch code imports downward to kernels/platform helpers, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:25 imports kernels.common._sdpa_seqlen and… |
| 5 | PASS | The platform installed-test entry uses the existing relative-test variable in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:867 and adds relative ignore paths in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:869. |
| 6 | PASS | Emitter-tree changes add explicit IR behavior rather than deleting unused handles, e.g. dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6154 and… |
| 7 | PASS | The routed-family process documentation moves with the SDPA routing change in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:588. |
| 8 | PASS | New docs use descriptive references such as the routed-family section in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:590 rather than code comments containing tracker IDs, reviewer handles, or PR numbers. |
| 9 | NEEDS-INFO | New post-review diff content includes changes outside the SDPA/rocKE scope, such as projects/composablekernel/include/ck_tile/ops/fused_moe/kernel/moe_sorting_kernel.hpp:1824 and projects/hipblaslt/clients/common/include/testing_matmul.hpp:1238. — no diff line supports this verdict |
| 10 | NEEDS-INFO | Not resolved since the previous review: the diff still does not include the requested merge-simulation record listing open PRs that touch the same helpers/tests/goldens, agreed land order, and affected-suite reruns on current develop plus those PRs. |
| 11 | PASS | The rocKE additions are deliverable docs, source, tests, and installed-test wiring, e.g. dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:588, dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1, and… |
| 12 | PASS | Both gfx942 and gfx950 are explicitly exercised by SDPA tests, e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_wiring.py:37 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_wiring.py:26. |
| 13 | PASS | The public rename away from private attention helpers updates definitions and callers in the same diff, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:753 defines attention_band and… |
| 14 | NEEDS-INFO | New post-review changes include omitted deletion hunks, so the diff provided is insufficient to verify that removed upstream code, comments, or KNOWN LIMITATION notes were intentionally ported after the rebase/port. |
| 15 | PASS | New helper parameters are used in their bodies, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:240 uses req_tensor/tensors/key and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:249 uses… |
| 16 | N/A | This item is a meta-instruction for reviewer grading rather than a property of the PR code. |
| 17 | PASS | no absolute throughput, bandwidth, or latency figure in added lines |
§F graded 11/11
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Branch users/tlakshma/hack/sdpa/st2-unified matches the recognized users/<user>/... prefix, and PR changes are in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:1. |
| 2 | PASS | title feat(hip-kernel-provider): route SDPA to unified and dense attention kernels matches type(scope): summary and is 76 characters |
| 3 | PASS | description is at least 30 characters and contains an issue reference |
| 4 | PASS | test file dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/attention/test_additive_registration.py is in the diff |
| 5 | PASS | All changed files in the diff (such as dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1) contain only code, docs, CMake, and test assets; no keys… |
| 6 | NEEDS-INFO | Still unresolved since previous review: local pre-commit execution status or CI lint check output is not visible in the diff. |
| 7 | NEEDS-INFO | Still unresolved since previous review: GitHub review-thread state and ticket IDs for any deferred review comments are not included in the provided PR context. |
| 8 | N/A | The PR targets Base: develop rather than a stacked feature branch, so the stacked-PR bot-skipping case does not apply. |
| 9 | N/A | There is no indication that the PR was artificially split or trimmed for a size cap (in fact, #14624 was folded into this PR). |
| 10 | PASS | The added and updated tests are substantive wiring, CPU data-movement, case-table, and GPU numeric tests rather than throwaway stubs, e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:82 and… |
| 11 | PASS | The PR metadata header in this prompt shows the updated title (feat(hip-kernel-provider): route SDPA to unified and dense attention kernels) and full multi-section body (covering dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1 and… |
§G graded 34/34
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Correctness is verified against an independent mathematical reference oracle rather than relying solely on IR byte-identity (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:61 and… |
| 2 | PASS | On-GPU numeric tests against the reference oracle are included and enrolled in the GPU test suite (dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1007 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:61). |
| 3 | NEEDS-INFO | In the new commits, attention_d128_swa.json is updated across all four LLVM flavors (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8) and attention_d256.json is updated (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8), but the tail of attention_d256.json is still… |
| 4 | PASS | New admission, binding, and length-validation rules have falsifiable unit tests checking exact decline prefixes, non-integer dtype rejection, data movement, and family entry points… |
| 5 | PASS | In test_an_out_of_range_block_table_physical_id_is_rejected_before_launch (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:83), the docstring explicitly documents why the test enters at validate_tuning_attention_tensors… |
| 6 | PASS | Prose claims in KERNEL_AUTHORING.md, docstrings, and code comments (including the 4-warp GQA clamping comments in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6363) match the adjacent code and map to concrete tests in… |
| 7 | PASS | KERNEL_AUTHORING.md specifies the reference oracle and four test layers as a first-class gate (dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:624), and the CPU and GPU binding tests validate against sdpa_reference… |
| 8 | PASS | Bindings and test harnesses perform strict dtype checks and loud-reject or skip unsupported dtypes/arches (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:267… |
| 9 | N/A | This PR does not introduce or claim any LDS size or occupancy numbers. |
| 10 | N/A | This PR is a functional routing and bug-fix change, not a performance optimization PR. |
| 11 | PASS | Kernel reachability through the SDPA routing gates and family entry points is verified in dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_wiring.py:480… |
| 12 | N/A | The changes in this PR are in Python-only dispatch layers and a gfx942 kernel builder… |
| 13 | N/A | No new LLVM intrinsic overload classes or mangled symbols are added or recommended. |
| 14 | N/A | No hardware bottleneck classification claims are made in this PR. |
| 15 | N/A | The PR diff is provided directly in the review prompt rather than inspected from a local git checkout. |
| 16 | PASS | Integration with existing attention dispatch symbols (ATTENTION_ROUTE_REGISTRY, UNIFIED_ALGORITHMS, DENSE_ALGORITHMS, attention_band, attention_problem, dense_scale_range, opt_in_probe) is verified by wiring and coverage tests… |
| 17 | N/A | No standalone CLI or benchmark scripts are shipped in this PR; new files are library modules and pytest suites enrolled in CMake. |
| 18 | N/A | This PR does not add statistical runtime guards or back-pressure detectors of the kind this item targets. |
| 19 | NEEDS-INFO | Fixtures include shapes/inputs that can express the named failures, including bad per-batch lengths in dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:139 and an out-of-range block-table physical id in dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:77. — no diff line supports this verdict |
| 20 | N/A | This PR does not add a claim-ban detector or regex sweep. |
| 21 | N/A | This PR does not introduce performance measurement or profiling deliverable code. |
| 22 | PASS | The SDPA numeric GPU tests are excluded from the host SDPA pytest shard but explicitly enrolled in the GPU CTest lane at dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1004 and dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1005, and… |
| 23 | N/A | This PR does not make speedup or routing A/B performance claims. |
| 24 | NEEDS-INFO | The diff shows re-blessed attention golden hashes, but the GPU validation hunk for the exact moved 4-warp GQA configurations is omitted; an untruncated diff or cited test lines proving coverage for each moved golden case are needed. |
| 25 | N/A | This PR is a functional routing and bounds/keyless-row correctness change, not a codegen micro-optimization report. |
| 26 | N/A | This PR does not make campaign-branch benchmark speedup claims. |
| 27 | N/A | This PR does not ship a performance knob or A/B toggle. |
| 28 | N/A | This PR does not present counterfactual performance tables. |
| 29 | N/A | This PR does not swap guarded math primitives for fast primitives. |
| 30 | N/A | This PR makes no register-spill or scratch-reduction claims. |
| 31 | N/A | This PR does not claim a performance win over a baseline defect. |
| 32 | N/A | The PR makes no timing, throughput, spill, or other measured performance-number claim requiring reproducible provenance. |
| 33 | N/A | The rocKE diff exposes and reuses attention routing helpers and adds SDPA gates, but does not modify an existing scalar… |
| 34 | PASS | The optional stats/LSE path is tested with the flag present on padded/keyless rows, asserting O is zero and stats is -inf at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:130 and… |
§H graded 5/5
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | sdpa_unified_bind allocates k_cache/v_cache sized to slots = batch * self.pages at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:352-358 and validates the launch tensors against plan.problem via… |
| 2 | PASS | _LaunchPlan builds an identity block_table in range(slots) at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:361-363, clamps logical page indices to k_last_page on device at… |
| 3 | NEEDS-INFO | Resolved since previous review: the new commits enforce that BS = spec.block_size is a power of two (if BS & (BS - 1): raise NotImplementedError(...)) and divides BN (if BN % BS: raise NotImplementedError(...)) in build_gfx942_4warp_gqa at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6203-6219, and compute k_last_page via b.lshr(..., b.const_i32(BS_LOG2)) at… |
| 4 | NEEDS-INFO | In build_gfx942_4warp_gqa, the new commits clamp page to k_last_page in phys_key at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6372 and 6391-6395 and zero O when context_off + q_inseq < 0 at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6159-6160 and 6749-6750, while bind_unified.py zeroes k_cache/v_cache at… |
| 5 | NEEDS-INFO | Resolved since previous review: seq_len_q and seq_len_kv are host-validated via _host_lengths and validate_lengths (preserving dtype to reject non-integers) at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:243-265 and dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_sdpa_seqlen.py:45, operand shapes/strides/ragged offsets are validated in Operand.__init__ at… |
§I graded 6/6
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR wires hipDNN SDPA forward requests to existing attention candidates and fixes gfx942 kernel edge cases without introducing… |
| 2 | PASS | Non-contiguous or strided layouts are repaired via copy buffers (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:276-284), unsupported features, bands, and non-power-of-two/non-divisible block sizes are loudly rejected… |
| 3 | PASS | Curated request and kernel case tables (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/cases.py:1361-1439) are combined with generated parameter sweeps and explicit algorithm-partition coverage checks… |
| 4 | N/A | This PR does not add or modify performance routing comparison tables or tuning reports. |
| 5 | N/A | This PR does not introduce baked tuning search results or static tuning snapshots. |
| 6 | N/A | This PR does not add any new DSL ops or hardware-instruction-named primitives. |
§J graded 7/7
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The only Markdown file in the PR (dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md) contains no mathematical… |
| 2 | N/A | The KERNEL_AUTHORING.md diff introduces no mathematical formula symbols, decay/log-decay domains, or log bases. |
| 3 | PASS | dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591-642 links to existing sections [E](#e-new-operator-with-a-hipdnn-graph-node), [F](#f-new-operand-or-output-on-an-existing-operator), and [A](#a-new-kernel-family) while stating only the delta for… |
| 4 | PASS | The novelty claims in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:321 and dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:638-641 name library/dispatch/sdpa/, bind_unified.py, bind_dense.py… |
| 5 | PASS | dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md makes no scaling/complexity claims, and the O(s_q) interval comparison claim in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:94-95 states the aggregate complexity across… |
| 6 | N/A | The KERNEL_AUTHORING.md diff does not introduce any model configuration tables or tuned band-edge constants. |
| 7 | N/A | This PR edits the existing dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md and does not add a new design or… |
§K graded 13/13
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR adds no multi-configuration candidate (SDPA_REGISTRY remains empty and SDPA delegates to existing attention candidates… |
| 2 | N/A | This PR does not add or modify any tunable kernel spec fields, KnobAxis declarations, or field classifications. |
| 3 | N/A | This PR does not change config_key payloads, canonical-knob rules, declared defaults, or knob meanings (unified_rules.py only… |
| 4 | N/A | This PR does not modify KnobSpace.defaults() or KnobSpace.recorded() in any KnobSpace subclass. |
| 5 | N/A | This PR does not touch knob drop/refuse rules, inert(), base_value(), or KNOWN_WRONG_KNOBS. |
| 6 | N/A | This PR does not touch KnobSpace.validate or library/dispatch/attention/tuning_specs.py. |
| 7 | N/A | This PR does not modify KnobSpace.base(req) or _tuning_problem (it only renames _problem to attention_problem in… |
| 8 | N/A | This PR does not modify normalize_knobs or request-boundary knob scalar coercion. |
| 9 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:231 (value_rejects) explicitly declines any pinned algorithm, spec_id, tuning_id, or tuning_knobs rather than adopting canonicalized ids or falling back to auto, tested in… |
| 10 | N/A | This PR does not modify KnobSpace.find or tuning_id resolution logic. |
| 11 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:119 (registered_sdpa_combos, sdpa_sweep_space, dispatch_sdpa_all) delegates to SDPA_REGISTRY and registered_attention_combos / dispatch_attention_all filtered by _serves(req)… |
| 12 | N/A | This PR does not reorder any production stack, change accept_default, or edit a tuning variant's base spec. |
| 13 | PASS | The benchmark changes in dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/attention_combo_sweep.py:44 and dnn-providers/hip-kernel-provider/rocke/library/benchmarks/gfx950/attention/decode/decode_dispatch_diagnose.py:13 only rename _problem to… |
§L graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This diff does not add or modify a rocKE gfx target in arch specs or ISA backend registries. |
| 2 | N/A | This diff does not add or modify an MMA op_id, MMA fragment-info row, or MMA lowering layout map. |
| 3 | N/A | This diff does not add or change MMA atom selection, MMA catalog fields, lane maps, or intrinsic-name emission. |
| 4 | N/A | This diff does not introduce or change wave-size handling or lane math derived from wave size. |
| 5 | PASS | The arch-specific codegen change is confined to the gfx942 4-warp kernel and masks fully masked rows without leaking gfx950 recipes at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6157. |
| 6 | N/A | This diff does not move pipeline vocabulary onto ArchTarget/core arch data, mutate ArchTarget from policies, or import CK Tile… |
| 7 | N/A | This diff only routes or declines SDPA dtype features and does not change rocKE logical dtype, bit-packing, scalar lowering, or… |
| 8 | N/A | This diff does not touch scaled-MMA scale contracts, retired scaled WMMA ids, or MfmaAtom.emit_scaled lowering. |
| 9 | N/A | This diff does not add a new rocKE architecture target. |
| 10 | N/A | This diff does not touch intrinsic_arch_domain..json, gen_arch_domain.py, or LLVM_FLAVORS. |
§M graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:5865, :6159, and :6749 use b.cmp_lt and b.select for device-time predicates rather than Python control flow. — no diff line supports this verdict |
| 2 | N/A | The PR diff does not touch legalize_for_target or target legalization logic. |
| 3 | N/A | The PR diff does not change print_ir, ck.dsl.ir/v1 serialization, wire versions, attrs, SSA ids, or parameter ordering. |
| 4 | N/A | The PR diff does not change portable IR schemas, importers, magic division constants, recipe VM code, or recorder coverage. |
| 5 | N/A | The PR diff does not touch AsyncTileLoader, async DRAM-to-LDS loading, LDS async layouts, or VMEM waits. |
| 6 | NEEDS-INFO | dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6017, :6067, :6382, and :6393 clamp the Q row, key, and KV page indices before loading, avoiding loads from out-of-bounds addresses. — no diff line supports this verdict |
| 7 | N/A | The PR diff does not add or modify buffer-resource word 3 handling or TemporalHint.STREAMING lowering. |
| 8 | N/A | The PR diff does not add LLVM intrinsics or change backend selection precedence. |
| 9 | N/A | The PR diff does not touch public C entry points, cpp/include/rocke/ headers, or C++ codegen launch behavior. |
| 10 | N/A | The PR diff does not emit !dbg metadata or change debug-info consumers. |
§N graded 18/18
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The diff does not add or promote a helper under platform/python/rocke/helpers/. |
| 2 | NEEDS-INFO | dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6576 uses the shared q_desc.offset TensorView mechanism for Q addressing after the clamped row/head selection, rather than hand-expanded Q offsets. — no diff line supports this verdict |
| 3 | N/A | The diff does not touch SoftwarePipeline, sync_before_issue, or overlap_vmcnt behavior. |
| 4 | N/A | The diff does not add or modify fused kernels, EpilogueOp, fusion schedules, or fusion validation paths. |
| 5 | N/A | The diff does not add or modify a public API under rocke.helpers.tiling. |
| 6 | N/A | The diff does not add or modify front-door/toolbox tiling helpers or LayoutStyle classes. |
| 7 | N/A | The diff does not touch tiling layout overrides, accumulator c_desc derivation, or MMA soundness gates. |
| 8 | N/A | The diff does not touch helpers/tiling LDS allocation, lds_store/lds_read, contraction padding, or wide-load tiling emit rules. |
| 9 | N/A | The diff does not use or modify N-D tiling axis roles, at_index/squeeze behavior, FlowStage labeling, or AB-swap routing. |
| 10 | N/A | The diff does not modify tiling emit, LDS staging, or MMA issue code in helpers/tiling. |
| 11 | N/A | The diff does not touch tiling MMA traits, TileMma target selection, lds_swizzle policy, or derive_sched_group_counts. |
| 12 | N/A | The PR diff makes no helpers/tiling LDS bank-conflict or coalescing optimization claim requiring simulator, counter, or ASM-backed… |
| 13 | N/A | The diff does not construct WarpDistributionEncoding, change make_tile_desc factoring, re-export RegisterMapper, or alter tiling… |
| 14 | N/A | The diff does not use or modify helpers/tiling verbs or add a new helpers/tiling LDS pipeline variant. |
| 15 | N/A | The diff does not touch helpers/tiling descriptor allocation or transformed-range bounds checks after swizzle, pad, or XOR. |
| 16 | N/A | The diff does not use or modify optimize_layout, recommend(), or evaluate_transform ranking behavior. |
| 17 | N/A | The diff does not modify tiling_recorder.py, round-trip analysis, or tiling visualization emission behavior. |
| 18 | N/A | The diff does not touch InterleavedStyle, helpers/tiling/kernels demos, or tiling numeric gates for interleaved layouts. |
§O graded 16/16
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The PR does not add or modify IR address spaces, required attributes, or core/verify.py. |
| 2 | N/A | The PR does not modify scf.for structure or terminator emission in core IR/lowerers. |
| 3 | N/A | The PR does not touch IRBuilder, intrinsic declaration tables, or metadata slot numbering. |
| 4 | N/A | The PR does not modify IR types or SmemType serialization/parsing. |
| 5 | N/A | The PR does not modify op purity tables or core/passes.py / cpp/core/passes.cpp. |
| 6 | N/A | The PR does not modify constant folding in Python or C++ passes. |
| 7 | N/A | The PR does not modify kernel attributes, codegen policy, or cluster dimension checks. |
| 8 | N/A | The PR does not set or modify scheduler strategies or CodegenPolicy. |
| 9 | N/A | The PR does not modify waitcnt encoding or barrier lowering. |
| 10 | N/A | The PR does not modify lowerer attribute validation or pointer parameter lowering. |
| 11 | PASS | dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_numeric.py:30 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:36 resolve the device architecture via… |
| 12 | N/A | The PR does not add hardcoded width/size caps in core/codegen or silent arch or "gfx950" defaults. |
| 13 | N/A | The PR does not touch any public C structs or enums under cpp/include/rocke/. |
| 14 | N/A | The PR does not touch C++ sources or add C++ compiler flags in platform/CMakeLists.txt. |
| 15 | N/A | The PR does not modify C++ engine fallback mechanisms or lowerer entry points. |
| 16 | N/A | The PR does not touch the C++ codegen library or gfx1250 builders/lowerers. |
§P graded 14/14
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The visible diff does not add or modify platform instance builder signatures or spec REQUIRED_FIELDS. |
| 2 | N/A | The visible diff does not touch platform kernarg signatures, manifest args_signature, runtime argument packing, or GEMM/conv ABI… |
| 3 | N/A | The visible diff does not touch GEMM, implicit-GEMM, or block-scaled GEMM launch geometry. |
| 4 | N/A | The visible diff does not modify LDS allocation gates or C++ validators. |
| 5 | N/A | The visible diff does not touch WMMA/RDNA GEMM, GemmMultiD, or GemmMultiAbd admission paths. |
| 6 | N/A | The visible diff does not touch Stream-K, MoE sorting, fused MoE, or split-K caller-zeroed workspaces in rocKE. |
| 7 | N/A | The visible diff does not touch Stream-K rejection rules, persistent GEMM, GroupedGemmLauncher, or SUPPORT_MATRIX.md. |
| 8 | N/A | The visible diff does not touch quantization helpers, MX GEMM, or MX scale decoding. |
| 9 | N/A | The visible diff does not add or modify numeric emulation or low-precision decode paths requiring edge-case emulation tests. |
| 10 | N/A | The visible diff does not add a manifest kind or touch helpers/manifest.py, run_manifest, or manifest allocate/verify consumers. |
| 11 | PASS | The new SDPA bindings allocate and retain launch buffers once in the bound plan and launch on resolve_stream: see dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:341-352 and… |
| 12 | N/A | The visible diff does not touch COMGR or HIP library loading logic. |
| 13 | N/A | The visible diff adds SDPA routing and binding logic rather than refactoring shared host-side grid, tile-count, workspace… |
| 14 | N/A | The visible diff does not touch examples/, helpers/ placement, or add new modules under platform instances/common family glue. |
§Q graded 7/7
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | Visible SDPA tests are registered in installed CTest entries at dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:867 and GPU numeric entries at dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1004, but the current-head hunks for additional new tests are omitted and the diff does not show ROCKE_ENGINE_test_categories_external.yaml entries or the CI lane that runs them. |
| 2 | N/A | No new kernel family or variant is added; the PR routes SDPA to existing attention candidates and edits existing gfx942 attention… |
| 3 | PASS | The representative golden tripwire is updated for the changed attention kernels at dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:12 and… |
| 4 | PASS | The affected representative goldens are updated across repeated LLVM-flavor blocks, including dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and… |
| 5 | NEEDS-INFO | Resolved for the visible SDPA tests: GPU numeric tests derive the active architecture and skip when unavailable at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:26 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_numeric.py:24, and torch-free length validation remains covered at… |
| 6 | N/A | This PR does not refactor a platform validation gate or differential harness with per-configuration timeout, hash fallback, or… |
| 7 | N/A | This PR does not change ROCKE_BACKEND, ROCKE_CPP_STRICT, binding-present fallback behavior, or engine rejection contracts. |
§R graded 11/11
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The diff routes SDPA over existing attention candidates and changes correctness/address guards in an existing gfx942 kernel, but… |
| 2 | NEEDS-INFO | The changed KV page clamp and Q-row hoist in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6364 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6374 need a runbook §6/§8 mapping, matching §13 attention checklist entry, and gfx942 arch recipe/ArchTarget rationale; those are not provided in the diff. |
| 3 | NEEDS-INFO | The global-load address changes at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6063 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6364 need utilities-skill evidence for prefetch/data-load or ISA inspection, especially because dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6368 makes a measured… |
| 4 | NEEDS-INFO | The intended primitive/address-computation change is stated at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6374, but the diff does not provide before/after probe_occupancy.py, probe_regs.py, or probe_intrinsic_counts.py/probe_isa_inspect.py evidence for the claimed movement. |
| 5 | NEEDS-INFO | Changes since the previous review do not resolve the prior NEEDS-INFO: the current diff still changes the Q load address at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6063 and the block-table/KV page address at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6364, but no… |
| 6 | N/A | The rocKE diff does not change an LDS layout, MMA layout, tiling pipeline identity/relabel/reorder, or bank-conflict lever. |
| 7 | N/A | The rocKE diff makes no bottleneck claim such as memory-bound, LDS-bound, barrier-bound, or latency-bound. |
| 8 | N/A | The rocKE diff does not add, remove, or modify waits, barriers, sched_barrier, sched_group_barrier, iglp_opt, or scheduling hints. |
| 9 | N/A | The rocKE diff is not presented as a performance-regression fix and does not add a shape-trading performance heuristic or dispatch… |
| 10 | NEEDS-INFO | The new comments make measured-performance claims at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6368 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6374, but the diff does not include runbook §2.3 benchmark-hygiene evidence for those claims. |
| 11 | N/A | The documentation added is an SDPA routed-family authoring checklist, not a new optimization failure mode, heuristic, case study… |
§S graded 8/8
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | The review-response change resolves part of the prior gap by adding an all public-surface test in dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/attention/test_attention.py:250 and SDPA host/GPU lane entries in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:867 and dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1004, but the visible diff still does not show the… |
| 2 | PASS | Existing behavior changes keep updated tests: missing hipDNN node type now raises in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/hipdnn.py:217 and is covered by dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_sdpa_from_hipdnn.py:454… |
| 3 | N/A | The visible rocKE diff does not remove or rename a public API; underscore-prefixed _problem and _validate_lengths migrations are… |
| 4 | PASS | The prior FAIL is resolved: dispatch.attention.all is edited in dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/init.py:561, and the new public-surface test imports and resolves every listed name in… |
| 5 | PASS | API tests assert behavior and negative contracts: torch_dtype resolves names and rejects an invalid dtype in dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:7658, device_tensor verifies values/dtype/device forwarding in… |
| 6 | PASS | CPU-checkable SDPA contracts are kept in the host SDPA pytest lane while GPU numeric files are explicitly ignored there in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:867, GPU numeric SDPA tests are separately added to the GPU pytest file set in… |
| 7 | N/A | The added and changed public APIs in this section are Python dispatch/runtime helpers or SDPA binding-layer APIs without a C++… |
| 8 | NEEDS-INFO | The prior NEEDS-INFO is not resolved: the visible diff does not contain a PR test-run statement showing the README entry points from rocke/platform were run, nor GPU result arch names for the new GPU SDPA coverage. |
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13115 — route SDPA to unified and dense attention kernels (ThruptiRajLakshmanaGowda, AICK-2613, AICK-2614) — 🟡 Follow-up: 14929 merge holds; nits remain
Utility: 🟢 Positive — routes hipDNN SDPA to unified and dense kernels; reuses 14929's clamps. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: delta since |
There was a problem hiding this comment.
🟡 Changes recommended
Critical SDPA binding defects break request semantics, ignore caller streams, and pass misaligned pointers to kernels, risking incorrect results and data races.
7 open issues
Preserve the SDPA contract in returned dispatch results · New Honor the requested stream for the full launch sequence · New Reject non-integer ragged offsets · New Route misaligned Q/O views through the workspace · New Reject non-positive scales or apply masking after scaling · New Check pointer alignment before using Q/O in place · New Honor the requested stream instead of discarding launch arguments · New
2 resolved since last review
Review effort: Max
…no-key-row O = 0, re-record attention_d128_swa and attention_d256 strided_decode: develop's _tiled_3d_signatures(arch, dtype, strided_kv=True) with this PR's attention_problem. gfx942 4-warp GQA bodies: develop's runtime num_query_heads (hq_p) in the O index, this PR's no-key-row O = 0 select. Golden shards re-blessed with the IR parity harness (llvm20/22/23/24). Co-Authored-By: Claude <noreply@anthropic.com>
…ic API Adds to the CI-run rocke_library_attention_pytest lane (tests/dispatch/ attention): __all__ must equal the listed public surface, holds no underscored name or submodule, and every exported *_ALGORITHM(S) names a registered candidate. Co-Authored-By: Claude <noreply@anthropic.com>
…st's offset dtype Operand cast ragged offsets with int(), so float32 offsets [0.5, 4.5, 8.5] bound like int32 [0, 4, 8]. It now refuses an offsets tensor whose dtype is not desc.ragged_offset.dtype (int32/int64) before reading it. Tested with float32 and int64 offsets against an int32 ragged_offset; both bound before. Co-Authored-By: Claude <noreply@anthropic.com>
…ied_bind The previous test called _lengths directly and passed with the old int() read. The new one binds a float32 seq_len_q of [1.5, 1.0] (truncates to the valid [1, 1]) through sdpa_unified_bind and asserts the integer-dtype refusal; it fails with the old read. The _lengths copy in the wiring tests is dropped. Co-Authored-By: Claude <noreply@anthropic.com>
|
@cgmillette thanks for the follow-up review. Addressed in head
|
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13115 — route SDPA to unified and dense attention kernels (ThruptiRajLakshmanaGowda, AICK-2613, AICK-2614) — 🔴 Follow-up: under-aligned Q/O give wrong O
Utility: 🟢 Positive — routes hipDNN SDPA to unified and dense kernels; fix alignment first. Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: delta since |
… reject foreign stream Copilot review (2026-11-05) findings on the unified/dense SDPA bindings: - Operand.packed_view (bind_shared.py) checked contiguity but not pointer alignment, so a contiguous Q/O view starting at an unaligned element could be handed in place to the gfx942/gfx950 tiled 2D kernels, which declare query_ptr/output_ptr with align=16. Now falls back to the workspace copy when the view's data_ptr() is not 16-byte aligned, matching _DensePlan's existing BSHD check. - sdpa_unified_admits (bind_unified.py) admitted non-positive and non-finite scales; the gfx942/gfx950 tiled 2D bodies mask with -1e30 before scaling in _online_softmax_rescale, so scale <= 0 lets masked keys regain or dominate softmax weight. Declined at admission with a new _scale_reject, mirroring dense's _scale_reject. - sdpa_unified_bind's and sdpa_dense_bind's launch() silently discarded a caller-supplied stream: pack()/unpack() always run on the ambient current stream while the bound workspace buffers are allocated once, at bind time, against whatever stream was then current, so honoring the kernel launch alone on a different stream would race on those buffers. launch() now rejects a non-current stream with NotImplementedError instead of silently ignoring it. Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
|
Follow-up on the 11-05 22:49Z review (Copilot's 6 remaining threads + N16), pushed as
All of Left open, by design, for your call rather than guessed at:
|
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max PR #13115 — route SDPA to unified and dense attention kernels (ThruptiRajLakshmanaGowda, AICK-2613, AICK-2614) — 🟠 Follow-up: stream refusal, CI test gaps
Utility: 🟢 Positive — three defects fixed; stream check and CI tests need one more pass Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: delta since |
|
I reviewed this PR again. The following review threads are still unresolved, or comments have no author reply:
For reassessment, the PR author should reply to top-level feedback or state that thread feedback is addressed. Before approval, top-level comments need explicit resolution from their author or that author's approval of the PR. |
yraparti
left a comment
There was a problem hiding this comment.
rocKE checklist re-review of the fixes since 6daf7e6b96b0
ROCm/rocm-libraries#13115 · 0d916301ba32 · model gemini-4-argon, fallback gpt-5.5
PASS · graded 238/238 · BLOCKER 0, HIGH 0, MEDIUM 0, LOW 0, NEEDS-INFO 12
The diff was truncated before grading. Items that needed a missing hunk are NEEDS-INFO.
Since the last review (6daf7e6b96b0)
Resolved 0 · still open 0 · new 0
Findings
- No FAIL rows.
Needs info
- §A.1 Unresolved since previous review:
dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6185and:6784modify_build_gfx942_4warp_gqa_leanandbuild_gfx942_4warp_gqawithout a C++ update in the PR diff, andknown_gaps.mdandcpp/instances/gfx942/are not in the PR diff to verify whether the 4-warp GQA body has a C++ twin or is documented as Python-only. - §D.14 The kernarg layout change for tiled attention (head counts and sliding window becoming kernargs) bumps ATTENTION_ABI_VERSION to v2 at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:80 and strided_decode's abi_version to v2 at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/strided_decode.py:222. — no diff line supports this verdict
- §E.10 Not resolved since the previous review: the PR description and diff (which touches shared files such as dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:894 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8) still do not provide a merge-simulation record listing open PRs touching the same tables/tests/goldens and suite reruns.
- §F.6 Still unresolved since previous review: local
pre-commitexecution status or CI lint check output is not visible in the diff. - §F.7 Still unresolved since previous review: Although the new commits add review-response fixes (e.g. stream rejection in
dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:364, 16-byte alignment check indnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:340, and positive scale check indnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:166), the GitHub review-thread list and any deferred-item ticket IDs are not included in the provided PR context. - §G.3 In the updated head,
attention_d128_swa.jsonis updated across all four LLVM flavors (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8) andattention_d256.jsonis updated (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8), but the tail ofattention_d256.jsonat line 98 is still truncated in the diff and local non-GPU gate sweep execution results cannot be verified from the diff alone. - §G.19 The new commits and current diff provide explicit fixtures and inline checks for the named failures: misaligned contiguous BSHD views at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:122 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_binding.py:150, bad per-batch sequence lengths at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:168 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_sdpa_seqlen.py:172, non-int32 ragged offsets at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:241, and an out-of-range block-table physical id…
- §G.24 The diff updates golden hashes in dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8 (including
attention_d256/gfx942/4warp_gqa_d192at line 16), while the SDPA unified gate admitsUNIFIED_HEAD_SIZES(64, 128, 256); confirmation of GPU validation coverage for every moved golden configuration (especially D192 4warp_gqa) is needed.
Section notes
- §A: Re-evaluated against current head 0d91630. Item 1 remains NEEDS-INFO because the PR modifies
_build_gfx942_4warp_gqa_leanandbuild_gfx942_4warp_gqainattention_tiled_2d.pywithout touching C++ orknown_gaps.mdin the PR diff. - §B: Re-reviewed against head 0d91630. New commits further strengthen trust-boundary rejections (ragged offset dtype check, scale positivity check, non-current stream rejection, and 16-byte alignment check on packed views).
- §E: Re-review of head 0d91630. The full PR diff is now clean of non-rocKE project changes (resolving item 9) and updated golden hashes are present for the gfx942 4-warp fix (resolving item 14).
- §G: Re-review of items 18-34 against head 0d91630.
- §H: Trace of device-side bounds/masks touched in this PR: 1.
Oepilogue store inbuild_gfx942_4warp_gqa(attention_tiled_2d.py:6187-6191,6786-6790): extent source is runtimeqlen(andcontext_off + q_inseq >= 0for rows with no visible key); visited range isq_inseq = qbase + wq + cacross the CTA's Q tile; guarded byscf_if(q_inseq < qlen)andselect(context_off + q_inseq < 0, 0.0, ...), agreeing with the valid row range. 2.k_cache/v_cacheandblock_tableinbind_unified.py:376-389: extent source isslots = batch * pageswherepages = max(1, ceil(s_kv / block)),cu_seqlens_qfromq_len, andseqused_kfromkv_len; validated before launch viavalidate_tuning_attention_tensors(bind_unified.py:466-467). - §J: Shapes: | Equation / Op | Operands | Inner Dims | Result vs Declared | |---|---|---|---| | (none in
KERNEL_AUTHORING.md) | — | — | — | Symbols: | Symbol | Defined at | Domain | |---|---|---| |AttentionRequest|dispatch/attention/common.py| Attention dispatch request | |SdpaRequest|dispatch/sdpa/common.py| hipDNN SDPA request contract | |ATTENTION_ROUTE_REGISTRY|dispatch/attention/__init__.py| Route candidate registry | |UNIFIED_ALGORITHMS|dispatch/attention/__init__.py:85|frozensetof unified algorithm names | |DENSE_ALGORITHMS|dispatch/attention/candidate.py:51|frozensetof dense algorithm names | |SDPA_<FAMILY>_FEATURES|dispatch/sdpa/bind_{unified,dense}.py|frozensetof carried SDPA feature tags | Novelty: | Claim | Symbol / Path | Status | |---|---|---| | Routed SDPA family & checklist H |library/dispatch/sdpa/,bind_unified.py,bind_dense.py| Added in this PR | | Unified route algorithm set |…
Graded with the rocKE PR checklist, one model pass per section. FAIL rows that sit on a diff line are inline comments.
§A graded 14/14
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | Unresolved since previous review: dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6185 and :6784 modify _build_gfx942_4warp_gqa_lean and build_gfx942_4warp_gqa without a C++ update in the PR diff, and known_gaps.md and cpp/instances/gfx942/ are not in the PR diff to verify whether the 4-warp GQA body has a C++ twin or is documented as Python-only. |
| 2 | N/A | The PR does not modify any core lowering ops or IR op-handlers under core/lower_* or core/ir. |
| 3 | PASS | New dispatch and routing logic is confined to dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1 (plus public exports in… |
| 4 | N/A | No C++ kernel twin or C++ IR emitter code is modified in this PR. |
| 5 | PASS | Both dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8 re-bless only the… |
| 6 | N/A | No C++ twin is touched in the PR diff and the PR does not claim cross-engine parity from a golden update. |
| 7 | PASS | The no-key output zeroing fix in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6185 and :6784 applies only to the 4warp_gqa builders, and the golden diffs in… |
| 8 | N/A | No rocKE C++ files are modified or added in the PR diff. |
| 9 | N/A | No shared op-handler or lowerer arch divergence fix is included in this PR. |
| 10 | N/A | No new kernel IR feature flag is introduced in this PR; it adds host-side SDPA dispatch bindings and fixes a bug in existing… |
| 11 | N/A | No new IR op is added in this PR. |
| 12 | N/A | core/backend.py and ROCKE_BACKEND handling are not touched in the PR diff. |
| 13 | N/A | LLVM .ll emission and debug-location handling are not modified in this PR. |
| 14 | N/A | No helper files under helpers/ or platform/cpp/helpers/ are modified in the PR diff. |
§B graded 6/6
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:176-191 classifies implementation-shortcut declines with NOT_YET_IMPLEMENTED: separately from ABI/kernel limits, and… |
| 2 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:211-214 probes each dense candidate with opt_in_probe, checks candidate.admits(probe), and builds the candidate spec only after admits succeeds… |
| 3 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:240-250 declines requests when no unified candidate covers them, and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:459-461 and… |
| 4 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/hipdnn.py:220-222 rejects a node missing its 'type' key instead of defaulting; dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:292-297 rejects ragged offsets whose dtype… |
| 5 | PASS | dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:49-57 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_numeric.py:43-51 run GPU numeric verification over every case admitted by sdpa_unified_admits… |
| 6 | N/A | The diff adds SDPA dispatch admission/binding logic and kernel bug fixes without adding or modifying an is_valid_spec or… |
§C graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Shared predicates, decline reasons, and translation logic across unified and dense SDPA bindings are factored into dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:46-246, and torch_dtype / device_tensor are consolidated in… |
| 2 | PASS | The SDPA bindings route to existing public entrypoints (bind_dense_attention_torch in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:345-347 and run_unified_attention_torch in… |
| 3 | N/A | No new kernel builder functions or kernel_name() definitions are introduced in this PR. |
| 4 | PASS | Arch-independent SDPA request/binding logic lives in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:1-348, while gfx942-specific kernel fixes stay in… |
| 5 | PASS | All touched files in the PR diff are directly scoped to SDPA dispatch/binding (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:1-152), SDPA tests (dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:894-895… |
| 6 | PASS | Shared SDPA request, band, decline-table, and operand abstractions are cleanly factored into dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:46-348, and single-source public attention-layer accessors (attention_band… |
| 7 | N/A | This PR is a feature routing and bug-fix change rather than a pure no-op cleanup. |
| 8 | PASS | After rebasing onto develop (which included #14929's clamp fixes), the PR reuses existing helpers including opt_in_probe, empty_workspace, and resolve_stream (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:31-32)… |
| 9 | PASS | The previous NEEDS-INFO is resolved in the current diff: the no_key zeroing fix for rows where qlen > klen in bottom-right causal mode is applied across both the 4-warp lean and 4-warp GQA bodies in… |
| 10 | PASS | Every feature in SDPA_FEATURES is explicitly partitioned into carried vs declined sets in both dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:85-131 and… |
§D graded 21/21
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR does not add or modify any KernelSpec knobs or environment variable reads. |
| 2 | PASS | Coverage limits and algorithm constants are imported from kernel/attention modules rather than transcribed (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:39 imports UNIFIED_ALGORITHMS and UNIFIED_HEAD_SIZES… |
| 3 | PASS | All candidate admission checks go through candidate.admits(probe) or ATTENTION_ROUTE_REGISTRY.supported(...) and never call _supports directly (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:212… |
| 4 | PASS | Selection and admission are pure functions of SdpaRequest taking req.arch directly without any device probes (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:74… |
| 5 | N/A | This PR does not introduce or modify any KernelSpec identity, spec_hash, compile_key, or kernel_name() methods. |
| 6 | N/A | This PR does not introduce or modify cohort routing, spec-builders, or kernel cache keys. |
| 7 | N/A | This PR does not add hand-pinned heuristic overrides. |
| 8 | N/A | This PR does not define or modify any kernel cohort predicates. |
| 9 | N/A | No new KernelCandidate or architecture is registered; the PR routes SdpaRequest onto existing ATTENTION_ROUTE_REGISTRY candidates. |
| 10 | PASS | Every SdpaRequest field and feature is translated into AttentionRequest and routed to the unified and dense bindings with dispatcher-level wiring, entry-point, and numeric tests (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:46… |
| 11 | PASS | Existing opt-in dense candidates remain opt-in and are queried explicitly via opt_in_probe(attn, candidate) in the SDPA dense gate (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:211). |
| 12 | PASS | Routing priority order is preserved and tested: dense SDPA asserts the first gfx950 variant in registry order is chosen at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_wiring.py:312, and unified SDPA asserts returned candidates are sorted… |
| 13 | PASS | Pinned SDPA requests are explicitly declined instead of falling back to auto at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:233, and tested for known and unknown pins at… |
| 14 | NEEDS-INFO | The kernarg layout change for tiled attention (head counts and sliding window becoming kernargs) bumps ATTENTION_ABI_VERSION to v2 at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:80 and strided_decode's abi_version to v2 at dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/strided_decode.py:222. — no diff line supports this verdict |
| 15 | N/A | No new KernelCandidate or arch registration module is added; SDPA routes through existing attention candidates. |
| 16 | PASS | registered_sdpa_combos, sdpa_sweep_space, and dispatch_sdpa_all short-circuit on _request_errors(req) before delegating at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/init.py:123, with coverage in… |
| 17 | N/A | No new KernelCandidate is registered in this PR, so per-candidate CPU routing and family fixture invariants do not apply. |
| 18 | N/A | This PR does not change registry ratchets, route-only exemptions, or candidate binding requirements. |
| 19 | PASS | The unified SDPA binding accurately documents and uses run_unified_attention_torch with no tuning spec (backend="auto") at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:467, while dense SDPA launches through… |
| 20 | N/A | The selection/routing changes are confined to the Python library dispatch layer (library/dispatch/), which has no C++ twin. |
| 21 | PASS | Band and scale mismatches are rejected in the chooser/admission gates rather than hacking the emitter, at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:152 and… |
§E graded 17/17
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Platform changes remain strictly one-way with no library imports, e.g. dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/torch_interop.py:96 and dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:7656. |
| 2 | PASS | Library SDPA bindings use rocke.runtime.torch_interop helpers rather than importing torch directly or querying devices with torch, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:32 and… |
| 3 | PASS | Torch remains optional: tests gate it via pytest.importorskip in dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_binding.py:21 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:22, while runtime… |
| 4 | PASS | Imports inside library/ flow downward from benchmarks to dispatch to kernels, e.g. dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/attention_combo_sweep.py:44 imports dispatch.attention.common and… |
| 5 | PASS | Installed CTest entries use relative paths via ${ROCKE_TEST_CTEST_REL} in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:894 and dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1031. |
| 6 | PASS | Emitter changes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6187 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6786 add explicit IR zero-select guards for masked rows rather than… |
| 7 | PASS | Authoring and routing documentation is updated alongside the code in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:321 and dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591. |
| 8 | PASS | Code comments and docstrings use descriptive references without Jira IDs, PR numbers, or user handles, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:166 and dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591. |
| 9 | PASS | Resolved since previous review: the current PR diff against develop is cleanly scoped to SDPA routing and its gfx942 4-warp fix (e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and… |
| 10 | NEEDS-INFO | Not resolved since the previous review: the PR description and diff (which touches shared files such as dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:894 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8) still do not provide a merge-simulation record listing open PRs touching the same tables/tests/goldens and suite reruns. |
| 11 | PASS | Committed files consist only of deliverable source, tests, CMake test wiring, and required authoring documentation, e.g. dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591, dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1, and… |
| 12 | PASS | Both gfx942 and gfx950 are covered in the SDPA routing gates and test suites, e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_wiring.py:39 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_wiring.py:43. |
| 13 | PASS | The rename of _problem, _attention_band, and _validate_lengths to public helpers updates definitions and callers across the tree, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:758… |
| 14 | PASS | Resolved since previous review: the rebase onto develop cleanly preserves upstream code while updating only the intended lines and golden hashes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6185 and… |
| 15 | PASS | New commits explicitly validate the stream parameter in launch(*, stream=None, **_kw) rather than ignoring it in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:350 and… |
| 16 | N/A | This item is a meta-instruction for reviewer grading rather than a property of the PR code. |
| 17 | PASS | No absolute or relative performance claims appear in the PR description, comments, or documentation such as dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591. |
§F graded 11/11
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Branch users/tlakshma/hack/sdpa/st2-unified matches the recognized users/<user>/... prefix, and PR changes are in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:1. |
| 2 | PASS | title feat(hip-kernel-provider): route SDPA to unified and dense attention kernels matches type(scope): summary and is 76 characters |
| 3 | PASS | description is at least 30 characters and contains an issue reference |
| 4 | PASS | test file dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/attention/test_additive_registration.py is in the diff |
| 5 | PASS | All changed files in the diff (such as dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1) contain only code, docs, CMake, and test assets; no keys… |
| 6 | NEEDS-INFO | Still unresolved since previous review: local pre-commit execution status or CI lint check output is not visible in the diff. |
| 7 | NEEDS-INFO | Still unresolved since previous review: Although the new commits add review-response fixes (e.g. stream rejection in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:364, 16-byte alignment check in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:340, and positive scale check in… |
| 8 | N/A | The PR targets Base: develop rather than a stacked feature branch, so the stacked-PR bot-skipping case does not apply. |
| 9 | N/A | There is no indication that the PR was artificially split or trimmed for a size cap (in fact, #14624 was folded into this PR). |
| 10 | PASS | The added and updated tests are substantive wiring, CPU data-movement, case-table, and GPU numeric tests rather than throwaway stubs, e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:82 and… |
| 11 | PASS | The PR metadata header in this prompt shows the updated title (feat(hip-kernel-provider): route SDPA to unified and dense attention kernels) and full multi-section body (covering dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1 and… |
§G graded 34/34
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Correctness is verified against an independent mathematical reference oracle (sdpa_reference) rather than relying solely on IR byte-identity (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:61 and… |
| 2 | PASS | On-GPU numeric tests against the reference oracle are included and enrolled in the GPU test suite (dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1031 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:61). |
| 3 | NEEDS-INFO | In the updated head, attention_d128_swa.json is updated across all four LLVM flavors (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8) and attention_d256.json is updated (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8), but the tail of attention_d256.json at line 98 is… |
| 4 | PASS | New admission, binding, alignment, stream-rejection, scale-rejection, and length/offset-validation rules have falsifiable unit tests (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:122… |
| 5 | PASS | In test_an_out_of_range_block_table_physical_id_is_rejected_before_launch (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:83), the docstring explicitly documents why the test enters at validate_tuning_attention_tensors… |
| 6 | PASS | Prose claims in KERNEL_AUTHORING.md, docstrings, and code comments (including the new alignment, scale, and stream rejection comments in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:322… |
| 7 | PASS | KERNEL_AUTHORING.md specifies the reference oracle and four test layers as a first-class gate (dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:624), and the CPU and GPU binding tests validate against sdpa_reference… |
| 8 | PASS | Bindings and test harnesses perform strict dtype checks on operands, lengths, and ragged offsets (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:293, dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:267… |
| 9 | N/A | This PR does not introduce or claim any LDS size or occupancy numbers. |
| 10 | N/A | This PR is a functional routing and bug-fix change, not a performance optimization PR. |
| 11 | PASS | Kernel reachability through the SDPA routing gates and family entry points is verified in dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_wiring.py:474… |
| 12 | N/A | The changes in this PR are in Python-only dispatch layers and a gfx942 4-warp GQA kernel builder… |
| 13 | N/A | No new LLVM intrinsic overload classes or mangled symbols are added or recommended. |
| 14 | N/A | No hardware bottleneck classification claims are made in this PR. |
| 15 | N/A | The PR diff is provided directly in the review prompt rather than inspected from a local git checkout. |
| 16 | PASS | Integration with existing attention dispatch symbols (ATTENTION_ROUTE_REGISTRY, UNIFIED_ALGORITHMS, DENSE_ALGORITHMS, attention_band, attention_problem, dense_scale_range, opt_in_probe) is verified by wiring and coverage tests… |
| 17 | N/A | No standalone CLI or benchmark scripts are shipped in this PR; new files are library modules and pytest suites enrolled in CMake. |
| 18 | N/A | This PR does not add statistical runtime guards or back-pressure detectors of the kind this item targets. |
| 19 | NEEDS-INFO | The new commits and current diff provide explicit fixtures and inline checks for the named failures: misaligned contiguous BSHD views at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:122 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_binding.py:150, bad per-batch sequence lengths at… |
| 20 | N/A | This PR does not add a claim-ban detector or regex sweep. |
| 21 | N/A | This PR does not introduce performance measurement or profiling deliverable code. |
| 22 | PASS | The SDPA numeric GPU tests are excluded from the host SDPA pytest shard at dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:894 and dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:895 and explicitly enrolled in the GPU CTest lane at… |
| 23 | N/A | This PR does not make speedup or routing A/B performance claims. |
| 24 | NEEDS-INFO | The diff updates golden hashes in dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8 (including attention_d256/gfx942/4warp_gqa_d192 at line 16), while the SDPA unified gate admits UNIFIED_HEAD_SIZES (64, 128, 256)… |
| 25 | N/A | This PR is a functional routing and bounds/keyless-row correctness change, not a codegen micro-optimization report. |
| 26 | N/A | This PR does not make campaign-branch benchmark speedup claims. |
| 27 | N/A | This PR does not ship a performance knob or A/B toggle. |
| 28 | N/A | This PR does not present counterfactual performance tables. |
| 29 | N/A | This PR does not swap guarded math primitives for fast primitives. |
| 30 | N/A | This PR makes no register-spill or scratch-reduction claims. |
| 31 | N/A | This PR does not claim a performance win over a baseline defect. |
| 32 | N/A | The PR makes no timing, throughput, spill, or other measured performance-number claim requiring reproducible provenance. |
| 33 | N/A | The rocKE diff exposes and reuses attention routing helpers and adds SDPA gates, but does not modify an existing scalar… |
| 34 | PASS | The optional stats/LSE path is tested on padded/keyless rows, asserting O is zero and stats is -inf at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:157 and… |
§H graded 5/5
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | sdpa_unified_bind allocates k_cache/v_cache sized to slots = batch * self.pages at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:376-383 and validates the launch tensors against plan.problem via… |
| 2 | PASS | _LaunchPlan builds an identity block_table in range(slots) at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:385-387, validates physical block IDs before launch via validate_tuning_attention_tensors at… |
| 3 | N/A | After rebasing onto develop (which already merged #14929), this PR's diff no longer adds or modifies any device-side div/mod or… |
| 4 | PASS | In build_gfx942_4warp_gqa, rows whose diagonal is left of key 0 (context_off + q_inseq < 0, where every score in the loop was masked) now explicitly select 0.0 for O before storing at… |
| 5 | NEEDS-INFO | Resolved since previous review: seq_len_q/seq_len_kv are host-validated via _host_lengths and validate_lengths at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:267-289 and dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_sdpa_seqlen.py:45, operand shapes/strides/ragged offsets (including offset dtype) and 16-byte pointer alignment are validated in Operand at… |
§I graded 6/6
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR wires hipDNN SDPA forward requests to existing attention candidates and fixes gfx942 kernel edge cases without introducing… |
| 2 | PASS | Non-contiguous or unaligned (<16-byte) layouts are repaired via workspace copy buffers (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:276-284, dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:340-348), and… |
| 3 | PASS | Curated request and kernel case tables (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/cases.py:1361-1439) are combined with generated parameter sweeps and explicit algorithm-partition coverage checks… |
| 4 | N/A | This PR does not add or modify performance routing comparison tables or tuning reports. |
| 5 | N/A | This PR does not introduce baked tuning search results or static tuning snapshots. |
| 6 | N/A | This PR does not add any new DSL ops or hardware-instruction-named primitives. |
§J graded 7/7
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The only Markdown file in the PR (dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md) contains no mathematical… |
| 2 | N/A | The KERNEL_AUTHORING.md diff introduces no mathematical formula symbols, decay/log-decay domains, or log bases. |
| 3 | PASS | dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591-642 links to existing sections [E](#e-new-operator-with-a-hipdnn-graph-node), [F](#f-new-operand-or-output-on-an-existing-operator), and [A](#a-new-kernel-family) while stating only the delta for… |
| 4 | PASS | The novelty claims in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:321 and dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:638-641 name library/dispatch/sdpa/, bind_unified.py, bind_dense.py… |
| 5 | PASS | dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md makes no scaling/complexity claims, and the O(s_q) interval comparison claim in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:96-97 states the aggregate complexity across… |
| 6 | N/A | The KERNEL_AUTHORING.md diff does not introduce any model configuration tables or tuned band-edge constants. |
| 7 | N/A | This PR edits the existing dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md and does not add a new design or… |
§K graded 13/13
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR adds no multi-configuration candidate (SDPA_REGISTRY remains empty and SDPA delegates to existing attention candidates… |
| 2 | N/A | This PR does not add or modify any tunable kernel spec fields, KnobAxis declarations, or field classifications. |
| 3 | N/A | This PR does not change config_key payloads, canonical-knob rules, declared defaults, or knob meanings (unified_rules.py only… |
| 4 | N/A | This PR does not modify KnobSpace.defaults() or KnobSpace.recorded() in any KnobSpace subclass. |
| 5 | N/A | This PR does not touch knob drop/refuse rules, inert(), base_value(), or KNOWN_WRONG_KNOBS. |
| 6 | N/A | This PR does not touch KnobSpace.validate or library/dispatch/attention/tuning_specs.py. |
| 7 | N/A | This PR does not modify KnobSpace.base(req) or _tuning_problem (it only renames _problem to attention_problem in… |
| 8 | N/A | This PR does not modify normalize_knobs or request-boundary knob scalar coercion. |
| 9 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:220 (value_rejects) explicitly declines any pinned algorithm, spec_id, tuning_id, or tuning_knobs rather than adopting canonicalized ids or falling back to auto, tested in… |
| 10 | N/A | This PR does not modify KnobSpace.find or tuning_id resolution logic. |
| 11 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:119 (registered_sdpa_combos, sdpa_sweep_space, dispatch_sdpa_all) delegates to SDPA_REGISTRY and registered_attention_combos / dispatch_attention_all filtered by _serves(req)… |
| 12 | N/A | This PR does not reorder any production stack, change accept_default, or edit a tuning variant's base spec. |
| 13 | PASS | The benchmark changes in dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/attention_combo_sweep.py:44 and dnn-providers/hip-kernel-provider/rocke/library/benchmarks/gfx950/attention/decode/decode_dispatch_diagnose.py:13 only rename _problem to… |
§L graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This diff does not add or modify a rocKE gfx target in arch specs or ISA backend registries. |
| 2 | N/A | This diff does not add or modify an MMA op_id, MMA fragment-info row, or MMA lowering layout map. |
| 3 | N/A | This diff does not add or change MMA atom selection, MMA catalog fields, lane maps, or intrinsic-name emission. |
| 4 | N/A | This diff does not introduce or change wave-size handling or lane math derived from wave size. |
| 5 | PASS | The arch-specific codegen change is confined to the gfx942 4-warp kernel and masks fully masked rows without leaking gfx950 recipes at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6187. |
| 6 | N/A | This diff does not move pipeline vocabulary onto ArchTarget/core arch data, mutate ArchTarget from policies, or import CK Tile… |
| 7 | N/A | This diff only routes or declines SDPA dtype features and does not change rocKE logical dtype, bit-packing, scalar lowering, or… |
| 8 | N/A | This diff does not touch scaled-MMA scale contracts, retired scaled WMMA ids, or MfmaAtom.emit_scaled lowering. |
| 9 | N/A | This diff does not add a new rocKE architecture target. |
| 10 | N/A | This diff does not touch intrinsic_arch_domain..json, gen_arch_domain.py, or LLVM_FLAVORS. |
§M graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Resolved in the current diff: dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6187 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6786 compute no_key with b.cmp_lt and select the… |
| 2 | N/A | The PR diff does not touch legalize_for_target or target legalization logic. |
| 3 | N/A | The PR diff does not change print_ir, ck.dsl.ir/v1 serialization, wire versions, attrs, SSA ids, or parameter ordering. |
| 4 | N/A | The PR diff does not change portable IR schemas, importers, magic division constants, recipe VM code, or recorder coverage. |
| 5 | N/A | The PR diff does not touch AsyncTileLoader, async DRAM-to-LDS loading, LDS async layouts, or VMEM waits. |
| 6 | PASS | Resolved in the current diff: the kernel changes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6189 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6788 only guard b.global_store… |
| 7 | N/A | The PR diff does not add or modify buffer-resource word 3 handling or TemporalHint.STREAMING lowering. |
| 8 | N/A | The PR diff does not add LLVM intrinsics or change backend selection precedence. |
| 9 | N/A | The PR diff does not touch public C entry points, cpp/include/rocke/ headers, or C++ codegen launch behavior. |
| 10 | N/A | The PR diff does not emit !dbg metadata or change debug-info consumers. |
§N graded 18/18
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The PR diff against develop does not add or modify any files under platform/python/rocke/helpers/. |
| 2 | N/A | In the current head diff against develop, the previously flagged Q-addressing clamp in attention_tiled_2d.py was dropped (reused… |
| 3 | N/A | The diff does not touch SoftwarePipeline, sync_before_issue, or overlap_vmcnt behavior. |
| 4 | N/A | The diff does not add or modify fused kernels, EpilogueOp, fusion schedules, or fusion validation paths. |
| 5 | N/A | The diff does not add or modify a public API under rocke.helpers.tiling. |
| 6 | N/A | The diff does not add or modify front-door or toolbox tiling helpers or LayoutStyle classes. |
| 7 | N/A | The diff does not touch tiling layout overrides, accumulator c_desc derivation, or MMA soundness gates. |
| 8 | N/A | The diff does not touch helpers/tiling LDS allocation, lds_store/lds_read, contraction padding, or wide-load tiling emit rules. |
| 9 | N/A | The diff does not use or modify N-D tiling axis roles, at_index/squeeze behavior, FlowStage labeling, or AB-swap routing. |
| 10 | N/A | The diff does not modify tiling emit, LDS staging, or MMA issue code in helpers/tiling. |
| 11 | N/A | The diff does not touch tiling MMA traits, TileMma target selection, lds_swizzle policy, or derive_sched_group_counts. |
| 12 | N/A | The PR diff makes no helpers/tiling LDS bank-conflict or coalescing optimization claim requiring simulator, counter, or ASM-backed… |
| 13 | N/A | The diff does not construct WarpDistributionEncoding, change make_tile_desc factoring, re-export RegisterMapper, or alter tiling… |
| 14 | N/A | The diff does not use or modify helpers/tiling verbs or add a new helpers/tiling LDS pipeline variant. |
| 15 | N/A | The diff does not touch helpers/tiling descriptor allocation or transformed-range bounds checks after swizzle, pad, or XOR. |
| 16 | N/A | The diff does not use or modify optimize_layout, recommend(), or evaluate_transform ranking behavior. |
| 17 | N/A | The diff does not modify tiling_recorder.py, round-trip analysis, or tiling visualization emission behavior. |
| 18 | N/A | The diff does not touch InterleavedStyle, helpers/tiling/kernels demos, or tiling numeric gates for interleaved layouts. |
§O graded 16/16
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The PR does not add or modify IR address spaces, required attributes, or core/verify.py. |
| 2 | N/A | The PR does not modify scf.for structure or terminator emission in core IR/lowerers. |
| 3 | N/A | The PR does not touch IRBuilder, intrinsic declaration tables, or metadata slot numbering. |
| 4 | N/A | The PR does not modify IR types or SmemType serialization/parsing. |
| 5 | N/A | The PR does not modify op purity tables or core/passes.py / cpp/core/passes.cpp. |
| 6 | N/A | The PR does not modify constant folding in Python or C++ passes. |
| 7 | N/A | The PR does not modify kernel attributes, codegen policy, or cluster dimension checks. |
| 8 | N/A | The PR does not set or modify scheduler strategies or CodegenPolicy. |
| 9 | N/A | The PR does not modify waitcnt encoding or barrier lowering. |
| 10 | N/A | The PR does not modify lowerer attribute validation or pointer parameter lowering. |
| 11 | PASS | dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_numeric.py:30 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:36 resolve the device architecture via… |
| 12 | N/A | The PR does not add hardcoded width/size caps in core/codegen or silent arch or "gfx950" defaults. |
| 13 | N/A | The PR does not touch any public C structs or enums under cpp/include/rocke/. |
| 14 | N/A | The PR does not touch C++ sources or add C++ compiler flags in platform/CMakeLists.txt. |
| 15 | N/A | The PR does not modify C++ engine fallback mechanisms or lowerer entry points. |
| 16 | N/A | The PR does not touch the C++ codegen library or gfx1250 builders/lowerers. |
§P graded 14/14
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The PR diff does not add or modify platform instance builder signatures or spec REQUIRED_FIELDS. |
| 2 | N/A | The PR diff does not touch platform kernarg signatures, manifest args_signature, runtime argument packing, or GEMM/conv ABI… |
| 3 | N/A | The PR diff does not touch GEMM, implicit-GEMM, or block-scaled GEMM launch geometry. |
| 4 | N/A | The PR diff does not modify LDS allocation gates or C++ validators. |
| 5 | N/A | The PR diff does not touch WMMA/RDNA GEMM, GemmMultiD, or GemmMultiAbd admission paths. |
| 6 | N/A | The PR diff does not touch Stream-K, MoE sorting, fused MoE, or split-K caller-zeroed workspaces in rocKE. |
| 7 | N/A | The PR diff does not touch Stream-K rejection rules, persistent GEMM, GroupedGemmLauncher, or SUPPORT_MATRIX.md. |
| 8 | N/A | The PR diff does not touch quantization helpers, MX GEMM, or MX scale decoding. |
| 9 | N/A | The PR diff does not add or modify numeric emulation or low-precision decode paths requiring edge-case emulation tests. |
| 10 | N/A | The PR diff does not add a manifest kind or touch helpers/manifest.py, run_manifest, or manifest allocate/verify consumers. |
| 11 | PASS | The SDPA bindings allocate and retain launch buffers once in the bound plan, launch on resolve_stream(0, device=device), and reject non-current streams to prevent cross-stream races on preallocated workspaces: see… |
| 12 | N/A | The PR diff does not touch COMGR or HIP library loading logic. |
| 13 | N/A | The PR diff adds SDPA routing and binding logic rather than refactoring shared host-side grid, tile-count, workspace, cache-key… |
| 14 | N/A | The PR diff does not touch examples/, helpers/ placement, or add new modules under platform instances/common family glue. |
§Q graded 7/7
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | In the full PR diff, the new SDPA CPU tests under tests/library/tests/sdpa run via the existing rocke_library_sdpa_pytest CTest target, with the two new GPU numeric files excluded from the CPU suite at… |
| 2 | N/A | This PR routes SDPA requests onto existing attention candidates and fixes an existing gfx942 4-warp attention kernel without… |
| 3 | PASS | The representative IR golden tripwires are updated for the modified gfx942 4-warp attention kernels at dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and… |
| 4 | PASS | The representative golden files are updated across all four LLVM flavor blocks in dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 through… |
| 5 | NEEDS-INFO | All new tests in the full PR diff are portable: GPU numeric suites dynamically query the active device arch and skip cleanly without torch or a supported GPU (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:17, dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:33… |
| 6 | N/A | This PR does not refactor any platform validation gate or differential harness. |
| 7 | N/A | This PR does not modify engine rejection contracts or ROCKE_BACKEND / ROCKE_CPP_STRICT fallback behavior. |
§R graded 11/11
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The diff routes SDPA over existing attention candidates and adds a correctness zero-select for fully-masked rows in… |
| 2 | N/A | Changes since the previous review dropped the KV page clamp and Q-row hoist hunks after rebasing onto develop (#14929), leaving… |
| 3 | N/A | Changes since the previous review removed the global-load address clamp and Q-row hoist hunks from… |
| 4 | N/A | Changes since the previous review removed the hoisted Q-row / KV clamp hunks and their performance comments from… |
| 5 | N/A | Changes since the previous review removed the Q load and block-table/KV page address clamp changes from… |
| 6 | N/A | The diff does not change an LDS layout, MMA layout, tiling pipeline identity/relabel/reorder, or bank-conflict lever. |
| 7 | N/A | The diff makes no bottleneck claim such as memory-bound, LDS-bound, barrier-bound, or latency-bound. |
| 8 | N/A | The diff does not add, remove, or modify waits, barriers, sched_barrier, sched_group_barrier, iglp_opt, or scheduling hints. |
| 9 | N/A | The diff is not a performance-regression fix and does not add a shape-trading performance heuristic or dispatch rule. |
| 10 | N/A | Changes since the previous review removed the measured-performance comments from… |
| 11 | N/A | The documentation added in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591 is an SDPA routed-family authoring… |
§S graded 8/8
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | The new commits since 6daf7e6 do not resolve the missing hunk: torch_dtype and device_tensor are added at dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/torch_interop.py:96 and tested in dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:7656, and SDPA tests are registered in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:892 and… |
| 2 | PASS | Existing API behavior changes update their tests in the same PR: from_hipdnn requiring 'type' in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/hipdnn.py:220 is tested at… |
| 3 | N/A | The diff does not remove or rename a public API; _problem, _attention_band, and _validate_lengths were private… |
| 4 | PASS | dispatch.attention.__all__ is updated in dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/__init__.py:563 and verified by TestAttentionPublicSurface in… |
| 5 | PASS | Tests assert behavioral outputs and exact fail-fast exceptions: torch_dtype and device_tensor in dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:7660, non-positive/non-finite unified scale rejection in… |
| 6 | PASS | CPU-checkable contracts run in the host lane rocke_library_sdpa_pytest while ignoring the GPU numeric files in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:892, GPU numeric tests are registered in the GPU test list at… |
| 7 | N/A | The public APIs added or modified in this PR are Python dispatch/routing and torch_interop helpers that are Python-only by… |
| 8 | NEEDS-INFO | The new commits do not resolve the prior NEEDS-INFO: the PR description is truncated at | Rows that see no key ... and does not show whether the README test entry points (python tests/test_rocke.py and python tests/run_all.py from rocke/platform/) were stated as run. |
…m, one dtype rule Follow-up to the 11-06 review of #13115 (N16-N18, the dispatch_sdpa_all thread, and the code-quality notes): - launch(*, stream=None) is keyword-only in both bindings and goes through bind_shared.launch_stream: the current stream is accepted unset, as 0, or named by handle (graph capture's form); any other stream raises, and an unknown keyword is a TypeError instead of being swallowed. - bind_shared.check_dtype is the one dtype rule: operands, stats, seq_len_* and ragged offsets must match their descriptor, so int64 seq_len_q and fp16 stats against an fp32 descriptor are refused. - bind_shared.usable_in_place (contiguous and 16-byte aligned) and nonfinite_scale_reject replace the copies in both bindings; the unified scale docstring names only the gfx942 4-warp bodies. - dispatch.sdpa exports sdpa_{unified,dense}_{admits,bind} and documents that dispatch_sdpa_all's results are attention results, not SDPA launches. - Tests: the unified scale declines move to the torch-free wiring test (CI runs it, both arches); a ragged twin of the misalignment test; stream, unknown-keyword, seq_len and stats dtype tests for both bindings; the strided test tensors live once in binding_harness. - KERNEL_AUTHORING §H lists the binding rules. Co-Authored-By: Claude <noreply@anthropic.com>
|
@cgmillette @CongMa13: follow-up on the 11-06 10:43Z review, pushed as
Each new test fails against its fix's mutant. Runs at
Threads: the 7 fixed Copilot threads have replies and can be resolved. AnaghaRaoAMD's K/V-copy thread stays deferred ("Not in this PR", §H), with a follow-up reply asking them to resolve it or ask for it here. |
yraparti
left a comment
There was a problem hiding this comment.
rocKE checklist re-review of the fixes since 0d916301ba32
ROCm/rocm-libraries#13115 · cf2b94adaae7 · model gemini-4-argon, fallback gpt-5.5
PASS · graded 238/238 · BLOCKER 0, HIGH 0, MEDIUM 0, LOW 0, NEEDS-INFO 11
The diff was truncated before grading. Items that needed a missing hunk are NEEDS-INFO.
Since the last review (0d916301ba32)
Resolved 0 · still open 0 · new 0
Findings
- No FAIL rows.
Needs info
- §A.1 Unresolved since previous review:
dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6258and:6857modify_build_gfx942_4warp_gqa_leanandbuild_gfx942_4warp_gqawithout a C++ update in the PR diff, andknown_gaps.mdandcpp/instances/gfx942/are not in the PR diff to verify whether the 4-warp GQA body has a C++ twin or is documented as Python-only. - §D.14 Previous NEEDS-INFO is not resolved: the visible diff does not include the cache-key use or ATTENTION_ABI_VERSION/abi_version hunks needed to verify compile-key identity and old-id refusal.
- §E.10 Not resolved since the previous review: the PR description and diff (which touches shared files such as dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:932 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8) still do not provide a merge-simulation record listing open PRs touching the same tables/tests/goldens and suite reruns.
- §F.6 Still unresolved since previous review: local
pre-commitexecution status or CI lint check output is not visible in the diff. - §F.7 Still unresolved since previous review: although the commits include review-response fixes (e.g.
launch_streamindnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:276,usable_in_place16-byte alignment check indnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:323, and_scale_rejectindnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:163), the GitHub review-thread list and any deferred-item ticket IDs are not included in the provided PR context. - §G.3 In the current head,
attention_d128_swa.json(dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8) andattention_d256.json(dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8) are updated, but the tail ofattention_d256.jsonat line 98 is still truncated in the diff and local non-GPU gate sweep execution results cannot be verified from the diff alone. - §G.24 The new commits did not resolve the previous question: the diff updates golden hashes in dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8 (including
attention_d256/gfx942/4warp_gqa_d192at dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:16) due to the shared keyless-row O store fix at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6258 and… - §H.5 Resolved since previous review: all launch inputs feeding bounds/masks are host-validated with negative tests —
seq_len_q/seq_len_kvvia_host_lengthsandvalidate_lengthsat dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:267-290 and dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_sdpa_seqlen.py:45 (tested at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:155-198, 317-330 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_sdpa_seqlen.py:172-196), operand dtypes, shapes, strides, ragged offsets, and 16-byte alignment incheck_dtype/usable_in_place/Operandat…
Section notes
- §C: Re-evaluated against cf2b94a; all review-response additions (dtype checks, alignment checks, stream checks) are factored cleanly into bind_shared.py and binding_harness.py.
- §F: Re-evaluated against head cf2b94a; items 6 and 7 remain NEEDS-INFO because pre-commit execution logs and GitHub review-thread states are not present in the diff context.
- §G: Re-grading items 18-34 of section G against head cf2b94a.
- §H: 1) gfx942 4-warp GQA epilogue empty-row mask (
no_key = context_off + q_inseq < 0) inattention_tiled_2d.py: extent source is runtimeqlenandklen(context_off = klen - qlen); compile-time range visited is the Q tile rows (q_inseq < qlen). Whenqlen > klenon bottom-right causal, rows withcontext_off + q_inseq < 0have all K positions masked out in the main loop;select(no_key, 0.0, acc * recip)zeroes their output before storing, so they agree. 2)sdpa_unified_bindpaged cache &block_tableinbind_unified.py: extent source isslots = batch * pageswherepages = max(1, ceil(s_kv / block)); compile-time/runtime visited range isseqused_k[b] <= s_kvacrosspagesblocks per batch;k_cache/v_cacheare zero-initialized to(slots, block, h, d),block_tableis identityrange(slots)shaped(batch, pages), andvalidate_tuning_attention_tensorsvalidates shapes/bounds before launch, so they agree. 3)cu_seqlens_q/seqused_kfrom… - §J: Shapes: | Equation | Operands | Inner dims | Result vs Declared | |---|---|---|---| | (none in Markdown diff) | - | - | - | Symbols: | Symbol | Defined at | Domain | |---|---|---| |
SDPA_<FAMILY>_FEATURES|KERNEL_AUTHORING.md:609| carried feature set per SDPA launch family | |UNIFIED_ALGORITHMS|KERNEL_AUTHORING.md:648| unified attention route algorithm names | |DENSE_ALGORITHMS|KERNEL_AUTHORING.md:649| dense attention route algorithm names | Novelty: | Claim | Symbol / Path | Status | |---|---|---| | Routed SDPA family overdispatch.attention|library/dispatch/sdpa/,bind_unified.py,bind_dense.py| Added in this PR | | Exported route algorithm sets |dispatch.attention.UNIFIED_ALGORITHMS,dispatch.attention.DENSE_ALGORITHMS| Added in this PR | - §K: Identity table: no tuned knob, axis, default, canonicalization rule, or builder behavior behind a tuning knob is changed by this PR (unified_rules.py only renames
_problemtoattention_problem).
Graded with the rocKE PR checklist, one model pass per section. FAIL rows that sit on a diff line are inline comments.
§A graded 14/14
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | Unresolved since previous review: dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6258 and :6857 modify _build_gfx942_4warp_gqa_lean and build_gfx942_4warp_gqa without a C++ update in the PR diff, and known_gaps.md and cpp/instances/gfx942/ are not in the PR diff to verify whether the 4-warp GQA body has a C++ twin or is documented as Python-only. |
| 2 | N/A | The PR does not modify any core lowering ops or IR op-handlers under core/lower_* or core/ir. |
| 3 | PASS | New dispatch and routing logic is confined to dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1 (plus public exports in… |
| 4 | N/A | No C++ kernel twin or C++ IR emitter code is modified in this PR. |
| 5 | PASS | Both dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8 re-bless only the… |
| 6 | N/A | No C++ twin is touched in the PR diff and the PR does not claim cross-engine parity from a golden update. |
| 7 | PASS | The no-key output zeroing fix in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6258 and :6857 applies only to the 4warp_gqa builders, and the golden diffs in… |
| 8 | N/A | No rocKE C++ files are modified or added in the PR diff. |
| 9 | N/A | No shared op-handler or lowerer arch divergence fix is included in this PR. |
| 10 | N/A | No new kernel IR feature flag is introduced in this PR; it adds host-side SDPA dispatch bindings and fixes a bug in existing… |
| 11 | N/A | No new IR op is added in this PR. |
| 12 | N/A | core/backend.py and ROCKE_BACKEND handling are not touched in the PR diff. |
| 13 | N/A | LLVM .ll emission and debug-location handling are not modified in this PR. |
| 14 | N/A | No helper files under helpers/ or platform/cpp/helpers/ are modified in the PR diff. |
§B graded 6/6
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:176-191 classifies shortcut feature declines with NOT_YET_IMPLEMENTED: separately from ABI/kernel limits, and… |
| 2 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:211-214 probes each dense candidate with opt_in_probe, checks candidate.admits(probe), and only then builds the candidate spec… |
| 3 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:239-255 declines requests when no unified candidate covers them, and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:459-461 plus… |
| 4 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/hipdnn.py:217-222 rejects a node missing its type key instead of defaulting it; dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:291-297 rejects tensors whose dtype differs… |
| 5 | PASS | dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:49-67 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_numeric.py:43-61 run GPU numeric verification over every case admitted by the unified and dense… |
| 6 | N/A | The visible diff adds SDPA dispatch admission/binding logic and kernel bug fixes without adding or modifying an is_valid_spec or… |
§C graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Shared predicates, decline reasons, dtype checks (check_dtype), 16-byte alignment checks (usable_in_place), and stream validation (launch_stream) are factored into dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:46-329, test tensor… |
| 2 | PASS | The SDPA bindings route to existing public entrypoints (bind_dense_attention_torch in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:345-347 and run_unified_attention_torch in… |
| 3 | N/A | No new kernel builder functions or kernel_name() definitions are introduced in this PR. |
| 4 | PASS | Arch-independent SDPA request/binding logic lives in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:1-408, while gfx942-specific kernel fixes stay in… |
| 5 | PASS | All touched files in the PR diff are directly scoped to SDPA dispatch/binding (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:1-192), SDPA tests (dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:930-935… |
| 6 | PASS | Shared SDPA request, band, decline-table, stream, dtype, and operand abstractions are cleanly factored into dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:62-408, and single-source public attention-layer accessors (attention_band… |
| 7 | N/A | This PR is a feature routing and bug-fix change rather than a pure no-op cleanup. |
| 8 | PASS | After rebasing onto develop, the PR reuses existing helpers including opt_in_probe, empty_workspace, and resolve_stream (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:30-31… |
| 9 | PASS | The no_key zeroing fix for rows where qlen > klen in bottom-right causal mode is applied across both the 4-warp lean and 4-warp GQA bodies in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6256-6260 and… |
| 10 | PASS | Every feature in SDPA_FEATURES is explicitly partitioned into carried vs declined sets in both dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:82-128 and… |
§D graded 21/21
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR does not add or modify any KernelSpec knobs or environment variable reads. |
| 2 | PASS | Coverage limits and algorithm constants are imported from kernel/attention modules rather than transcribed (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:35 imports UNIFIED_ALGORITHMS and UNIFIED_HEAD_SIZES… |
| 3 | PASS | All candidate admission checks go through candidate.admits(probe) or ATTENTION_ROUTE_REGISTRY.supported(...) and never call _supports directly (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:215… |
| 4 | PASS | Selection and admission are pure functions of SdpaRequest taking req.arch directly without any device probes (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:90… |
| 5 | N/A | This PR does not introduce or modify any KernelSpec identity, spec_hash, compile_key, or kernel_name() methods. |
| 6 | N/A | This PR does not introduce or modify cohort routing, spec-builders, or kernel cache keys. |
| 7 | N/A | This PR does not add hand-pinned heuristic overrides. |
| 8 | N/A | This PR does not define or modify any kernel cohort predicates. |
| 9 | N/A | No new KernelCandidate or architecture is registered; the PR routes SdpaRequest onto existing ATTENTION_ROUTE_REGISTRY candidates. |
| 10 | PASS | Every SdpaRequest field and feature is translated into AttentionRequest and routed to the unified and dense bindings with dispatcher-level wiring, entry-point, and numeric tests (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:62… |
| 11 | PASS | Existing opt-in dense candidates remain opt-in and are queried explicitly via opt_in_probe(attn, candidate) in the SDPA dense gate (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:214). |
| 12 | PASS | Dense SDPA asserts the first serving candidate is the first variant in registry order at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_wiring.py:312, and unified SDPA asserts candidate priorities remain sorted at… |
| 13 | PASS | Pinned SDPA dispatch choices are declined before routing rather than retried as auto at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:234, with tests covering known and missing pins at… |
| 14 | NEEDS-INFO | Previous NEEDS-INFO is not resolved: the visible diff does not include the cache-key use or ATTENTION_ABI_VERSION/abi_version hunks needed to verify compile-key identity and old-id refusal. |
| 15 | N/A | No new KernelCandidate or arch registration module is added; SDPA routes through existing attention candidates. |
| 16 | PASS | The SDPA family entry points short-circuit on _request_errors(req) before delegating or adding attention combos at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/init.py:112, with family-entry coverage at… |
| 17 | N/A | No new KernelCandidate is registered in this PR, so per-candidate CPU routing and family fixture invariants do not apply. |
| 18 | N/A | This PR does not change registry ratchets, route-only exemptions, or candidate binding requirements. |
| 19 | PASS | The SDPA docs state the returned dispatch results are attention results and that SDPA requests must launch through sdpa_unified_bind or sdpa_dense_bind at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/init.py:17, while unified binding uses… |
| 20 | N/A | The visible selection/routing changes are in the Python SDPA dispatch/binding layer over existing attention candidates, not in a… |
| 21 | PASS | Band mismatches are rejected in the SDPA admission chooser at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:135, and dense candidates are accepted only when their selected spec keeps the requested keys at… |
§E graded 17/17
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Platform changes remain strictly one-way with no library imports, e.g. dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/torch_interop.py:96 and dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:7658. |
| 2 | PASS | Library SDPA bindings use rocke.runtime.torch_interop helpers rather than importing torch directly or querying devices with torch, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:31 and… |
| 3 | PASS | Torch remains optional: tests gate it via pytest.importorskip in dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_binding.py:21 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:22, while runtime… |
| 4 | PASS | Imports inside library/ flow downward from benchmarks to dispatch to kernels, e.g. dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/attention_combo_sweep.py:44 imports dispatch.attention.common and… |
| 5 | PASS | Installed CTest entries use relative paths via ${ROCKE_TEST_CTEST_REL} in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:932 and dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1075. |
| 6 | PASS | Emitter changes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6258 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6857 add explicit IR zero-select guards for masked rows rather than… |
| 7 | PASS | Authoring and routing documentation is updated alongside the code in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:321 and dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591. |
| 8 | PASS | Code comments and docstrings use descriptive references without Jira IDs, PR numbers, or user handles, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:167 and dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591. |
| 9 | PASS | The current PR diff against develop is cleanly scoped to SDPA routing and its gfx942 4-warp fix, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and… |
| 10 | NEEDS-INFO | Not resolved since the previous review: the PR description and diff (which touches shared files such as dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:932 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8) still do not provide a merge-simulation record listing open PRs touching the same tables/tests/goldens and suite reruns. |
| 11 | PASS | Committed files consist only of deliverable source, tests, CMake test wiring, and required authoring documentation, e.g. dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591, dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1, and… |
| 12 | PASS | Both gfx942 and gfx950 are covered in the SDPA routing gates and test suites, e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_wiring.py:39 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_wiring.py:43. |
| 13 | PASS | The rename of _problem, _attention_band, and _validate_lengths to public helpers updates definitions and callers across the tree, e.g. dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/common.py:758… |
| 14 | PASS | The rebase onto develop cleanly preserves upstream code while updating only the intended lines and golden hashes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6256 and… |
| 15 | PASS | Every parameter in the bindings is used and validated, including stream in launch(*, stream: Any = None) in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:350 and… |
| 16 | N/A | This item is a meta-instruction for reviewer grading rather than a property of the PR code. |
| 17 | PASS | No absolute or relative performance claims appear in the PR description, comments, or documentation such as dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591. |
§F graded 11/11
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Branch users/tlakshma/hack/sdpa/st2-unified matches the recognized users/<user>/... prefix, and PR changes are in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:1. |
| 2 | PASS | title feat(hip-kernel-provider): route SDPA to unified and dense attention kernels matches type(scope): summary and is 76 characters |
| 3 | PASS | description is at least 30 characters and contains an issue reference |
| 4 | PASS | test file dnn-providers/hip-kernel-provider/rocke/library/tests/dispatch/attention/test_additive_registration.py is in the diff |
| 5 | PASS | All changed files in the diff (such as dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:1 and dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1) contain only code, docs, CMake, and test assets; no keys… |
| 6 | NEEDS-INFO | Still unresolved since previous review: local pre-commit execution status or CI lint check output is not visible in the diff. |
| 7 | NEEDS-INFO | Still unresolved since previous review: although the commits include review-response fixes (e.g. launch_stream in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:276, usable_in_place 16-byte alignment check in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:323, and _scale_reject in… |
| 8 | N/A | The PR targets Base: develop rather than a stacked feature branch, so the stacked-PR bot-skipping case does not apply. |
| 9 | N/A | There is no indication that the PR was artificially split or trimmed for a size cap (in fact, #14624 was folded into this PR). |
| 10 | PASS | The added and updated tests are substantive wiring, CPU data-movement, case-table, and GPU numeric tests rather than throwaway stubs, e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:82… |
| 11 | PASS | The PR metadata header in this prompt shows the updated title (feat(hip-kernel-provider): route SDPA to unified and dense attention kernels) and full multi-section body (covering dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:1 and… |
§G graded 34/34
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | Correctness is verified against an independent mathematical reference oracle (sdpa_reference) rather than relying solely on IR byte-identity (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:61 and… |
| 2 | PASS | On-GPU numeric tests against the reference oracle are included and enrolled in the GPU test suite (dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:1075, dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:61, and… |
| 3 | NEEDS-INFO | In the current head, attention_d128_swa.json (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8) and attention_d256.json (dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8) are updated, but the tail of attention_d256.json at line 98 is still truncated in the diff and local… |
| 4 | PASS | New admission, binding, alignment, stream-rejection, scale-rejection, dtype-rejection, and length/offset-validation rules have falsifiable unit tests (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:106… |
| 5 | PASS | In test_an_out_of_range_block_table_physical_id_is_rejected_before_launch (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:83), the docstring explicitly documents why the test enters at validate_tuning_attention_tensors… |
| 6 | PASS | Prose claims in KERNEL_AUTHORING.md (dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591), docstrings, and code comments (including the alignment, scale, dtype, and stream rejection comments in… |
| 7 | PASS | KERNEL_AUTHORING.md specifies the reference oracle and four test layers as a first-class gate (dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:633), and the CPU and GPU binding tests validate against sdpa_reference… |
| 8 | PASS | Bindings and test harnesses perform strict dtype checks on operands, stats, lengths, and ragged offsets (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:305… |
| 9 | N/A | This PR does not introduce or claim any LDS size or occupancy numbers. |
| 10 | N/A | This PR is a functional routing and bug-fix change, not a performance optimization PR. |
| 11 | PASS | Kernel reachability through the SDPA routing gates and family entry points is verified in dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_wiring.py:493… |
| 12 | N/A | The changes in this PR are in Python-only dispatch layers and a gfx942 4-warp GQA kernel builder… |
| 13 | N/A | No new LLVM intrinsic overload classes or mangled symbols are added or recommended. |
| 14 | N/A | No hardware bottleneck classification claims are made in this PR. |
| 15 | N/A | The PR diff is provided directly in the review prompt rather than inspected from a local git checkout. |
| 16 | PASS | Integration with existing attention dispatch symbols (ATTENTION_ROUTE_REGISTRY, UNIFIED_ALGORITHMS, DENSE_ALGORITHMS, attention_band, attention_problem, dense_scale_range, opt_in_probe) is verified by wiring and coverage tests… |
| 17 | N/A | No standalone CLI or benchmark scripts are shipped in this PR; new files are library modules and pytest suites enrolled in CMake. |
| 18 | N/A | This PR does not add statistical runtime guards or back-pressure detectors of the kind this item targets. |
| 19 | PASS | The fixtures and inline checks explicitly construct and catch the named failure modes across the binding and wiring gates: misaligned contiguous BSHD/ragged views at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/binding_harness.py:17… |
| 20 | N/A | This PR does not add a claim-ban detector or regex sweep. |
| 21 | N/A | This PR does not introduce performance measurement or profiling deliverable code. |
| 22 | PASS | The SDPA numeric GPU tests are excluded from the host SDPA pytest shard at dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:933 and dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:934 and explicitly enrolled in the GPU CTest lane at… |
| 23 | N/A | This PR does not make speedup or routing A/B performance claims. |
| 24 | NEEDS-INFO | The new commits did not resolve the previous question: the diff updates golden hashes in dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d256.json:8 (including attention_d256/gfx942/4warp_gqa_d192 at… |
| 25 | N/A | This PR is a functional routing and bounds/keyless-row correctness change, not a codegen micro-optimization report. |
| 26 | N/A | This PR does not make campaign-branch benchmark speedup claims. |
| 27 | N/A | This PR does not ship a performance knob or A/B toggle. |
| 28 | N/A | This PR does not present counterfactual performance tables. |
| 29 | N/A | This PR does not swap guarded math primitives for fast primitives. |
| 30 | N/A | This PR makes no register-spill or scratch-reduction claims. |
| 31 | N/A | This PR does not claim a performance win over a baseline defect. |
| 32 | N/A | The PR makes no timing, throughput, spill, or other measured performance-number claim requiring reproducible provenance. |
| 33 | N/A | The rocKE diff exposes and reuses attention routing helpers and adds SDPA gates, but does not modify an existing scalar… |
| 34 | PASS | The optional stats/LSE path is tested on padded/keyless rows, asserting O is zero and stats is -inf at dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_binding.py:144 and… |
§H graded 5/5
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | sdpa_unified_bind allocates k_cache/v_cache sized to slots = batch * self.pages at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:376-382 and validates the launch tensors against plan.problem via… |
| 2 | PASS | _LaunchPlan builds an identity block_table in range(slots) at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:385-387, validates physical block IDs before launch via validate_tuning_attention_tensors at… |
| 3 | N/A | This PR's diff does not add or modify any device-side div/mod or shift/mask page-indexing logic (only epilogue zeroing for empty… |
| 4 | PASS | In build_gfx942_4warp_gqa, rows whose diagonal is left of key 0 (context_off + q_inseq < 0, where every score in the loop was masked) explicitly select 0.0 for O before storing at… |
| 5 | NEEDS-INFO | Resolved since previous review: all launch inputs feeding bounds/masks are host-validated with negative tests — seq_len_q/seq_len_kv via _host_lengths and validate_lengths at dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_unified.py:267-290 and dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_sdpa_seqlen.py:45 (tested at… |
§I graded 6/6
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR wires hipDNN SDPA forward requests to existing attention candidates and fixes gfx942 kernel edge cases without introducing… |
| 2 | PASS | Non-contiguous or unaligned (<16-byte) layouts are repaired via workspace copy buffers (dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_dense.py:276-284, dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:323-328), and… |
| 3 | PASS | Curated request and kernel case tables (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/cases.py:1361-1439) are combined with generated parameter sweeps and explicit algorithm-partition coverage checks… |
| 4 | N/A | This PR does not add or modify performance routing comparison tables or tuning reports. |
| 5 | N/A | This PR does not introduce baked tuning search results or static tuning snapshots. |
| 6 | N/A | This PR does not add any new DSL ops or hardware-instruction-named primitives. |
§J graded 7/7
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The only Markdown file in the PR (dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md) contains no mathematical… |
| 2 | N/A | The KERNEL_AUTHORING.md diff introduces no mathematical formula symbols, decay/log-decay domains, or log bases. |
| 3 | PASS | dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591-651 links to existing sections [E](#e-new-operator-with-a-hipdnn-graph-node), [F](#f-new-operand-or-output-on-an-existing-operator), and [A](#a-new-kernel-family) while stating only the delta for… |
| 4 | PASS | The novelty claims in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:321 and dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:647-650 name library/dispatch/sdpa/, bind_unified.py, bind_dense.py… |
| 5 | PASS | dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md makes no scaling/complexity claims, and the O(s_q) interval comparison claim in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:112-113 states the aggregate complexity… |
| 6 | N/A | The KERNEL_AUTHORING.md diff does not introduce any model configuration tables or tuned band-edge constants. |
| 7 | N/A | This PR edits the existing dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md and does not add a new design or… |
§K graded 13/13
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This PR adds no multi-configuration candidate (SDPA_REGISTRY remains empty and SDPA delegates to existing attention candidates… |
| 2 | N/A | This PR does not add or modify any tunable kernel spec fields, KnobAxis declarations, or field classifications. |
| 3 | N/A | This PR does not change config_key payloads, canonical-knob rules, declared defaults, or knob meanings (unified_rules.py only… |
| 4 | N/A | This PR does not modify KnobSpace.defaults() or KnobSpace.recorded() in any KnobSpace subclass. |
| 5 | N/A | This PR does not touch knob drop/refuse rules, inert(), base_value(), or KNOWN_WRONG_KNOBS. |
| 6 | N/A | This PR does not touch KnobSpace.validate or library/dispatch/attention/tuning_specs.py. |
| 7 | N/A | This PR does not modify KnobSpace.base(req) or _tuning_problem (it only renames _problem to attention_problem in… |
| 8 | N/A | This PR does not modify normalize_knobs or request-boundary knob scalar coercion. |
| 9 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/bind_shared.py:244 (value_rejects) explicitly declines any pinned algorithm, spec_id, tuning_id, or tuning_knobs rather than adopting canonicalized ids or falling back to auto, tested in… |
| 10 | N/A | This PR does not modify KnobSpace.find or tuning_id resolution logic. |
| 11 | PASS | dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/__init__.py:130 (registered_sdpa_combos, sdpa_sweep_space, dispatch_sdpa_all) delegates to SDPA_REGISTRY and registered_attention_combos / dispatch_attention_all filtered by _serves(req)… |
| 12 | N/A | This PR does not reorder any production stack, change accept_default, or edit a tuning variant's base spec. |
| 13 | PASS | The benchmark changes in dnn-providers/hip-kernel-provider/rocke/library/benchmarks/common/attention_combo_sweep.py:44 and dnn-providers/hip-kernel-provider/rocke/library/benchmarks/gfx950/attention/decode/decode_dispatch_diagnose.py:13 only rename _problem to… |
§L graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | This diff does not add or modify a rocKE gfx target in arch specs or ISA backend registries. |
| 2 | N/A | This diff does not add or modify an MMA op_id, MMA fragment-info row, or MMA lowering layout map. |
| 3 | N/A | This diff does not add or change MMA atom selection, MMA catalog fields, lane maps, or intrinsic-name emission. |
| 4 | N/A | This diff does not introduce or change wave-size handling or lane math derived from wave size. |
| 5 | PASS | The arch-specific codegen change is confined to the gfx942 4-warp kernel and masks fully masked rows without leaking gfx950 recipes at dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6258. |
| 6 | N/A | This diff does not move pipeline vocabulary onto ArchTarget/core arch data, mutate ArchTarget from policies, or import CK Tile… |
| 7 | N/A | This diff only routes or declines SDPA dtype features and does not change rocKE logical dtype, bit-packing, scalar lowering, or… |
| 8 | N/A | This diff does not touch scaled-MMA scale contracts, retired scaled WMMA ids, or MfmaAtom.emit_scaled lowering. |
| 9 | N/A | This diff does not add a new rocKE architecture target. |
| 10 | N/A | This PR diff against base does not touch intrinsic_arch_domain..json, gen_arch_domain.py, or LLVM_FLAVORS. |
§M graded 10/10
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6258 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6857 compute no_key with b.cmp_lt and select the output with b.select and… |
| 2 | N/A | The PR diff does not touch legalize_for_target or target legalization logic. |
| 3 | N/A | The PR diff does not change print_ir, ck.dsl.ir/v1 serialization, wire versions, attrs, SSA ids, or parameter ordering. |
| 4 | N/A | The PR diff does not change portable IR schemas, importers, magic division constants, recipe VM code, or recorder coverage. |
| 5 | N/A | The PR diff does not touch AsyncTileLoader, async DRAM-to-LDS loading, LDS async layouts, or VMEM waits. |
| 6 | PASS | The kernel changes in dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6261 and dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tiled_2d.py:6860 only guard b.global_store inside… |
| 7 | N/A | The PR diff does not add or modify buffer-resource word 3 handling or TemporalHint.STREAMING lowering. |
| 8 | N/A | The PR diff does not add LLVM intrinsics or change backend selection precedence. |
| 9 | N/A | The PR diff does not touch public C entry points, cpp/include/rocke/ headers, or C++ codegen launch behavior. |
| 10 | N/A | The PR diff does not emit !dbg metadata or change debug-info consumers. |
§N graded 18/18
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The PR diff against develop does not add or modify any files under platform/python/rocke/helpers/. |
| 2 | N/A | The PR diff against develop only adds the no-key select guard in library/kernels/gfx942/attention_tiled_2d.py and does not add or… |
| 3 | N/A | The diff does not touch SoftwarePipeline, sync_before_issue, or overlap_vmcnt behavior. |
| 4 | N/A | The diff does not add or modify fused kernels, EpilogueOp, fusion schedules, or fusion validation paths. |
| 5 | N/A | The diff against develop does not add or modify a public API under rocke.helpers.tiling. |
| 6 | N/A | The diff against develop does not add or modify front-door or toolbox tiling helpers or LayoutStyle classes. |
| 7 | N/A | The diff against develop does not touch tiling layout overrides, accumulator c_desc derivation, or MMA soundness gates. |
| 8 | N/A | The diff against develop does not touch helpers/tiling LDS allocation, lds_store/lds_read, contraction padding, or wide-load… |
| 9 | N/A | The diff against develop does not use or modify N-D tiling axis roles, at_index/squeeze behavior, FlowStage labeling, or AB-swap… |
| 10 | N/A | The diff against develop does not modify tiling emit, LDS staging, or MMA issue code in helpers/tiling. |
| 11 | N/A | The diff against develop does not touch tiling MMA traits, TileMma target selection, lds_swizzle policy, or… |
| 12 | N/A | The PR diff makes no helpers/tiling LDS bank-conflict or coalescing optimization claim requiring simulator, counter, or ASM-backed… |
| 13 | N/A | The diff against develop does not construct WarpDistributionEncoding, change make_tile_desc factoring, re-export RegisterMapper… |
| 14 | N/A | The diff against develop does not use or modify helpers/tiling verbs or add a new helpers/tiling LDS pipeline variant. |
| 15 | N/A | The diff against develop does not touch helpers/tiling descriptor allocation or transformed-range bounds checks after swizzle… |
| 16 | N/A | The diff against develop does not use or modify optimize_layout, recommend(), or evaluate_transform ranking behavior. |
| 17 | N/A | The diff against develop does not modify tiling_recorder.py, round-trip analysis, or tiling visualization emission behavior. |
| 18 | N/A | The diff against develop does not touch InterleavedStyle, helpers/tiling/kernels demos, or tiling numeric gates for interleaved… |
§O graded 16/16
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The PR does not add or modify IR address spaces, required attributes, or core/verify.py. |
| 2 | N/A | The PR does not modify scf.for structure or terminator emission in core IR/lowerers. |
| 3 | N/A | The PR does not touch IRBuilder, intrinsic declaration tables, or metadata slot numbering. |
| 4 | N/A | The PR does not modify IR types or SmemType serialization/parsing. |
| 5 | N/A | The PR does not modify op purity tables or core/passes.py / cpp/core/passes.cpp. |
| 6 | N/A | The PR does not modify constant folding in Python or C++ passes. |
| 7 | N/A | The PR does not modify kernel attributes, codegen policy, or cluster dimension checks. |
| 8 | N/A | The PR does not set or modify scheduler strategies or CodegenPolicy. |
| 9 | N/A | The PR does not modify waitcnt encoding or barrier lowering. |
| 10 | N/A | The PR does not modify lowerer attribute validation or pointer parameter lowering. |
| 11 | PASS | dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_dense_sdpa_numeric.py:30 and dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:36 resolve the device architecture via… |
| 12 | N/A | The PR does not add hardcoded width/size caps in core/codegen or silent arch or "gfx950" defaults. |
| 13 | N/A | The PR does not touch any public C structs or enums under cpp/include/rocke/. |
| 14 | N/A | The PR does not touch C++ sources or add C++ compiler flags in platform/CMakeLists.txt. |
| 15 | N/A | The PR does not modify C++ engine fallback mechanisms or lowerer entry points. |
| 16 | N/A | The PR does not touch the C++ codegen library or gfx1250 builders/lowerers. |
§P graded 14/14
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The PR diff does not add or modify platform instance builder signatures or spec REQUIRED_FIELDS. |
| 2 | N/A | The PR diff does not touch platform kernarg signatures, manifest args_signature, runtime argument packing, or GEMM/conv ABI… |
| 3 | N/A | The PR diff does not touch GEMM, implicit-GEMM, or block-scaled GEMM launch geometry. |
| 4 | N/A | The PR diff does not modify LDS allocation gates or C++ validators. |
| 5 | N/A | The PR diff does not touch WMMA/RDNA GEMM, GemmMultiD, or GemmMultiAbd admission paths. |
| 6 | N/A | The PR diff does not touch Stream-K, MoE sorting, fused MoE, or split-K caller-zeroed workspaces in rocKE. |
| 7 | N/A | The PR diff does not touch Stream-K rejection rules, persistent GEMM, GroupedGemmLauncher, or SUPPORT_MATRIX.md. |
| 8 | N/A | The PR diff does not touch quantization helpers, MX GEMM, or MX scale decoding. |
| 9 | N/A | The PR diff does not add or modify numeric emulation or low-precision decode paths requiring edge-case emulation tests. |
| 10 | N/A | The PR diff does not add a manifest kind or touch helpers/manifest.py, run_manifest, or manifest allocate/verify consumers. |
| 11 | PASS | The SDPA bindings allocate and retain launch buffers once in the bound plan via empty_workspace/device_tensor, launch on resolve_stream(0, device=device), and reject non-current streams to prevent cross-stream races on preallocated workspaces: see… |
| 12 | N/A | The PR diff does not touch COMGR or HIP library loading logic. |
| 13 | N/A | The PR diff adds SDPA routing and binding logic rather than refactoring shared host-side grid, tile-count, workspace, cache-key… |
| 14 | N/A | The PR diff does not touch examples/, helpers/ placement, or add new modules under platform instances/common family glue. |
§Q graded 7/7
| # | Verdict | Evidence |
|---|---|---|
| 1 | PASS | The new SDPA CPU tests under tests/library/tests/sdpa run via the existing rocke_library_sdpa_pytest CTest target, with the two new GPU numeric files excluded from the CPU suite at dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:933-934 and… |
| 2 | N/A | This PR routes SDPA requests onto existing attention candidates and fixes an existing gfx942 4-warp attention kernel without… |
| 3 | PASS | The representative IR golden tripwires are updated for the modified gfx942 4-warp attention kernels at dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8 and… |
| 4 | PASS | The representative golden files are updated across all four LLVM flavor blocks in dnn-providers/hip-kernel-provider/rocke/platform/tests/golden/rocke_representative_ir_sha256/attention_d128_swa.json:8-50 and… |
| 5 | NEEDS-INFO | The full PR diff is now visible and all tests are portable: GPU numeric suites dynamically query the active device arch and skip cleanly without torch or a supported GPU (dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:17, dnn-providers/hip-kernel-provider/rocke/library/tests/sdpa/test_unified_sdpa_numeric.py:33… |
| 6 | N/A | This PR does not refactor any platform validation gate or differential harness. |
| 7 | N/A | This PR does not modify engine rejection contracts or ROCKE_BACKEND / ROCKE_CPP_STRICT fallback behavior. |
§R graded 11/11
| # | Verdict | Evidence |
|---|---|---|
| 1 | N/A | The diff routes SDPA over existing attention candidates and adds a correctness zero-select for fully-masked rows in… |
| 2 | N/A | The only kernel change is the correctness no_key zero-select at `dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/a… |
| 3 | N/A | The diff only adds the epilogue no_key correctness select in `dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/atte… |
| 4 | N/A | The diff adds only the no_key correctness guard in `dnn-providers/hip-kernel-provider/rocke/library/kernels/gfx942/attention_tile… |
| 5 | N/A | The diff does not add or change any global load or store address expression or access width (global_store at… |
| 6 | N/A | The diff does not change an LDS layout, MMA layout, tiling pipeline identity/relabel/reorder, or bank-conflict lever. |
| 7 | N/A | The diff makes no bottleneck claim such as memory-bound, LDS-bound, barrier-bound, or latency-bound. |
| 8 | N/A | The diff does not add, remove, or modify waits, barriers, sched_barrier, sched_group_barrier, iglp_opt, or scheduling hints. |
| 9 | N/A | The diff is not a performance-regression fix and does not add a shape-trading performance heuristic or dispatch rule. |
| 10 | N/A | The diff makes no codegen optimization, assembly-parity, occupancy, or benchmark-timing claims. |
| 11 | N/A | The documentation added in dnn-providers/hip-kernel-provider/rocke/KERNEL_AUTHORING.md:591 is an SDPA routed-family authoring… |
§S graded 8/8
| # | Verdict | Evidence |
|---|---|---|
| 1 | NEEDS-INFO | The new commits (rebase onto develop) do not resolve the missing hunk: torch_dtype and device_tensor are added at dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/runtime/torch_interop.py:96 and tested in dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:7658, and SDPA tests are registered in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:931 and… |
| 2 | PASS | Existing API behavior changes update their tests in the same PR: from_hipdnn requiring 'type' in dnn-providers/hip-kernel-provider/rocke/library/dispatch/sdpa/hipdnn.py:220 is tested at… |
| 3 | N/A | The diff does not remove or rename a public API; _problem, _attention_band, and _validate_lengths were private… |
| 4 | PASS | dispatch.attention.__all__ is updated in dnn-providers/hip-kernel-provider/rocke/library/dispatch/attention/__init__.py:563 and verified by TestAttentionPublicSurface in… |
| 5 | PASS | Tests assert behavioral outputs and exact fail-fast exceptions: torch_dtype and device_tensor in dnn-providers/hip-kernel-provider/rocke/platform/tests/test_rocke.py:7662, non-positive/non-finite unified scale rejection in… |
| 6 | PASS | CPU-checkable contracts run in the host lane rocke_library_sdpa_pytest while ignoring the GPU numeric files in dnn-providers/hip-kernel-provider/rocke/platform/CMakeLists.txt:931, GPU numeric tests are registered in the GPU test list at… |
| 7 | N/A | The public APIs added or modified in this PR are Python dispatch/routing and torch_interop helpers that are Python-only by… |
| 8 | NEEDS-INFO | The new commits do not resolve the prior NEEDS-INFO: the PR description is truncated at LSE is written in place into a contiguous FP32 (B and does not show whether the README test entry points (python tests/test_rocke.py and python tests/run_all.py from rocke/platform/) were stated as run. |
|
Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: high PR #13115 — route SDPA to unified and dense attention kernels (ThruptiRajLakshmanaGowda, AICK-2613, AICK-2614) — 🟠 Follow-up: develop conflict, scale range
Utility: 🟢 Positive — stream, dtype and CI test gaps fixed; merging develop needs two ports Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted Follow-up review: first pass by qwen3.6 (local), verified by the orchestrator (Claude Opus 5.5). Checks: delta since |


Summary
Routes hipDNN SDPA forward requests to the unified attention kernels (AICK-2613) and the dense attention kernels (AICK-2614): for each family an admission gate that never admits a request and then drops an operand, and a binding that launches admitted requests; plus a fix for one gfx942 kernel bug the unified binding's GPU test found (the other two it found are fixed on
developby #14929, which this PR now reuses). Supersedes #14624, which is folded in here.The translation, band comparison, value checks and operand views both families use live in
dispatch/sdpa/bind_shared.py.Admission (
dispatch/sdpa/bind_unified.py)to_attention_request(bind_shared.py) translates anSdpaRequest(from feat(rocke): add hipDNN SDPA request contract and reference oracle #13596) into theAttentionRequestthe unified candidates already consume: batch, heads, sequence lengths, head dims, causal/window bounds (sliding_window = left_bound: since fix(hipdnn): follow cuDNN semantics for the SDPA left diagonal-band bound #14483 hipDNN'sleft_bound = Lkeeps the L keys ending at the diagonal, the countsliding_windowuses), sinks and LSE.SDPA_FEATURESis classified once.SDPA_UNIFIED_FEATURESis what the unified launch path carries. Every other feature is declined with a reason that starts with its name:kv_heads_differ,dtype_mixed,compute_dtype,scale_tensor,bias,bias_broadcast,alibi,sinks(no reference oracle yet),dropout,paged,block_mask,fp8,fp8_s_o,max_sum_exp. Declines for kernel features not implemented yet (scale_tensor,bias,bias_broadcast,alibi,sinks,dropout,paged,block_mask,fp8,fp8_s_o) are taggedNOT_YET_IMPLEMENTED:; the ABI limits are not.mma_core_modemust equal Q's dtype,implementation="composite"is declined,padding_maskwith no seq_len tensor is declined, a pinnedalgorithm,spec_id,tuning_idortuning_knobsis declined (dispatch_pin:), because the unified launch takes no pin, and a non-finite or non-positive scale is declined (scale:): the gfx942 4-warp bodies mask with-1e30before scaling, so a scale <= 0 would bring masked keys back.attention_band, whatattention_problemhands every unified candidate — both exported fromdispatch.attentionfor callers outside the package, rather than aliasing a private name) keeps different keys from the requested one. On gfx942 and gfx1250 full attention comes back bottom-right causal, so it is declined except at decode; this relaxes once those arches gain the full band (feat(rocke): attention band coverage on gfx942 #13211). Outside the band arches (gfx1151, gfx1201, gfx90a, gfx11/gfx12 generics) the problem is always bottom-right causal, so top-left bands with s_q != s_kv are declined there. Bands are compared as per-row kept-key intervals, O(s_q).unified_2d,unified_3d,d256_decode,d256_gfx950,dense_pipe), named fromdispatch.attention.UNIFIED_ALGORITHMS. Both the tiled 2D and split-KV (3D) paths write LSE on gfx942/gfx950 (_reject_lseinkernels/common/attention_unified.pygates each on its own arch/feature rule); admission trusts each candidate's own gate rather than a separate, narrower algorithm allowlist, so an LSE request reaches whichever path the registry actually admits it to.admits()decide.Family entry points (
dispatch/sdpa/__init__.py)registered_sdpa_combos/sdpa_sweep_space/dispatch_sdpa_allkeep #13596'sSDPA_REGISTRYdelegation and add thedispatch.attentioncombos of the family an SDPA gate admits the request to: the unified tuning variants when the unified gate admits it, and the dense candidates the dense gate names. A request both gates decline (bias, dropout, block mask, paged, ...) gets nothing. These results describe attention kernels for ranking and triage: a result'sbind_torch()takes attention-format operands and the attention default scale, so an SDPA request is launched throughsdpa_unified_bind/sdpa_dense_bind, whichdispatch.sdpaexports.Binding (
sdpa_unified_bind(req, tensors) -> TorchBinding)Launches an admitted request through
run_unified_attention_torchwith no tuning spec andbackend="auto": both the tiled 2D and split-KV (3D) paths can write LSE on gfx942/gfx950, so the launch defers to the kernel's own path selection rather than forcing 3D whenever LSE is requested. The unified kernels take packed[total_q, H, D]Q/O and a paged K/V cache, so the binding adapts the hipDNN operands:seq_len_q/seq_len_kv, else the ragged offsets (element start =offset * ragged_offset_multiplier, as hipDNN reads them), else the full extent.align=16Q/O pointers). O is handled the same way and scattered back otherwise.seq_len_*, ragged offsets) must have the dtype its request descriptor names (bind_shared.check_dtype); nothing is truncated or cast on the way in.ceil(S_kv / block)wide, which the scalar fallback requires.(B, H_q, S_q, 1)stats. Rows past a batch's query length get O = 0 and LSE = -inf, as hipDNN's reference writes them.1/sqrt(d).Lengths are snapshotted and buffers allocated once at bind time, so every launch hands the kernel the same tensors (the 2D path replays graphs keyed on them).
launch(*, stream=None)runs pack, kernel and unpack on Torch's current stream; it accepts that stream unset or named by handle and refuses any other (bind_shared.launch_stream), because the copies and the workspace belong to it. The paged ABI is checked withvalidate_tuning_attention_tensors, because the route path checks no shapes itself.Dense (
dispatch/sdpa/bind_dense.py, AICK-2614)sdpa_dense_admits/sdpa_dense_candidates/sdpa_dense_bindroute a request to the gfx942 / gfx950 dense candidates (grid, qb-minor grid, persistent, persistent wide-DMA) with the same rules as the unified side:SDPA_DENSE_FEATURESis what the dense launch carries: GQA, every layout (operands that are not contiguous[B, S, H, D]are copied through buffers allocated once at bind time), top-left and bottom-right causal, windows, lookahead bands and LSE. Declined with a reason:kv_heads_differ,d_v_differs(one head size),dtype_mixed,compute_dtype,scale_tensor,bias,bias_broadcast,alibi,sinks(no reference oracle yet),paddedandragged(the dense candidates take one length per batch),dropout,paged,block_mask,fp8,fp8_s_o,max_sum_exp; and the same value checks as unified.algorithm/spec_idpinned (opt_in_probe); their ownadmits()decide every shape, dtype, head-size and band rule (tile multiples on gfx942, bottom-right only on the gfx950 grid body, windows anchored top-left).causal,sliding_window,right_bound,causal_bottom_right) keeps the keys the request asks for. The gfx950 kernel's scale range[2^-64, 16]is checked here, read from the kernel module, because anAttentionRequestdoes not carry the scale.sdpa_dense_bindlaunches the first serving candidate in registry order (the grid body on gfx950) at its default spec throughbind_dense_attention_torch, with the request's scale, under the same dtype, alignment and stream rules as the unified binding. LSE is written in place into a contiguous FP32(B, H_q, S_q, 1)stats tensor, else through a buffer. Rows with no visible key come back O = 0, LSE = -inf from the kernel.NO_MASKgap (feat(rocke): attention band coverage on gfx942 #13211), not by this module.gfx942 4-warp kernel fix
build_gfx942_4warp_gqa(the D128 sliding-window cohort and the D256 lean body) assumed callers over-allocate Q and the block table, which hipDNN's exactly sized operands don't. On MI325X the unified binding's GPU test found three bugs there. Two of them, the unbounded Q load of a 128-row q-block and the last KV tile reading block-table entries pastceil(klen / BS), are fixed ondevelopby #14929; this PR takes #14929's clamps unchanged and adds no clamp of its own. It fixes the third:qlen > klen) weren't special-casedx_mha_padded_br_win_none_gt,x_gqa4_ragged_br_win_none_gt)Against
developthe kernel diff is that one select in each of the two bodies, so this intentionally changes the IR of exactly the two kernels: only theattention_d128_swaandattention_d256shards of the representative IR golden are re-blessed (--write … --only attention_d128_swa,attention_d256, llvm20/22/23/24, re-recorded on top of #14929 and #15024), and nothing else drifts. MI325X timings of the cohort are unchanged within noise (relative to a 5-round alternating median, 200 attempts each; measured before the #14929 merge):Tests
tests/sdpa/test_unified_sdpa_wiring.py: every feature classified exactly once, one decline case per declined feature, value checks including dispatch pins, band checks on gfx942, gfx950 and four non-band arches, the interval comparison against a dense mask, a 128k band check under 64 MiB, layouts, padded, ragged, LSE path.tests/sdpa/test_unified_sdpa_case_table.py: for everycases.pycase on gfx942 and gfx950, an admitted request reports only carried features and a unified candidate's ownadmits()accepts it; a declined one has a reason. gfx950 admits 405 cases, gfx942 283 (up from 261/202: the N8 fix below reaches the LSE cases a narrower algorithm allowlist used to drop).tests/sdpa/test_unified_sdpa_numeric.py(GPU, gfx942/gfx950): binds every admitted case, all tiers, launches it twice, and compares withsdpa_referenceon the valid rows, fully masked rows included.--ignored byrocke_library_sdpa_pytest, enrolled in_ROCKE_LIBRARY_GPU_TEST_RELPATHS.tests/sdpa/test_unified_sdpa_binding.py(CPU, needs torch): replaces the kernel withsdpa_referenceover exactly the packed and paged tensors the binding produces, across every layout and length mode;sdpa_unified_bindrefuses badseq_len_*values before launch, including a float32 or int64seq_len_q, ragged offsets whose dtype is not the request'sragged_offsetdtype, and FP16 stats against an FP32 descriptor; misaligned BSHD and ragged Q/O go through the workspace;launchruns on the current stream (unset,0or named), refuses another stream and an unknown keyword.tests/sdpa/test_dense_sdpa_wiring.py,test_dense_sdpa_case_table.py,test_dense_sdpa_binding.py(CPU, needs torch) andtest_dense_sdpa_numeric.py(GPU): the same four layers for the dense binding. The case-table test also checks, independently of the gate, that each named candidate's spec keeps exactly the requested keys. The GPU test is--ignored byrocke_library_sdpa_pytestand enrolled in the GPU lane.test_coverage_source_of_truth.py: the unified and dense twins, every SDPA admit is a candidate admit.tests/sdpa/cases.py: 21 tile-multipledense_tile_*cases (S_q a multiple of 256: the lengths gfx942's dense kernel takes), including one with fully masked rows and LSE.tests/sdpa/request_bridge.py,tests/sdpa/binding_harness.py: build theSdpaRequestand the binding's tensors for any case. The bridge lifts the case table's FlashAttention windowWto hipDNN'sleft_bound = W + 1, so every case keeps the same keys; the oracles convert back withhipdnn_left_to_window.tests/sdpa/test_sdpa_family_entry_points.py: the entry points return exactly the gate-admitted families' combos on gfx942 and gfx950 (pinned on named candidates, e.g. gfx942 full attention reaches dense but not the unified tuning variants, gfx950 padded reaches unified but not dense — not just re-derived from the same_serveslogic the entry points use), both the tiled 2D and split-KV unified paths with LSE, and nothing for bias, dropout, block-mask or paged requests.tests/sdpa/test_unified_sdpa_wiring.py::test_sliding_window_keeps_exactly_left_bound_keys:to_attention_request(...).sliding_windowkeeps the L keys the reference oracle keeps, at left_bound = 1, 4, 63, 64 — a CPU pin that fails if a+ 1returns to the translation.tests/sdpa/test_unified_sdpa_wiring.py(torch-free, so CI runs it): the scale decline (0, -1, NaN, ±inf) on gfx942 and gfx950;seq_len_q/seq_len_kvbind-time validation (negative, too large, wrong count, past the ragged lengths) and__all__resolution for thedispatch/sdpamodules.tests/dispatch/attention/test_attention.py::TestAttentionPublicSurface(run byrocke_library_attention_pytest): everydispatch.attention.__all__name resolves and star-imports;__all__equals the listed public surface, so a new name fails until it is added on purpose; no underscored name or submodule is exported; every exported*_ALGORITHM(S)names a registered candidate.Results
On this branch merged with
developat10805993013(including #13596, #14483, #14929 and #15024). Admission: 283 unified / 32 dense on gfx942, 405 unified / 192 dense on gfx950 — unified counts up from 202 (gfx942) / 261 (gfx950) after this round's fix to the unified LSE admission, which was wrongly narrowed to split-KV only (dense is unchanged); see "Admission" and the entry-points test above.gfx942 (MI325X) / gfx950 (MI355X) at
cf2b94adaae, one process each: unified and dense numeric suites (every admitted case, all tiers), both binding tests, both wiring tests and the entry points: 676 / 958 passed, 0 failed, 0 skipped.Host at
cf2b94adaae(CPU torch 2.15):tests/sdpa+tests/dispatch/attention6509 passed, 257 skipped (the skips are the GPU suites). Earlier, after the perf(rocke): runtime heads, window and batch for gfx942/gfx950 tiled attention #15024 merge:tests/sdpa5676 passed, 4 skipped (unified wiring 122, dense wiring 78, unified case table 1433, dense case table 1434, entry points 19, unified binding 74);tests/dispatch/attention805 passed, 253 skipped; coverage source of truth 11 passed;test_sdpa_seqlen.py+test_strided_kv_decode.py+test_attention_tiled_launch_signature.py168 passed, 12 skipped; platformtest_sdpa_reference.py46 passed.IR golden parity OK at llvm20/22/23/24 (
rocke_ir_parity_harness.py --check).pre-commit clean.
Not in this PR
s_kv % block == 0would avoid the copy, but it drops the zero-filled tail the copy gives rows past each sequence's length, so it needs a GPU numeric check first.New public API
dispatch.attentiongains, for callers outside the package (SDPA's gates, in particular) that must not depend on a private symbol:attention_band(req) -> AttentionBand/attention_problem(req) -> UnifiedAttentionProblem: the single home for each (no more_attention_band/_problemalongside a public wrapper).UNIFIED_2D_ALGORITHM,UNIFIED_3D_ALGORITHM,D256_DECODE_ALGORITHM,D256_GFX950_ALGORITHM,GFX942_DENSE_PIPE_ALGORITHM, and the setUNIFIED_ALGORITHMSbuilt from them.dense_scale_range(arch) -> tuple[float, float] | None: the dense kernel's scale range where its module declares one (gfx950:[2^-64, 16]), so the dense gate does not read a private_MIN_SCALE/_MAX_SCALE.DENSE_ALGORITHMS(alongsideUNIFIED_ALGORITHMS): every algorithm a routed family's gate must account for so a new candidate cannot go unreachable by omission;bind_dense.py's_DENSE_ALGORITHMSand the attention route-coverage test both build from it now instead of re-listing the three dense algorithm constants.dispatch.sdpaexportssdpa_unified_admits/sdpa_unified_bindandsdpa_dense_admits/sdpa_dense_bind, the launch path for an SDPA request.rocke.runtime.torch_interopgainstorch_dtypeanddevice_tensor(empty_workspaceandresolve_streamalready existed) — the torch launch-integration glue bothbind_unified.pyandbind_dense.pyuse instead of importingtorchdirectly indispatch/sdpa/.See
KERNEL_AUTHORING.mdAppendix §H ("Route a family overdispatch.attention") for the binding-layer checklist this package's gates follow.JIRA ID : AICK-2613, AICK-2614
🤖 Generated with Claude Code