Skip to content

fix(rocke): take the HIP buffer descriptor word from the target - #13886

Merged
cgmillette merged 3 commits into
developfrom
users/cgmillette/rocke/hip-buffer-descriptor
Oct 23, 2026
Merged

cgmillette merged 3 commits into
developfrom
users/cgmillette/rocke/hip-buffer-descriptor

Conversation

@cgmillette

@cgmillette cgmillette commented Oct 17, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

The HIP lowering of a buffer resource (tile.buffer_rsrc) passed the gfx9 descriptor word 0x00027000 as the flags argument of __builtin_amdgcn_make_buffer_rsrc on every target. That word's format and out-of-bounds encoding is ISA-specific: on gfx11 and gfx12 it makes every raw buffer load read 0 and every store drop. The LLVM lowering already takes the word from the target.

ISSUE ID : AICK-2573

Technical Details

  • Both engines' HIP lowerings now read the descriptor word from the same per-target source as their LLVM lowering: Python backend_for(arch).buffer_rsrc_word3, C++ rocke_ll_backend_for(...)->buffer_rsrc_word3.
  • The C++ backend lookup is now read-only: rocke_ll_backend_for returns the target's immutable table row instead of overwriting one process-wide copy, and each LLVM lowering fills the triple and its flavor's datalayout in its own copy. Concurrent HIP and LLVM lowerings for different targets or flavors can no longer change each other's backend. The emitted IR is unchanged.
  • gfx9 targets and the default arch keep 0x00027000; gfx11 and gfx12 get 0x31014000. gfx1250 also gets 0x31014000 today, because its backend inherits the gfx12 word (unchanged by this PR).
  • No LLVM IR changes. HIP backend docs describe the per-target word.
  • New native test rocke_buffer_rsrc_hip (both engines' words agree per target), registered in the CTest tiers.

Test Plan

  • The HIP word pinned per target family; HIP word equal to the LLVM word per target in both engines; Python and C++ HIP text byte-identical for every target (new native test binary).
  • gfx11 / gfx12 HIP source compiled with hipcc; the word checked in the device IR.
  • A bounds-checked buffer copy through the HIP route on the current GPU, lowered for its target (or its family's generic target): lanes inside both bounds copy, lanes past the load bound read 0, lanes past the store bound are untouched.

Test Result

  • Buffer copy through the HIP route, before and after this change:
GPU develop this PR
gfx1201 wrong: 112 of 128 elements untouched correct
gfx1100 wrong: 112 of 128 elements untouched correct
gfx90a correct correct
  • Byte identity green at llvm20, llvm22 and llvm23; ctest passes; strict both-engines tests pass.
  • Full platform suite: no new failures against develop.
  • pre-commit passes.

Submission Checklist

🤖 Generated with Claude Code

@therock-pr-bot

therock-pr-bot Bot commented Oct 17, 2026 •

Copy link
Copy Markdown

✅ 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

therock-pr-bot Bot commented Oct 17, 2026 •

Copy link
Copy Markdown

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

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

🟡 Changes recommended

Shared mutable backend state can corrupt concurrent cross-target lowering, and the GPU test can compile for a different device than it executes on.

2 open issues
What changed in this PR

This PR makes HIP buffer descriptors target-specific, aligning both engines with LLVM lowering while preserving gfx9 behavior.

Changes:

  • Read descriptor flags from each target’s ISA backend.
  • Add descriptor, engine-parity, compilation, and GPU checks.
  • Register native tests and update lowering documentation.
File Description
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​tools/​test_run_all.py Verify native fixture discovery and forwarding.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​run_all.py Register the descriptor test fixture.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​test_buffer_rsrc_hip.py Add descriptor and HIP execution checks.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​test_buffer_rsrc_hip.cpp Add native descriptor checks and parity output.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​CMakeLists.txt Build, register, and install the native test.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​core/​lower_llvm.py Clarify descriptor documentation only.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​core/​lower_hip.py Emit target-specific descriptor flags.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​ir_lowering/​lowering_pipeline.md Document target-derived flags.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​ir_lowering/​backend_details.md Document HIP/LLVM descriptor agreement.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​core/​lower_hip/​lower_hip_mem.cpp Obtain flags from the C++ ISA backend.
dnn-providers/​hip-kernel-provider/​ROCKE_ENGINE_test_categories_external.yaml Include the native test in test tiers.

Review effort: Balanced

@cgmillette
cgmillette enabled auto-merge (squash) October 19, 2026 20:42
@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/hip-buffer-descriptor branch 4 times, most recently from e5ce600 to 299941e Compare October 20, 2026 22:48
@cgmillette

Copy link
Copy Markdown
Contributor Author

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

PR #13886 — take the HIP buffer descriptor word from the target — 🟢 First-round: no closer look needed — bug and fix reproduced on gfx11 hardware

Prior item Status
Copilot (10-17) HIGH: rocke_ll_backend_for() returns shared mutable LL_BACKEND_RESOLVED → cross-target race 🟢 FIXED de92827d — the static is gone; rocke_ll_backend_for returns the target's const table row, each LLVM lowering copies it into L->backend_resolved and fills triple / datalayout there, and the new HIP call only reads it (read)
Copilot (10-17) MEDIUM: GPU test compiles for device 0, launches on the current device 🟢 FIXED 299941e0 — compiles for current_device_arch(), falling back to the family's generic target; ran on gfx1151 (ran)
Bots / CI 🟢 Policy bot, pre-commit, clang-tidy pass; TheRock CI hipkernelprovider green on Linux gfx94X and Windows gfx110X; hipDNN superbuild green. No human reviews yet (4 requested)
Our checks 🟢 On a gfx1151 GPU, the PR's bounds-checked HIP-route copy passes with this PR and fails on develop (112 of 128 elements wrong) — the bug and the fix both reproduced on real gfx11 hardware (ran). Merged with current develop: native build clean, new rocke_buffer_rsrc_hip ctest passes (ran); the file's other 27 tests pass, including Python ↔ C++ HIP text byte-identical for all 8 targets via the native binary, and hipcc device IR carrying 0x31014000 on gfx1151 / gfx1201 (ran). Processors without a backend (gfx1100, gfx1102, gfx1200) are rejected identically on develop and the PR — no new failure mode (ran). No conflicts; rule_scan clean (ran)
# Sev file:line Finding
N1 🟡 tests/core/test_buffer_rsrc_hip.py:15 The module docstring still says the copy runs "through the HIP route on a gfx9 GPU"; since 299941e0 it runs on the current device, gfx11 / gfx12 included → update the sentence.

Utility: 🟢 Positive — fixes silent data loss on the HIP route for gfx11 / gfx12 (raw buffer loads read 0, stores dropped) and removes a cross-lowering race; small, and now verified on the hardware it affects.


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

First-round review: a local qwen3.6 model made the independent first pass; the orchestrator (Claude Opus 5.5) verified it on hardware — native build + ctest, engine-parity and hipcc tests, and the HIP-route copy on a gfx1151 GPU against develop — and owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

@cgmillette
cgmillette disabled auto-merge October 21, 2026 02:52

@CongMa13 CongMa13 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

No new issues found. I agree with the following existing comments:

  • Update module docstring in test_buffer_rsrc_hip.py to reflect multi-device target execution: cgmillette (#13886 (comment))

@CongMa13 CongMa13 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

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

@cgmillette
cgmillette enabled auto-merge (squash) October 22, 2026 02:50
@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/hip-buffer-descriptor branch from 299941e to 8569306 Compare October 22, 2026 15:09
@cgmillette

Copy link
Copy Markdown
Contributor Author
  • Copilot, lower_hip_mem.cpp (shared mutable backend state): already fixed in "fix(rocke): make the C++ ISA backend lookup read-only" (now f8bd159). rocke_ll_backend_for returns the target's read-only table row; there is no process-wide copy left, each LLVM lowering fills its own backend_resolved, and the HIP lowering only reads buffer_rsrc_word3. Re-checked after the rebase: no other caller writes through the returned pointer.
  • Copilot, test_buffer_rsrc_hip.py (compile for device 0, launch on the current device): already fixed in "test(rocke): run the HIP buffer copy on the current GPU's target" (now 8569306). The test compiles for current_device_arch(), which reads HIP's current device, the same one Runtime loads and launches on.
  • N1 / CongMa13 (module docstring still says "on a gfx9 GPU"): fixed in 8569306. The docstring now says the copy runs on the current GPU, lowered for its target or its family's generic target, so gfx9, gfx11 and gfx12 GPUs all run it.
  • Rebased onto current develop; the only conflict was the fixture setup in tests/tools/test_run_all.py, where both develop's new fixtures and this PR's buffer descriptor fixture are kept.

@cgmillette

Copy link
Copy Markdown
Contributor Author

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

PR #13886 — take the HIP buffer descriptor word from the target — 🟢 Follow-up: no closer look needed — rebase clean, our nit fixed, GPU fix re-verified on the merge

Prior item Status
Ours (10-20) N1 🟡: the module docstring says the copy runs "on a gfx9 GPU" 🟢 FIXED — it now says the current GPU, lowered for its target or for its family's generic target (read)
Copilot (10-17) HIGH / MEDIUM: shared LL_BACKEND_RESOLVED race; compiling for device 0 🟢 Still fixed — the rebase leaves those changes intact (range-diff)
Rebase (tier-YAML conflict with #13964) 🟢 rocke_qk_scale_spec_validate and rocke_buffer_rsrc_hip are both kept in quick and ffm-quick, and the test_run_all.py fixture merge keeps develop's new variables (read). Merges cleanly with today's develop (27 behind) (ran)
CongMa13 APPROVED (10-21) at 299941e0 🟡 Given before the rebase; the conflict resolution and the docstring have changed since
Bots / CI 🟠 Re-run at head 8569306e (10-22) and still running: pre-commit, clang-tidy, the ASAN summary, Component CI and the hipDNN Windows superbuild pass; TheRock multi-arch (compiler-runtime stage) and the hipDNN Linux superbuild are running
Our checks 🟢 On a gfx1151 GPU, the HIP-route copy passes on the merge and still fails on develop (112 of 128 elements wrong) (ran). test_buffer_rsrc_hip.py passes 28 / 28 with both engines, the native binary and the GPU (ran). Native build and 28 ctests pass; the representative golden passes with no drift; rule_scan is clean (ran)
# Sev file:line Finding
N1 🟡 tests/run_all.py:148; #14050's platform/CMakeLists.txt list Land order with #14050, which adds an installed-lane copy of NATIVE_PARITY_TESTS. The two PRs conflict in run_all.py and test_run_all.py. Unless ROCKE_BUFFER_RSRC_HIP_TEST=rocke_buffer_rsrc_hip joins that CMake list, 8 native-parity cases here skip in the installed lane; with the pair added, they run and pass (ran, in a scratch merge of both). This PR, #13885 and #13954 also insert at the same tier-YAML line, so the later ones need a rebase that keeps both entries (ran).

Utility: 🟢 Positive — fixes silent data loss on the HIP route for gfx11 / gfx12 and removes a cross-lowering race; re-verified on the hardware it affects.


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

Follow-up review: a local qwen3.6 model made the independent first pass. The orchestrator (Claude Opus 5.5) verified it with a range-diff against the reviewed head, the HIP-route copy on a gfx1151 GPU against develop, the native build and tests, the golden gate, and scratch merges with #14050, #13885 and #13954, and owns this merge-worthiness call. It flags whether a closer look is needed; it is not an approval or a rejection.

@yraparti yraparti left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

rocKE checklist review

ROCm/rocm-libraries#13886 · 8569306e8015 · model gemini-4-argon, fallback gpt-5.5

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

Findings

  • HIGH §D · dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:113 · Put the knob on the KernelSpec and delete the environment read.
  • HIGH §Q · dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:1 · Register tests/core/test_buffer_rsrc_hip.py in the installed rocke_pytest list in platform/CMakeLists.txt (with ROCKE_BUFFER_RSRC_HIP_TEST wired if needed) and in ROCKE_ENGINE_test_categories_external.yaml.

Needs info

  • §F.6 Local pre-commit execution status is not verifiable from the diff alone (though formatting in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:1 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1035 appears clean).
  • §F.7 Bot, Copilot, and human review threads are not included in the provided PR context.
  • §F.11 The PR title and body are populated in the prompt metadata, but live GitHub API verification (gh pr view --json body,title) is not a path:line in the diff.

Section notes

  • §A: Both Python and C++ HIP lowerings of tile.buffer_rsrc are updated in lockstep, with C++ backend lookup made thread-safe and cross-engine HIP parity verified by a new native/Python test pair.
  • §D: This PR only modifies platform IR/HIP/LLVM lowering and platform tests; no dispatcher registration, routing, or KernelSpec changes are made.
  • §D: Platform-only PR modifying HIP/LLVM buffer resource lowering and tests; does not touch dispatcher registration, routing, candidates, or selectors.
  • §G: Items 1-17 of 34
  • §G: Items 18–34 of Section G graded for PR #13886.
  • §H: Platform IR-lowering PR (touches platform/cpp/core/lower_hip/, platform/cpp/core/lower_llvm/, platform/python/rocke/core/lower_hip.py, and unit tests). No library kernel, spec builder, dispatcher, or paged/indirection logic is added or modified. In the synthetic unit-test copy kernel (test_buffer_rsrc_hip.py:56), src_rsrc bound comes from runtime param load_bytes (160B in test) and dst_rsrc bound comes from store_bytes (224B in test) while the 64-lane wave visits [0, 256B) intentionally to verify hardware OOB zero-on-load and drop-on-store semantics.

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 N/A This PR does not modify or add any kernel-body emitters.
2 PASS The tile.buffer_rsrc HIP lowering change is mirrored in both Python (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243) and C++ (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036).
3 N/A This PR does not touch kernel-level selectors, spec builders, or codegen knobs.
4 PASS In the new parity test dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:52, C++ IR builder calls bind operands to temporaries in separate statements to match Python's evaluation order in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:61.
5 PASS LLVM lowering is unchanged (only rocke_ll_backend_for was made non-mutating in dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:287), and no representative-IR golden files are modified in the diff.
6 PASS Cross-engine parity is directly verified by comparing C++ and Python lowered HIP text in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:112 via ROCKE_BUFFER_RSRC_HIP_TEST (dnn-providers/hip-kernel-provider/rocke/platform/tests/run_all.py:148).
7 PASS For gfx9 targets and the default arch, backend_for(self.arch).buffer_rsrc_word3 resolves to the same 0x00027000 word as before (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243, tested at dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:105).
8 PASS The C++ changes in dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1035 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2319 are exercised by the native C++ test dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:154.
9 PASS The shared op-handler fix is proven via the mini-pair dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:45 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:56.
10 N/A This PR fixes a bug in existing HIP lowering of tile.buffer_rsrc and does not introduce a new feature flag or change LLVM IR.
11 N/A This PR does not add a new IR opcode.
12 N/A This PR does not touch ROCKE_BACKEND=both fallback handling or core/backend.py::CPP_UNPORTED_ARCHES.
13 PASS The changes in dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2319 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1041 do not add any build id, version, timestamp, or content hash to emitted output.
14 N/A This PR does not modify any classic helpers under helpers/.
§B graded 6/6
# Verdict Evidence
1 PASS In dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1037-1040, the new check rejects unknown architectures with ROCKE_ERR_KEY when no ISA backend is registered, matching the Python backend_for(self.arch) lookup in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243.
2 N/A This PR does not modify or introduce any preflight gates or kernel spec constructors.
3 N/A This PR fixes HIP buffer resource lowering and does not touch model-zoo dispatch or fallback routing.
4 PASS Both HIP lowerers reject unknown target arch strings at the lowering boundary rather than normalizing or falling back to a default word (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036-1040 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243).
5 N/A No dispatch or spec gates are opened or modified in this PR.
6 N/A This PR does not touch any is_valid_spec or supports_* functions.
§C graded 10/10
# Verdict Evidence
1 PASS Instead of duplicating target-specific word3 logic, both HIP lowerers query the existing ISA backend helpers (backend_for in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 and rocke_ll_backend_for in dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036).
2 N/A This PR modifies platform HIP/LLVM IR lowering and adds a unit test, rather than authoring a kernel family builder.
3 N/A This PR does not add or rename any production kernel builders; the only kernel is a unit test fixture (buffer_rsrc_copy).
4 PASS Target-specific descriptor words are looked up via the ISA backend table rather than branching on arch in the shared lowering code (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243, dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036).
5 PASS All 13 modified/added files are directly tied to the HIP buffer descriptor word fix across Python and C++ engines, docs (dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/backend_details.md:150), and test suites (dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:171).
6 PASS To safely share rocke_ll_backend_for with the C++ HIP lowerer without mutating global state, LL_BACKEND_RESOLVED was refactored out of rocke_ll_backend_for into per-lowerer L->backend_resolved (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:287, dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2319).
7 N/A This PR is a targeted bug fix rather than a standalone cleanup refactor.
8 PASS The PR reuses the existing backend_for (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:43) and rocke_ll_backend_for (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:20) rather than inventing a new arch-lookup helper.
9 PASS The hardcoded 0x00027000 bug was fixed across both the Python HIP lowerer (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2249), its C++ twin (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1045), and the prose docs (dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/backend_details.md:150…
10 PASS Widening rocke_ll_backend_for to be called by the HIP lowerer was accompanied by making its return value an immutable table row (dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h:223) and updating the LLVM lowerer consumer to populate datalayout and triple on its own L->backend_resolved copy (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2319).
§D graded 21/21
# Verdict Evidence
1 FAIL dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:113 reads the environment
2 N/A This PR does not add or modify any dispatcher candidates or Capability declarations.
3 N/A This PR does not add or modify any candidate _supports predicates or admits() calls.
4 N/A This PR does not touch dispatcher selectors, spec builders, or the selection path.
5 N/A This PR does not modify KernelSpec identity, cache keys, or kernel_name() generation.
6 N/A This PR does not touch cohort routing, spec-builders, or launch metadata.
7 N/A This PR does not add or modify any dispatcher heuristics or hand-pinned overrides.
8 N/A This PR does not define or modify any dispatcher cohort predicates.
9 N/A This PR does not register a new kernel or arch candidate in the dispatcher registries.
10 N/A This PR does not add any new spec fields or features.
11 N/A This PR does not register, promote, or extend arch coverage for any dispatcher candidate.
12 N/A This PR only modifies platform HIP/LLVM lowering and tests and does not add or modify any dispatcher candidates, priorities, or rankers.
13 N/A This PR does not touch dispatcher selection, pin handling, or entry points.
14 N/A This PR does not modify kernel identity keys, cache keys, or candidate kernarg layouts/ABI versions.
15 N/A This PR does not add or modify any dispatcher candidate or arch module in the dispatch registry.
16 N/A This PR does not introduce a new operator family in the dispatcher.
17 N/A This PR does not add or change dispatcher candidates or CPU routing tests.
18 N/A This PR does not modify any CandidateRegistry settings, binding requirements, or exempt sets.
19 N/A This PR does not touch attention dispatch or make any claims about run_unified_attention_torch registration.
20 N/A This PR is a platform lowering change (mirrored in Python and C++) and does not modify any kernel-level selector, spec builder, or registry/tuning layer.
21 N/A This PR fixes an ISA-specific HIP lowering bug for tile.buffer_rsrc rather than working around a spec selection issue in the emitter.
§E graded 17/17
# Verdict Evidence
1 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:33-35 imports only from rocke.core, rocke.helpers, rocke.runtime, and rocke.tests with no library dependency.
2 N/A This PR only touches files under platform/ and does not modify library/.
3 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:31-32 uses numpy and pytest without importing torch.
4 N/A No files under library/ or KNOWN_BAD in check_ir_validity.py are touched in this diff.
5 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:29,116 uses pathlib.Path and contains no hardcoded absolute or repo-escaping paths.
6 PASS dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 only adds word3 = backend_for(self.arch).buffer_rsrc_word3 without deleting any DSL emitter handles.
7 PASS dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/backend_details.md:150 and dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/lowering_pipeline.md:186 update the reference lowering docs alongside the code change (the new ROCKE_BUFFER_RSRC_HIP_TEST env var is a test-fixture wiring variable in run_all.py:148, matching sibling ROCKE_*_TEST entries).
8 PASS dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1032-1034 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2234-2241 use descriptive comments with no Jira IDs, PR numbers, or handles in code, while the bare ticket ID is kept in the PR description.
9 PASS dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:287-291 are tightly scoped to reading buffer_rsrc_word3 from the target backend and making rocke_ll_backend_for safe for HIP lowering to call.
10 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:169-183 and dnn-providers/hip-kernel-provider/ROCKE_ENGINE_test_categories_external.yaml:40,80 cleanly append the new test target without conflicting struct or table edits.
11 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:1 and dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/backend_details.md:150 commit only deliverable implementation, tests, and required lowering docs.
12 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:170-175 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:42-51 cover gfx9 (gfx90a, gfx942, gfx950) as well as gfx11/gfx12 (gfx1151, gfx11-generic, gfx1201, gfx12-generic, gfx1250).
13 PASS dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:229-231,291,2319-2322 replaces the static LL_BACKEND_RESOLVED with per-lowerer L->backend_resolved (lower_llvm_internal.h:402) and updates the surrounding comments in the same commit.
14 PASS dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2234-2249 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2298-2322 only remove the obsolete hardcoded 0x00027000 comments/code and LL_BACKEND_RESOLVED static scratch copy.
15 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:65-115 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:74-156 use every parameter in every newly added or modified function.
16 N/A This item applies to review grading conduct rather than code in the pull request diff.
17 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:1-210 and the PR description contain only functional correctness results and no performance numbers or superiority claims.
§F graded 11/11
# Verdict Evidence
1 PASS branch users/cgmillette/rocke/hip-buffer-descriptor uses the users// prefix
2 PASS title fix(rocke): take the HIP buffer descriptor word from the target matches type(scope): summary and is 63 characters
3 PASS description is at least 30 characters and contains an issue reference
4 PASS test file dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp is in the diff
5 PASS All 13 modified/added paths in the diff (e.g., dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:20 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:43) are source, test, build, or doc files with no keys, certs, or secrets.
6 NEEDS-INFO Local pre-commit execution status is not verifiable from the diff alone (though formatting in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:1 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1035 appears clean).
7 NEEDS-INFO Bot, Copilot, and human review threads are not included in the provided PR context.
8 N/A This PR targets develop directly rather than being stacked on a feature branch.
9 PASS The PR keeps the full logical change together across Python (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243), C++ (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036), docs, and tests without artificial splitting.
10 PASS The added tests in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:126 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:91 are substantive parity, compilation, and GPU execution tests rather than throwaway stubs.
11 NEEDS-INFO The PR title and body are populated in the prompt metadata, but live GitHub API verification (gh pr view --json body,title) is not a path:line in the diff.
§G graded 34/34
# Verdict Evidence
1 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:112 tests byte-identity between Python and C++ HIP output, while lines 91, 98, 129, and 158 separately verify descriptor values, hipcc device IR, and on-GPU execution correctness.
2 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:158 adds on-silicon execution (test_hip_route_copy_is_bounds_checked_on_the_gpu) and the PR body reports silicon verification across gfx1201, gfx1100, and gfx90a.
3 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:171 and dnn-providers/hip-kernel-provider/rocke/platform/tests/run_all.py:147 wire the new test into CTest and the runner; PR body attests byte-identity at llvm20/22/23, ctest, and full platform suite pass.
4 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:91 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:174 assert 0x31014000 for gfx11/gfx12 targets, which failed before the change when 0x00027000 was hardcoded.
5 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:92 enters through the public lower_kernel_to_hip and compile_kernel_via_hipcc APIs, and reverting dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 makes the test fail.
6 PASS Prose claims in docs (dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/backend_details.md:150) and comments (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py:5921) match the implementation and tests (dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:37), and the PR body's '112 of 128 elements untouched' matches 2 * _WAVE = 128 and `2 * store_lanes =…
7 N/A This PR does not add or modify a design document or fp32 numerical oracle.
8 N/A No dtype fallback logic is introduced; dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:172 explicitly skips only when no ISA backend or hipcc is available.
9 N/A This PR does not make LDS or occupancy claims.
10 N/A This is a correctness bug fix with no performance claims or benchmarks.
11 N/A This PR fixes platform HIP lowering of tile.buffer_rsrc rather than adding or optimizing a dispatched library kernel.
12 PASS Both engines are updated in lockstep in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036, and cross-checked in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:112.
13 N/A No new LLVM intrinsic symbols or overload classes are introduced or recommended.
14 N/A No performance bottleneck classification is made in this PR.
15 N/A The PR diff is provided directly in the review prompt from head commit 8569306.
16 PASS dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036 and dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 reuse the existing rocke_ll_backend_for and backend_for symbols and buffer_rsrc_word3 field.
17 PASS The modified runner script dnn-providers/hip-kernel-provider/rocke/platform/tests/run_all.py:147 and test dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:158 are exercised end-to-end by dnn-providers/hip-kernel-provider/rocke/platform/tests/tools/test_run_all.py:119 and on-GPU execution.
18 N/A This PR does not add a statistical guard or back-pressure/jitter probe; it fixes the target buffer descriptor word in HIP lowering and adds deterministic unit/GPU copy tests.
19 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:175 sets distinct bounds (load_lanes = 40, store_lanes = 56 out of _WAVE = 64) and pre-fills dst with sentinel = -7.0 so in-bounds copy, load-OOB zeroing, and store-OOB preservation are all distinguishable.
20 N/A This PR does not add or modify a text detector or linter for banned performance claims.
21 N/A This PR does not introduce performance measurement or profiling code.
22 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/run_all.py:147 registers ROCKE_BUFFER_RSRC_HIP_TEST (rocke_buffer_rsrc_hip) so the runner builds and passes the native binary to pytest, and dnn-providers/hip-kernel-provider/rocke/platform/tests/tools/test_run_all.py:119 verifies the environment variable is wired.
23 N/A This PR makes no performance speedup or dispatcher routing claims.
24 N/A This PR does not re-bless any golden fixtures (LLVM IR emission is unchanged).
25 N/A This PR is a correctness fix for the HIP buffer descriptor flags word and not a codegen micro-optimization.
26 N/A This PR does not contain performance benchmarks.
27 N/A This PR does not present a performance A/B experiment.
28 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:105 (test_default_arch_keeps_the_gfx950_word) and the PR body explicitly state that gfx9 targets and the default arch are 0-delta (0x00027000 before and after).
29 N/A This PR does not swap any math primitives such as exp2 to exp2_fast.
30 N/A This PR makes no register spill or scratch memory claims.
31 PASS dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1045 show the defect (0x00027000 hardcoded on all targets) existed in develop and was verified on gfx1201/gfx1100/gfx90a before and after the fix.
32 N/A This PR does not report measured performance or resource-usage numbers.
33 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:42 (_WORD_BY_ARCH) and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:170 exhaustively evaluate every supported target arch (gfx90a, gfx942, gfx950, gfx1151, gfx11-generic, gfx1201, gfx12-generic, gfx1250) across both Python and C++ HIP and LLVM lowerers.
34 N/A This PR does not add or modify an opt-in kernel output flag.
§H graded 5/5
# Verdict Evidence
1 N/A This PR modifies platform HIP/LLVM lowering for tile.buffer_rsrc descriptor flags (DWORD3) and does not add or modify any kernel spec, host launch validation, or spec-derived buffer bounds.
2 N/A This PR does not touch kernel indirection tables, page/block tables, or indexed memory lookups.
3 N/A This PR does not introduce or modify any kernel hot-loop div or mod operations or tile/page size specs.
4 N/A This PR does not add or modify any kernel device-side masks or bounds derived from runtime sequence length inputs (kv_lens, cu_seqlens, etc.).
5 N/A This PR is a platform lowering fix and does not add or modify kernel launch inputs feeding device-side bounds or masks.
§I graded 6/6
# Verdict Evidence
1 N/A This PR fixes HIP buffer descriptor lowering and does not modify kernel routing, dispatch selection, or search spaces.
2 PASS In dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036-1040, an unknown target architecture in _op_tile_buffer_rsrc is loudly rejected via rocke_h_fail(lw, ROCKE_ERR_KEY, "no ISA backend registered for '%s'", lw->arch.gfx) instead of silently emitting an invalid gfx9 descriptor word.
3 N/A This PR does not add or modify hypothesis/property-based input generators or shape coverage suites.
4 N/A This PR does not include routing benchmark tables or kernel selection reporting.
5 N/A This PR does not introduce or update any baked tuning or search snapshots.
6 PASS The generic tile.buffer_rsrc op selects the target-specific descriptor word3 at lowering time via backend_for(self.arch).buffer_rsrc_word3 in dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 and rocke_ll_backend_for(lw->arch.gfx, &backend_status)->buffer_rsrc_word3 in dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036-1049.
§J graded 7/7
# Verdict Evidence
1 N/A The Markdown diff only updates two bullet lines describing __builtin_amdgcn_make_buffer_rsrc arguments and introduces no mathematical equations or tensor shapes.
2 PASS The symbol word3 introduced in dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/backend_details.md:150 and dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/lowering_pipeline.md:186 is defined in the same sentence as the target ISA backend's buffer_rsrc_word3.
3 N/A The PR does not add or modify any algorithm derivations or design docs (ALGORITHM.md / DESIGN.md).
4 PASS In dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/backend_details.md:150 and dnn-providers/hip-kernel-provider/rocke/platform/dsl_docs/ir_lowering/lowering_pipeline.md:186 (and docstring at dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py:5921), buffer_rsrc_word3 is accurately described as the existing word the LLVM route already emits, now shared by the HIP lowering.
5 N/A No scaling or complexity claims are made in the diff.
6 N/A No model config tables or tuned band-edge constants are introduced in the documentation.
7 N/A No new documentation files are added in this PR.
§K graded 13/13
# Verdict Evidence
1 N/A This PR only modifies HIP/LLVM buffer resource lowering and platform tests, with no tuned candidates or KnobSpace changes.
2 N/A No kernel spec fields or KnobAxis declarations are added or modified.
3 N/A No tuning identity payload, canonical knob rules, declared defaults, or builder behavior behind a knob are touched.
4 N/A No KnobSpace.defaults() or KnobSpace.recorded() implementations are touched.
5 N/A No knob canonicalization, inertness, or refusal rules are modified.
6 N/A No KnobSpace.validate or tuning spec construction code is touched.
7 N/A No KnobSpace.base(req) or tuning problem normalization is touched.
8 N/A No knob normalization or request boundary validation is touched.
9 N/A No benchmark recording, configuration replay, or PinRefused handling is touched.
10 N/A No tuning_id resolution or KnobSpace.find logic is touched.
11 N/A No sweep entry points or sweep harnesses are touched.
12 N/A No production stacks, accept_default rules, or base tuned specs are touched.
13 N/A No benchmark harnesses or auto-selection claims are touched.
§L graded 10/10
# Verdict Evidence
1 N/A This PR does not add or change any gfx target or ISA backend table row; it only makes rocke_ll_backend_for return the existing immutable row and makes HIP lowering read buffer_rsrc_word3 from the existing backend.
2 N/A This PR does not add or modify any MMA/WMMA op_id or layout map.
3 N/A This PR does not touch MMA selection, MfmaAtom catalogs, or MMA lowering.
4 N/A This PR does not modify target wave size handling or lane math in production code.
5 N/A This PR does not add or modify arch-specific recipes such as transpose-LDS reads, K-packed atoms, or AsyncTileLoader.
6 N/A This PR does not touch pipeline vocabulary, policy modules, ArchTarget, or dispatcher imports.
7 N/A This PR does not add or change any logical dtype, bit packing format, or instruction fragment contract.
8 N/A This PR does not touch scaled-MMA contracts, operation IDs, or MfmaAtom.emit_scaled().
9 N/A This PR does not add a new gfx architecture target.
10 N/A This PR does not modify intrinsic arch-domain artifacts or flavor membership logic.
§M graded 10/10
# Verdict Evidence
1 PASS The new test kernel in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:56 (_copy_kernel) and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:45 (build) emits straight-line buffer operations without using SSA values in host control flow.
2 N/A The diff does not touch legalize_for_target, tile.mma, or scf.for legalization.
3 N/A The diff does not modify print_ir or ck.dsl.ir/v1 serialization in core/ir_serialize.py or cpp/core/ir/serialize.cpp.
4 N/A The diff does not touch portable IR schemas, importers, recorders, or magic division helpers.
5 N/A The diff does not modify or use AsyncTileLoader or async DRAM-to-LDS operations.
6 PASS The test kernel in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:68 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:57 uses bounds-checked buffer descriptors (buffer_rsrc with buffer_load_vN_f16 / buffer_store_vN_f16) for OOB safety and emits no barriers or cross-lane shuffles.
7 PASS Both Python (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243) and C++ (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036) HIP lowerings now take buffer_rsrc_word3 from the target ISA backend (0x00027000 on CDNA, 0x31014000 on RDNA/gfx1250) instead of hardcoding 0x00027000.
8 N/A The diff does not add new LLVM intrinsic declarations or modify engine selection precedence.
9 PASS Only the internal header dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h:228 (rocke_ll_backend_for comment and rocke_lower struct field at line 402) is updated; no public C entry points are changed and no headers are renamed.
10 N/A The diff does not modify !dbg metadata emission or DWARF debug info handling.
§N graded 18/18
# Verdict Evidence
1 N/A The PR does not add or modify any files under platform/python/rocke/helpers/.
2 N/A The PR does not add or modify kernel builders, atoms, TensorView, transforms, loaders, epilogues, or pipelines.
3 N/A The PR does not use or modify SoftwarePipeline.
4 N/A The PR does not involve kernel fusion or EpilogueOp.
5 N/A The PR does not touch rocke.helpers.tiling or add any tiling API.
6 N/A The PR does not add or modify tiling helpers or LayoutStyle.
7 N/A The PR does not involve MMA layout overrides or tiling drivers.
8 N/A The PR does not touch tiling memory emission, LDS staging, or global tile loads/stores.
9 N/A The PR does not touch N-D tensor axis roles, label flow, or tiling transforms.
10 N/A The PR does not change tiling emit, LDS staging, or MMA issue.
11 N/A The PR does not touch tiling traits, TileMma, lds_swizzle, or scheduling helpers.
12 N/A The PR makes no LDS bank-conflict or memory coalescing claims.
13 N/A The PR does not touch WarpDistributionEncoding, make_tile_desc, or RegisterMapper.
14 N/A The PR does not modify tiling verbs or LDS pipeline synchronization.
15 N/A The PR does not touch tiling allocation or swizzled/padded descriptor bounds checks.
16 N/A The PR does not touch optimize_layout or evaluate_transform.
17 N/A The PR does not touch the tiling recorder, roundtrip analysis, or visualization.
18 N/A The PR does not touch interleaved layouts or tiling demo kernels.
§O graded 16/16
# Verdict Evidence
1 PASS Hand-built test IR uses standard PtrType(..., 'global') and rocke_ptr_type(..., 'global') with no verifier or address-space changes (dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:62, dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:47).
2 N/A The diff does not touch scf.for, scf.if, region scoping, or block terminator emission rules.
3 PASS Emission order and SSA naming remain deterministic; rocke_ll_backend_for now returns the read-only per-arch table row while rocke_lower copies it into L->backend_resolved (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:287, dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2319).
4 PASS Test code uses module singletons F16 and I32 in Python and rocke_f16() / rocke_i32() in C++ (dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:62, dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:47).
5 N/A The diff does not modify optimize_kernel, canonicalize_region, core/passes.py, cpp/core/passes.cpp, or op purity tables.
6 N/A The diff does not modify constant folding in Python or C++.
7 N/A The diff does not touch kernel attributes (amdgpu-flat-work-group-size, waves_per_eu, agpr_alloc, or cluster_dims).
8 N/A The diff does not touch CodegenPolicy, SchedulerStrategy, or apply_codegen_policy.
9 N/A The diff does not touch waitcnt encoding, barriers, or optimization_barrier.
10 N/A The diff does not change op attribute validation or pointer parameter lowering (_param_llvm_type).
11 PASS Target buffer descriptor word3 is queried from the ISA backend via backend_for(self.arch).buffer_rsrc_word3 in Python and rocke_ll_backend_for(lw->arch.gfx, &backend_status) in C++ (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243, dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036).
12 PASS No new arch defaults or hard-coded byte/dword caps are added in core; rocke_ll_backend_for still rejects a NULL arch (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2304).
13 PASS Only the internal struct struct rocke_lower in lower_llvm_internal.h is modified (backend_resolved field added), leaving public C ABI structs and enums untouched (dnn-providers/hip-kernel-provider/rocke/platform/cpp/include/rocke/lower_llvm_internal.h:402).
14 PASS C++ changes follow rocke_* naming and C style in dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1035 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_llvm/core.cpp:2301, with no new -Wno-* flags in dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:171.
15 PASS No fallback paths are modified, and lower_hip.py and lower_hip_mem.cpp are updated to match lower_llvm's per-target buffer_rsrc_word3 (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243, dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036).
16 PASS known_gaps.md debt is not widened and gfx1250's inherited backend behavior is explicitly noted without overclaiming (dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:172, dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:41).
§P graded 14/14
# Verdict Evidence
1 N/A The diff does not touch any kernel spec or instance builder signatures.
2 N/A The diff does not touch kernel argument signatures, ABI definitions, or manifest argument packing.
3 N/A The diff does not touch GEMM, convolution, or block-scaled GEMM launch geometry or specs.
4 N/A The diff does not touch LDS capacity gates or validators.
5 N/A The diff does not touch WMMA/RDNA GEMM, GemmMultiD, or GemmMultiAbd specs or validators.
6 N/A The diff does not touch Stream-K, MoE, or split-K workspace sizing or initialization.
7 N/A The diff does not touch Stream-K, persistent GEMM, GroupedGemmLauncher, or SUPPORT_MATRIX.md.
8 N/A The diff does not touch quantization helpers or MX GEMM scale decoding.
9 N/A The diff does not touch numeric emulation, type conversion, or decode paths.
10 N/A The diff does not add or modify manifest kinds, emitters, or runners.
11 N/A The diff does not touch runtime/ launch utilities, KernelLauncher, PipelineLauncher, or WorkspacePool.
12 N/A The diff does not modify COMGR or HIP library loading or fix a hipError launch failure.
13 N/A The diff does not refactor host-side grid, tile-count, workspace math, or kernel-cache admission guards.
14 N/A The diff does not touch examples/, instances/, or helpers/.
§Q graded 7/7
# Verdict Evidence
1 FAIL dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:1 adds a new pytest module, while dnn-providers/hip-kernel-provider/rocke/platform/tests/CMakeLists.txt:169 and dnn-providers/hip-kernel-provider/ROCKE_ENGINE_test_categories_external.yaml:40 register only the C++ executable rocke_buffer_rsrc_hip; platform/CMakeLists.txt is not updated to register test_buffer_rsrc_hip.py in rocke_pytest.
2 N/A This PR does not add a new kernel family or variant; it fixes the target buffer descriptor word in tile.buffer_rsrc HIP lowering.
3 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:90-158 separates string/IR parity checks, hipcc device-IR compilation checks, and actual GPU execution bounds-checking.
4 N/A No golden files are modified and the valid spec space of builders or lowerers is not widened.
5 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:113-125 and 169-173 resolve the GPU target from current_device_arch() and skip cleanly when hipcc, the native binary, or a supported GPU is absent, and subprocess text=True normalizes CRLF on Windows.
6 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/run_all.py:147-151 only adds the ROCKE_BUFFER_RSRC_HIP_TEST fixture mapping without altering per-configuration timeout or hashing semantics.
7 N/A This PR does not change any rejection contracts or exception types across backends.
§R graded 11/11
# Verdict Evidence
1 N/A This PR fixes functional buffer descriptor word3 encoding in HIP lowering on gfx11/gfx12 and does not introduce a new kernel body, tiling strategy, pipeline variant, or perf lever.
2 N/A No performance levers are modified by this correctness and thread-safety fix.
3 N/A The diff does not touch any optimization levers covered by the utilities skills.
4 N/A This PR makes no performance or codegen policy optimization claims and does not change emitted LLVM IR or gfx9 HIP output.
5 N/A No production kernel global memory access patterns or tiling layouts are added or modified.
6 N/A No LDS or MMA layouts are added or modified in this diff.
7 N/A The PR makes no bottleneck claims and is purely a functional correctness fix for gfx11/gfx12 HIP buffer descriptors.
8 N/A No waits, barriers, or scheduling hints are changed in this PR.
9 N/A This PR is not a performance regression fix and does not trade performance across shapes.
10 N/A No benchmarking, occupancy conclusions, or optimization codegen changes are made in this PR.
11 N/A No new optimization lesson, probe, or analysis tool is introduced.
§S graded 8/8
# Verdict Evidence
1 N/A The diff fixes tile.buffer_rsrc lowering behavior in existing HIP lowerers and adds no new public Python symbols or C entry points.
2 PASS dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243 and dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036 update tile.buffer_rsrc lowering for non-gfx9 targets while keeping default/gfx9 behavior (dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:105), and add tests in…
3 N/A No public API is removed or renamed in this diff.
4 N/A No __all__ list or public module export list is modified in this diff.
5 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:91 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:126 assert exact emitted HIP descriptor flags, parity with LLVM IR flags, byte-identical HIP output between Python and C++, and GPU numeric behavior (dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:210).
6 PASS dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:91 and dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.cpp:169 test HIP/LLVM lowering and cross-engine parity on CPU without torch or a GPU, while dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:158 separately tests GPU launch and bounds-checked numerics.
7 PASS Both Python (dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_hip.py:2243) and C++ (dnn-providers/hip-kernel-provider/rocke/platform/cpp/core/lower_hip/lower_hip_mem.cpp:1036) lowerers are updated and compared for byte-identity across all targets in dnn-providers/hip-kernel-provider/rocke/platform/tests/core/test_buffer_rsrc_hip.py:112 and wired into…
8 PASS The PR body (Test Result) reports running byte-identity across llvm20/22/23, ctest, strict both-engines tests, the full platform suite, and GPU tests naming gfx1201, gfx1100, and gfx90a, while dnn-providers/hip-kernel-provider/rocke/platform/tests/run_all.py:147 and dnn-providers/hip-kernel-provider/rocke/platform/tests/tools/test_run_all.py:119 keep run_all.py updated.

@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/hip-buffer-descriptor branch from 8569306 to f7e51ef Compare October 23, 2026 02:09
@cgmillette
cgmillette requested review from CongMa13 and removed request for CongMa13 October 23, 2026 02:10
cgmillette and others added 3 commits October 23, 2026 00:20
The HIP lowering of tile.buffer_rsrc passed the gfx9 descriptor word
0x00027000 as the flags argument of __builtin_amdgcn_make_buffer_rsrc for
every target. That word's format and out-of-bounds encoding is
ISA-specific: on gfx11 / gfx12 it makes every raw buffer load read 0 and
every store drop. The LLVM lowering already takes the word from the
target's ISA backend.

Both HIP lowerings now read the same per-target source as the LLVM
lowering of their engine: the Python one asks backend_for(arch) for
buffer_rsrc_word3, the C++ one asks rocke_ll_backend_for. gfx9 targets
and the default arch keep 0x00027000; gfx11 and gfx12 targets get
0x31014000, as does gfx1250, whose backend currently inherits the gfx12
word. The LLVM IR does not change.

Tests pin the HIP word per target, check it equals the LLVM word for the
same target in both engines, byte-compare the Python and C++ HIP text
through a new native test binary, compile the gfx11 / gfx12 HIP source
with hipcc and check the word in the device IR, and run a bounds-checked
buffer copy through the HIP route on a gfx9 GPU. The native test binary is
resolved for the source-tree run and set in the installed rocke_pytest
entry, so its parity cases run in both. The HIP backend docs
describe the per-target word.

Co-Authored-By: Claude <noreply@anthropic.com>
rocke_ll_backend_for copied the target's backend into one process-wide
scratch struct and returned it, and the LLVM lowering then rewrote that
struct's datalayout for its flavor. Every caller shared the copy, so a
lowering for one target or flavor could change it under another running at
the same time, now including the HIP lowering's buffer descriptor lookup.

The lookup returns the target's immutable table row. Each LLVM lowering
keeps its own copy in its lowering context and fills in the triple and the
flavor's datalayout there. The emitted IR is unchanged.

Co-Authored-By: Claude <noreply@anthropic.com>
The bounds-checked buffer copy through the HIP route compiled for device 0
and ran only on gfx9 GPUs, while the runtime launches on the current
device. It now compiles for the current device, lowering for its family's
generic target when the arch has no backend of its own (gfx11-generic on a
gfx1100), so it also covers the gfx11 and gfx12 descriptor word.

Co-Authored-By: Claude <noreply@anthropic.com>
@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/hip-buffer-descriptor branch from f7e51ef to 20f7942 Compare October 23, 2026 04:41
@cgmillette
cgmillette merged commit 07394ac into develop Oct 23, 2026
96 checks passed
@cgmillette
cgmillette deleted the users/cgmillette/rocke/hip-buffer-descriptor branch October 23, 2026 15:17
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants