Skip to content

fix(rocke): refuse buffer_rsrc operands its lowerings cannot take - #14842

Merged
cgmillette merged 8 commits into
developfrom
users/cgmillette/rocke/buffer-rsrc-refusals
Nov 5, 2026
Merged

cgmillette merged 8 commits into
developfrom
users/cgmillette/rocke/buffer-rsrc-refusals

Conversation

@cgmillette

@cgmillette cgmillette commented Nov 1, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

buffer_rsrc(ptr, num_bytes) lowers to @llvm.amdgcn.make.buffer.rsrc. Two operand forms reached
the LLVM lowering of both engines unchecked, and the module they produced was then rejected by llc
with an error that did not point at the op:

  • a pointer parameter declared addr_space="constant": it is ptr addrspace(4) in the kernel
    header, but the call names its pointer ptr addrspace(1) (LLVM 20 llc: '%X' defined with type 'ptr addrspace(4)' but expected 'ptr addrspace(1)');
  • an i64 num_bytes at llvm20: that flavor's intrinsic takes an i32 num_records, and the lowering
    printed any operand as i32 (LLVM 20 llc: '%n' defined with type 'i64' but expected 'i32').

The HIP lowering has the same kind of bug for the byte count: it calls
__builtin_amdgcn_make_buffer_rsrc with num_bytes cast to (int), so an i64 byte count was
silently truncated to 32 bits (hipcc emits trunc i64 %n to i32 before the call), and the
descriptor bounded a smaller range than asked. ROCm 7.1's builtin takes an int num_records
and later clang a 64-bit one (int64_t in LLVM 22 and 23, uint64_t after); the truncation comes from
the cast the lowering emits.

The constant-space mismatch was not specific to buffer_rsrc. The global loads, stores, atomics,
flag ops, cooperative store and global/LDS copies also named a constant-space parameter
ptr addrspace(1), and LLVM rejected the module.

This PR refuses these forms up front, in both engines, with a message that says what is wrong, and
lets the global loads read through a constant-space parameter.

Technical Details

  • Both checks need the lowering's context (the parameter's header type, the LLVM flavor), so they
    sit in the LLVM lowering of each engine, at the same point: Python _op_tile_buffer_rsrc
    (core/lower_llvm.py), C++ op_tile_buffer_rsrc (cpp/core/lower_llvm/mem.cpp). The
    global-pointer ops (global_ptr_add, global_ptr_offset, global_ptr_to_i64) already refuse a
    constant-space parameter there, and buffer_rsrc now uses the same helper
    (_require_global_header_ptr / rocke_ll_require_global_header_ptr) and so the same text:

    buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) (its parameter's addr_space)
    
  • At llvm20 a num_bytes that is not i32 is refused:

    buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20 make.buffer.rsrc takes an
    i32 num_records, so pass the byte count as i32
    

    llvm22 and later still take an i32 (zero-extended) or an i64, unchanged.

  • The HIP lowering of both engines (Python _op_tile_buffer_rsrc in core/lower_hip.py, C++
    _op_tile_buffer_rsrc in cpp/core/lower_hip/lower_hip_mem.cpp) refuses a num_bytes that is not
    i32, at every LLVM version, because it passes num_records cast to (int):

    buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; the HIP lowering casts
    num_records to (int), so a wider byte count would be truncated; pass it as i32
    

    An i64 byte count still lowers through LLVM at llvm22 and later.

  • The HIP lowering keeps a constant-space pointer: its kernel signature declares every pointer
    parameter the same way (float* __restrict__ X) whatever its addr_space, and the builtin takes a
    void*. A hipcc probe of that source compiles, and the descriptor is built from the kernel
    argument's address.

  • Other global-memory ops on a constant-space parameter, in the LLVM lowering of both engines:

    • global_load_f16, global_load and global_load_vN name the pointer's own type (a table of
      the accepted spaces, _GLOBAL_LOAD_PTR_TYPES / ROCKE_LL_GLOBAL_LOAD_PTR_TYPES, resolved by
      the existing pointer-space resolver). They read a constant-space parameter as
      load addrspace(4) and refuse a flat or LDS pointer, which they also used to name
      ptr addrspace(1). The resolver's message now says "op" rather than "intrinsic".
    • The stores, atomics, flag ops, cooperative_global_store, global_load_lds and
      global_load_async_to_lds refuse a constant-space parameter with the global_ptr_* text.
    • One check per engine: every op that names its pointer in one fixed space (the ops above,
      global_prefetch, flat_prefetch, global_store_async_from_lds, global_load_tr16_b128
      and both cluster loads) calls _require_header_ptr / rocke_ll_require_header_ptr, whose
      global-pointer form is _require_global_header_ptr / rocke_ll_require_global_header_ptr.
      The C++ global_store_async_from_lds and global_load_tr16_b128 lowerings used to check only
      the IR type; they now check the header type, as the Python engine already did.
    • Every one of these refusals, and the pointer-space resolver's, is a TypeError in Python and
      ROCKE_ERR_TYPE in C++. The text names the op's builder (store_f16, global_store,
      global_load_f16, global_load) and the builder's name for the operand (src_ptr for the
      copies into LDS), and says "(its parameter's addr_space)" only when the parameter's override
      moved its header type. The buffer_rsrc byte-count refusals also name the builder.
  • The portable IR route lowers through the C++ engine, so it refuses the LLVM forms too.

  • No emission changes for a kernel whose IR LLVM accepted before: byte identity is unchanged, no
    golden is re-blessed, and the HIP source of every parity kernel is unchanged.

  • dsl_docs/reference/op_vocabulary.md states the rules next to the buffer ops.

Test Plan

  • tests/core/test_buffer_rsrc_operands.py:
    • both LLVM lowerings refuse a constant-space pointer at llvm20, llvm22, llvm23 and llvm24, and an
      i64 num_bytes at llvm20, with the exact text;
    • controls: an i32 num_bytes at every flavor and an i64 one at llvm22 and later lower to
      identical IR in both engines;
    • both HIP lowerings refuse an i64 num_bytes with the exact text, and lower an i32 num_bytes
      over a global or a constant-space pointer to the same source.
  • tests/core/test_constant_space_global_ops.py: each LLVM lowering reads through a
    constant-space parameter with each global load at every flavor, the result compiles with the
    host clang, and a flat or LDS pointer is refused; every other op listed above refuses a
    constant-space parameter with the exact text and type, an LDS parameter is refused without the
    addr_space clause, and flat_prefetch refuses a parameter moved to constant space.
  • Both test files run each case per engine (rocke/tests/engines.py gains per-engine helpers that
    also check the C++ status), so the Python cases run in lanes that do not build the
    rocke_engine binding; a C++ case there skips before either engine runs. Develop's
    available_engines() returns the same ENGINES list, or the Python engine alone when the
    binding is not built.
  • tests/core/future_intrinsic_lowering.cpp (ctest rocke_future_intrinsic_lowering): the C++
    LLVM and HIP lowerings refuse the buffer_rsrc forms, with their status codes, and still lower
    the accepted ones; a constant_space_global_ops case checks a sample of the constant-space
    loads and refusals, so the C++ side also runs in those lanes.
  • HIP source of every parity-family config (95 families) for gfx942, gfx950 and gfx1201, before and
    after the change.
  • Build, ctest, byte identity at llvm20 / llvm22 / llvm23 / llvm24, golden checks, portable IR
    against a freshly built shared engine, the touched tests with both engines in strict mode, and
    the platform suite compared by failure name against develop.

Test Result

  • Before the change (develop), both buffer_rsrc LLVM forms lowered without error and LLVM 20
    llc rejected the result with the errors above; the HIP lowering of an i64 num_bytes compiled,
    and hipcc truncated the byte count to 32 bits. A global load through a constant-space parameter
    produced a module the host clang rejects.
  • HIP source (first revision): the Python HIP lowering's 3772 (config, arch) texts across the 95
    parity families are identical to develop's; the later commits change the HIP lowering only in the
    refusal's text.
  • Byte identity GREEN at llvm20, llvm22, llvm23 and llvm24.
  • Golden checks with both engines at all four flavors: no blessed config changed. The report also
    lists the Python-only barrier_fence family as absent from the both-engine run, and configs not
    yet in the golden file. No golden file changed.
  • ctest: 45 run, all pass except rocke_arch_domain_reproducibility, which skips; this includes
    rocke_future_intrinsic_lowering.
  • The two new test files: 107 cases with both engines in strict mode; without the binding, 55
    Python cases pass and the 52 C++ cases skip.
  • The 35 test files that use rocke/tests/engines.py, with both engines in strict mode: 5266
    passed, 134 skipped.
  • Platform suite (tests and python/rocke/tests) with both engines in strict mode and no GPU
    visible: 11255 passed, 14 failed. All 14 also fail on develop:
    • test_the_cpp_fallback_does_not_warn_twice stubs the engine and fails only under
      ROCKE_CPP_STRICT=1;
    • four test_installed_ctest_entries_get_scoped_environment_contracts cases fail the same way
      on develop;
    • the nine interleaved-GEMM demo tests need a GPU.
  • Library golden tests: 93 passed.
  • pre-commit passes.

Submission Checklist

ISSUE ID : AICK-2900

🤖 Generated with Claude Code

@therock-pr-bot

therock-pr-bot Bot commented Nov 1, 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 Nov 1, 2026 •

Copy link
Copy Markdown

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

@cgmillette

Copy link
Copy Markdown
Contributor Author

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

PR #14842 — refuse buffer_rsrc operands its lowerings cannot take (cgmillette, AICK-2900) — 🟡 First-round: refusals sound; builtin claim inaccurate

Prior item Status
Reviewers 🟢 None yet
• no reviews or threads; PR bot checks pass (read)
Author: refused forms broke before; kept forms unchanged 🟢 Verified
• develop: LLVM parser rejects both; (int) cast truncates i64 (ran)
• kept forms byte-identical to develop, LLVM and HIP (ran)
• 2221 parity configs: all i32, global; none reaches a refusal (ran)
Author: both engines refuse with the same text 🟢 Verified
• 15 ✓ with the native engine, strict mode (ran)
• C++ ctest 90 cases ✓, new case included (ran)
• CI collects the file but skips all 15: no engine (read)
Bots / CI 🟡 Red: known flake
• fails: gfx94X hipkernelprovider shard 1/2 — #14733 golden sha256 timeout (read)
• hit 3 of 4 shard runs; 49 ✓, Windows gfx110X shards incl.
• required Summary red: needs a green re-run or #14820
Our checks 🟢 Clean
• merge vs 61b3ac0e: clean, 0 behind (ran)
• build + ctest 34 ✓ / 1 skip; byte identity GREEN ×4 flavors (ran)
• golden PASS ×4; differential 0 changed; rules 1–3 clean (ran)
# Sev file:line Finding
N1 🟡 core/lower_hip.py:2305, lower_hip_mem.cpp:1089, op_vocabulary.md:161 HIP refusal blames the builtin; LLVM 24's takes a 64-bit count
• impact: text says int at every version; clang 24: unsigned long (ran)
• fix: blame the (int) cast instead; keep the refusal

Utility: 🟢 Positive — same-text refusals replace late llc errors and silent truncation; no emission change.


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

First-round review: first pass by a local qwen3.6 model, verified by the orchestrator (Claude Opus 5.5). Checks: merge with develop, build + ctest, byte identity, goldens, both-engine tests, develop before/after, parity operand sweep, CI logs at head. Flags whether a closer look is needed; not an approval or rejection.

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

  • HIP refusal rationale overstates the builtin's int signature — cgmillette, #14842 (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.

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

Remaining global-memory handlers in both engines hard-code address space 1 for constant-space parameters declared in address space 4, producing invalid LLVM IR.

1 open issue
What changed in this PR

Adds early buffer_rsrc operand validation to rocKE’s Python and C++ lowering engines, preventing invalid LLVM IR and silently truncated HIP descriptor bounds.

Changes:

  • Reject unsupported pointer address spaces and byte-count widths.
  • Add regression tests for rejected operands and supported forms across both engines.
  • Document the LLVM and HIP operand constraints.
File Description
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​test_buffer_rsrc_operands.py Adds rejection, acceptance, and engine-parity tests.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​tests/​core/​future_intrinsic_lowering.cpp Adds C++ LLVM and HIP regression cases.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​core/​lower_llvm.py Validates pointer address space and LLVM 20 byte-count width.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​python/​rocke/​core/​lower_hip.py Rejects non-i32 byte counts before HIP emission.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​dsl_docs/​reference/​op_vocabulary.md Documents backend-specific operand rules.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​core/​lower_llvm/​mem.cpp Mirrors the Python LLVM operand checks.
dnn-providers/​hip-kernel-provider/​rocke/​platform/​cpp/​core/​lower_hip/​lower_hip_mem.cpp Mirrors the HIP byte-count restriction.

Review effort: Max

@cgmillette

Copy link
Copy Markdown
Contributor Author

Feedback addressed at b54101d:

  • N1 (HIP refusal blamed the builtin): fixed in 8be45df. The HIP refusal, its comments in both engines and the op reference now blame the (int) cast the lowering emits. ROCm 7's builtin takes an int num_records, and newer clang takes a uint64_t. The new message is tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; the HIP lowering casts num_records to (int), so a wider byte count would be truncated; pass it as i32. The refusal itself is unchanged. The Python test and the C++ ctest assert the new text.
  • Copilot thread (other global-memory ops on a constant-space parameter): fixed in 8c6f652 and b54101d. Details are in the thread reply.
  • PR description: updated. It no longer says the builtin takes an int at every LLVM version, and it now covers the two new commits.

@cgmillette

Copy link
Copy Markdown
Contributor Author

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

PR #14842 — refuse buffer_rsrc operands its lowerings cannot take (cgmillette, AICK-2900) — 🟠 Follow-up: copied check; CI-skipped tests

Prior item Status
Our 11-02 N1: HIP refusal blamed the builtin 🟢 FIXED 8be45df8
• message, op reference blame the (int) cast; both tests assert it (ran)
• left: ROCm 7.2, LLVM 22–23 builtin takes int64_t; comments, body say int/uint64_t (read)
CongMa13 (11-02) APPROVED 🟢 APPROVED 792fafa8
• older head; still counts on develop
• predates 8c6f6523: load typing, 14 new refusals
Copilot (11-02) Rule 9: other global ops emit addrspace(1) 🟢 FIXED 8c6f6523
• 3 loads name addrspace(4) and compile; 14 ops refuse; engines agree (ran)
• left: select over it still names addrspace(1); LLVM rejects (ran)
Bots / CI 🟡 Running
• running: Windows gfx110X tests, hipDNN integration; 38 ✓ · 5 running
• gfx94X hipkernelprovider ×2 ✓, ctest incl.; last round's flake gone (read)
• merge: waits on Multi-Arch Summary (read)
Our checks 🟢 Clean
• merge vs a007e937: clean, 72 behind (ran)
• build + ctest 34 ✓ / 1 skip; byte identity GREEN ×4 (ran)
• goldens PASS, 0 changed; no parity config newly refused (ran)
Code quality 🟠 Should-fix
• N2: one space rule, 6 copies per engine; statuses split (ran)
• N3: the new cases run in no CI lane (ran)
• else reuses resolver and helper; flavor tuple restated twice (read)
# Sev file:line Finding
N2 🟠 lower_llvm/mem.cpp:1794, :1842, :1897, :1979; control.cpp:1386; lower_llvm.py:1860, :5069–:5701 Space check copied, not called; copies diverge
• impact: TYPE reached 3 of 6 C++ copies; 3 still VALUE (ran)
• impact: LDS pointer: loads raise ValueError, stores TypeError (ran)
• fix: one helper per engine at every site; one status, one type
N3 🟠 test_constant_space_global_ops.py:38, test_buffer_rsrc_operands.py:34 No CI lane runs the 50 new cases
• impact: CI has no rocke_engine: 50 of 50 skip, both gfx94X shards (read)
• fix: assert the Python half first; add load, store, atomic ctest cases
N4 🟡 lower_llvm.py:3410, :3437, :1913 Refusals name ops the caller didn't call
• impact: store_f16 says "global_store ptr …"; LDS param told "its parameter's addr_space" (ran)
• fix: builder names; that clause only when addr_space is set

Utility: 🟢 Positive — same-text refusals replace late LLVM errors; constant-space loads now compile.


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: merge with develop, build + ctest, byte identity, goldens, both-engine op probe, LLVM verifier, parity census, CI logs at head, code quality. Flags whether a closer look is needed; not an approval or rejection.

@cgmillette

Copy link
Copy Markdown
Contributor Author

Re our review (#14842 (comment)):

All items are addressed in 7425d42.

N2 (one check, copies diverge): fixed in both engines.

  • Every LLVM-lowering op that names its pointer in one fixed space now calls one helper per engine: _require_header_ptr in Python, rocke_ll_require_header_ptr in C++. A global-pointer wrapper covers the 21 sites that need ptr addrspace(1).
  • No inline copy remains. That includes global_prefetch, global_store_async_from_lds, global_load_tr16_b128, both cluster loads, and flat_prefetch (the same check against ptr).
  • Every one of these refusals is now a TypeError in Python and ROCKE_ERR_TYPE in C++. The multi-space resolver the global loads, av_*_b128 and s_prefetch_* use raises the same type, so an LDS pointer gets a TypeError from a load and from a store alike.
  • The per-engine test helpers now check the C++ status against the expected exception type. That makes the type claim a tested fact wherever the binding is built.

N3 (cases no CI lane runs): fixed.

  • rocke/tests/engines.py gains per-engine helpers: lower_with_engine, assert_engine_rejects, and their HIP forms. The existing both-engine helpers share their internals. They still compare only the C++ message, because one older refusal (the scale-encoding buffer access) does not yet share its status with Python.
  • test_constant_space_global_ops.py and test_buffer_rsrc_operands.py are parametrized over the two engines. Without the binding, all 55 Python cases run and only the 52 C++ cases skip.
  • A new future_intrinsic_lowering CTest case, constant_space_global_ops, covers part of the C++ side in those lanes:
    • the three loads through a constant-space parameter;
    • the store_f16, global_store and global_atomic_add refusals;
    • an LDS-pointer store and load.
  • buffer_rsrc_refusals already covered the buffer_rsrc forms.

N4 (names and clause): fixed.

  • Refusals name the op's builder: store_f16, global_store, global_load_f16, global_load.
  • The "(its parameter's addr_space)" clause appears only when the parameter's override moved its header type. An LDS parameter now reads global_store ptr must be a global pointer, got ptr addrspace(3). Both forms are tested.

Open notes in the table:

  • Builtin signature: corrected. LLVM 21 declares int num_records, LLVM 22 int64_t, and current clang uint64_t. The comments now say "a 64-bit num_records (int64_t in LLVM 22, uint64_t after)".
  • select over a constant-space parameter: not fixed here. It still names ptr addrspace(1). This is older than this PR and is about an op that takes a pointer value rather than addressing memory through it, so it needs its own fix.
  • Flavor tuple: fixed. Both test files now take LLVM_FLAVORS from rocke.core.lower_llvm.

Gates on 7425d42:

  • build and ctest 35 / 35;
  • byte identity GREEN at llvm20, llvm22, llvm23 and llvm24;
  • the touched tests and every user of the shared helpers, with both engines in strict mode, with and without the binding;
  • the platform suite compared by failure name against the PR base (no new failures);
  • pre-commit.

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.

🔵 Needs a closer look

The cross-engine global-memory lowering and public refusal-contract changes require maintainer CI/toolchain confirmation of byte identity, parity, and llc/hipcc acceptance.

No open issues 🎉

1 resolved since last review

Review effort: Max

@cgmillette

Copy link
Copy Markdown
Contributor Author

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

PR #14842 — refuse buffer_rsrc operands its lowerings cannot take (cgmillette, AICK-2900) — 🟠 Follow-up: rebase with one engine list

Prior item Status
Our 11-02 N1: HIP refusal blamed the builtin 🟢 FIXED 8be45df8
• 7425d423 fixes the ROCm 7.2 / LLVM 22 wording (read)
• left: LLVM 23 declares int64_t too; comments, body say uint64_t (read)
CongMa13 (11-02) APPROVED 🟢 APPROVED 792fafa8
• older head; still counts on develop
• predates 8c6f6523–7425d423: load typing, new refusals, ValueError→TypeError
Copilot (11-02) Rule 9: other global ops emit addrspace(1) 🟢 FIXED 8c6f6523, 7425d423
• 21 global-pointer sites refuse via one helper; 3 loads read addrspace(4) (ran)
• select still names addrspace(1): deferred as older than the PR; fair, develop same (ran)
• re-review at head (11-04): no open issues (read)
Our 11-03 N2: space check copied; copies diverge 🟢 FIXED 7425d423
• one helper per engine at all 22 fixed-space doors; 17 probed: TypeError ↔ status 2, same text (ran)
• LDS pointer: loads and stores both TypeError (ran)
Our 11-03 N3: no CI lane runs the 50 new cases 🟢 FIXED 7425d423
• installed rocke_pytest, no binding: 52 Python cases ✓, 3 skip (ran)
• new ctest case: C++ loads, refusals, statuses (ran)
• CI has not run them yet: head conflicts (N5)
Our 11-03 N4: refusals name ops the caller didn't call 🟡 OPEN, narrowed 7425d423
• builder names, addr_space blame fixed in both engines (ran)
• left: ptr for src_ptr at global_load_lds, global_load_async_to_lds; tile.buffer_rsrc in the width refusals (read)
• fix: pass "src_ptr" (lower_llvm.py:6446, :5570, C++ twins); name the builder
Bots / CI 🟠 Not run
• no build, tests or pre-commit at 7425d423: conflicting since #14328 landed
• 4 ✓ + 2 statuses ✓ (bots, Math CI Summary)
• merge: CONFLICTING (dirty)
Our checks 🟠 Conflicts
• merge vs 98e132f8: engines.py conflicts, 130 behind; resolved locally (ran)
• local merge: build, ctest 45 ✓ / 1 skip; byte identity GREEN ×4, every family (ran)
• goldens PASS ×4; differential 0 changed, NEW as on develop (ran)
• 34 test files: 4438 ✓ with engine, 1515 ✓ without, installed 1509 ✓; rules clean (ran)
Code quality 🟠 Should-fix
• N2, N3 fixed: one check per engine; Python half runs (ran)
• engine list duplicates develop's available_engines(): N5 (read)
• N4's names; both-engine helpers skip the status: 2844 of 2848 agree, rest develop's scale encoding (ran)
# Sev file:line Finding
N5 🟠 rocke/tests/engines.py:13, :41, :53, :66-67 Conflicts with develop's #14328 in engines.py; ENGINES repeats its available_engines()
• impact: no build, tests or pre-commit at this head; cannot merge; keeping both sides leaves two engine lists (API rule 5)
• fix: rebase with one list: available_engines() returns ENGINES, documented for loops only (ran: local merge)

Utility: 🟢 Positive — one space check per engine; Python cases now run without the binding


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: merge-tree with develop; local merge with develop 98e132f8, its engines.py conflict resolved by hand (build, ctest, byte identity llvm20–24, both goldens, 34 test files with and without the engine, installed lanes); saved probes N2 and N3 with negative controls; select probe; status parity of the both-engine refusals; CI at head; code quality. Flags whether a closer look is needed; not an approval or rejection.

cgmillette and others added 7 commits November 4, 2026 11:35
buffer_rsrc lowers to the p1 overload of llvm.amdgcn.make.buffer.rsrc,
which names its pointer ptr addrspace(1). A pointer parameter declared
addr_space="constant" is ptr addrspace(4) in the kernel header, so the
call did not match the header and llc rejected the module. The LLVM
lowering of both engines now refuses it with the text the global-pointer
ops already use:

  buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) (its
  parameter's addr_space)

At llvm20 the intrinsic takes an i32 num_records, and the lowering
printed any num_bytes operand as i32, so an i64 operand did not match
its own definition. Both engines now refuse a num_bytes that is not i32
at llvm20:

  tile.buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20
  make.buffer.rsrc takes an i32 num_records, so pass the byte count as i32

llvm22 and later still take an i32 (zero-extended) or an i64.

The HIP lowering calls __builtin_amdgcn_make_buffer_rsrc, whose
num_records is an int, and cast num_bytes to (int): an i64 byte count was
silently truncated to 32 bits, at every LLVM version. The HIP lowering of
both engines now refuses a num_bytes that is not i32:

  tile.buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; the
  HIP builtin takes an int num_records, so the byte count would be
  truncated at every LLVM version; pass it as i32

The HIP lowering keeps a constant-space pointer: its kernel signature
declares every pointer parameter the same way whatever its addr_space,
and the builtin takes a void*, so the source compiles and the descriptor
is built from the right address.

No IR or HIP source changes for a kernel that lowered before.

Tests: tests/core/test_buffer_rsrc_operands.py checks every refusal in
both engines with the exact text, that the accepted forms lower to the
same IR in both engines at every flavor, and that both HIP lowerings keep
an i32 num_bytes and a constant-space pointer; the C++ test
future_intrinsic_lowering gains the same cases. op_vocabulary.md states
the rules.

Co-Authored-By: Claude <noreply@anthropic.com>
The HIP lowering refuses a buffer_rsrc num_bytes wider than i32. Its
message and comments said the builtin takes an int num_records at every
LLVM version, which is not true: newer clang declares num_records as a
uint64_t. What truncates a wider byte count is the (int) cast the
lowering emits, so the message, the comments and the op reference now
say that, in both engines. The refusal itself is unchanged.

Co-Authored-By: Claude <noreply@anthropic.com>
A pointer parameter declared addr_space="constant" keeps its global IR
type, but the kernel header types it ptr addrspace(4). The global loads,
stores, atomics, flag ops, cooperative store and global/LDS copies named
it ptr addrspace(1) anyway, so LLVM rejected the module.

The global loads now name the pointer's own type, so a kernel can read
through a constant-space parameter; a flat or LDS pointer, which they
also named ptr addrspace(1), is refused. Every other op that addresses
global memory through the pointer refuses a constant-space parameter in
the LLVM lowering of both engines, as buffer_rsrc and the global_ptr ops
already do. The C++ global_store_async_from_lds and
global_load_tr16_b128 lowerings now check the header type as the Python
engine does. The shared pointer-space resolver's message says "op"
rather than "intrinsic", since plain loads use it too.

Co-Authored-By: Claude <noreply@anthropic.com>
The Python engine refuses a global-pointer op on a parameter whose header
type is not ptr addrspace(1) with a TypeError. The C++ engine failed the
same refusal with ROCKE_ERR_VALUE, and so did its
global_store_async_from_lds and global_load_tr16_b128 pointer checks. It
now fails them with ROCKE_ERR_TYPE, the status that matches TypeError.

Co-Authored-By: Claude <noreply@anthropic.com>
… engine

Every LLVM-lowering op that names its pointer in one fixed space now refuses
any other header type through one helper per engine
(_require_header_ptr / rocke_ll_require_header_ptr, with a global-pointer
form), so the store, atomic, flag, prefetch, transposed-load, cluster-load
and LDS-copy refusals share one text and one error type: TypeError in
Python, ROCKE_ERR_TYPE in C++. The multi-space resolver the global loads,
av_*_b128 and s_prefetch_* use refuses with the same type.

The refusal names the op's builder (store_f16 rather than global_store,
global_store rather than global_store_typed) and the operand, and blames
the parameter's addr_space only when the override moved its header type.

The engine-parity test helpers gain per-engine forms, which also check the
C++ status against the expected exception type. The constant-space and
buffer_rsrc tests run per engine, so their Python cases run in lanes that
do not build the rocke_engine binding, and a new future_intrinsic_lowering
CTest case checks a sample of the C++ side there.

The comments on the HIP builtin now say that clang after ROCm 7.1 takes a
64-bit num_records (int64_t in LLVM 22, uint64_t after).

Co-Authored-By: Claude <noreply@anthropic.com>
…refusals

The buffer_rsrc byte-count refusals now name the builder (buffer_rsrc)
rather than the op (tile.buffer_rsrc), and global_load_lds and
global_load_async_to_lds name their pointer operand src_ptr, as their
builders do. Both engines change together. The HIP lowering's comment
now says the builtin takes an int64_t num_records in LLVM 22 and 23.

Co-Authored-By: Claude <noreply@anthropic.com>
The per-engine helpers now check for the rocke_engine binding as soon as
they are asked for the C++ engine, so a "cpp" case skips with the
binding's reason in a lane that does not build it, rather than lowering
with the Python engine first. available_engines() says when to use it,
and the helpers' own argument errors name the bad value and what they
expected.

Co-Authored-By: Claude <noreply@anthropic.com>
@cgmillette
cgmillette force-pushed the users/cgmillette/rocke/buffer-rsrc-refusals branch from 7425d42 to 6efc600 Compare November 4, 2026 16:59
@cgmillette

Copy link
Copy Markdown
Contributor Author

Rebased onto develop and addressed the 11-04 follow-up review.

N5 (conflict in rocke/tests/engines.py): resolved during the rebase. There is now one engine list. available_engines(), which develop added, returns ENGINES, or ("python",) when the rocke_engine binding is not importable. Its docstring says it is for tests that loop over the engines themselves; pytest tests parametrize over ENGINES. all_engines_backend() is unchanged. Every develop caller of both functions passes with both engines in strict mode.

N4 (refusals name what the caller called):

  • The buffer_rsrc byte-count refusals (LLVM at llvm20, LLVM width check, HIP) now name the builder, buffer_rsrc, rather than the op tile.buffer_rsrc.
  • global_load_lds and global_load_async_to_lds name their pointer src_ptr, as their builders do.
  • Both engines change together, and the tests assert the new text.

N1 leftover: the HIP lowering's comment and the PR description now say the builtin takes an int64_t num_records in LLVM 22 and 23 and a uint64_t after that. This matches BuiltinsAMDGPU in the LLVM release/22.x, release/23.x and main branches.

Both-engine helpers and the status code: they still compare only the message. The module docstring says why: some older refusals do not yet report the same status as the Python exception type. The per-engine helpers this PR adds do check the status.

Our review gate also changed the per-engine helpers in one way: a "cpp" case now checks for the binding first, so in a lane without it the case skips with the binding's reason before the Python engine runs.

Checks:

  • build and ctest: 45 run, all pass except rocke_arch_domain_reproducibility, which skips;
  • byte identity GREEN at llvm20, llvm22, llvm23 and llvm24;
  • the new test files: 107 cases pass with both engines in strict mode; without the binding, 55 pass and 52 skip;
  • the 35 test files that use rocke.tests.engines pass with both engines in strict mode;
  • platform suite: every failure also fails on develop;
  • pre-commit passes.

@cgmillette

Copy link
Copy Markdown
Contributor Author

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

PR #14842 — refuse buffer_rsrc operands its lowerings cannot take (cgmillette, AICK-2900) — 🟡 Follow-up: items fixed; CI needs push

Prior item Status
Our 11-02 N1: HIP refusal blamed the builtin 🟢 FIXED 61a805c5
• f99650e0, body: int64_t in LLVM 22 and 23, uint64_t after; LLVM agrees (read)
• left: test docstring test_buffer_rsrc_operands.py:25 still says "LLVM 22, uint64_t after" (read)
CongMa13 (11-02) APPROVED 🟢 APPROVED 792fafa8
• older head; still counts on develop
• predates 61a805c5–6efc6001: load typing, TypeError contract, per-engine test helpers
Copilot (11-02) Rule 9: other global ops emit addrspace(1) 🟢 FIXED 738c9095, 6196bec3
• rebase changed context only; N2 probe FIXED at this head (ran)
• select still names addrspace(1), as on develop (ran); tracked in AICK-2977 (read)
Our 11-03 N2: space check copied; copies diverge 🟢 FIXED 6196bec3
• saved probe at this head: 17 ops, one Python site, same C++ status and text (ran)
Our 11-03 N3: no CI lane runs the 50 new cases 🟢 FIXED 6196bec3
• installed rocke_pytest: 52 Python cases ✓; C++, compile cases skip (ran)
• CI: C++ ctest ✓ in the Linux superbuild; Python shards not yet run (read)
Our 11-03 N4: refusals name ops the caller didn't call 🟢 FIXED f99650e0
• 3 byte-count texts say buffer_rsrc; LDS copies say src_ptr; engines agree (ran)
Our 11-04 N5: engines.py conflict; two engine lists 🟢 FIXED 6196bec3
• one list: available_engines() returns ENGINES or Python alone (read)
• 35 files using the helpers pass with and without engine, bar host hipcc (ran)
Bots / CI 🟡 Red: develop's break
• fails: 4 math-libs stages — develop's hipCCL configure (libhipcxx), 14976 (read)
• ref 6116d1d6 predates revert 15089; re-runs reuse it: rocKE shards need a push
• 24 ✓ incl. pre-commit, both hipDNN superbuilds; Math CI Summary not posted
Our checks 🟢 Clean
• merge vs 0c4056ee: clean, 28 behind (ran)
• build, ctest 46 ✓ / 1 skip; byte identity GREEN ×4; goldens PASS, differential 0 changed (ran)
• changed tests 563 ✓ at llvm22; installed 389 ✓ / 197 skip; rules clean (ran)
Code quality 🟡 Minor
• N5, N4 fixed; cpp cases skip first, one gate (ran)
• both-engine status gap: develop's, tracked in AICK-2977 (read)
• overlap with PR 15036: N6; this PR owns _assert_python_rejects (ran)
# Sev file:line Finding
N6 🟡 rocke/tests/engines.py:129 Overlap: PR 15036 adds its own _assert_python_rejects here
• impact: copies at its engines.py:86, :99; heads conflict in 4 hunks, both defs outside them (ran)
• fix: none here: owner this PR (same module; lands first, wave 1); PR 15036 drops its copies

Utility: 🟢 Positive — refusals now name the builder and operand; one engine list


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 and replay of the rebase; merge with develop (build, ctest, byte identity llvm20–24, both goldens, changed tests at llvm22, installed lane); 35 files using the engine helpers, unforced, with and without the engine; saved probes N2 and N3; both engines on the third byte-count text; select probe; LLVM's builtin declarations; CI logs of both attempts; trial merge with PR 15036; code quality. Flags whether a closer look is needed; not an approval or rejection.

@cgmillette
cgmillette enabled auto-merge (squash) November 5, 2026 03:10
@cgmillette
cgmillette merged commit f28ed8f into develop Nov 5, 2026
98 checks passed
@cgmillette
cgmillette deleted the users/cgmillette/rocke/buffer-rsrc-refusals branch November 5, 2026 05:25
dlangbe added a commit that referenced this pull request Nov 5, 2026
…e-gfx942-i64-cfvst-loader

Resolve conflicts with #15024 (runtime kv head count) and #14842
(buffer_rsrc operand checks). The i64 cfvst V page offset and its C++ twin
now use the runtime token stride. Re-bless gfx942_attention_tiled_2d
cfg18/19 in the differential golden.

Co-Authored-By: Claude <noreply@anthropic.com>
cycheng pushed a commit that referenced this pull request Nov 6, 2026
…4842)

## Motivation

`buffer_rsrc(ptr, num_bytes)` lowers to `@llvm.amdgcn.make.buffer.rsrc`.
Two operand forms reached
the LLVM lowering of both engines unchecked, and the module they
produced was then rejected by `llc`
with an error that did not point at the op:

- a pointer parameter declared `addr_space="constant"`: it is `ptr
addrspace(4)` in the kernel
header, but the call names its pointer `ptr addrspace(1)` (LLVM 20
`llc`: `'%X' defined with type
  'ptr addrspace(4)' but expected 'ptr addrspace(1)'`);
- an i64 `num_bytes` at llvm20: that flavor's intrinsic takes an i32
`num_records`, and the lowering
printed any operand as `i32` (LLVM 20 `llc`: `'%n' defined with type
'i64' but expected 'i32'`).

The HIP lowering has the same kind of bug for the byte count: it calls
`__builtin_amdgcn_make_buffer_rsrc` with `num_bytes` cast to `(int)`, so
an i64 byte count was
silently truncated to 32 bits (hipcc emits `trunc i64 %n to i32` before
the call), and the
descriptor bounded a smaller range than asked. ROCm 7.1's builtin takes
an `int` `num_records`
and later clang a 64-bit one (`int64_t` in LLVM 22 and 23, `uint64_t`
after); the truncation comes from
the cast the lowering emits.

The constant-space mismatch was not specific to `buffer_rsrc`. The
global loads, stores, atomics,
flag ops, cooperative store and global/LDS copies also named a
constant-space parameter
`ptr addrspace(1)`, and LLVM rejected the module.

This PR refuses these forms up front, in both engines, with a message
that says what is wrong, and
lets the global loads read through a constant-space parameter.

## Technical Details

- Both checks need the lowering's context (the parameter's header type,
the LLVM flavor), so they
sit in the LLVM lowering of each engine, at the same point: Python
`_op_tile_buffer_rsrc`
(`core/lower_llvm.py`), C++ `op_tile_buffer_rsrc`
(`cpp/core/lower_llvm/mem.cpp`). The
global-pointer ops (`global_ptr_add`, `global_ptr_offset`,
`global_ptr_to_i64`) already refuse a
constant-space parameter there, and `buffer_rsrc` now uses the same
helper
(`_require_global_header_ptr` / `rocke_ll_require_global_header_ptr`)
and so the same text:

  ```
buffer_rsrc ptr must be a global pointer, got ptr addrspace(4) (its
parameter's addr_space)
  ```

- At llvm20 a `num_bytes` that is not i32 is refused:

  ```
buffer_rsrc num_bytes must be i32 at llvm20, got i64; the llvm20
make.buffer.rsrc takes an
  i32 num_records, so pass the byte count as i32
  ```

llvm22 and later still take an i32 (zero-extended) or an i64, unchanged.
- The HIP lowering of both engines (Python `_op_tile_buffer_rsrc` in
`core/lower_hip.py`, C++
`_op_tile_buffer_rsrc` in `cpp/core/lower_hip/lower_hip_mem.cpp`)
refuses a `num_bytes` that is not
i32, at every LLVM version, because it passes `num_records` cast to
`(int)`:

  ```
buffer_rsrc num_bytes must be i32 in the HIP lowering, got i64; the HIP
lowering casts
num_records to (int), so a wider byte count would be truncated; pass it
as i32
  ```

  An i64 byte count still lowers through LLVM at llvm22 and later.
- The HIP lowering keeps a constant-space pointer: its kernel signature
declares every pointer
parameter the same way (`float* __restrict__ X`) whatever its
`addr_space`, and the builtin takes a
`void*`. A hipcc probe of that source compiles, and the descriptor is
built from the kernel
  argument's address.
- Other global-memory ops on a constant-space parameter, in the LLVM
lowering of both engines:
- `global_load_f16`, `global_load` and `global_load_vN` name the
pointer's own type (a table of
the accepted spaces, `_GLOBAL_LOAD_PTR_TYPES` /
`ROCKE_LL_GLOBAL_LOAD_PTR_TYPES`, resolved by
the existing pointer-space resolver). They read a constant-space
parameter as
`load addrspace(4)` and refuse a flat or LDS pointer, which they also
used to name
`ptr addrspace(1)`. The resolver's message now says "op" rather than
"intrinsic".
- The stores, atomics, flag ops, `cooperative_global_store`,
`global_load_lds` and
`global_load_async_to_lds` refuse a constant-space parameter with the
`global_ptr_*` text.
- One check per engine: every op that names its pointer in one fixed
space (the ops above,
`global_prefetch`, `flat_prefetch`, `global_store_async_from_lds`,
`global_load_tr16_b128`
and both cluster loads) calls `_require_header_ptr` /
`rocke_ll_require_header_ptr`, whose
global-pointer form is `_require_global_header_ptr` /
`rocke_ll_require_global_header_ptr`.
The C++ `global_store_async_from_lds` and `global_load_tr16_b128`
lowerings used to check only
the IR type; they now check the header type, as the Python engine
already did.
- Every one of these refusals, and the pointer-space resolver's, is a
`TypeError` in Python and
`ROCKE_ERR_TYPE` in C++. The text names the op's builder (`store_f16`,
`global_store`,
`global_load_f16`, `global_load`) and the builder's name for the operand
(`src_ptr` for the
copies into LDS), and says "(its parameter's addr_space)" only when the
parameter's override
moved its header type. The `buffer_rsrc` byte-count refusals also name
the builder.
- The portable IR route lowers through the C++ engine, so it refuses the
LLVM forms too.
- No emission changes for a kernel whose IR LLVM accepted before: byte
identity is unchanged, no
golden is re-blessed, and the HIP source of every parity kernel is
unchanged.
- `dsl_docs/reference/op_vocabulary.md` states the rules next to the
buffer ops.

## Test Plan

- `tests/core/test_buffer_rsrc_operands.py`:
- both LLVM lowerings refuse a constant-space pointer at llvm20, llvm22,
llvm23 and llvm24, and an
    i64 `num_bytes` at llvm20, with the exact text;
- controls: an i32 `num_bytes` at every flavor and an i64 one at llvm22
and later lower to
    identical IR in both engines;
- both HIP lowerings refuse an i64 `num_bytes` with the exact text, and
lower an i32 `num_bytes`
    over a global or a constant-space pointer to the same source.
- `tests/core/test_constant_space_global_ops.py`: each LLVM lowering
reads through a
constant-space parameter with each global load at every flavor, the
result compiles with the
host clang, and a flat or LDS pointer is refused; every other op listed
above refuses a
constant-space parameter with the exact text and type, an LDS parameter
is refused without the
addr_space clause, and `flat_prefetch` refuses a parameter moved to
constant space.
- Both test files run each case per engine (`rocke/tests/engines.py`
gains per-engine helpers that
also check the C++ status), so the Python cases run in lanes that do not
build the
`rocke_engine` binding; a C++ case there skips before either engine
runs. Develop's
`available_engines()` returns the same `ENGINES` list, or the Python
engine alone when the
  binding is not built.
- `tests/core/future_intrinsic_lowering.cpp` (ctest
`rocke_future_intrinsic_lowering`): the C++
LLVM and HIP lowerings refuse the `buffer_rsrc` forms, with their status
codes, and still lower
the accepted ones; a `constant_space_global_ops` case checks a sample of
the constant-space
  loads and refusals, so the C++ side also runs in those lanes.
- HIP source of every parity-family config (95 families) for gfx942,
gfx950 and gfx1201, before and
  after the change.
- Build, ctest, byte identity at llvm20 / llvm22 / llvm23 / llvm24,
golden checks, portable IR
against a freshly built shared engine, the touched tests with both
engines in strict mode, and
  the platform suite compared by failure name against develop.

## Test Result

- Before the change (develop), both `buffer_rsrc` LLVM forms lowered
without error and LLVM 20
`llc` rejected the result with the errors above; the HIP lowering of an
i64 `num_bytes` compiled,
and hipcc truncated the byte count to 32 bits. A global load through a
constant-space parameter
  produced a module the host clang rejects.
- HIP source (first revision): the Python HIP lowering's 3772 (config,
arch) texts across the 95
parity families are identical to develop's; the later commits change the
HIP lowering only in the
  refusal's text.
- Byte identity GREEN at llvm20, llvm22, llvm23 and llvm24.
- Golden checks with both engines at all four flavors: no blessed config
changed. The report also
lists the Python-only `barrier_fence` family as absent from the
both-engine run, and configs not
  yet in the golden file. No golden file changed.
- ctest: 45 run, all pass except `rocke_arch_domain_reproducibility`,
which skips; this includes
  `rocke_future_intrinsic_lowering`.
- The two new test files: 107 cases with both engines in strict mode;
without the binding, 55
  Python cases pass and the 52 C++ cases skip.
- The 35 test files that use `rocke/tests/engines.py`, with both engines
in strict mode: 5266
  passed, 134 skipped.
- Platform suite (`tests` and `python/rocke/tests`) with both engines in
strict mode and no GPU
  visible: 11255 passed, 14 failed. All 14 also fail on develop:
- `test_the_cpp_fallback_does_not_warn_twice` stubs the engine and fails
only under
    `ROCKE_CPP_STRICT=1`;
- four `test_installed_ctest_entries_get_scoped_environment_contracts`
cases fail the same way
    on develop;
  - the nine interleaved-GEMM demo tests need a GPU.
- Library golden tests: 93 passed.
- pre-commit passes.

## Submission Checklist

- [x] Look over the contributing guidelines at
  https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

ISSUE ID : AICK-2900

🤖 Generated with [Claude Code](https://claude.com/claude-code)

---------

Co-authored-by: Claude <noreply@anthropic.com>
@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.

3 participants