[afagaj]

JIRA ID : AICK-2762

Summary: the SDPA backward kernels of rocKE can only be built from Python today; this PR adds C++ twins of every backward kernel builder that emit the same LLVM IR byte for byte, and a CI check that fails when the two drift apart.

Motivation

rocKE's C++ engine has no backward kernel builder, so a C++ caller (such as the hip-kernel-provider) cannot build the attention backward without Python. The engine's dual-engine rule asks every family the C++ engine builds to emit byte-identical IR from both engines, and CI to hold it.

What this does

Ports every attention-backward kernel builder from Python to C++ ("C++ twins"): the aux kernels, the general main kernel, the split dQ kernel, the dense kernels and the unified paged kernel, with all their helpers. Each twin emits LLVM IR that is byte-identical to Python's for every config of its parity pair, and CI now checks that (commit 15).

This PR is a stack of 16 commits, one per layer of the port, meant to be read commit by commit. Review fixes are folded into the commit that owns the code (fixup + rebase), so the stack stays clean. Squashing would lose the layering.

The commits

#	Commit	What it adds
1	test	C++/Python lowerer byte-identity gate for the backward kernels (test_attention_bwd_ir_cpp_parity.py; local only, see below)
2	test	ctest rocke_buffer_load_hip: the full C++ HIP lowering of tile.buffer_load, operands included (the op itself came to develop with #14394); portable-IR comment fix
3	feat	aux kernels + addr / worklist / caps helpers
4	feat	band helpers (shared platform helper + backward band)
5	feat	main spec, caps, tiles, LDS planner (+ a parity config for every LAYOUT_SPILLING_COMBOS row)
6	feat	fragment + block-GEMM helpers
7	feat	counted, valu, sinks, paged helpers
8	refactor	remove helpers the parallel ports duplicated; DMA slab layout in 64-bit
9	feat	stage scheduler
10	feat	tile-step body + prefetch ring
11	feat	main-kernel builder
12	feat	split dQ kernel builder (+ the DQ_SPILLING_COMBOS / DQ_SCHED_SPILLING refusals)
13	feat	dense backward kernels (+ _kv_tile dense-shell hooks)
14	feat	unified paged backward main kernel (+ ctest rocke_attention_bwd_spec_reinit)
15	test	CI check of the twins: the 20 backward parity families blessed in the differential golden; their C++ emitters built and installed; new ctest rocke_native_golden_ll
16	docs	the design doc, dispatcher notes and barrier-fence parity files describe the twins

Every twin has a parity pair under library/tests/parity/ (a Python emit script and a C driver that build the same configs); check_byte_identity.py finds them automatically. The twins' sources are globbed, so they need no CMake edits; commit 15 builds the backward parity emitters so CI can run them.

Not ported, on purpose

• Routing, plans, tuning. The dispatcher (library/dispatch/attention_bwd/) stays Python-only. No C++ caller of these builders exists yet.
• The recorded resource-gate refusals RESOURCE_REFUSED and DQ_RESOURCE_REFUSED, both keyed by spec_digest. The C++ main and dQ builders build those specs. A future C++ caller must ask Python first (resource_refusal for the main kernel, validate_attn_bwd_dq_spec for the dQ kernel) until the tables are ported. Follow-up: AICK-3015 (draft fix feat(rocke): port the SDPA backward resource refusals to C++ (AICK-3015) #15046).
• LOG2E of attention_fwd_ext is defined locally in attention_bwd_aux.cpp, because that forward helper has no C twin. The nearest shared constant is ROCKE_LSE_LN2 in instance_lse_store_internal.h.

Testing

Every commit on its own: pre-commit clean, the engine builds, the HIP ctest passes (from commit 2), and that commit's own parity families are byte-identical at llvm22 (commits 1-2: the parity pytest).

The whole stack, on develop 0e689f1:

Check	Result
pre-commit, whole stack	clean
CI's Python golden check (run_diff.py --python-only --check-golden), ll at llvm20/22/23/24 and canonical IR	GOLDEN OK
Engine build + ctest	50/50 passed
rocke_native_golden_ll (20 families × llvm20/22/23/24)	passed (in the ctest run)
check_byte_identity.py --only attention_bwd,attention_dense_bwd, llvm22 and llvm20	GREEN
rocke_library_attention_cpu_pytest lane	422 passed, 25 skipped
differential pytests (test_check_native_golden.py, test_gate_verdict.py, test_golden_store.py)	77 passed

The golden has its own ll.llvm24 section (no aliases). For these 20 families its hashes equal llvm23's: LLVM 24 changes only the named-barrier intrinsics, which these kernels don't emit.

What CI runs

CI runs ctest by the names in dnn-providers/hip-kernel-provider/ROCKE_ENGINE_test_categories_external.yaml, from an installed tree with no engine source and no rocke_engine module (ROCKE_BUILD_PYBIND is off). So:

• check_byte_identity.py can't run there (it builds the engine from source).
• test_attention_bwd_ir_cpp_parity.py is in no ctest list, so it never runs in CI. It is a local gate.
• The lower_hip_with_both_engines pytest checks skip their C++ side.

This PR's CI checks are therefore ctests, all in the quick and ffm-quick tiers:

• rocke_buffer_load_hip (commit 2): the C++ HIP lowering of tile.buffer_load.
• rocke_native_golden_ll (commit 15): the installed C++ emitters against the golden hashes that rocke_differential_golden_ll checks Python against. Both passing means C++ == Python, so a Python change to a backward builder fails CI until its twin follows.
• rocke_attention_bwd_spec_reinit (commit 14) and rocke_attention_bwd_dma_slab_int64 (commit 15): the spec validators re-run init; the DMA slab layout keeps 64-bit offsets.
• In rocke_library_attention_cpu_pytest: test_attention_bwd_dq_refusal_rows.py and test_attention_bwd_layout_spill_rows.py, which fail if a Python refusal row has no parity config.

Known gap, not from this PR

The THD token base tok0 = cu_seqlens_q[b] is an unbounded memory offset for the Q / O / dO reads and the dQ convert store, in Python and now in the C++ twin. Follow-up: AICK-3016 (draft fix #15045).

🤖 Generated with Claude Code

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

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

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

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

🙋 Wish to Override Policy?

[therock-pr-bot]

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

[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 #14365 — C++ twins of the SDPA backward kernel builders (afagaj, AICK-2762) — 🔴 First-round: blocking — the 20 new parity families push CI's differential LL golden check past its time limit on all four hipkernelprovider shards (N1), and the rebase must drop the C++ tile.buffer_load that develop's #14394 already added (N2); the canonical golden failure was develop's, and the 18 attention_bwd_* families were byte-identical in our 10-27 run

Prior item	Status
CongMa13 (10-27, 10-28): "No new issues found"; APPROVED at 664745b5	🟢 Current — the approval is at head, but it was given while all four hipkernelprovider shards were red (N1 is this PR's), before AnaghaRaoAMD's two threads, and the rebase that N2 requires will leave it on an old commit (read)
AnaghaRaoAMD (10-29) ROCKE_ENGINE_test_categories_external.yaml:62: a quick-lane ctest that runs check_byte_identity.py for the backward families	🟡 OPEN, valid — no reply. Confirmed, and wider than stated (read): no ctest or workflow runs check_byte_identity.py; test_attention_bwd_ir_cpp_parity.py is in no CMake lane and in none of the four shard logs, so it would not run even with rocke_engine built; no library/tests/test_attention_bwd* file runs in CI. CI compiles the twins (the engine globs them) but runs none of them, and the differential golden lists all 20 new families as NEW, so it checks none of their IR. Her fallback, a tracked follow-up plus a line in the body, is acceptable; blessing the families (N1's fix) at least makes a Python-side change fail CI
AnaghaRaoAMD (10-29) test_attention_bwd_ir_cpp_parity.py:34: the knob cases pick the tuned-table key by index	🟡 OPEN, valid — no reply. _plan takes list(GENERAL_TUNED_TABLE)[ki] (:99), so reordering or retuning the table moves the case to a key where ws_bounds=buffer is inert, and the case still passes without a tile.buffer_load. Her saw_buffer_load assert is the right guard (read). Until a lane collects this file (row above), the guard only helps local runs
Bots / CI	🔴 • Policy bot pass (10-25). At head 664745b5 (CI of 10-27), pre-commit, clang-tidy, Component CI, both hipDNN superbuilds, the Multi-Arch ASAN summary and base freshness pass. In TheRock Multi-Arch, every math-libs build (gfx94X, gfx950, gfx125X, Windows gfx110X) passes, and so do the gfx94X and gfx110X sanity, hipdnn-samples and integration tests; no gfx950 or gfx125X test job runs. • All four hipkernelprovider shards (Linux gfx94X 1/2 and 2/2, Windows gfx110X 1/2 and 2/2) fail the same 2 of 45 ctests, in both attempts (read: all 8 logs). GitHub tested merge 3f9c7c09: this head on its own base, develop 6ca11545.   – rocke_differential_golden_canonical: "UNBLESSED CHANGE: 4", target_intrinsics[7,9,10,11]. These are develop's stale anchors after #13723 (#14474); develop was red on them that day, and #14501 re-blessed them on 10-28. Not this PR's (its Python edits are docstrings only), and the rebase clears it.   – rocke_differential_golden_ll: killed at the test's 300 s limit every time. The develop-based run that #14474 cites failed it in about 20 s with a golden mismatch, and a develop-based run today (#14706's shards) passes it in about 30 s at four flavors. This one is the PR's (N1). The other 43 pass, including the new rocke_buffer_load_hip; the red Multi-Arch CI Summary is these shards. • GitHub: dirty, so no CI has run on a merge since; auto-merge off. CongMa13's is the only approval; amd-khushbu and vidyasagar-amd are requested. Next: afagaj rebases (N2), fixes N1 and answers both threads; the approval then needs renewing
Our checks	🟠 • Merge (ran): no clean merge with develop 6a27c440 (151 behind, 14 ahead): ir.h (1 hunk) and core_types.cpp (2 hunks), the opcode enum, name and purity tables. git merge-tree gives the same with today's 80cc53d9 (#14379) and with develop plus #14263's head 0d1172b6, and in both, three auto-merged files would not compile (N2). With no merged tree, the check script skipped the build, goldens, byte identity and tests. Rules 1–3 clean (0/0/0). • Our 10-27 run at this same head (ran 10-27, on a scratch merge with that day's develop, before #14394): native build and 34/34 ctests, including rocke_buffer_load_hip. Byte identity GREEN at llvm20/22/23 for the 18 attention_bwd_* families (2,568 configs). The PR's cpp-parity test, the backward goldens and the representative golden: 160 passed; installed golden gate PASS ×3. Not covered: the two attention_dense_bwd_* families (the run used --only attention_bwd). The Python-suite comparison did not finish, but the PR's Python edits are docstrings only (read). That run used this PR's tile.buffer_load lowering, so the rebase needs a fresh one (N2). • Twins (read, sampled): the six static caches are mutex-guarded (the kernarg-name store is filled once, so the names it hands out stay valid), and the thread-local fragment cache keys on static catalog storage. The C++ main and dq builders skip Python's recorded resource-gate refusals on purpose (instance_attention_bwd.h:46, instance_attention_bwd_dq.h:31). Nothing in C++ calls them yet, so a future C++ dispatch must consult resource_refusal first. The parity configs are mostly gfx942 and gfx950; the wave32 targets (gfx1100, gfx1151, gfx1201) appear in 14 of the 20 emitters with a few entries each, and gfx1250 only in sched_params

#	Sev	file:line	Finding
N1	🔴	platform/CMakeLists.txt:535; the 20 new library/tests/parity/attention_bwd_*_emit.py and attention_dense_bwd_*_emit.py families	Problem: the differential harness collects every parity pair, so CI's rocke_differential_golden_ll (Python-only, three flavors at this head) also emits this PR's 20 families. That takes it from 83 to 103 families, with 1,846 new configs, none blessed (canonical: NEW 1,847 against develop's 1). With them, it hits its TIMEOUT 300 on all four shards in both attempts, while develop runs it in about 20–30 s (read: CI logs). One of the 20, attention_bwd_unified_main (236 configs), takes about 13 s per flavor on one local core (ran), so it is unverified whether the overrun is the total cost or one slow family under CI's installed tree. Consequence: the rebase clears the canonical failure, not this one, and today's develop CI already runs the ll check over 90 families at four flavors (llvm24 joined with #14415). CI stays red on every hipkernelprovider shard, and merged as is, every later PR's shards would fail too. Fix: time each new family with CI's command on the installed tree (run_diff.py --mode ll --python-only --check-golden --flavor llvm20,llvm22,llvm23,llvm24 --only <family>) and fix any outlier. Then raise this test's TIMEOUT with margin, or give the twins' families their own ll ctest. Bless the 20 families in both modes (--record-golden --only attention_bwd,attention_dense_bwd, ll at every flavor plus canonical), so CI checks their Python IR
N2	🟠	include/rocke/ir.h:480, :1558; core/ir/core_types.cpp:942, :1230; core/ir/ir_flow.cpp:1242; core/lower_llvm/mem.cpp:1134, :2035; core/lower_hip/lower_hip_mem.cpp:1061, :1768	Problem: develop's #14394 (10-29) already added the C++ tile.buffer_load that commit 2 adds: the same enumerator right after ROCKE_OP_ARITH_FDOT2, the same name, not pure, the same rocke_b_buffer_load signature, and LLVM and HIP handlers that emit the same text (they differ only in their error paths). Git marks the three table hunks but merges the rest silently. The merged ir_flow.cpp, mem.cpp and lower_hip_mem.cpp each define the function twice (rocke_b_buffer_load, op_tile_buffer_load, _op_tile_buffer_load), and the last two register it twice (ran: merge-tree with develop 80cc53d9, and with develop plus #14263's head; read: the merged files). Consequence: a rebase that only resolves the marked hunks does not compile. Nothing is renumbered: both sides append the opcode at the same index. Fix: drop commit 2's engine half. Take develop's side of all three hunks. Once #14263 lands, that side is develop's TILE_BUFFER_LOAD followed by #14263's TILE_BUFFER_STORE_VN, TILE_GLOBAL_PTR_OFFSET and TILE_TRAP; this PR adds no opcode. Delete this PR's builder, both lowerings, their two registrations and the second ir.h declaration. Keep rocke_buffer_load_hip: it checks only the emitted HIP text, which both versions share. Then re-run byte identity for all 20 families on the rebased tree

Utility: 🟢 Positive — C++ twins of every backward builder; byte-identical in our run, but CI-ungated.

---

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; today's build, goldens, byte identity and test runs wait for the rebase, and this review reuses our 10-27 run at the same head on that day's develop. The orchestrator simulated the merge with git merge-tree against today's develop and against develop plus #14263, and compared the duplicated tile.buffer_load code on both sides. It read the eight hipkernelprovider shard logs at this head (both attempts), the develop-based run #14474 cites and today's #14706 shards, and timed one new family's Python-only LL golden check. It sampled the twins' caches, refusal split and parity coverage, and checked both open threads. The orchestrator owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

[afagaj]

Re: N1, the rocke_differential_golden_ll timeout

the 20 new parity families push CI's differential LL golden check past its time limit on all four hipkernelprovider shards (N1)
[…] it is unverified whether the overrun is the total cost or one slow family

One slow family. Fixed in db6a697.

Python reference, llvm22, 1 core	Time
attention_bwd_layout	164 s
each of the other 19 families	≤ 18 s

• Cause. The 18 choose|R,C cases of attention_bwd_layout_emit.py took 152 of the 164 s: 720 XT swizzle searches that re-predict identical instructions. planner_xt_swizzle repeats choose_xt_swizzle, and the reads do not depend on the writer (_predict_max: 348 M calls for choose|128,256 alone). One family runs in one worker and the flavors run in series, so this family alone needed about 650 s.
• Fix. The emitter memoizes _predict_max for its process. Test-only: _attention_bwd_lds.py and the C++ twin are unchanged, and the output is byte-identical.
• Result. CI's command (run_diff.py --mode ll --python-only --check-golden --flavor llvm20,llvm22,llvm23,llvm24, 8 jobs) takes 180 s locally with GOLDEN OK at all four flavors and 0 unblessed changes; it was killed at 300 s before. The choose cases take 35 s (152 s before). TIMEOUT 300 is unchanged.
• Blessing. The 20 families are blessed in 35415c2 (.ll at llvm20/22/23, llvm24 via the alias; canonical IR), and the new rocke_native_golden_ll checks the C++ emitters against the same hashes.

I'll confirm both golden ctests on the CI shards.

[afagaj]

Re: N2, the duplicate tile.buffer_load

develop's #14394 (10-29) already added the C++ tile.buffer_load that commit 2 adds
[…] Keep rocke_buffer_load_hip: it checks only the emitted HIP text, which both versions share.

Done. The engine copy is dropped and rocke_buffer_load_hip is kept, in 6ff9fc4.

• Engine. Rebased onto develop 7cd54ff and took develop's opcode, builder and both lowerings as they are. This PR adds no opcode.
• Ctest kept. rocke_buffer_load_hip (f16, bf16, f32, i32 on gfx942 and gfx950) is the only CI check of the C++ HIP lowering. test_buffer_load_scalar.py skips its two-engine cases without rocke_engine (SKIPPED on feat(rocke): buffer address space in the shared C++ tensor view #14706's gfx94X shards, 10-30). Registered in tests/cmake/CoreTests.cmake, in the quick and ffm-quick tiers.
• LLVM side. 15 backward emitter configs emit the op (14 in attention_bwd_main, 1 in attention_bwd_body), so rocke_native_golden_ll checks its C++ LLVM lowering against the golden.
• Verified. Builds with no warnings under gcc and clang; ctest 37/37; byte identity GREEN at llvm20 and llvm22 after the rebase.

[cgmillette]

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

PR #14365 — C++ twins of the SDPA backward kernel builders (afagaj, AICK-2762) — 🔴 Follow-up: new CI lane red

Prior item	Status
CongMa13 (10-28) APPROVED at 664745b5	🟢 APPROVED 664745b5 • older head; still counts on develop
AnaghaRaoAMD (10-29) CI gate for the backward twins	🟠 Mostly 7f9df6fb • 20 families blessed; C++ lane ✓ on gfx94X shards (read) • red on Windows; breaks rocke_pytest everywhere: N1, N2 • reply's "rocke_buffer_load_hip is gone" is stale: kept, ✓ (read)
AnaghaRaoAMD (10-29) knob cases pick the table key by index	🟢 FIXED bd7d4f33 • assert passes at key 1, fires at inert key 0 (ran) • 89 ✓ at head (ran); no CI lane collects the file (read)
Our 10-30 N1 🔴: 20 new families time out the ll golden lane	🟢 FIXED f917487c, 7f9df6fb • ll lane ✓ on all 4 shards, both attempts (read) • 20 families blessed; no new unblessed configs (ran) • 167–210 s of 300 s; develop takes 31 s (read)
Our 10-30 N2 🟠: duplicate C++ tile.buffer_load vs develop	🟢 FIXED c30393eb • engine half dropped; PR adds no opcode (ran: range-diff) • builds clean; rocke_buffer_load_hip ✓ here and in CI (ran) • byte identity GREEN ×3 for all 20 families (ran)
Bots / CI	🔴 Red • 37 ✓ · 49 skip · 7 ✗; attempt 1 failed the same way; none running • ✗ Windows shards + HipDNN Windows: native golden 80/80 differ (N1) • ✗ all 4 shards: registration test (N2); gfx94X also develop's 14733 • merge: MERGEABLE; blocked: these checks, 2 open threads
Our checks	🔴 ctest red • head has develop 7cd54ff9; +1 hipBLASLt commit since: clean (ran) • ctest 35 ✓ / 1 skip / 1 ✗: body emitter aborts (N3) (ran) • byte identity GREEN ×3; parity 2 ✓ ×3; goldens ✓; 89 ✓; rules 1–3 clean (ran) • compile 6 of 20 families: 2 dq spills, knob-only specs, not dispatch picks (ran)

#	Sev	file:line	Finding
N1	🔴	platform/tests/instances/differential/check_native_golden.py:46	Windows: all 80 native golden checks differ • impact: emitters fputs the .ll to text-mode stdout (LF→CRLF), raw bytes are hashed (read); both Windows shards and HipDNN Windows red, develop's Windows CI too on merge • fix: normalize CRLF before hashing (Linux CRLF simulation fails as CI does, passes normalized: ran) or binary stdout; rerun a Windows shard
N2	🔴	platform/tests/CMakeLists.txt:14	new include breaks the registration contract test • impact: test_cmake_test_registration_is_concern_owned fails on all 4 shards and here (ran); develop passes it; rocke_pytest red on merge • fix: add NativeGoldenTests.cmake to _FRAGMENTS in tests/tools/test_cmake_test_registration.py
N3	🟠	library/tests/parity/attention_bwd_body_emit.c:1262	memset clears a Probe that holds a std::string (UB) • impact: gcc 15 -O3 emitter aborts ("double free"), so the lane fails here; UBSan flags it (ran); CI's clang warns too, -Wno-error hides it (read) • fix: p = Probe{}; gives 246/246 ✓ at all 4 flavors, UBSan clean (ran)

Utility: 🟢 Positive — byte-identical C++ twins of every backward builder; its CI gate needs fixes.

---

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: range-diff since our last review; CI logs at head (both attempts) against develop's at its base; the new lane on this host, a CRLF simulation of the Windows failure, and -O3 and UBSan rebuilds of the crashing emitter with the one-line fix; build, ctest, byte identity, parity, goldens, the PR's tests and the registration test at head; the knob assert on an inert key; the two spilling dq configs; a merge-tree with today's develop. Flags whether a closer look is needed; not an approval or rejection.

[afagaj]

rocke_library_attention_cpu_pytest: 12 failed, 410 passed (Linux gfx94X and Windows gfx110X)
FileNotFoundError: ... tests/library/tests/parity/attention_bwd_spec_emit.c

TL;DR: fixed in 46ceded. The installed test tree now ships the one C file this test reads, so it runs in full in CI.

Cause: test_attention_bwd_layout_spill_rows.py checks each LAYOUT_SPILLING_COMBOS row twice: a Python parity config is refused at that row, and the same config is in the C table, which it reads from attention_bwd_spec_emit.c. CI runs the tests from the installed tree, and the library-tests install skips every *.c file. I had only run the test from the source tree, where the file exists.

Fix: one install(FILES ...) rule in platform/CMakeLists.txt, with a comment, installs attention_bwd_spec_emit.c next to its .py twin.

Why not skip the C half in CI instead: the golden checks treat an unblessed new config as "informational", so a new spill config that is neither blessed nor in the C table would pass every CI check. Catching that is this test's job.

Checked (provider-style install, the lane's own command):

Run	Result
Without the rule	12 failed, same FileNotFoundError
With the rule	spill-rows test 12 passed; whole rocke_library_attention_cpu_pytest lane 422 passed, 25 skipped

[amd-khushbu]

Re-reviewed at d83c2d6, rebased on develop 0e689f1. All 17 threads from my last review are fixed or tracked (AICK-3014, AICK-3015), so I've resolved them. The non-inline points are handled too: the stale prose, the llvm24 wording, "never runs in CI", the resource_refusal wording, LOG2E, and the THD follow-up (AICK-3016).

Re-run locally at the PR head:

Check	Result
Engine and test build	clean
ctest (50, including rocke_native_golden_ll, rocke_attention_bwd_spec_reinit, rocke_attention_bwd_dma_slab_int64, rocke_buffer_load_hip)	50/50 pass (1 skipped: arch_domain_reproducibility)
rocke_native_golden_ll, 20 families × llvm20/22/23/24	OK
Python golden, ll at llvm20/22/23/24 and ir --canonical	GOLDEN OK
test_attention_bwd_dq_refusal_rows.py, test_attention_bwd_layout_spill_rows.py, test_check_native_golden.py	26 passed

One new item (inline): the two new ctests are in no tier list, so the repo's tier guard fails.

CI: the math-libs stages failed in the hipCCL configure step (hipcub find_package), which this PR doesn't touch. Because of that, no test jobs ran, including the hipkernelprovider shards. rocke_native_golden_ll hasn't run in CI yet. Please re-run once that's fixed.

[afagaj]

TL;DR: rebased onto develop a6b0022cc5c to clear merge conflicts. No code changed, and all checks pass.

Conflicts: two files, both where develop and this PR added entries at the same spot. I kept both sides.

File	Develop added	This PR added
ROCKE_ENGINE_test_categories_external.yaml	rocke_attention_dense_spec	rocke_buffer_load_hip
llvm_gfx_all.json (golden)	attention_dense_lse and other families	the 20 backward families

The golden file was merged entry by entry. No entry was changed on both sides, so nothing had to be chosen. No backward hash moved after the rebase.

Checked at 94ea1c10621:

Check	Result
pre-commit (whole stack)	clean
Python golden, ll at llvm20/22/23/24 and ir --canonical	GOLDEN OK
Build + ctest	62/62 pass, including rocke_native_golden_ll
C++ vs Python byte identity, backward + dense, llvm22 and llvm20	GREEN
Attention CPU tests	726 passed
tests/tools (includes the tier check)	249 passed

CI note: the Windows hipkernelprovider shards have been hitting their 30-minute limit. I'm looking at that separately with the CI owners.

[Copilot]

🟡 Changes recommended

The parity checker can falsely pass when trailing C++ refusal cases are missing because it checks blessed hashes without enforcing the full Python config count, including refusals.

1 open issue

• Reject truncated C++ config tables after the refused tail · New

What changed in this PR

This PR adds C++ counterparts for rocKE’s SDPA backward kernel construction and installed-artifact parity checks against the Python reference. Dispatch, tuning and recorded resource-refusal tables remain Python-only.

Changes:

• Adds backward builder, spec and helper interfaces, plus the dense-main builder wrapper.
• Builds and installs native parity emitters for per-LLVM-flavor golden checks in CI.
• Adds parity and regression tests and updates documentation to describe the dual-engine support.

File	Description
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​tools/​test_cmake_test_registration.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​differential/​test_check_native_golden.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​differential/​run_diff.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​differential/​golden/​README.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​differential/​check_native_golden.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​attention_bwd_spec_reinit.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​instances/​attention_bwd_dma_slab_int64.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​test_buffer_load_hip.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​CMakeLists.txt	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​cmake/​NativeGoldenTests.cmake	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​cmake/​InstanceTests.cmake	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​cmake/​CoreTests.cmake	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​helpers/​attention_band.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​instances/​attention_bwd.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​portable_ir/​recipe_vm.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​portable_ir/​ir_import_json.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​instances/​common/​attention_dense_bwd.cpp	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_attention_dense_bwd_spec.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_attention_dense_bwd_dq.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_attention_bwd_unified.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_attention_bwd_dq.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​instance_attention_bwd_aux.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_rocke.instances.common._attention_dense_bwd_shell.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_rocke.instances.common._attention_bwd_worklist.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_rocke.instances.common._attention_bwd_valu.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_rocke.instances.common._attention_bwd_tiles.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_rocke.instances.common._attention_bwd_sinks.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_rocke.instances.common._attention_bwd_counted.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​include/​rocke/​helper_rocke.helpers.attention_band.h	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​CMakeLists.txt	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_barrier_fence_parity_emit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_attention_bwd_layout_spill_rows.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_attention_bwd_ir_cpp_parity.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​test_attention_bwd_dq_refusal_rows.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​barrier_fence_emit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​attention_bwd_valu_emit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​attention_bwd_sched_ir_emit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​attention_bwd_lds_emit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​attention_bwd_lds_emit.c	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​attention_bwd_counted_emit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​tests/​parity/​attention_bwd_band_emit.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​attention_bwd_aux.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_dense_bwd_shell.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_worklist.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_valu.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_sinks.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_ring.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_paged.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_dqbody.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_counted.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_body.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_band.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​kernels/​common/​_attention_bwd_addr.py	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​rocke/​library/​dispatch/​AGENTS.md	Updated as part of this pull request.
dnn-providers/​hip-kernel-provider/​ROCKE_ENGINE_test_categories_external.yaml	Updated as part of this pull request.

_{Review effort: Max}

[cgmillette]

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

PR #14365 — C++ twins of the SDPA backward kernel builders (afagaj, AICK-2762) — 🔴 Follow-up: does not build on develop

Prior item	Status
CongMa13 (10-28) APPROVED at 664745b5	🟢 APPROVED 664745b5 • older head; still counts on develop • does not cover the stack rewrite or the N1 port
AnaghaRaoAMD (11-02) APPROVED at 35415c25	🟢 APPROVED 35415c25 • older head; still counts on develop • does not cover the N1 port
amd-khushbu (11-06) APPROVED at bd912643	🟢 APPROVED bd912643 • older head; still counts on develop • since then a rebase only: two conflict files resolved (ran)
AnaghaRaoAMD (10-29) CI gate for the backward twins	🟢 FIXED c3ebbed6 • rocke_native_golden_ll ✓ at head build (ran) • not yet run in CI: no test job ran at head (read)
AnaghaRaoAMD (10-29) knob cases pick the table key by index	🟢 FIXED ead24391 • fix carried through the rebuilt stack (ran: range-diff)
Our 10-30 N1 🔴: 20 new families time out the ll golden lane	🟢 FIXED c4c0dc0a, c3ebbed6 • carried; not re-timed: no CI test job at head (read)
Our 10-30 N2 🟠: duplicate C++ tile.buffer_load vs develop	🟢 FIXED cd688404 • rocke_buffer_load_hip ✓ at head build (ran)
Our 10-31 N1 🔴: Windows: all 80 native golden checks differ (CRLF)	🟢 FIXED c3ebbed6 • saved probe: simulated Windows run ✓ at head build (ran)
Our 10-31 N2 🔴: new include breaks the registration contract test	🟢 FIXED c3ebbed6 • saved probe: contract test 5 ✓ on the merge (ran)
Our 10-31 N3 🟠: memset clears a Probe holding a std::string	🟢 FIXED e7f6005e • saved probe: -O3 and UBSan emit equal bytes at head (ran)
amd-khushbu (11-03) re-bless after 14328; 17 inline items	🟢 FIXED • threads resolved by amd-khushbu after her re-run (read) • native golden, spec_reinit, dma_slab ctests ✓ at head (ran)
amd-khushbu (11-03) one home for salted name; C++ refusal tables	🟡 DEFERRED • AICK-3014, AICK-3015; draft PRs 15044, 15046 stacked here (read) • "no C++ caller exists": latent today, fair to defer
amd-khushbu (11-03) THD tok0 offset unbounded	🟡 DEFERRED • AICK-3016, draft PR 15045; gap shared by both engines (read) • fair to defer; listed under "Known gap" in the body
amd-khushbu (11-04) two new ctests in no tier; out-of-date dQ comment	🟢 FIXED c3ebbed6 • both ctests in quick and ffm-quick; ✓ at head build (ran)
Copilot (11-07) truncated refusal tail passes the check	🟡 OPEN • reproduced: dQ emitter cut before config 124 reads OK (ran) • golden blesses 0–106; refusal configs 107–125 go unchecked • afagaj to fix or reply in the thread
Bots / CI	🔴 Failing • fails: 4 math-libs stages, both HipDNN superbuilds — N1's compile error (read) • CI merged onto a8445ead (11642 itself); 2 Summaries follow • 48 skipped: no rocKE test shard ran at head • merge: MERGEABLE; BLOCKED
Our checks	🔴 Merge does not build • merge vs 1702865a: no conflict; build ✗, 52 errors in 10 files (ran) • head alone: build ✓, ctest 62 ✓ / 1 skip (ran) • merge: byte identity, parity, differential golden not run: no build (ran) • representative golden ✓ ×4; 313 tests ✓; no ABI change; rules clean (ran)
Code quality	🟠 Should-fix • twins bypass develop's division and repr helpers: N2, N4 (read) • LDS predictor and arch list in a kernel file: N3 (read) • macro shared with open PR 14464: N5; MMA ops reuse develop's catalog (read)

#	Sev	file:line	Finding
N1	🔴	platform/cpp/instances/common/attention_bwd.cpp:1738	Twins use rocke_mma_op_t members develop's 11642 removed • impact: merge fails, e.g. _attention_bwd_caps.cpp:315:51: error: no member named 'a_frag_len' in 'rocke_mma_op' (CI, every build stage) • impact: 52 errors in 10 files; attention_bwd_valu_emit.py imports cast_c_frag, now cast_dst_frag: differential golden lanes fail • fix: rebase onto a8445ead or later; port to srcs[i] / dst, cast_dst_frag; re-run byte identity
N2	🟠	platform/cpp/instances/common/_attention_bwd_sched.cpp:176	11 private integer-division copies beside develop's rocke_floor_div • impact: fdiv, cdiv, floordiv, py_floordiv, align_up … in 9 files restate helper_rocke.helpers.spec.h:373-375 • impact: draft 15044 adds one more shared py_floordiv, not develop's • fix: call rocke_floor_div / rocke_ceil_div / rocke_align_up; one checked wrapper
N3	🟠	platform/cpp/instances/common/_attention_bwd_lds.cpp:364	LDS conflict predictor and its arch list ported into a kernel file • impact: wave64 targets hard-coded; Python reads them from the profile registry • impact: a new or changed profile moves Python's LDS plan, not C++'s • fix: move it to a C++ twin of rocke.analysis.lds; targets from its profile table
N4	🟠	platform/cpp/instances/common/_attention_bwd_caps.cpp:67	Second Python str repr beside develop's rocke_py_str_repr • impact: py_repr plus 7 wrappers; ir.h:831 already owns this • impact: the two spell U+0080–U+00FF differently in error texts • fix: move the Latin-1-exact algorithm into rocke_py_str_repr; keep one wrapper
N5	🟠	platform/cpp/include/rocke/helper_rocke.instances.common._attention_bwd_sched.h:110	ROCKE_COMGR_FRESH_PROCESS_ATTR also defined by open PR 14464 • impact: two homes, same value: here and 14464's compile_options.h:31 • owner: 14464 (its header is the twin of helpers/compile.py) • fix: keep the guarded copy until 14464 lands; whichever lands second deletes :108-111

Utility: 🟡 Mixed — byte-identical twins at head; must be ported to develop's indexed MMA operands.

---

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: range-diff since our last review and since the last approval; the CI logs of every failing job and their merge ref; build of the head and of a merge with today's develop (all errors listed); ctest at head; our three saved probes at head; a reproduction of Copilot's truncated-tail case; goldens, the PR's tests, rules and ABI on the merge; code quality. Flags whether a closer look is needed; not an approval or rejection.

[cgmillette]

Automated overlap notice (rocKE merge monitor)

Hi @afagaj, this PR and #14464 (vidyasagar-amd) both define the comgr attribute name ROCKE_COMGR_FRESH_PROCESS_ATTR. Our follow-up review of 7 November (N5) records the owner; this notice says how the copy goes away.

The copies, here at 94ea1c1062 and in #14464 at 61501dd46f (paths under dnn-providers/hip-kernel-provider/rocke/platform/cpp/):

• here: include/rocke/helper_rocke.instances.common._attention_bwd_sched.h:108-111, a comment and an #ifndef guard around #define ROCKE_COMGR_FRESH_PROCESS_ATTR "comgr_fresh_process", used at instances/common/_attention_bwd_sched.cpp:2231 and :2457;
• in feat(hipTensor/rocKE): Add a new rocKE backend to hipTensor and add optimized kernel support #14464: include/rocke/compile_options.h:31, the same name and value.

Owner: #14464: compile_options.h is the header that holds the comgr compile options. Resolution: this PR uses that definition and drops its own.

Landing order: the merge plan (develop 1702865a2f) has this PR in wave 2. #14464 is not in the plan yet: it conflicts with develop. So this PR will likely land first, with its guarded copy: keep the copy until #14464 lands. When #14464 rebases onto a develop that has this PR, it makes the step below. If #14464 lands first, this PR makes the step at its rebase.

Rebase (the PR that lands second), in include/rocke/helper_rocke.instances.common._attention_bwd_sched.h:

1. Replace the comment and the #ifndef ROCKE_COMGR_FRESH_PROCESS_ATTR block (:108-111) with #include "rocke/compile_options.h", beside the header's other rocke/ includes.
2. Leave the two uses in _attention_bwd_sched.cpp as they are: the name and value do not change.

Checked: both definitions, read at the heads above, have the same name and value, and no other file in either PR defines it. compile_options.h is a C header with extern "C" guards, so this C header can include it. Nothing was built.

If you think this PR should own it instead, reply here and on #14464, 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.

[yraparti]

rocKE checklist review

ROCm/rocm-libraries#14365 · 94ea1c106219 · model gemini-4-argon, fallback gpt-5.5

PASS · graded 238/238 · BLOCKER 0, HIGH 3, MEDIUM 3, LOW 0, NEEDS-INFO 54

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

Findings

• HIGH §D · dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/check_native_golden.py:149 · Put the knob on the KernelSpec and delete the environment read.
• HIGH §D · dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:19 · Port the RESOURCE_REFUSED resource-gate refusals to the C++ spec validation path, or make the C++ builder refuse the same specs with the same reasons before building.
• HIGH §O · dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:42 · Resolve the arch once through ArchTarget and gate dense-builder support on target/catalog capabilities instead of comparing gfx strings directly.
• MEDIUM §E · dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:432 · Replace the internal https://amd-hub.atlassian.net/browse/AICK-... links in the PR description with bare ticket IDs (AICK-3015, AICK-3016).
• MEDIUM §O · dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:37 · Replace the fixed AUX_MIN_WG_ATTR constant with a value derived from the resolved target wave size and max workgroup size before setting kernel attributes.
• MEDIUM §O · dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_dense_bwd_shell.cpp:41 · Use a dtype-size helper or target/catalog dtype metadata for element-byte calculations instead of the ELEM_BYTES literal.

Needs info

• §A.4 Full C++ hunks for several new twins are omitted or unavailable, so operand/order identity cannot be verified from the provided diff.
• §A.5 The representative golden file hunk for platform/tests/instances/differential/golden/llvm_gfx_all.json is omitted, so it cannot be verified whether golden hashes were unchanged or intentionally re-blessed.
• §A.7 The hot-path logic is spread across large omitted C++ hunks, so it cannot be verified from the provided diff that all new logic is guarded and covered shapes emit identical IR.
• §A.8 The C++ implementation hunks and CMake/native-golden gate hunks are omitted or unavailable, so the C++ review/gate cannot be verified from the provided diff.
• §A.9 The relevant test_buffer_load_hip.cpp and test-registration hunks are omitted, so it cannot be verified whether any shared op-handler/lowerer change is proven by the required mini-pair or equivalent parity test.
• §A.13 The native-golden and run_diff.py hunks are omitted, so it cannot be verified that .ll output remains unstamped and gates run with debug locations disabled.
• §A.14 platform/python/rocke/helpers/attention_band.py:29 documents a C++ twin, but the actual platform/cpp/helpers/attention_band.cpp and header hunks are unavailable, so the twin update cannot be verified.
• §E.1 Diffs for platform tests and CMake files (such as platform/tests/cmake/NativeGoldenTests.cmake, platform/tests/instances/differential/check_native_golden.py, and platform/tests/instances/differential/run_diff.py) are omitted from the diff.

Section notes

• §A: The provided diff is truncated and several relevant hunks are omitted or unavailable, so items that require inspecting those hunks are marked NEEDS-INFO.
• §B: C++ twins of SDPA backward kernel builders; dispatcher/routing remains untouched and Python-only.
• §D: PR ports existing Python attention backward kernel builders and specs to C++ twins under platform/cpp/ without adding new spec knobs or modifying Python dispatcher registration, routing, or candidate selection.
• §F: PR metadata and file list satisfy policy checks; truncated diff hunks and external CI/GitHub thread states require NEEDS-INFO where line citations are mandatory.
• §G: Items 1-17 of 34 in Section G. Many test and doc diffs are omitted in the truncated patch.
• §G: Items 18-34 of Section G graded against PR #14365.

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	The PR adds C++ twins for attention backward builders, e.g. platform/cpp/instances/common/attention_bwd_aux.cpp:4 identifies the C port of library/kernels/common/attention_bwd_aux.py, and library/kernels/common/attention_bwd_aux.py:6 now documents the byte-identical C++ twin.
2	N/A	No core lowering-op or op-handler change under core/lower_* or core/ir is shown in the provided diff.
3	PASS	Kernel-level attention backward builders are mirrored in C++, e.g. platform/cpp/instances/common/attention_bwd_main.cpp:4 ports the Python main-kernel builder and platform/cpp/instances/common/attention_bwd_unified.cpp:4 ports the unified builder.
4	NEEDS-INFO	Full C++ hunks for several new twins are omitted or unavailable, so operand/order identity cannot be verified from the provided diff.
5	NEEDS-INFO	The representative golden file hunk for platform/tests/instances/differential/golden/llvm_gfx_all.json is omitted, so it cannot be verified whether golden hashes were unchanged or intentionally re-blessed.
6	PASS	The PR includes explicit cross-engine parity pairs rather than relying only on goldens; library/tests/parity/attention_bwd_band_emit.py:8 says the probe is matched by a C twin and library/tests/parity/attention_bwd_band_emit.py:12 says the two .ll outputs must match.
7	NEEDS-INFO	The hot-path logic is spread across large omitted C++ hunks, so it cannot be verified from the provided diff that all new logic is guarded and covered shapes emit identical IR.
8	NEEDS-INFO	The C++ implementation hunks and CMake/native-golden gate hunks are omitted or unavailable, so the C++ review/gate cannot be verified from the provided diff.
9	NEEDS-INFO	The relevant test_buffer_load_hip.cpp and test-registration hunks are omitted, so it cannot be verified whether any shared op-handler/lowerer change is proven by the required mini-pair or equivalent parity test.
10	N/A	The shown change ports existing attention backward IR emission to C++ rather than adding a new feature with separate FLAG-ON IR.
11	N/A	No new IR opcode is shown in the provided diff.
12	N/A	The provided diff does not touch ROCKE_BACKEND=both fallback logic or core/backend.py::CPP_UNPORTED_ARCHES.
13	NEEDS-INFO	The native-golden and run_diff.py hunks are omitted, so it cannot be verified that .ll output remains unstamped and gates run with debug locations disabled.
14	NEEDS-INFO	platform/python/rocke/helpers/attention_band.py:29 documents a C++ twin, but the actual platform/cpp/helpers/attention_band.cpp and header hunks are unavailable, so the twin update cannot be verified.

§B graded 6/6

#	Verdict	Evidence
1	PASS	Implementation shortcuts vs hardware/functional limits are explicitly distinguished with NotImplError (ROCKE_ERR_NOTIMPL) vs ValueError (ROCKE_ERR_VALUE), e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:478 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:67.
2	N/A	This PR does not add or modify dispatcher preflight gates; routing and dispatch remain Python-only and untouched.
3	N/A	This PR only adds C++ twins of backward kernel builders and parity tests without changing model-zoo routing or fallback selection.
4	PASS	Last-gate validators re-run spec initialization rather than normalizing malformed C struct fields into valid alternatives (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:150-156 re-runs post_init so a misspelled page_mode is rejected rather than silently building as 'direct', and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:69-101 strictly…
5	N/A	No dispatch gate scopes are opened or modified in this PR; the C++ builders mirror existing Python spec validation.
6	PASS	Spec validators enforce only functional, hardware, and emit-safety invariants (tile divisibility, atom legality, LDS capacity, register estimates, compiler scratch-spill refusals) with hardware/algorithm reasons, e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:482-598.

§C graded 10/10

#	Verdict	Evidence
1	PASS	Shared logic such as the per-kv-tile driver and epilogues is factored into shared helpers (main_kernel::kv_tile) reused across the general, dense (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_dense_bwd_shell.cpp:34), and unified (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:27) builders.
2	PASS	Each family exposes a single spec-driven builder entry point (e.g., rocke_build_attn_bwd_unified_main_new in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:903 and rocke_build_attn_dense_bwd_main_new in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:26).
3	PASS	Kernel names are derived from operation, stage, and codegen knobs rather than workload dimensions (e.g., rocke_attn_bwd_aux_kernel_name in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:114 and dq_kernel_name in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:375).
4	PASS	All C++ backward builders live under platform/cpp/instances/common/ and dispatch arch differences via target facts (arch_facts(arch)) rather than structural if arch == ... branches (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_dense_bwd_shell.cpp:133, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:234).
5	PASS	All touched files are directly part of the C++ backward port, its parity emitters (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:1), docstrings (dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7), and CI registration.
6	PASS	The C++ builders are structured cleanly around shared internal headers and helpers (instance_attention_bwd_main_internal.h in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:44) rather than ad-hoc patches.
7	N/A	This PR ports existing Python attention backward builders to C++ rather than performing a standalone cleanup refactor.
8	PASS	The port reuses existing C builder primitives and tile.buffer_load from develop (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_dense_bwd_shell.cpp:544, dnn-providers/hip-kernel-provider/rocke/platform/cpp/portable_ir/ir_import_json.cpp:1084).
9	PASS	The portable-IR comment fix and spec re-validation pattern were applied consistently across sibling sites (dnn-providers/hip-kernel-provider/rocke/platform/cpp/portable_ir/ir_import_json.cpp:1082 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/portable_ir/recipe_vm.cpp:128; dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:151).
10	PASS	Adding C++ twins for the backward modules updated all corresponding Python module docstrings (dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7, dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/attention_band.py:30) and parity/golden gates.

§D graded 21/21

#	Verdict	Evidence
1	FAIL	dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/check_native_golden.py:149 reads the environment
2	N/A	This PR does not add or modify any dispatcher KernelCandidate or Capability declarations.
3	N/A	This PR does not add or modify any candidate _supports predicates or dispatcher admission checks.
4	N/A	This PR does not modify dispatcher selectors or spec builders on the request selection path.
5	N/A	This PR does not add or modify Python KernelSpec identity, compile_key, or dispatcher cache-key logic.
6	N/A	This PR does not add or modify cohort routing, launch metadata, or dispatcher cache keys.
7	N/A	This PR does not touch dispatcher routing heuristics or hand-pinned selector overrides.
8	N/A	This PR does not define or modify any dispatcher routing cohorts.
9	N/A	This PR ports existing backward kernel builders to C++ and does not register a new kernel candidate or architecture in the dispatcher.
10	N/A	This PR does not introduce any new spec fields or features that need request-level reachability.
11	N/A	This PR does not register, promote, or change the opt-in status or arch coverage of any dispatcher candidate.
12	N/A	This PR does not add or reorder dispatcher candidates or custom rankers, so priority/ranker routing does not apply.
13	N/A	This PR does not add a CandidateRegistry wrapper or dispatcher entry point that could catch PinRefused or retry pinned requests as auto.
14	N/A	This PR does not change dispatcher cache keys, KernelId identity, or candidate abi_version handling.
15	N/A	This PR does not add a dispatcher candidate or arch registry module, so additive candidate registration rules do not apply.
16	N/A	This PR adds C++ twins for an existing attention-backward family rather than introducing a new dispatcher operator family.
17	N/A	This PR does not add a dispatcher candidate whose routing must be covered by CPU admits() tests.
18	N/A	This PR does not modify dispatcher registry ratchets, binding requirements, route-only registries, or exempt sets.
19	N/A	This PR does not claim or change production unified attention auto-routing through run_unified_attention_torch.
20	FAIL	The new C++ main builder documents a Python/C++ spec-validation divergence: it mirrors _emit_main(..., main_kernel_refusals=False) and says recorded resource-gate refusals stay Python-only, so C accepts RESOURCE_REFUSED specs that Python refuses at dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:19.
21	N/A	This PR is a C++ parity port of existing emitters, not a selection-choice fix that should be made in a chooser instead of an artifact.

§E graded 17/17

#	Verdict	Evidence
1	NEEDS-INFO	Diffs for platform tests and CMake files (such as platform/tests/cmake/NativeGoldenTests.cmake, platform/tests/instances/differential/check_native_golden.py, and platform/tests/instances/differential/run_diff.py) are omitted from the diff.
2	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:22 imports only _emit_common, kernels.common, and rocke, with no torch import.
3	NEEDS-INFO	Diffs for the added pytest files (test_attention_bwd_dq_refusal_rows.py, test_attention_bwd_ir_cpp_parity.py, test_attention_bwd_layout_spill_rows.py, test_barrier_fence_parity_emit.py, test_check_native_golden.py) are omitted.
4	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:22 imports downward from kernels.common and rocke; neither test_library_layering.py nor check_ir_validity.py is modified.
5	NEEDS-INFO	Diffs for added files under platform/tests/ (including check_native_golden.py, test_check_native_golden.py, and NativeGoldenTests.cmake) are omitted.
6	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7 and the other emitter files under library/kernels/common/ and platform/python/rocke/helpers/attention_band.py:30 only update module docstrings without deleting IR handles.
7	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7 updates docstrings for C++ twins alongside updates to platform/dsl_docs/instances/attention_bwd.md and library/dispatch/AGENTS.md.
8	FAIL	PR description body under 'Not ported, on purpose' and 'Known gap, not from this PR' contains internal Atlassian tracker URLs (https://amd-hub.atlassian.net/browse/AICK-3015 and https://amd-hub.atlassian.net/browse/AICK-3016), and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:432 has an unused parameter live_. (Note: item 8 specifically forbids internal tracker URLs in public PR threads; also see dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7 for code reference).
9	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:1 and all other changed files in the PR are scoped to the SDPA backward C++ builder twins and their parity/golden tests.
10	NEEDS-INFO	Cannot verify open PR collisions and merge simulation from the truncated diff alone.
11	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:1 and the file list show only deliverable C++ twin sources, parity emitters, tests, and required reference docs are committed.
12	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:38 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:45 cover both gfx942 and gfx950.
13	N/A	This PR ports existing Python attention backward builders to C++ twins and does not remove dead code or rename concepts.
14	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/portable_ir/ir_import_json.cpp:1082 show removed lines are limited to updating 'Python-only, no C++ twin' docstrings and a portable-IR comment.
15	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_bwd_valu.cpp:41 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:74 use all accepted parameters in function bodies (only callback adapter signatures in attention_bwd_unified.cpp:253,432 intentionally ignore unused hook arguments to match KvEpilogue / row_source_t function types).
16	N/A	Severity honesty applies to how review findings are stated rather than the PR code itself.
17	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7 and the PR description contain no absolute or relative performance claims.

§F graded 11/11

#	Verdict	Evidence
1	PASS	Branch users/jafagani/rocke-bwd-cpp-stack matches users/<user>/... (see dnn-providers/hip-kernel-provider/rocke/library/kernels/common/attention_bwd_aux.py:6).
2	PASS	title feat(rocke): C++ twins of the SDPA backward kernel builders matches type(scope): summary and is 59 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/test_attention_bwd_dq_refusal_rows.py is in the diff
5	PASS	No keys, certificates, or secrets are present in the changed files (see dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:1).
6	NEEDS-INFO	PR body reports pre-commit, whole stack: clean, but local pre-commit execution status cannot be independently verified from the truncated diff alone.
7	NEEDS-INFO	GitHub review threads and bot/Copilot comments are not included in the prompt payload; deferred items in the PR body cite AICK-3015 and AICK-3016.
8	N/A	This PR targets develop directly (Base: develop), not a feature branch.
9	PASS	The PR keeps the full 16-commit port intact across 108 files without artificial splitting for size caps (see dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:1).
10	NEEDS-INFO	Parity emitter dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:86 is a real test probe, but the diff bodies of the test_* unit test files are omitted and needed to verify their contents.
11	PASS	The PR title (feat(rocke): C++ twins of the SDPA backward kernel builders) and body are populated and present in the PR metadata (see dnn-providers/hip-kernel-provider/rocke/library/kernels/common/attention_bwd_aux.py:6).

§G graded 34/34

#	Verdict	Evidence
1	PASS	dnn-providers/hip-kernel-provider/rocke/library/kernels/common/attention_bwd_aux.py:6 documents that the C++ twins provide byte-identity parity with the already-shipped Python backward builders rather than claiming new kernel math correctness.
2	N/A	This PR ports existing Python SDPA backward builders to emit byte-identical LLVM IR in C++ without changing the emitted kernel IR or numerical behavior.
3	NEEDS-INFO	The PR body lists gate runs, but the diffs for the test files and golden files (e.g., check_native_golden.py, NativeGoldenTests.cmake, llvm_gfx_all.json) are omitted and full local gate output (tools/run_checks.py, check_ir_validity.py, rocke_installed_golden_test.py) is not verifiable from the diff alone.
4	NEEDS-INFO	The diffs for all newly added unit and parity tests (e.g., test_buffer_load_hip.cpp, attention_bwd_spec_reinit.cpp, test_attention_bwd_dq_refusal_rows.py, test_check_native_golden.py) are omitted from the prompt budget.
5	NEEDS-INFO	Cannot verify falsifiability and reachability of the new tests because the diffs for test_attention_bwd_dq_refusal_rows.py, test_attention_bwd_layout_spill_rows.py, attention_bwd_spec_reinit.cpp, and attention_bwd_dma_slab_int64.cpp are omitted.
6	NEEDS-INFO	While code comments indnn-providers/hip-kernel-provider/rocke/platform/cpp/portable_ir/ir_import_json.cpp:1082 match the surrounding logic, the diffs for platform/dsl_docs/instances/attention_bwd.md, library/dispatch/AGENTS.md, and the refusal/spill row tests are omitted.
7	NEEDS-INFO	The diff for dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/instances/attention_bwd.md is omitted, so any changes to the design doc cannot be inspected.
8	N/A	This PR does not add or modify numerical execution harnesses or fp8 dtype handling (only fp16 and bf16 builder specs are ported).
9	N/A	This PR does not introduce new LDS or occupancy claims; it ports the existing Python LDS planner and register model to C++ for byte-identity.
10	N/A	This PR is a C++ builder port for byte-identity and makes no performance or benchmarking claims.
11	N/A	This PR does not add a new optimized kernel or change dispatch routing; it ports existing kernel builders to C++.
12	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:15 explicitly documents that the recorded resource-gate refusals (RESOURCE_REFUSED) stay Python-only because no C++ caller exists yet and the dispatcher remains in Python.
13	N/A	This PR does not introduce new LLVM intrinsic manglings or .td intrinsic recommendations.
14	N/A	This PR does not perform hardware bottleneck classification or kernel performance optimization.
15	N/A	This is a review process meta-rule for the reviewer rather than a code property of the PR.
16	PASS	dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:32 verifies the absence of a C twin for attention_fwd_ext.LOG2E and cites the exact nearest shared constant ROCKE_LSE_LN2 in rocke/instance_lse_store_internal.h.
17	PASS	dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:171 defines a runnable CLI emitter entry point exercised end-to-end by the parity and golden checks.
18	NEEDS-INFO	The PR adds parity and native-golden verification gates (check_native_golden.py, test_check_native_golden.py, test_attention_bwd_dq_refusal_rows.py, test_attention_bwd_layout_spill_rows.py), whose diffs are omitted from the prompt.
19	NEEDS-INFO	The diffs for the new gate fixtures and unit tests (test_check_native_golden.py, test_attention_bwd_dq_refusal_rows.py, test_attention_bwd_layout_spill_rows.py, attention_bwd_spec_reinit.cpp) are omitted.
20	N/A	This PR does not add or modify a claims-ban or text-sweep detector.
21	N/A	This PR ports kernel builders to C++ and adds IR/LL byte-identity checks, with no performance or statistical measurement code.
22	NEEDS-INFO	The PR adds 20 backward parity families to platform/tests/instances/differential/golden/llvm_gfx_all.json and adds check_native_golden.py, run_diff.py, and test_attention_bwd_ir_cpp_parity.py, but their diffs are omitted.
23	N/A	This PR makes no performance or routing claims; it is a byte-identical C++ port of Python kernel builders.
24	NEEDS-INFO	platform/tests/instances/differential/golden/llvm_gfx_all.json is modified in this PR, but its diff is omitted so it cannot be verified whether only new backward parity families were added or existing hashes were re-blessed.
25	N/A	This PR is a byte-identical C++ port of existing Python builders and does not change emitted IR or perform a codegen micro-optimization.
26	N/A	This PR does not include performance benchmarks.
27	N/A	This PR does not perform A/B performance experiments or toggle optimization knobs.
28	N/A	This PR does not report performance deltas or counterfactual comparisons.
29	PASS	In dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_bwd_valu.cpp:103-127, rocke_attn_bwd_softmax_p ports the existing bwd_exp2 lever (rocke_b_exp2_fast vs rocke_b_exp2) solely for softmax probability sv * scale_log2 - lse2 where the non-positive precondition holds.
30	N/A	This PR makes no new no-spill claims and emits byte-identical LLVM IR to existing Python builders.
31	N/A	This PR is a C++ port for IR byte-identity and makes no baseline vs. after defect/performance claims.
32	N/A	This PR reports no hardware performance or resource measurements.
33	N/A	This PR does not change any existing pure decision function, heuristic, or spec builder; it ports existing Python builders to C++ with byte-identity gates.
34	N/A	This PR does not add or modify any opt-in kernel output flags.

§H graded 5/5

#	Verdict	Evidence
1	NEEDS-INFO	Missing launch-side host-validation hunks showing tensors backing spec-derived buffer/resource bounds are shape/size-checked against the spec before compile or launch.
2	NEEDS-INFO	Missing paged-helper implementation and validation/test hunks showing BLOCK_TABLE physical values are range-checked or proven contained by an effective hardware bound.
3	NEEDS-INFO	Missing complete codegen/spec-validation hunks for all hot-loop div/mod sites, especially paged page/sub_page/block_n divisors, so power-of-two enforcement or an overhead justification cannot be confirmed.
4	NEEDS-INFO	Missing host-validation or proof hunks showing runtime sequence extents such as CU_SEQLENS_Q and SEQ_LENS_KV cover the full compile-time tile range or that the softmax masks any gap.
5	NEEDS-INFO	Missing complete host-validation and negative-test hunks for every launch input feeding device masks/bounds, including sequence lengths, block tables, worklist counts, page parameters, tile counts and strides.

§I graded 6/6

#	Verdict	Evidence
1	N/A	This PR ports SDPA backward kernel builders from Python to C++ for IR byte-identity and does not perform routing or tuning space measurements.
2	PASS	The C++ spec validators and builders loudly reject invalid, unsupported, or spilling configurations with explicit status codes and diagnostic messages instead of failing cryptically or running silently slow (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:419, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:74…
3	N/A	This PR adds deterministic C++/Python parity emitters for existing backward kernel families rather than modifying problem-space test coverage or hypothesis generators.
4	N/A	This PR does not add or modify performance benchmark tables comparing shipped versus achievable kernel selections.
5	PASS	The golden hash store is treated as an explicit blessed snapshot checked by CI, paired with live parity emitters (*_emit.py / *_emit.c) and differential scripts (run_diff.py, check_native_golden.py) that provide the regeneration and verification path (dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:171).
6	PASS	The C++ builders use general IR primitives such as rocke_b_warp_shuffle_xor (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:477) and enumerate all lanes/slots on a scratch builder to prove C fragment layout properties rather than asserting them (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_bwd_valu.cpp:224).

§J graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	The Markdown hunks for platform/dsl_docs/instances/attention_bwd.md and other Markdown files are omitted, so equations, operand annotations, and declared result shapes cannot be checked.
2	NEEDS-INFO	The Markdown hunks are omitted, so the notation table, symbol domains, and any log-base/domain bridges cannot be verified.
3	NEEDS-INFO	The relevant Markdown content is omitted, so it is not possible to tell whether the docs restate an algorithm owned by another document instead of linking and stating only the delta.
4	NEEDS-INFO	The Markdown hunks containing documentation claims are omitted, so novelty claims cannot be checked against develop or classified as inherited, added, or family-only.
5	NEEDS-INFO	The Markdown hunks are omitted, so any scaling or complexity prose and its repeated sites cannot be inspected.
6	NEEDS-INFO	The Markdown hunks are omitted, so constants, model-table provenance, tuned band edges, and stated assumptions cannot be verified.
7	N/A	No new Markdown document is shown as added, so the new-doc sibling filename and entry-point convention does not apply.

§K graded 13/13

#	Verdict	Evidence
1	N/A	This PR adds C++ kernel builder twins and parity tests without modifying or adding dispatcher candidates or KnobSpace subclasses.
2	N/A	No tunable kernel spec fields, KnobAxis declarations, or field classifications are added or changed in the Python tuning layer.
3	N/A	No knob definitions, defaults, canonicalization rules, or Python builder behaviors behind a knob are changed (C++ twins emit byte-identical IR to existing Python builders).
4	N/A	This PR does not touch KnobSpace.defaults() or KnobSpace.recorded().
5	N/A	This PR does not touch KnobSpace drop/refuse logic, inert(), base_value(), or KNOWN_WRONG_KNOBS.
6	N/A	This PR does not modify KnobSpace.validate or policy-free tuning spec construction.
7	N/A	This PR does not touch KnobSpace.base(req) or request normalization in the dispatcher.
8	N/A	This PR does not modify normalize_knobs or knob value coercion at the request boundary.
9	N/A	This PR does not touch benchmark recording, replay logic, spec_hash verification, or PinRefused handling.
10	N/A	This PR does not touch KnobSpace.find or bare tuning_id resolution.
11	N/A	This PR does not touch sweep entry points, pin_to_spec, or sweep harnesses.
12	N/A	This PR does not alter production stacks, accept_default, or base specs for auto selection.
13	N/A	This PR does not touch benchmark harnesses or make claims about sweep results driving auto selection.

§L graded 10/10

#	Verdict	Evidence
1	N/A	This PR ports SDPA backward kernel builders to C++ and does not add a new gfx target.
2	N/A	This PR does not add or modify any MMA op_id or catalog entry.
3	NEEDS-INFO	The C++ block-GEMM and fragment helpers (_attention_bwd_gemm.cpp and _attention_bwd_frag.cpp) that handle MMA selection are omitted from the diff, though _attention_bwd_valu.cpp:212-246 properly queries op->c_layout from the catalog.
4	PASS	Wave size is queried from the target via rocke_arch_target_from_gfx(arch)->wave_size in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:178 and arch_facts(arch).wave_size in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:299.
5	PASS	Arch-specific capabilities such as transpose-LDS reads (tr_read) and wide-K atoms (16x16x32) are gated on target capability flags (facts.has_tr_read and facts.wide_k_atom) and validated per arch in dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:265-272 and 515-524.
6	PASS	The PR does not touch core/arch/ or ArchTarget, nor does it import CK Tile's dispatcher; all kernel-family policy and spec vocabulary lives under platform/cpp/instances/common/ (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:47-59).
7	N/A	This PR does not add or modify logical dtypes, bit packing formats, or FP8/FP4/FP6 support.
8	N/A	This PR does not add, modify, or use scaled-MMA operations.
9	N/A	This PR does not add a new gfx architecture target.
10	N/A	This PR does not touch intrinsic_arch_domain..json artifacts or LLVM flavor membership lists.

§M graded 10/10

#	Verdict	Evidence
1	NEEDS-INFO	Need the untruncated Python and C++ builder hunks to verify no SSA/device-time Value is used in host-time control flow and that values defined inside scf regions do not escape improperly.
2	N/A	No target legalizer or legalize_for_target path is modified by this diff.
3	N/A	No ck.dsl.ir serialization, print_ir interchange, wire-version, attr-tagging, SSA-id, or parameter-order code is changed.
4	PASS	The visible portable-IR changes are comment-only clarifications of opcode alias resolution in dnn-providers/hip-kernel-provider/rocke/platform/cpp/portable_ir/ir_import_json.cpp:1079 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/portable_ir/recipe_vm.cpp:125, with no schema/version/importer behavior changed.
5	N/A	No AsyncTileLoader or async DRAM-to-LDS loader path is introduced or modified in the visible diff.
6	NEEDS-INFO	Need the untruncated kernel-body hunks, especially the omitted attention backward builders, to verify all barriers are outside divergent control flow and all potentially invalid memory accesses are masked or sentinel-protected before load.
7	NEEDS-INFO	Need the untruncated C++ builder hunks to verify every buffer resource uses backend-owned construction and no kernel hard-codes buffer-resource word 3 or misuses temporal-hint constants.
8	N/A	No intrinsic declaration table, llvm22/llvm23 override table, flavor check, or engine-precedence code is modified in the visible diff.
9	PASS	New visible C entry points are declared inside extern "C" and return kernel definitions rather than fabricating HIP launch behavior, for example dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:23 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:25.
10	N/A	The diff does not emit !dbg metadata or add consumers that read debug information.

§N graded 18/18

#	Verdict	Evidence
1	N/A	No new helper is promoted into platform/python/rocke/helpers/ in the shown diff; the existing attention_band helper only receives documentation about its C++ twin.
2	N/A	The shown diff does not add or change tiling output mapping, TensorView addressing, transform DAG code, loaders, epilogues, pipelines, or helpers/transforms.py.
3	N/A	The shown diff does not add or change SoftwarePipeline or overlap_vmcnt behavior.
4	N/A	The shown diff does not add or change fused kernels, EpilogueOp composition, fusion validation, or fusion legalizers.
5	N/A	The shown diff does not introduce a new public tiling API or change rocke.helpers.tiling exports.
6	N/A	The shown diff does not add front-door/toolbox tiling helpers or a new LayoutStyle.
7	N/A	The shown diff does not change tiling layout overrides, accumulator descriptor derivation, or MMA soundness validation.
8	N/A	The shown diff does not change tiling LDS allocation, lds_store/lds_read, contraction-axis padding, or wide global-load tiling rules.
9	N/A	The shown diff does not change N-D tiling axis-role declarations, label flow, at_index, or squeeze behavior.
10	N/A	The shown diff does not change tiling emit, LDS staging, MMA issue, or tiling analysis gates.
11	N/A	The shown diff does not change tiling atom traits, LDS swizzle policy, scheduling group counts, or iglp tiling interactions.
12	N/A	The shown diff does not make LDS bank-conflict or coalescing claims in the tiling analysis code.
13	N/A	The shown diff does not change WarpDistributionEncoding, make_tile_desc factoring, RegisterMapper exports, or tiling store casts.
14	N/A	The shown diff does not change tiling verbs or add a new tiling pipeline variant.
15	N/A	The shown diff does not change tiling allocation or bounds checks for transformed descriptor address ranges.
16	N/A	The shown diff does not change optimize_layout, recommend(), evaluate_transform, or tiling layout ranking.
17	N/A	The shown diff does not change the tiling recorder, tiling witness flags, round-trip analysis, or visualization emission behavior.
18	N/A	The shown diff does not change interleaved tiling layouts, tiling demo kernels, or numeric gates for exact tiling operations.

§O graded 16/16

#	Verdict	Evidence
1	NEEDS-INFO	Most new C++ builder/header hunks are omitted or unavailable; need the full IR-building and verification-related diffs to confirm no new attrs/address spaces and that hand-built IR is verified consistently.
2	PASS	The new worklist loop declares two iter-args and yields two values, preserving scf.for result arity: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_bwd_worklist.cpp:91 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_bwd_worklist.cpp:112.
3	PASS	New kernel-name knob ordering is made deterministic with a stable sort by field name: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:365.
4	PASS	The visible new type comparison uses value equality via rocke_type_eq rather than pointer identity: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_bwd_valu.cpp:82.
5	N/A	The visible diff does not edit optimize_kernel, canonicalize_region, pass purity tables, or C++ pass implementations.
6	N/A	The visible diff does not edit constant-folding logic in either engine.
7	FAIL	The new aux builder hard-codes the flat-work-group minimum as a fixed 64 instead of deriving the minimum from the target wave/workgroup contract: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:37.
8	PASS	The split-dQ spec rejects unsupported scheduler_strategy strings with rocke_scheduler_strategy_is_valid before emission: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:172.
9	PASS	The dense backward shell emits synchronization through the fenced barrier helper rather than raw/named barrier text: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_dense_bwd_shell.cpp:62.
10	N/A	The visible diff does not change lowerer attr validation or pointer parameter lowering code.
11	FAIL	The dense backward C++ entry point makes an architecture decision by direct gfx string comparison: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:42.
12	FAIL	The dense backward shell hard-codes the tensor element byte size instead of deriving it from the IR dtype size: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_dense_bwd_shell.cpp:41.
13	NEEDS-INFO	Public headers under cpp/include/rocke are omitted or unavailable, so it cannot be verified whether public C structs/enums grew or reordered and whether ABI/version/static_assert updates were made.
14	NEEDS-INFO	Many new .cpp and header hunks plus platform CMake hunks are omitted or unavailable, so style guards, recursive CMake inclusion, warning flags, and minimal binding changes cannot be verified.
15	N/A	The visible diff does not change engine fallback handling, pybind fallback records, or backend flavor resolution paths.
16	PASS	The dense backward C++ entry point is codegen-only and returns a rocke_kernel_def_t builder result rather than adding HIP launch or KernelLauncher behavior: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:26.

§P graded 14/14

#	Verdict	Evidence
1	PASS	C++ builders take (b, spec, arch) matching Python's (spec, *, arch) contract and validate arch-specific spec classes (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:26 and :48; dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:912).
2	PASS	The C++ ports declare kernel arguments in the exact existing Python ABI order without modifying parameter order or ABI versions (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:362, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:200, dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:283).
3	N/A	This PR does not touch GEMM, implicit-GEMM, or gfx950 block-scaled GEMM launch geometry or host launch harnesses.
4	PASS	The C++ LDS planner and validator mirror Python's phase-based LDS sizing and check plan.total_bytes > plan.capacity_bytes against arch_facts(arch).lds_capacity_bytes (dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_dq.cpp:415 and :484).
5	N/A	This PR does not touch WMMA/RDNA GEMM, GemmMultiD, or GemmMultiAbd paths.
6	N/A	This PR ports SDPA backward kernel builders to C++ and does not touch Stream-K, MoE, split-K, or host-side workspace allocation/zeroing.
7	N/A	This PR does not touch Stream-K, persistent GEMM, GroupedGemmLauncher, or SUPPORT_MATRIX.md.
8	N/A	This PR does not touch quantization helpers or MX GEMM kernels.
9	N/A	This PR ports existing Python attention backward builders (rne_noguard_bf16, softmax_p) to C++ for byte-identical IR parity without introducing new numeric emulation algorithms or error bounds.
10	N/A	This PR does not add or modify manifest kinds, emitters, or runtime manifest consumers.
11	N/A	This PR does not touch kernel launching, stream synchronization, WorkspacePool, or timing utilities.
12	N/A	This PR does not touch COMGR or HIP library loading.
13	N/A	This PR does not refactor host-side grid/workspace math, kernel-cache keys, or admission guards.
14	PASS	No expected.json files are touched, and C++ twins of family glue are placed under platform/cpp/instances/common/ and platform/cpp/helpers/ matching their Python counterparts (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_bwd_valu.cpp:1, dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/helpers/attention_band.py:30).

§Q graded 7/7

#	Verdict	Evidence
1	NEEDS-INFO	The diffs for ROCKE_ENGINE_test_categories_external.yaml, platform/CMakeLists.txt, platform/tests/CMakeLists.txt, platform/tests/cmake/{CoreTests,InstanceTests,NativeGoldenTests}.cmake, and test_cmake_test_registration.py are omitted.
2	NEEDS-INFO	rocke_ir_parity_harness.py is not in the file list and the diffs for test_attention_bwd_ir_cpp_parity.py, NativeGoldenTests.cmake, and check_native_golden.py are omitted; need those diffs or confirmation of whether rocke_ir_parity_harness.py should be updated alongside the parity pairs.
3	NEEDS-INFO	The diffs for platform/tests/instances/differential/run_diff.py, check_native_golden.py, and test_check_native_golden.py are omitted.
4	NEEDS-INFO	The diffs for platform/tests/instances/differential/golden/llvm_gfx_all.json, golden/README.md, and run_diff.py are omitted.
5	NEEDS-INFO	The diffs for all new test files (test_buffer_load_hip.cpp, attention_bwd_dma_slab_int64.cpp, attention_bwd_spec_reinit.cpp, check_native_golden.py, test_check_native_golden.py, and library/tests/test_attention_bwd_*.py) are omitted.
6	NEEDS-INFO	The diffs for platform/tests/instances/differential/run_diff.py and check_native_golden.py are omitted.
7	NEEDS-INFO	The diffs for library/tests/test_attention_bwd_dq_refusal_rows.py, test_attention_bwd_layout_spill_rows.py, and attention_bwd_spec_reinit.cpp are omitted.

§R graded 11/11

#	Verdict	Evidence
1	N/A	This PR ports existing Python attention backward builders to byte-identical C++ twins without introducing a new kernel algorithm, tiling strategy, pipeline variant, or structural performance claim.
2	N/A	This PR changes no performance lever; it only creates C++ twins that emit byte-identical LLVM IR to the existing Python builders.
3	N/A	This PR touches no optimization levers, tile sizes, LDS layouts, or launch configurations, as it is a byte-identical C++ port of existing Python builders.
4	N/A	This PR makes no performance or codegen changes—the C++ twins are verified to emit byte-identical LLVM IR to the existing Python builders.
5	N/A	This PR does not add or change any global memory access patterns; it ports existing Python builders to emit byte-identical IR in C++.
6	N/A	This PR does not change any LDS or MMA layouts; it ports existing Python builders to C++ with byte-identical IR output.
7	N/A	This PR makes no bottleneck claims or performance optimizations.
8	N/A	This PR does not change any waits, barriers, or scheduling hints; the C++ port emits byte-identical IR to the existing Python implementation.
9	N/A	This PR is not a performance regression fix and does not trade performance across shapes.
10	N/A	This PR makes no codegen policy changes, assembly parity claims, occupancy conclusions, or benchmark comparisons.
11	N/A	This PR introduces no new optimization lessons, probes, or analysis tools.

§S graded 8/8

#	Verdict	Evidence
1	NEEDS-INFO	The PR adds new rocke_* C entry points (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:801-946, attention_bwd_aux.cpp:46-214, attention_dense_bwd.cpp:26) and headers under platform/cpp/include/rocke/, plus CMake/YAML test registrations (platform/tests/cmake/*.cmake, ROCKE_ENGINE_test_categories_external.yaml), but the header diffs, emitter diffs, and test registration diffs are omitted; test_attention_bwd_ir_cpp_parity.py is also noted in the PR body as not in any ctest list.
2	NEEDS-INFO	Visible Python changes in dnn-providers/hip-kernel-provider/rocke/library/kernels/common/*.py and platform/python/rocke/helpers/attention_band.py:26-32 are docstring-only, but diffs for platform/tests/instances/differential/run_diff.py, platform/tests/tools/test_cmake_test_registration.py, and ROCKE_ENGINE_test_categories_external.yaml are omitted, so it cannot be verified from the diff that no existing API behavior or test expectation was altered or narrowed.
3	N/A	The PR only adds C++ twins and tests for attention backward builders and does not remove or rename any public API.
4	N/A	No __all__ lists or Python package exports are modified; the only file touched under platform/python/rocke/ is a docstring update in…
5	NEEDS-INFO	New C entry points include fail-fast validation errors (e.g. dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:69-105, attention_bwd_unified.cpp:808-899, attention_dense_bwd.cpp:40-70), and although library/tests/parity/attention_bwd_band_emit.py:86-168 is visible, the diffs for the C parity drivers (library/tests/parity/*_emit.c) and unit tests (platform/tests/instances/attention_bwd_spec_reinit.cpp, attention_bwd_dma_slab_int64.cpp, platform/tests/core/test_buffer_load_hip.cpp) are omitted.
6	NEEDS-INFO	The new tests (test_buffer_load_hip.cpp, attention_bwd_dma_slab_int64.cpp, attention_bwd_spec_reinit.cpp, check_native_golden.py, test_attention_bwd_dq_refusal_rows.py, test_attention_bwd_layout_spill_rows.py) and ROCKE_ENGINE_test_categories_external.yaml appear to target CPU lanes, but their diffs are omitted so CPU-only execution without torch/GPU cannot be confirmed from the diff.
7	NEEDS-INFO	Each new C++ twin has a parity emitter pair (e.g. dnn-providers/hip-kernel-provider/rocke/library/tests/parity/attention_bwd_band_emit.py:1-180 and attention_bwd_band_emit.c) and golden harness (platform/tests/instances/differential/check_native_golden.py, NativeGoldenTests.cmake), but the diffs for check_native_golden.py, NativeGoldenTests.cmake, and the .c emitters are omitted.
8	NEEDS-INFO	The PR body lists run_diff.py --python-only --check-golden, ctest (50/50), check_byte_identity.py, and targeted pytests, but does not state that PYTHONDONTWRITEBYTECODE=1 PYTHONPATH=python:../library python tests/test_rocke.py and python tests/run_all.py were run from rocke/platform/; please confirm those README entry points were executed.

[yraparti]

[HIGH] §D.1 ★ ★ Knobs/features on the KernelSpec (validator + selector + name tag + cache-key), NOT os.environ (env only as a selector-level diagnostic override folded onto the spec). Reviewers may reject env vars outright — even a reproducibility/pin one (#9739: "please dont use env variables", on a…

dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/differential/check_native_golden.py:149 reads the environment

Fix: Put the knob on the KernelSpec and delete the environment read.

[yraparti]

[HIGH] §D.20 Selection-layer fix in a kernel-level selector or spec builder mirrored in C++ (shipped provider consumes C++ selectors) — now a LIVE break. A fix confined to the Python registry/tuning layer has no C++ consumer; the PR says so instead of claiming parity.

The new C++ main builder documents a Python/C++ spec-validation divergence: it mirrors _emit_main(..., main_kernel_refusals=False) and says recorded resource-gate refusals stay Python-only, so C accepts RESOURCE_REFUSED specs that Python refuses at dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_main.cpp:19.

Fix: Port the RESOURCE_REFUSED resource-gate refusals to the C++ spec validation path, or make the C++ builder refuse the same specs with the same reasons before building.

[yraparti]

[HIGH] §O.11 ★ ★ Hardware decisions query ArchTarget (from_gfx, fits_lds, lds_capacity_bytes, supports_dtype_combo, async_lds_max_dwords, has_*, the MMA catalog), never gfx.startswith(…) or string family compares. Resolve the gfx string once, never mutate the lru_cached ArchTarget or…

The dense backward C++ entry point makes an architecture decision by direct gfx string comparison: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_dense_bwd.cpp:42.

Fix: Resolve the arch once through ArchTarget and gate dense-builder support on target/catalog capabilities instead of comparing gfx strings directly.

[yraparti]

[MEDIUM] §E.8 No Jira/tracking IDs, reviewer usernames/GitHub handles, or PR numbers in code (comments, docstrings, strings, kernel names, tests). Use a descriptive reference; tracker + commit message + PR description only. In public PR threads and replies use the bare ticket ID, never an internal tracker URL.…

PR description body under 'Not ported, on purpose' and 'Known gap, not from this PR' contains internal Atlassian tracker URLs (https://amd-hub.atlassian.net/browse/AICK-3015 and https://amd-hub.atlassian.net/browse/AICK-3016), and dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_unified.cpp:432 has an unused parameter live_. (Note: item 8 specifically forbids internal tracker URLs in public PR threads; also see dnn-providers/hip-kernel-provider/rocke/library/kernels/common/_attention_bwd_addr.py:7 for code reference).

Fix: Replace the internal https://amd-hub.atlassian.net/browse/AICK-... links in the PR description with bare ticket IDs (AICK-3015, AICK-3016).

[yraparti]

[MEDIUM] §O.7 Kernel attributes keep their contracts and order: amdgpu-flat-work-group-size min is min(wave_size, max_workgroup_size) (an inverted range is silently discarded by LLVM); waves_per_eu is an int or (lo, hi); agpr_alloc is (min, max) with min ≤ max, and (0, 0) stays coupled to the…

The new aux builder hard-codes the flat-work-group minimum as a fixed 64 instead of deriving the minimum from the target wave/workgroup contract: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/attention_bwd_aux.cpp:37.

Fix: Replace the fixed AUX_MIN_WG_ATTR constant with a value derived from the resolved target wave size and max workgroup size before setting kernel attributes.

[yraparti]

[MEDIUM] §O.12 Widths and size caps come from the target catalog or the dtype size, per arch, in both engines (vector-load bytes, async_lds_max_dwords, element size): no literal 16 bytes, 4 dwords or * 4, and non-positive values are rejected. No new silent arch or "gfx950" defaults: a forgotten arch=…

The dense backward shell hard-codes the tensor element byte size instead of deriving it from the IR dtype size: dnn-providers/hip-kernel-provider/rocke/platform/cpp/instances/common/_attention_dense_bwd_shell.cpp:41.

Fix: Use a dtype-size helper or target/catalog dtype metadata for element-byte calculations instead of the ELEM_BYTES literal.