[amd-khushbu]

JIRA ID : AICK-2712

Summary

This is PR 6 of a stack of 9 that brings the rocKE C++ unified-attention engine to parity with Python (see the stack below). It sits on PR 5.

Python routes gfx942 D256 bf16 prefill, and D128 prefill with a sliding window, to build_gfx942_4warp_gqa. C++ had no port, so the C++ tiled-2D gate rejected that cohort as not ported. This PR ports the builder and its host predicates, and adds a byte-identity family for it.

The builder has no C++ dispatcher caller yet. Only the unified tiled-2D gate and the bindings reach it, as with the other C++ builders.

Changes

All paths are under dnn-providers/hip-kernel-provider/rocke/ unless noted.

• Builder: the new platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp ports four bodies and the online-softmax rescale they share:
  • the D256 lean natural-QK body;
  • the D128 double-buffered pipe, with and without the GQA head-fold;
  • the D128 block_size=64 single-buffer body.
• Host predicates: gfx942_gqa_fold_eligible, gfx942_4warp_launch_grid, and the spec-shaped supports_tiled_2d_spec / supports_gfx942_4warp_gqa_spec wrappers, exposed through rocke_engine.
• Gate: the unified tiled-2D gate admits the gfx942 4-warp cohort, ahead of the ring and mask-limit checks the 4-warp builder does not read. The 4-warp cohort excludes i64 caches, so it never overlaps the i64 + ring reject from PR 5.

Stack

1. feat(hip-kernel-provider): bring the rocKE C++ unified-attention engine to parity with Python #14522 Stopgaps: stop C++ silently misbuilding specs
2. feat(hip-kernel-provider): port rocKE unified-attention selectors, routing and support gates to C++ #14523 Selectors, routing and support gates
3. feat(hip-kernel-provider): assemble rocKE unified-attention specs from a problem in C++ and run selector parity in CTest #14524 Spec from problem (gfx942 2D, 3D), scalar fallback, selector parity in CTest
4. feat(hip-kernel-provider): port the rocKE gfx950 tiled-2D V double-buffer, D256 fast path and softmax interleave to C++ #14525 gfx950 tiled 2D: V double-buffer, D256 fast path, softmax interleave
5. feat(hip-kernel-provider): port the rocKE gfx942 tiled-2D mask-limit, sliced-K ring, phase split and HBM-direct K/V to C++ #14526 gfx942 tiled 2D: mask-limit, sliced-K ring, phase split, HBM-direct K/V
6. feat(hip-kernel-provider): port the rocKE gfx942 4-warp GQA attention builder to C++ #14527 gfx942 4-warp GQA builder ← this PR
7. fix(hip-kernel-provider): rocKE unified-attention C++ parity follow-ups #14528 Follow-ups: gfx942 3D i64 gate, reject messages, emit_lse, CTest labels
8. feat(hip-kernel-provider): close the remaining rocKE unified-attention C++ builder gaps #14758 i64 sliced-K ring, shared arch gate, 3D kv_layout, LDS resolver, gfx950 Q re-read and depth-3 ring, parity coverage
9. feat(hip-kernel-provider): port the rocKE unified-attention mask band and split-KV LSE to C++ #14759 Mask band (top-left / right-bound) and split-KV LSE

Each PR is based on the one before it. Review from 1 upward; each diff shows only its own slice.

Test plan

• New byte-identity family gfx942_attention_4warp_gqa: 9 configs (D256 lean, D128 pipe with and without fold, D128 bs64, D128 causal), built with both engines
• New test_attention_4warp_gqa_cpp_parity.py. It compares with Python every dispatched 4-warp spec over a prefill grid, the builder's rejects, and the host predicates. It is added to the rocke_attention_cpp_parity_pytest CTest entry.
• test_attention_routing_cpp_parity.py and attention_unified_routing_test.cpp: the 4-warp cohort is admitted and builds, and at least one admitted problem goes through the 4-warp builder
• CTest, pybind build: 39/39
• Byte-identity gate GREEN at llvm20, llvm22, llvm23 and llvm24
• Golden: gfx942_attention_4warp_gqa 0–8 blessed; no other attention entry moves
• pre-commit's black (25.12.0) and clang-format (18) hooks, run by hand on every file this PR changes: clean (pre-commit itself is not installed locally)
• GPU numeric lanes not run for this PR; the byte-identity gate is the parity check

🤖 Generated with Claude Code

[CongMa13]

No new issues found.

[CongMa13]

No new P0 or P1 issues were identified during this review. The changes are approved.

[cgmillette]

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max

PR #14527 — port the rocKE gfx942 4-warp GQA attention builder to C++ (amd-khushbu, AICK-2712) — 🟠 First-round: not merge-ready — rebase needed — conflicts inherited from #14523/#14524/#14526; the port matches Python, but no CI lane runs its C++

Prior item	Status
CongMa13 (10-28): "No new issues found", then APPROVED at b79fcfca	🟢 Current — given at head (the PR's only commit); there are no findings or threads to verify. It covers the diff against the stacked base, not the develop conflict below (read)
Stack items: CongMa13's #14523 threads (10-28); our #14522–#14525 first rounds (10-29)	🟠 • #14150 drift (our #14523 N1): carried. The same two attention_unified_selectors.cpp hunks conflict here, and attention_unified_selectors_test.cpp:171 still pins T=32, nw=2, mw=16. • CI coverage (our #14524 N1): extended. This PR adds test_attention_4warp_gqa_cpp_parity.py to rocke_attention_cpp_parity_pytest, which is registered only under ROCKE_BUILD_PYBIND and has no tier label. It also extends the native rocke_attention_unified_routing, which is not in the tier YAML. Neither appears in any shard log at head. #14528 adds both, and four more, to ROCKE_ENGINE_test_categories_external.yaml; the pytest entry still needs a CI build of rocke_engine. • emit_lse (CongMa13): widened. The problem conversion still drops emit_lse, and the C++ gate now admits the 4-warp cohort (attention_unified_gates.cpp:160). So for an emit_lse problem in that cohort, C++ answers "supported" where Python refuses ("LSE output is not implemented on the tiled 2D attention path"); before this PR both refused. No C++ caller uses the gate yet, and #14528 makes every problem binding refuse emit_lse (read). • #13889 copy (our #14522 N1): unchanged. This head conflicts with #13889's head in the same 8 files as #14526's head, with the same 8 golden hunks; this PR adds none (ran: merge-tree). • Not affected here: the Rule 3 getenv knobs (this PR adds no reads), the scalar binding's fp8/mask refusal (our #14524 N2), the gfx942 ALiBi x8 twin (our #14523 N2; the 4-warp cohort excludes score bias), and CongMa13's early-V + staggered-wait race and our K-slab pad item (both gfx950, #14525)
Bots / CI	🟡 • At head b79fcfca, TheRock Multi-Arch CI (10-28) fails only Test hipkernelprovider, on all four shards; an earlier run at this head was cancelled. Linux gfx94X 1/2 and 2/2 fail rocke_differential_golden_ll and _canonical: per LLVM flavor, the same 13 moe_fused_mega_fp8/target_intrinsics configs show unblessed changes (4 in canonical IR), plus 1 new config. Windows gfx110X 1/2 and 2/2 fail those two plus the gfx1100–1103 census entries of TestGfx1151WmmaAttentionDecodePacks, 9 of 9 cases each ("descriptor set 'hipkernel:Gfx1151WmmaAttentionDecode' did not load"). • None of this is the PR's. #14526's head fails the same tests (its Windows 2/2 shard timed out), as did the lower PRs. Both failures are develop breaks in the stack base 0e1cfd3c, which develop fixed in #14501 (golden re-bless) and #14553 (census descriptor root). • The new family passes where CI runs it. CI's golden lanes find every family by its _emit.{c,py} pair but run only the Python side. They see 84 families here against 83 at #14526's head, with 9 more blessed configs at llvm20, llvm22 and llvm23 and in canonical IR (1533 vs 1524 at llvm20), and no new failure (read: CI logs). • Coverage: CI compiles the new .cpp (the cpp/*.cpp glob) but never runs it. The C emitter runs only in the local byte-identity gate, the new pytest needs the pybind module, and the routing ctest is outside the tier YAML. The gfx950 test jobs were skipped, and no GPU numeric lane ran, as the author's checklist says. ASAN, the hipDNN superbuild (Linux, Windows), clang-tidy and Component CI pass; no pre-commit run exists at this head. GitHub reports MERGEABLE (UNSTABLE) against #14526's branch, not against develop (read)
Our checks	🟠 • Merge: git merge-tree with develop 050acb2a (56 behind; 6 ahead: the five lower PRs' commits and this one) leaves 3 conflicted files with 7 hunks, all inherited:   – attention_unified_selectors.cpp: 2 hunks, #14523 against #14150;   – prove_parity_binding.py: 1 hunk, #14524 against #14150;   – llvm_gfx_all.json: 4 hunks, one per section, all #14526's gfx942_attention_tiled_2d 11–19 anchors. #14526's head fb4e2463 gives the identical set and counts, and this PR's new family (9 configs in each of the 4 sections) merges cleanly outside the hunks (ran). With no merged tree, the check script skipped the build, byte identity, goldens and tests. • Drift: since the base, develop has not touched build_gfx942_4warp_gqa (gfx942/attention_tiled_2d.py is unchanged), the cohort predicates, gfx942_gqa_fold_eligible, gfx942_4warp_launch_grid or _get_2d_launcher. #14150, develop's only change to attention_unified.py, does not reach the 4-warp route. So there is nothing to mirror (ran: git). • Goldens: the 9 new anchors reproduce at llvm20, llvm22 and llvm23 and in canonical IR on head's Python. They also reproduce with develop's rocke core in place of head's (run_diff.py --python-only --check-golden --pyroot), so the rebase leaves them valid (ran). • Parity (read): the C++ follows Python op for op. That covers:   – the D256 lean body: V staging, direct-K QK, the causal and varlen masks, PV and the epilogue;   – the general body: the D128 double-buffer prefetch (knext, nbuf_off, loads issued before compute and stored after), the head-fold Q map, mask and inverse output map, the D128 bs64 single-buffer path, the windowed KV start, the three-loop interior split with its clamps, and the swizzle hoist;   – _online_softmax_rescale, including the bperm cross-lane max and sum;   – the reject order and messages (spec validation first, as at Python construction, then arch, i64, dtype, head size, and the lean body's bf16 and window checks) and the LDS guard;   – gfx942_gqa_fold_eligible and gfx942_4warp_launch_grid; the two spec wrappers forward the same 21 fields as _supports_tiled_2d_spec. The gate claim holds. The C++ cohort predicate excludes i64 caches (attention_unified_selectors.cpp:290), so the early admit never meets the i64 + ring reject. It also runs only after the per-arch support check, which both engines make on the same branch. Since #14526 that reject is the only gfx942 "not ported" check left; the mask-limit check the body mentions is gone. Neither builder reads sinks, softcap, bias, fp8 K/V or mask fields; in both engines the cohort predicate and the gate keep such problems away from it. • Evaluation order: none of the 619 calls in the new file passes two or more IR-emitting arguments (ran: scan). • Tests: test_attention_4warp_gqa_cpp_parity.py on head's Python gives 1 passed and 12 skipped without rocke_engine. The pass shows that the routed grid reaches D256, the D128 pipe with and without the fold, and D128 bs64 (ran). No pybind or native build of this head exists here, so the body's byte-identity and CTest results are not re-run. • Rules: 1–3 clean on all 2276 added lines, including the C++: there is no getenv (only a test's monkeypatch.setenv of the LLVM flavor), and no links or tracker IDs (ran)

Net-new: none.

Utility: 🟢 Positive — C++ gfx942 4-warp builder matches Python by reading; CI never runs it.

---

Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

First-round review: a local qwen3.6 model made the independent first pass, and the orchestrator (Claude Opus 5.5) verified it. The PR conflicts with develop, so there is no merged tree yet; the build, byte identity, goldens and test runs wait for the rebase. The orchestrator simulated the merge with git merge-tree for this head and #14526's head, against develop and against #13889's head. It read the ported builder, host predicates, spec wrappers and gate change against their Python twins, and scanned the new file for unsequenced IR-emitting calls. It re-checked the new golden anchors on head's Python and with develop's core, ran the new pytest file on head's Python, and compared the four hipkernelprovider CI logs at head with #14526's. The orchestrator owns this merge-worthiness call. This review flags whether a closer look is needed; it is not an approval or a rejection.

[Copilot]

🟡 Changes recommended

Unresolved critical issues permit out-of-bounds query and K/V reads on partial tiles and division by zero for zero block sizes, risking invalid memory access and process termination.

4 open issues

• Out-of-bounds page-table lookups on ragged key tiles · New
• Out-of-bounds Q loads on ragged final query tiles · New
• Reject nonpositive block sizes before geometry arithmetic · New
• Unordered spec names invalidate unfolded-pipe coverage assertion · New

What changed in this PR

Ports the gfx942 four-warp grouped-query attention builder to rocKE’s C++ engine, advancing Python/C++ parity in the unified-attention stack. C++ dispatcher integration remains separate.

Changes:

• Adds D256 lean and D128 pipelined, head-folded, and single-buffer kernel variants.
• Exposes support and launch predicates through bindings and admits the cohort through the tiled-2D gate.
• Adds byte-identity emitters, golden hashes, routing checks, and CTest coverage.

Paths are relative to dnn-providers/hip-kernel-provider/rocke/.

File	Description
platform/​tests/​instances/​differential/​golden/​llvm_gfx_all.json	Adds golden hashes for nine kernel configurations.
platform/​tests/​instances/​attention_unified_routing_test.cpp	Checks cohort admission and launch grids.
platform/​tests/​CMakeLists.txt	Registers the new Python parity tests.
platform/​cpp/​instances/​gfx942/​attention_gfx942_4warp_gqa.cpp	Implements the builder and host predicates.
platform/​cpp/​instances/​common/​attention_unified_gates.cpp	Removes the cohort’s not-ported rejection.
platform/​cpp/​include/​rocke/​instance_gfx942_attention_4warp_gqa.h	Declares the public C interface.
platform/​cpp/​bindings/​rocke_engine_attention.cpp	Exposes building, lowering, verification, and predicates.
library/​tests/​test_attention_routing_cpp_parity.py	Uses the dedicated builder for admitted four-warp problems.
library/​tests/​test_attention_4warp_gqa_cpp_parity.py	Compares builder output, rejects, and predicates with Python.
library/​tests/​parity/​gfx942_attention_4warp_gqa_emit.py	Adds the Python reference emitter.
library/​tests/​parity/​gfx942_attention_4warp_gqa_emit.c	Adds the matching C parity emitter.

_{Review effort: Max}

[tenpercent]

Review by Codex / GPT-6.1 Sol / high

Assessment: REQUEST CHANGES. Findings rechecked against head 0d2548302b59a507781c094c2a41cea6abc42864, based on stacked base f87527fa33ade872ca46a93d0ee94c7f8725eaf3. The finding-related builder, support wrapper, and bindings are unchanged from the initially reviewed head b79fcfca04b8564514efd0b452c1191fe4f9a22d; intervening selector changes do not prevent these triggers.

1. High, confirmed by source proof — partial query tiles load beyond Q.

   The Q load at line 920 is unconditional. The prologue rejects only wholly empty CTAs; masking output stores does not protect the earlier loads.

   A bf16 D128 problem with 33 query tokens, 32 query heads, 8 KV heads, block size 16, and sliding window 1024 is admitted and builds. Its grid is (8, 2, 1). In the final CTA, folded row 4 loads token 33 at element offset 135168, exactly beyond an exact-sized [33,32,128] Q allocation. The D256 lean body has the same unguarded pattern at line 664.

   This already exists in Python; the PR exposes it through the new public native builder. Guard invalid query loads and supply zero values in both engines, then validate ragged tails with memory-access checks. This is an address proof, not a reproduced GPU fault.

2. Medium, confirmed — the direct builder silently ignores requested attention features.

   The builder validation at line 1323 accepts use_sinks, has_softcap, use_alibi, and use_qq_bias, although the body never consumes those inputs.

   Enabling each flag separately returns supported, construction succeeds, and emitted LLVM is identical to the feature-off kernel after replacing its name. Direct callers therefore receive a kernel that omits the requested semantics. Production selection excludes these cases, so this affects the new public builder and bindings. Python shares the omission.

   Reject these features at the builder and support boundary unless implemented, and mirror the checks in Python.

3. Low, confirmed — D256 sliding-window specs pass support but fail construction.

   The new support wrapper at line 187 reports D256 bf16 sliding-window specs as supported, while the lean builder rejects them as causal-only. This reproduces with use_register_pv=False, allowing shared spec validation to pass. Make the wrapper enforce the builder's window restriction.

The previously reported emit_lse=True support mismatch also remains at this head: Python rejects the tiled-2D request while C++ reports supported. The #14528 follow-up adds rejection of this unrepresented problem field. No C++ dispatcher currently demonstrates runtime LSE loss.

Validation performed on the initially reviewed head b79fcfca04b8:

• Fresh native engine and pybind build; native routing CTest passed.
• Attention C++ parity suite: 78 passed, 1 skipped.
• All nine new configurations passed byte identity and golden checks at LLVM 20, 22, and 23, plus canonical IR.
• 36 additional gfx942/gfx950 combinations passed byte identity and IR verification.
• The checked Linux CI shard had the same golden failures as the stacked base.

The current-head refresh traced the affected source; it did not rerun the build or parity suite after the rebase. No GPU numerical or memory-safety validation was run.

[cgmillette]

Automated review: Model - qwen3.6:35b-a3b-coding-mtp-q4_K_M, effort: max

PR #14527 — port the rocKE gfx942 4-warp GQA attention builder to C++ (amd-khushbu, AICK-2712) — 🟠 Follow-up: merges clean; reviewer findings confirmed

Prior item	Status
CongMa13 (10-28) APPROVED at b79fcfca	🟢 Still counts. The rebase left the PR's patch identical (range-diff =; only the golden-JSON hunk offsets moved) (ran). The approval predates the findings below
Ours (10-29 13:44Z) 🟠: rebase needed; conflicts inherited from #14523 / #14524 / #14526	🟢 FIXED by the stack rebase onto develop 4a3fc208, which already has #14150 and #13889. The merge with develop 7e3d78a3 is clean (7 behind, 6 ahead), and so is the merge with today's 49a6d20f (ran). The #14150 drift is gone: attention_unified_selectors_test.cpp:171 now pins #14150's T=64, and the ctest passes. The #13889 copy dropped out: the merge changes no Python kernel file (ran)
Ours (10-29): no CI lane runs this C++ (extends our #14524 N1)	🟡 OPEN, unchanged. The pybind pytest entry, which now includes the 4-warp file, needs ROCKE_BUILD_PYBIND, which CI never sets. Neither that entry nor rocke_attention_unified_routing is in ROCKE_ENGINE_test_categories_external.yaml; #14528 adds both (ran: grep). Locally the new file passes 13 of 13 on the merge's engine (ran)
CongMa13 (10-28, #14523) / ours (10-29): the C++ gate admits emit_lse	🟡 OPEN, as tenpercent also notes. For a D128 sliding-window or D256 4-warp problem with emit_lse, C++ answers "supported" and Python answers "LSE output is not implemented on the tiled 2D attention path" (ran). No C++ caller uses the gate yet, and #14528 makes the problem bindings refuse emit_lse (read)
Copilot (10-29) attention_gfx942_4warp_gqa.cpp:255, :664, :920; tenpercent (10-29) #1: ragged final tiles read past Q and past the sequence's block-table entries	🟠 OPEN, valid, inherited. No row guard exists. We checked the emitted IR for tenpercent's 33-token D128 case (C++ IR = Python IR; grid (8, 2, 1) in both engines). There, qlen feeds only the early exit, klen − qlen and the 64 store guards; klen feeds only the tile bounds and the 32 score masks. No guard covers any of the 48 Q loads or the 8 block-table loads (ran: use-def). The last CTA reads token rows 33–63 (read). Develop's Python builder emits the same IR (the 9 parity configs hash-equal) (ran), and develop's dispatcher ships it for D256 bf16 causal and D128 sliding-window prefill (read). The C++ copy has no dispatcher caller, so this PR adds a public entry point but no production path → guard the Q rows and the page lookups in both engines (byte identity makes that a single change), and add a ragged-tail test
Copilot (10-29) :1324: reject nonpositive block_size	🟠 OPEN, valid; only C++ crashes. Through the binding, block_size=0 kills the process with SIGFPE, both without a tile size and with tile_size=32. Python raises ZeroDivisionError instead. block_size=-16 builds a _b-16_ kernel in both engines (ran: child processes). The faults are fwg_geometry's BN / BS and, with a tile size, the stack validator's tile_size % block_size (attention_tiled_2d.cpp:668, from #14522) (read) → reject block_size <= 0 first in the shared validator, in both engines
Copilot (10-29) test_attention_4warp_gqa_cpp_parity.py:119: the coverage assertion zips a set with a list	🟡 OPEN, valid. With all 56 unfolded-pipe specs removed from the 112 routed ones, the test still passes, 3 of 3 runs (ran: mutant). The coverage itself exists today, and the routed byte-identity test compares all 112 specs (read)
tenpercent (10-29) #2: the builder ignores use_sinks, has_softcap, use_alibi and use_qq_bias	🟠 OPEN, valid, inherited. In both engines and both bodies (D128 fold, D256), each flag is reported "supported" and builds IR identical to the flag-off kernel. Only the name differs, and it claims the feature (_sinks_, _softcap_, …) (ran). Dispatch never sends these flags (_gfx942_4warp_eligible), so only direct callers are affected (read) → reject the flags in the builder and the support wrapper, in both engines
tenpercent (10-29) #3: D256 sliding-window specs pass the support wrapper, then fail to build	🟡 OPEN, valid, inherited. Both engines answer "supported", then raise "causal-only" (ran). Dispatch never routes D256 with a window, so only pack-time checks are misled
Bots / CI	🟢 None of the red checks is this PR's. At 0d254830, every workflow ran twice at 18:19Z, and the first copy of each was cancelled within 2 s. The four red checks are those copies' summaries: Multi-Arch CI, its ASAN twin, clang-tidy and the hipDNN superbuild. No job under them ran (ran: gh jobs). The live copies pass: ASAN, clang-tidy, both hipDNN superbuilds and Component CI. TheRock Multi-Arch CI is still building: Linux compiler-runtime started at 19:26Z, and Windows gfx110X math-libs has been running since 18:52Z. So no hipkernelprovider test has run at this head yet. The old head's develop reds were fixed by #14501 and #14553, which are now below the base. GitHub reports MERGEABLE (UNSTABLE) against #14526's branch. tenpercent's REQUEST CHANGES is a conversation comment, not a review, so the review decision stays APPROVED. No author reply yet on any thread or comment (ran: gh)
Our checks	🟢 • Build: native and pybind builds OK. 37 ctests pass, with the usual rocke_arch_domain_reproducibility skip (ran). • IR: byte identity is GREEN at llvm20 / 22 / 23 for all 10 families the gate covers, including gfx942_attention_4warp_gqa (9 of 9). The representative golden passes ×3. The differential golden is OK, with nothing changed or missing; its unblessed configs are develop's (ran). • Tests: the PR's 6 test files on the merge give 178 passed, 1 skipped and 2 failed. Both failures fail on develop too: test_selector_parity.py's graph-replay case needs torch, and TestKvAddrLimit is #14651 (ran). • Compile (ROCm clang 24, -O3, on the Python-emitted IR): gfx942_attention_4warp_gqa compiles 9 of 9, and 6 spill. Develop's Python emits the same IR for all nine (ran: hashes), so the spills belong to the shipped kernels, not to this port. The 1 spill in gfx950_attention_tiled_2d is in a config added lower in the stack (#14524's checks show it). The gfx950 Python builder is unchanged, and this PR touches no gfx950 file. gfx942 3D cfg 2 fails on develop too (ran). • Rules 1–3: clean (ran). • Land order: this is PR 6 of 7. It lands after #14526 and before #14528, which closes the coverage and emit_lse rows

Net-new: none.

Utility: 🟢 Positive — byte-identical C++ port of the shipped 4-warp builder, inherited bounds gaps included.

---

Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good · ⚪ retracted

Follow-up review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it against our 10-29 13:44Z review at b79fcfca, the rebase to 0d254830, Copilot's four threads and tenpercent's comment. On the merge with develop 7e3d78a3, the checks covered the native and pybind builds and the ctests, byte identity at three LLVM flavors, the representative and differential goldens, the PR's test files, a ROCm clang 24 compile of the touched families, and a scan of the diff. Host probes on the merge's engine and Python covered zero and negative block sizes (in child processes), the four ignored flags, the D256 window wrapper, emit_lse, a mutant of the coverage test, the uses of qlen and klen in the emitted IR, and the 4-warp IR compared with develop's Python. No GPU was available, so the out-of-bounds reads rest on the IR and the address arithmetic, not on an observed fault. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

[yraparti]

rocKE checklist review

ROCm/rocm-libraries#14527 · 0d2548302b59 · model gemini-4-argon, fallback gpt-5.5

PASS · graded 238/238 · BLOCKER 0, HIGH 1, MEDIUM 2, LOW 0, NEEDS-INFO 30

The diff was truncated before grading. Items that needed a missing hunk are NEEDS-INFO.

Findings

• HIGH §S · dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:91 · Add a CTest unit test (e.g. in platform/tests/instances/) or route t942_4w_lower / the C parity emitter through rocke_gfx942_attention_4warp_gqa_lower_to_llvm so the convenience C entry point is exercised in CI.
• MEDIUM §F · summary · Retitle as type(scope): summary, 10–80 characters, scope [a-z0-9-]+.
• MEDIUM §O · dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:544 · Remove the trailing rocke_b_ret(b); at the end of fwg_epilogue in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:544 (and its Python counterpart if present).

Needs info

• §B.1 dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 811 before rocke_build_gfx942_4warp_gqa, where the builder's rejections (ROCKE_ERR_NOTIMPL vs ROCKE_ERR_VALUE) are implemented.
• §B.4 dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 811 before rocke_build_gfx942_4warp_gqa, so the C ABI entrypoint's input validation cannot be inspected.
• §C.2 The definition of rocke_build_gfx942_4warp_gqa in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is beyond the diff truncation at line 825 (1419 total lines).
• §C.3 The builder entry point rocke_build_gfx942_4warp_gqa and its kernel_name construction in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp are beyond the diff truncation at line 825.
• §E.8 No tracking IDs, handles, or PR numbers appear in the visible diff (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:1-102), but dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 732 of 1419.
• §E.15 All parameters in visible functions are used (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:113-199), but lines 733-1419 of attention_gfx942_4warp_gqa.cpp are truncated from the diff.
• §F.6 The PR body notes black and clang-format were run by hand (pre-commit itself is not installed locally), and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 811, so full formatting compliance cannot be verified from the diff alone.
• §F.7 Bot and Copilot review threads and their resolution status are not included in the provided diff context.

Section notes

• §B: attention_gfx942_4warp_gqa.cpp is truncated before the builder entrypoint rocke_build_gfx942_4warp_gqa, so builder-side rejection and validation hunks are not visible in the diff.
• §C: attention_gfx942_4warp_gqa.cpp is truncated at line 825 (total 1419 lines), hiding the general builder body and public entry point rocke_build_gfx942_4warp_gqa.
• §D: This PR ports an existing Python kernel builder (build_gfx942_4warp_gqa) and its C99 host predicates to C++ for engine parity, without adding or modifying Python dispatcher candidates, specs, or routing.
• §D: C++ builder and host-predicate port for gfx942 4-warp GQA; Python dispatcher registry is not modified.
• §E: attention_gfx942_4warp_gqa.cpp is truncated after line 732; items requiring inspection of the remaining function bodies are marked NEEDS-INFO.
• §F: Merge-readiness / PR-bot policy

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	PASS	Ports build_gfx942_4warp_gqa to C++ in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:1 and pairs it with parity emitters in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:1 and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.py:1.
2	N/A	No core IR or lowering op-handlers (core/lower_*, core/ir) are modified in this PR.
3	PASS	Host predicates and spec support wrappers (rocke_gfx942_gqa_fold_eligible, rocke_gfx942_4warp_launch_grid, rocke_gfx942_attention_tiled_2d_supports_spec, rocke_gfx942_attention_4warp_gqa_supports_spec) are ported in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:107 and tested against Python in…
4	PASS	In dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:217 (and throughout, e.g., lines 522-569), C++ expressions are decomposed into sequential statements binding operands left-to-right before combining, and byte-identity with Python is verified in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:123.
5	PASS	Existing attention golden hashes are unchanged and only the new gfx942_attention_4warp_gqa family (indices 0-8) is added in dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222.
6	PASS	Both the cross-engine parity emitter pair (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:109 and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.py:148) and direct Python-vs-C++ lowering assertions (dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:127) are included alongside the golden hashes.
7	PASS	Python kernel emitters are untouched, and in C++ the gate change is guarded by rocke_unified_attn_gfx942_4warp_fast(p, arch) in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:160.
8	PASS	C++ parity is wired into CTest in dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:691 and dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:368, and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:1 provides the C byte-identity harness.
9	N/A	This PR ports an instance kernel builder rather than fixing a shared core op-handler or arch lowerer divergence.
10	PASS	Golden hash entries for gfx942_attention_4warp_gqa configs 0-8 across LLVM flavors are added in dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222 (and lines 2925, 4627).
11	N/A	No new IR opcodes are introduced in this PR.
12	N/A	core/backend.py and ROCKE_BACKEND fallback logic are not touched in this PR.
13	PASS	The C emitter lowers via rocke_lower_kernel_to_llvm without stamps or debug locations in dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:123.
14	N/A	No classic helpers under helpers/ or platform/cpp/helpers/ are modified in this PR.

§B graded 6/6

#	Verdict	Evidence
1	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 811 before rocke_build_gfx942_4warp_gqa, where the builder's rejections (ROCKE_ERR_NOTIMPL vs ROCKE_ERR_VALUE) are implemented.
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:351 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:184 verify that unified_attention_supports_tiled_2d and gfx942_attention_4warp_gqa_supports_spec match the spec and builder across gfx942 and gfx90a.
3	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:160 removes the C++ not-ported rejection for rocke_unified_attn_gfx942_4warp_fast(p, arch) so the 4-warp GQA cohort is admitted rather than demoted.
4	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 811 before rocke_build_gfx942_4warp_gqa, so the C ABI entrypoint's input validation cannot be inspected.
5	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:160 opens the rocke_unified_attn_gfx942_4warp_fast cohort, and dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:123 plus dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:362 test all routed specs in that cohort.
6	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:139 (fwg_supports_spec) delegates directly to the functional rocke_gfx942_attention_tiled_2d_supports check over spec geometry/capability fields, while routing policy stays in rocke_unified_attn_gfx942_4warp_fast.

§C graded 10/10

#	Verdict	Evidence
1	PASS	Shared predicates and sub-operations are factored into single named helpers, e.g. rocke_gfx942_gqa_fold_eligible called by rocke_gfx942_4warp_launch_grid at dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:112-127, fwg_supports_spec at :141-192, and fwg_online_softmax_rescale at :312-358.
2	NEEDS-INFO	The definition of rocke_build_gfx942_4warp_gqa in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is beyond the diff truncation at line 825 (1419 total lines).
3	NEEDS-INFO	The builder entry point rocke_build_gfx942_4warp_gqa and its kernel_name construction in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp are beyond the diff truncation at line 825.
4	PASS	The gfx942-specific 4-warp GQA builder lives in its own arch module dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:1, while common gate logic in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:156-162 only evaluates target predicates.
5	PASS	All 11 touched files are directly scoped to porting the gfx942 4-warp GQA C++ builder, its bindings, gates, parity emitters, tests, and golden hashes (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:1 and dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:691).
6	PASS	The PR is a structured slice (PR 6 of 7) cleanly porting the C++ builder and parity tests with shared helpers like t942_4w_build in dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:1021-1035 and fwg_supports_spec in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:141-192.
7	N/A	This PR ports a missing C++ kernel builder rather than performing a no-op cleanup or refactoring.
8	PASS	The C++ builder reuses existing platform helpers such as rocke_binary_search_seq_idx, rocke_decode_mfma_lanes, and rocke_transforms_descriptor_offset in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:25-28, :298, :433, :449.
9	N/A	This PR is a feature port of an existing Python builder to C++, not a bug fix.
10	PASS	Widening tiled-2D C++ support to admit the gfx942 4-warp GQA cohort in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:158-161 is updated consistently across all consumers in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:254-256, :362-366 and dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:273.

§D graded 21/21

#	Verdict	Evidence
1	N/A	This PR ports an existing kernel builder to C++ and does not introduce any new KernelSpec fields, knobs, or environment variable reads.
2	N/A	This PR does not declare or modify any Python dispatcher Capability objects.
3	N/A	This PR does not add or modify any Python dispatcher candidate _supports predicates; it only ports C++ instance support wrappers.
4	PASS	The C++ support wrappers and builder take arch explicitly as a parameter (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:128 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:82).
5	N/A	This PR does not modify Python KernelSpec identity, spec_hash, compile_key, or kernel_name() logic on the Python side (and the C++ kernel_name suffix construction in attention_gfx942_4warp_gqa.cpp is in the truncated…
6	PASS	In C++, the builder and launch grid share rocke_gfx942_gqa_fold_eligible (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:96-119) to keep launch-meta and kernel body in agreement, tested in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:161-180.
7	N/A	This PR does not add or modify hand-pinned overrides or machine-regenerable heuristics in the dispatcher.
8	PASS	The C++ unified tiled-2D gate reuses rocke_unified_attn_gfx942_4warp_fast (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:160) and rocke_gfx942_gqa_fold_eligible is shared between the launch grid and builder (dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:48-53).
9	N/A	This PR ports an existing Python builder to C++ and does not register a new KernelCandidate in the Python dispatcher registry.
10	N/A	This PR does not add any new spec fields or features.
11	N/A	This PR does not register or promote any dispatcher candidates in the Python routing registries.
12	N/A	This PR ports the C++ gfx942 4-warp GQA builder and does not register, rename, or reprioritize any Python dispatcher candidate.
13	N/A	This PR does not touch CandidateRegistry.select, pinned spec_id handling, or any Python dispatcher wrapper.
14	N/A	This PR does not modify compiled-binary caches, KernelId, or candidate abi_version.
15	N/A	This PR does not add or modify a candidate in the Python dispatcher registry.
16	N/A	This PR does not introduce a new operator family in the Python dispatcher.
17	N/A	This PR does not add a Python dispatcher candidate, so Python admits() CPU routing tests do not apply (C++ gate/predicate parity is tested instead).
18	N/A	This PR does not touch Python CandidateRegistry ratchets or exempt sets.
19	PASS	The PR body explicitly states that no C++ dispatcher caller exists yet and dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:94 tests against _tiled_spec_from_problem(p, _ARCH) without over-claiming dispatcher registration.
20	PASS	The C++ tiled-2D gate and host predicates (rocke_gfx942_gqa_fold_eligible, rocke_gfx942_4warp_launch_grid, supports_*_spec) are mirrored in C++ at dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:160 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:113 and tested against Python in…
21	N/A	This PR ports an existing Python kernel builder to C++ for engine parity rather than working around a spec-selection bug in an emitter.

§E graded 17/17

#	Verdict	Evidence
1	PASS	Platform files do not import library; platform/tests/CMakeLists.txt:691 only registers a pytest runner invoking library tests with WORKING_DIRECTORY set to the library tests directory.
2	PASS	No torch imports appear in the added library test or parity files (dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:19-35).
3	PASS	Torch is not imported anywhere in the diff, and optional rocke_engine is gated with pytest.importorskip at dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:43.
4	PASS	Library test files only import downward from kernels.* and rocke.core.* (dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:23-35, dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.py:15-19).
5	PASS	CMakeLists.txt uses relative CMake variables (${_rocke_library_tests}) at dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:691 and no shell scripts or absolute paths are added.
6	N/A	This PR does not modify Python emitter trees under core/, helpers/, instances/, or library/kernels/.
7	N/A	This PR ports an existing Python gfx942 4-warp GQA builder to C++ for engine parity without adding new env vars, dtypes, ops, kernel families, manifest kinds, or perf levers.
8	NEEDS-INFO	No tracking IDs, handles, or PR numbers appear in the visible diff (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:1-102), but dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 732 of 1419.
9	PASS	All touched files are strictly scoped to porting the gfx942 4-warp GQA builder, its predicates/bindings, and its parity tests/goldens (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:1-102).
10	PASS	This PR is explicitly part of a 7-PR stack with a defined land order and modifies only the 4-warp GQA slice in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:158-161.
11	PASS	Only deliverable C++/Python source, tests, CMake, and golden hashes are committed (e.g. dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222-1232).
12	PASS	The kernel is explicitly gfx942-specific and bothgfx942/gfx950 routing and non-gfx942 rejection are covered in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:357 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:147.
13	N/A	This PR does not remove dead code or rename any existing concept or symbol.
14	PASS	The only removed lines in the diff are the obsolete 'not ported to C++ yet' stub checks in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:159-162, dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:255, and dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:272.
15	NEEDS-INFO	All parameters in visible functions are used (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:113-199), but lines 733-1419 of attention_gfx942_4warp_gqa.cpp are truncated from the diff.
16	N/A	This item is a meta-rule for review grading honesty and no FAILs were softened.
17	PASS	No performance claims or numbers appear in the PR description or visible code (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:1-32).

§F graded 11/11

#	Verdict	Evidence
1	PASS	branch users/khuagarw/rocke/aick-2712-cpp-parity-6-gfx942-4warp-gqa uses the users// prefix
2	FAIL	title feat(hip-kernel-provider): port the rocKE gfx942 4-warp GQA attention builder to C++ is 84 characters and does not match type(scope): summary with a 10–80 character length
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/test_attention_4warp_gqa_cpp_parity.py is in the diff
5	PASS	No keys, certificates, or secrets appear in the diff (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:1 through dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:4640 are all source/test/golden files).
6	NEEDS-INFO	The PR body notes black and clang-format were run by hand (pre-commit itself is not installed locally), and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 811, so full formatting compliance cannot be verified from the diff alone.
7	NEEDS-INFO	Bot and Copilot review threads and their resolution status are not included in the provided diff context.
8	NEEDS-INFO	This is a stacked PR targeting users/khuagarw/rocke/aick-2712-cpp-parity-5-gfx942-kernel-body, and the PR title in metadata is 84 chars (>80 limit in item 2), but PR metadata has no diff path:line to cite.
9	N/A	This item is guidance that PR-size caps are not enforced; no artificial splitting or trimming for a size cap is applicable to grade here.
10	PASS	The added tests are substantive parity and routing tests rather than throwaway placeholders (dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:111, dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:337).
11	NEEDS-INFO	The PR title and body are populated in the prompt metadata, but there is no diff path:line for GitHub API metadata verification.

§G graded 34/34

#	Verdict	Evidence
1	PASS	The PR ports an existing Python builder to C++ and frames byte-identity strictly as a parity check against the existing Python emitter (dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:3-15, dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:4-11).
2	N/A	This PR ports an existing Python kernel builder to C++ with byte-identical LLVM IR emission rather than authoring or modifying a GPU kernel algorithm.
3	PASS	The PR updates the golden hashes across LLVM flavors in dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222-1232, registers the new parity test in dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:691, and reports CTest 39/39 and byte-identity green across llvm20/22/23/24.
4	PASS	New tests assert exact LLVM IR equality, exact reject messages, and exact grid tuples that would fail without the C++ implementation (dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:122-158, dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:273).
5	PASS	Each test in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:99-158 is falsifiable by a single change and explicitly separates _routed_specs() from _off_route_specs() (with docstring explaining why off-route specs are tested and how _reject_cases() bypasses upstream register-PV validation to reach builder rejects).
6	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 656, so code comments and control-flow claims in the remainder of the builder cannot be verified against the implementation.
7	N/A	This PR does not add or modify a design document.
8	N/A	No GPU execution or numerical benchmark harnesses are added or modified in this PR.
9	N/A	This PR is a byte-identical C++ port of an existing Python builder and does not introduce new LDS or occupancy claims.
10	N/A	This PR makes no performance measurements or speedup claims.
11	N/A	This PR is a C++ parity port of an existing Python builder rather than a kernel optimization PR, and explicitly notes the C++ dispatcher caller is not wired yet.
12	PASS	The PR explicitly documents who calls the new C++ builder now (only the unified tiled-2D gate and pybind bindings) and tests both surfaces in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:351-372 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:1848.
13	N/A	No new LLVM intrinsic overloads or mangled intrinsic symbols are introduced or recommended in this PR.
14	N/A	This PR does not classify hardware bottlenecks or make performance claims.
15	N/A	This is a review workflow instruction rather than a code property of the PR diff.
16	N/A	This PR implements code and tests directly rather than proposing an unverified plan.
17	PASS	The new parity emitter script dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.py:155-165 was executed end-to-end across configs 0-8 to produce and bless the hashes in dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222-1232.
18	N/A	This PR ports an existing Python attention builder to C++ for byte-identity parity and does not add a statistical or back-pressure guard.
19	PASS	In dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:132-148 and :209-229, fixtures construct reject shapes and spec-field variations, and dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:372 asserts that at least one 4-warp problem was admitted.
20	N/A	This PR does not add or modify a text detector for banning performance claims.
21	N/A	This PR does not add measurement or profiling code.
22	PASS	New golden entries in dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222-1232, :2925-2935, and :4627-4637 record valid sha256 hashes (no build errors), and PR description confirms CTest pybind build ran 39/39.
23	N/A	This PR is a C++ parity port of an existing Python kernel builder and makes no performance or new routing speedup claims.
24	N/A	No existing golden hashes were re-blessed; only the new gfx942_attention_4warp_gqa family was added in dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222.
25	N/A	This PR ports an existing Python kernel to C++ with byte-identical LLVM IR rather than introducing a codegen micro-optimization.
26	N/A	This PR does not include benchmarks or performance claims.
27	N/A	This PR does not include A/B performance comparisons.
28	N/A	This PR does not report performance deltas or counterfactual A/B tables.
29	PASS	In dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:324 and :331, rocke_b_exp2_fast is used only in online softmax rescale on (m_old - m_new) and (s - m_new), where m_new is the running maximum so the argument is provably <= 0.
30	N/A	This PR makes no register-spill or scratch-reduction claims.
31	N/A	This PR does not claim to fix a performance defect or report before/after performance baselines.
32	N/A	This PR does not report measured hardware performance or resource numbers.
33	PASS	The ported predicates and routing gates are differentially tested across their input domains (including unrouted arches like gfx90a and gfx950) against the Python reference in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:161-206 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:351-372.
34	N/A	This PR does not add or modify an opt-in kernel output flag such as emit_lse.

§H graded 5/5

#	Verdict	Evidence
1	N/A	This PR ports a C++ IR builder and static spec/grid predicates with no runtime tensor launch path or AMDGPU buffer-SRD num_records setup.
2	NEEDS-INFO	In dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:240-255, fwg_phys_key and fwg_kv_elem index BT, Kp, and Vp via unguarded global pointer loads (rocke_b_global_load_i32 / rocke_b_global_load_vN), while the file is truncated after line 860 and the PR has no C++ launch/dispatcher caller yet to confirm host-side block_tables value validation.
3	PASS	Hot-loop divisors/moduli on BS, HD, FWG_FOLD_HEADS (4), 8, 16, 32, 64 in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:237-244 and 675-691 are powers of two, enforced by rocke_gfx942_attention_4warp_gqa_supports_spec (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:176-181) which delegates to…
4	NEEDS-INFO	The lean body masks every visited tile with m_causal || m_varlen (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:598-615), but fwg_body in the general body has an unmasked branch (if(!masked), lines 817-821) whose caller and loop bounds are in the truncated portion of dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp after line 860.
5	N/A	This PR only adds the C++ IR builder and static problem/spec helpers (rocke_build_gfx942_4warp_gqa, rocke_gfx942_4warp_launch_grid, rocke_gfx942_gqa_fold_eligible) and does not add or modify a kernel launch…

§I graded 6/6

#	Verdict	Evidence
1	N/A	This PR is a C++ parity port of an existing Python kernel builder and does not modify routing search space or benchmarking measurements.
2	PASS	The C++ builder and bindings loudly reject unsupported specs and architectures with descriptive error messages matching Python (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:69, dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:152).
3	PASS	Parity test coverage combines curated pin configs in the byte-identity harness (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.py:30) and generated Cartesian sweeps over prefill shapes and off-route specs (dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:68).
4	N/A	This PR does not include performance benchmark tables or routing evaluation reports.
5	N/A	This PR does not add or modify baked tuning snapshots or search results.
6	N/A	This PR ports an instance-level kernel builder using existing IR builder primitives and introduces no new hardware-named IR ops or DSL primitives.

§J graded 7/7

#	Verdict	Evidence
1	N/A	The diff contains no Markdown or design/algorithm documents.
2	N/A	The diff contains no Markdown or design/algorithm documents.
3	N/A	The diff contains no Markdown or design/algorithm documents.
4	N/A	The diff contains no Markdown or design/algorithm documents.
5	N/A	The diff contains no Markdown or design/algorithm documents.
6	N/A	The diff contains no Markdown or design/algorithm documents.
7	N/A	The diff contains no Markdown or design/algorithm documents.

§K graded 13/13

#	Verdict	Evidence
1	N/A	The diff does not add or modify a tuned dispatch candidate or KnobSpace-based candidate construction.
2	N/A	The diff does not add or modify tunable kernel fields, KnobAxis declarations, or tuning-space classification tests.
3	N/A	The diff does not change tuning identity payloads, canonical knob rules, declared defaults, or TUNING_ID_VERSION.
4	N/A	The diff does not add or modify KnobSpace.defaults() or recorded tuning fields.
5	N/A	The diff does not modify tuning-space drop, inert, base_value, refuse, or known-wrong knob handling.
6	N/A	The diff does not modify KnobSpace.validate or policy-free tuning spec construction for sweeps and pins.
7	N/A	The diff does not add or modify a KnobSpace base(req) implementation.
8	N/A	The diff does not modify request-boundary knob normalization or JSON scalar knob validation.
9	N/A	The diff does not modify benchmark replay, tuning_id recording, spec_hash matching, or PinRefused handling.
10	N/A	The diff does not modify tuning_id lookup or KnobSpace.find behavior.
11	N/A	The diff does not modify sweep entry points, sweep sampling, sweep levels, or sweep harness iteration.
12	N/A	The diff does not reorder a production tuning stack, change accept_default, or edit a variant base spec used for tuning auto selection.
13	N/A	The diff does not modify benchmark legality decisions, sweep filtering, crash recording, or priority ranking.

§L graded 10/10

#	Verdict	Evidence
1	N/A	This PR does not add or modify a gfx target.
2	N/A	This PR does not add or modify an MMA op_id or fragment layout.
3	NEEDS-INFO	attention_gfx942_4warp_gqa.cpp is truncated after line 868; while rocke_mfma_atom("bf16", 32, 32, 8) is used at line 552, fwg_lane_to_output at line 284 manually computes lane-to-output coordinates and the general body's atom selection is in the truncated portion of the diff.
4	NEEDS-INFO	attention_gfx942_4warp_gqa.cpp is truncated after line 868, and the visible part (fwg_lanes at line 424) uses hardcoded ci(c, 64) for gfx942 wave64 math without showing the remainder of the file.
5	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:77 gates arch capabilities properly via fwg_require_tiled_attention_arch and does not leak gfx950 ds_read_tr* or wide-K atoms onto gfx942.
6	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:1 places the instance builder under instances/gfx942/ without mutating ArchTarget or importing CK Tile dispatcher code.
7	N/A	This PR does not add or modify dtypes, bit packing formats, or FP8/FP4/FP6 dialect handling.
8	N/A	This PR does not use or modify scaled MMA operations or scale contracts.
9	N/A	This PR does not add a new architecture target.
10	N/A	This PR does not touch intrinsic arch-domain artifacts or toolchain flavor membership lists.

§M graded 10/10

#	Verdict	Evidence
1	PASS	Runtime conditions use rocke_b_cmp_* with rocke_b_scf_if / rocke_b_select and loop state is carried via rocke_b_scf_for_iter / rocke_b_scf_yield in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:478 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:604.
2	N/A	This PR does not modify legalize_for_target or target legalization passes.
3	N/A	This PR does not modify ck.dsl.ir/v1 serialization or print_ir.
4	N/A	This PR does not touch portable IR schemas, recipe_expand.py, recipe_vm.cpp, or record_coverage.py.
5	N/A	The ported 4-warp GQA kernel stages tiles via rocke_b_global_load_vN and rocke_b_smem_store_vN rather than AsyncTileLoader.
6	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 889 before the rest of fwg_body and rocke_build_gfx942_4warp_gqa, so barrier placement and pointer clamping/loading in the second half of the file cannot be verified.
7	N/A	This PR does not touch ISABackend.buffer_rsrc_word3 or TemporalHint usage.
8	N/A	This PR does not add LLVM intrinsic declarations or modify backend selection precedence.
9	PASS	The new public C header dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:45 uses extern "C" declarations for all C entry points (rocke_build_gfx942_4warp_gqa, rocke_gfx942_attention_4warp_gqa_lower_to_llvm, etc.) and does not rename existing headers or fabricate HIP launches.
10	N/A	This PR does not touch !dbg emission, Debug Info Version module flags, or DWARF consumers.

§N graded 18/18

#	Verdict	Evidence
1	N/A	This PR does not add or modify any helpers under platform/python/rocke/helpers/.
2	PASS	The C++ builder port uses shared C++ helper twins (rocke_mfma_atom, rocke_decode_mfma_lanes, rocke_binary_search_seq_idx, rocke_tensor_descriptor_naive, rocke_transforms_descriptor_offset) in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:25 and :286, :367, :383.
3	N/A	This PR does not use or modify SoftwarePipeline.
4	N/A	This PR does not add or modify fused kernels or EpilogueOp.
5	N/A	This PR does not add or modify any public tiling API under helpers/tiling/.
6	N/A	This PR does not add or modify front-door/toolbox tiling helpers or LayoutStyle.
7	N/A	This PR does not use helpers/tiling/ or layout overrides.
8	N/A	This PR does not use or modify helpers/tiling/ memory verbs (lds_tile_alloc, lds_store, lds_read).
9	N/A	This PR does not use or modify helpers/tiling/ N-D tensor axis roles or FlowStage.
10	N/A	This PR does not modify tiling emit, LDS staging, or MMA issue in helpers/tiling/.
11	N/A	This PR does not touch helpers/tiling/ traits, lds_swizzle policies, or derive_sched_group_counts.
12	N/A	This PR makes no LDS bank-conflict or coalescing claims and does not touch helpers/tiling/.
13	N/A	This PR does not touch WarpDistributionEncoding, make_tile_desc, or RegisterMapper.
14	N/A	This PR does not use helpers/tiling/ verbs or add a new tiling pipeline variant.
15	N/A	This PR does not modify helpers/tiling/ allocation or bounds checks.
16	N/A	This PR does not use or modify optimize_layout or evaluate_transform.
17	N/A	This PR does not touch the tiling recorder, round-trip analysis, or visualization.
18	N/A	This PR does not touch InterleavedStyle or demo kernels under helpers/tiling/kernels/.

§O graded 16/16

#	Verdict	Evidence
1	PASS	Pointer parameters use valid address space "global" (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:365-367), and verification is exposed and wired into the parity emitter (dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:1057-1064, dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:146-161).
2	FAIL	fwg_epilogue appends a trailing rocke_b_ret(b); at the end of the kernel body (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:544), violating the rule not to append a cf.return to terminate a kernel because the lowerer adds ret void.
3	PASS	C++ builder statements are sequenced explicitly to preserve deterministic SSA emission order matching Python (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:13-15, 237-242).
4	PASS	IR types are constructed via singleton/helper functions rocke_i32(), rocke_f32(), and rocke_attention_tiled_2d_spec_dtype_ir(spec) (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:271-274, 366-367, 557).
5	N/A	This PR does not modify core/passes.py, cpp/core/passes.cpp, or op purity tables.
6	N/A	This PR does not modify constant folding in the Python or C++ optimization passes.
7	PASS	Kernel attributes max_workgroup_size (256) and optional waves_per_eu are set cleanly via fwg_kernel_attrs (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:458-463).
8	N/A	This PR does not set or modify CodegenPolicy or SchedulerStrategy.
9	N/A	This PR does not touch waitcnt encoding, gfx950 split barriers, or optimization_barrier.
10	N/A	This PR does not modify the LLVM lowerer or parameter/attribute validation in lower_llvm.
11	PASS	Target capabilities are queried via rocke_arch_target_from_gfx(arch) and rocke_mma_catalog_has_shape (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:69-85).
12	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp is truncated at line 882, so the remainder of the general builder body and entry point rocke_build_gfx942_4warp_gqa (including LDS capacity checks and default arch handling) cannot be inspected.
13	PASS	The new public header dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:1-102 declares only new functions and reuses existing structs without growing or reordering any public C ABI struct or enum.
14	PASS	The new header and source follow style/CPP_STYLE.md with #ifndef ROCKE_INSTANCE_GFX942_ATTENTION_4WARP_GQA_H guards, no #pragma once, a Python-to-C correspondence table, and rocke_* snake_case symbols (dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:8-34, 52-96), and bindings are added in the attention family file…
15	PASS	No quiet fallback paths are added, and the C++ binding resolves the Python LLVM flavor via resolve_python_llvm_flavor() (dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:1044-1045).
16	PASS	The C++ builder stays codegen-only, mirrors Python's NotImplementedError rejections via ROCKE_ERR_NOTIMPL (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:74, 102), and verifies exact reject parity in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:151-158.

§P graded 14/14

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:82 declares rocke_build_gfx942_4warp_gqa(b, spec, arch) taking only the builder, spec, and arch.
2	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:324 emits the exact 18-parameter production paged-attention ABI matching Python.
3	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:113 ports rocke_gfx942_4warp_launch_grid using rocke_gfx942_gqa_fold_eligible as the single source of truth, verified against Python in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:174.
4	NEEDS-INFO	The LDS-capacity check inside rocke_build_gfx942_4warp_gqa is in the truncated tail of dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp.
5	N/A	This PR ports a gfx942 CDNA3 MFMA attention builder and does not touch WMMA/RDNA GEMM or GemmMultiD/GemmMultiAbd.
6	N/A	The gfx942 4-warp GQA attention kernel does not use caller-zeroed global workspaces (such as Stream-K, split-K, or MoE workspaces).
7	N/A	This PR does not touch Stream-K, persistent GEMM, GroupedGemmLauncher, or SUPPORT_MATRIX.md.
8	N/A	The 4-warp GQA attention builder only supports bf16 and fp16 and does not perform quantization or MX scale decoding.
9	N/A	This PR does not add or modify numeric emulation or quantization decode paths.
10	N/A	This PR does not add or modify manifest kinds, emitters, or runtime manifest consumers.
11	N/A	This PR only adds C++ IR builder and parity tests and does not perform GPU kernel launches or stream management.
12	N/A	This PR does not touch COMGR or HIP module loading.
13	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:161 and :174 prove exact parity with Python for host predicates gfx942_gqa_fold_eligible, gfx942_4warp_launch_grid, and support wrappers without refactoring existing Python host math.
14	PASS	New C++ instance code is placed under instances/ in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:1 and no examples or helpers/ modules are modified.

§Q graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	test_attention_4warp_gqa_cpp_parity.py is registered in dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:691, but ROCKE_ENGINE_test_categories_external.yaml and the full CMake hunk showing how gfx942_attention_4warp_gqa_emit.c is registered are not in the diff.
2	NEEDS-INFO	The PR body names the covered family gfx942_attention_4warp_gqa and dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:56 references rocke_ir_parity_harness pins, but platform/tests/instances/rocke_ir_parity_harness.py is not in the diff to verify cases() already contains or adds these variants.
3	PASS	The PR body Test plan explicitly scopes the byte-identity gate (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.py:1) as a cross-engine parity check and notes that GPU numeric lanes were not run rather than claiming numerics from IR parity.
4	NEEDS-INFO	dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/golden/llvm_gfx_all.json:1222 (and lines 2925, 4627) adds golden hashes for gfx942_attention_4warp_gqa, but the diff is truncated after line 4637 so the fourth LLVM flavor section cannot be verified.
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:42 uses pytest.importorskip("rocke_engine"), pins ROCKE_LLVM_FLAVOR via monkeypatch, requires no GPU or torch, and tests negative cases with inputs that genuinely raise NotImplementedError/RuntimeError/ValueError (lines 154-156, 179).
6	N/A	This PR does not refactor any gate or test harness infrastructure.
7	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:154 and :179 assert exact exception types (NotImplementedError, RuntimeError, ValueError) per engine call rather than a tuple of unrelated types, and the PR does not alter ROCKE_BACKEND or ROCKE_CPP_STRICT fallback dispatch.

§R graded 11/11

#	Verdict	Evidence
1	NEEDS-INFO	Missing Step 0 lever-sweep evidence: enumerate existing §12.1/spec levers, correctness-prune against an fp32 reference, and name the best existing baseline for the target shapes before the new C++ 4-warp GQA kernel body is admitted.
2	NEEDS-INFO	Missing checked-in mapping from each changed lever to the runbook sections, the §13 attention checklist, the gfx942 §21 recipe, and any §14/§16 anti-pattern safety rationale.
3	NEEDS-INFO	Missing evidence that the PR followed the relevant utilities skills for MFMA/GEMM, LDS layout/swizzle, prefetch/load pipeline, ISA inspection, and launch shape, or measured justification for any contradiction.
4	NEEDS-INFO	Missing static probe outputs: before/after occupancy and registers, intrinsic or ISA inspection showing the intended primitives, and lowering comparison evidence for the C++/Python lowering-impacting builder port.
5	NEEDS-INFO	Missing tiling analysis for the added or changed global accesses: analyze_coalescing with vector_transactions/addr_fn_from_strides, CoalescingReport.summary, and assert_asm_backed against emitted access width.
6	NEEDS-INFO	Missing LDS/MMA layout proof: verify_lds_roundtrip, verify_mma_soundness or an MMA-free record_build declaration, optimize_layout recommendation evidence, or a bit-exact layout test for the non-tiling kernel path.
7	N/A	The diff does not make a bottleneck claim such as memory-bound, LDS-bound, barrier-bound, or latency-bound.
8	NEEDS-INFO	Missing runbook §8.3/§8.4 and scheduler-policy.md rationale plus ISA evidence for the added waits/barriers/scheduling behavior in the new builder.
9	N/A	This PR is not presented as a performance regression fix and does not report a shape-trading optimization result.
10	NEEDS-INFO	Missing methodology evidence that codegen/compiler caveats were reviewed, especially DSL compiler hazards and alias semantics for the new C++ emitted-IR path; no benchmark-hygiene evidence is needed absent benchmark claims.
11	N/A	The diff adds a parity port and tests, not a new optimization lesson, probe, analysis tool, or measured-results update.

§S graded 8/8

#	Verdict	Evidence
1	FAIL	dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:91 declares the new public C entry point rocke_gfx942_attention_4warp_gqa_lower_to_llvm, but neither dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:1036 (which calls rocke_build_gfx942_4warp_gqa + rocke_lower_kernel_to_llvm_ex directly), dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:109, nor dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:337 exercises rocke_gfx942_attention_4warp_gqa_lower_to_llvm.
2	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_unified_gates.cpp:160 changes the C++ tiled-2D support gate to admit the gfx942 4-warp GQA cohort, and both dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:273 and dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_routing_cpp_parity.py:256 (:362) are updated in the same PR to expect…
3	N/A	This PR does not remove or rename any public API.
4	N/A	This PR does not touch any Python __all__ list or add public symbols to Python modules.
5	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:123 asserts exact LLVM IR output equality against Python, :152 (test_cpp_builder_rejects_like_python) negatively tests reject cases (i64, head_size, lean_fp16, lean_window, arch) asserting RuntimeError with matching message, and :179 negatively tests gfx942_4warp_launch_grid raising ValueError on invalid num_kv_heads=3 (also…
6	PASS	All new tests in dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:22 and dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:337 check IR lowering, spec validation, and host predicates on CPU without requiring torch or a GPU, and are registered in CTest at dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:691.
7	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/test_attention_4warp_gqa_cpp_parity.py:123 (test_cpp_builder_matches_python), :161 (test_fold_eligible_matches_python), :174 (test_launch_grid_matches_python), and :184 (test_spec_wrappers_match_python) test the C++ implementations via rocke_engine bindings directly against their Python twins, and…
8	NEEDS-INFO	The PR description (Test plan section) states CTest, pybind build: 39/39 and Byte-identity gate GREEN at llvm20, llvm22, llvm23 and llvm24, and dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:691 is in the diff, but the PR body does not explicitly state whether PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py were run from rocke/platform/.

[yraparti]

[HIGH] §S.1 ★ ★ Every new public symbol has a test that exercises it and runs in a CI lane: a name added to __all__ of rocke, rocke.core, rocke.runtime, rocke.analysis or the rocke.helpers.tiling package root; a new <Op>Spec, build_<op>, signature or grid helper; a new parameter of…

dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/instance_gfx942_attention_4warp_gqa.h:91 declares the new public C entry point rocke_gfx942_attention_4warp_gqa_lower_to_llvm, but neither dnn-providers/hip-kernel-provider/rocke/platform/cpp/bindings/rocke_engine_attention.cpp:1036 (which calls rocke_build_gfx942_4warp_gqa + rocke_lower_kernel_to_llvm_ex directly), dnn-providers/hip-kernel-provider/rocke/library/tests/parity/gfx942_attention_4warp_gqa_emit.c:109, nor dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/attention_unified_routing_test.cpp:337 exercises rocke_gfx942_attention_4warp_gqa_lower_to_llvm.

Fix: Add a CTest unit test (e.g. in platform/tests/instances/) or route t942_4w_lower / the C parity emitter through rocke_gfx942_attention_4warp_gqa_lower_to_llvm so the convenience C entry point is exercised in CI.

[yraparti]

[MEDIUM] §O.2 scf.for keeps its structure: lower, upper and step share a type; iter-args, inits and results match in count and type; a body with iter-args ends in scf.yield of the same arity. A loop's or if's own result is not visible inside its regions, and sibling regions don't share block-defined values…

fwg_epilogue appends a trailing rocke_b_ret(b); at the end of the kernel body (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:544), violating the rule not to append a cf.return to terminate a kernel because the lowerer adds ret void.

Fix: Remove the trailing rocke_b_ret(b); at the end of fwg_epilogue in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/gfx942/attention_gfx942_4warp_gqa.cpp:544 (and its Python counterpart if present).

[amd-khushbu]

This stack is now combined into one PR, #14522, so it can be reviewed and merged as one change. GitHub shows this PR as merged because its commit is now on #14522's branch; nothing has gone into develop. This PR's change is one commit there, unchanged. Open review comments from here are listed in #14522's description and will be answered there.