fix(rocke): add address spaces p10-p15 to the LLVM 23 datalayout by CongMa13 · Pull Request #12672 · ROCm/rocm-libraries

CongMa13 · 2026-09-28T22:26:25Z

fix(rocke): add address spaces p10-p15 to the LLVM 23 datalayout

AICK-2416

Summary

rocKE hardcodes a per-flavor AMDGPU target datalayout string. The llvm23 constant was
stale — upstream 5bf967cb132b added address spaces p10–p15 and rocKE never picked
them up. This PR adds the six fields to all three copies of the constant and teaches the
drift guard to handle mixed-toolchain boxes.

Every attention kernel currently fails to build on a current amd-staging toolchain.

Why it surfaced only now

The bug has been latent since 5bf967cb132b. It became fatal at fc6829a3 ("clang: Do
not overwrite a module's DataLayout in the backend").

	Before `fc6829a3`	After `fc6829a3`
Backend behaviour	silently replaces the module's DataLayout with the target's	validates the module's DataLayout against the target's
Effect of the stale string	corrected in flight — nothing fails	hard codegen error

So the stale constant was being quietly repaired on the way through the backend. Once
clang stopped doing that, a module carrying the short form fails outright:

Can't create a MachineFunction using a Module with a Target-incompatible
DataLayout attached

What changed

Three copies of the constant, kept byte-identical per the #1 invariant (the Python and
C++ engines must emit identical LLVM-IR bytes):

File	Constant
`platform/python/rocke/core/lower_llvm.py`	`_DATALAYOUT_LLVM23`
`platform/cpp/core/lower_llvm/data.cpp`	`ROCKE_LL_DATALAYOUT_LLVM23`
`platform/cpp/core/isa/backend.cpp`	`ROCKE_DATALAYOUT_LLVM23`

The added fields, inserted between p9 and i64:64:

-p10:32:32-p11:32:32-p12:32:32-p13:32:32-p14:32:32-p15:32:32

Confirmed gfx-invariant — these address spaces are AMDGPU-generic, not per-arch —
across every wired arch (gfx90a / 942 / 950 / 1100 / 1151 / 1201 / 1250), so one value
covers all of them. Nothing else in the string drifts.

Plus a narrow skip arm in test_datalayout_matches_hipcc_emitted_ir; see below.

Why the superset is safe on older toolchains

A box can legitimately pair a new-numbered ROCm with an older LLVM, so one hardcoded
constant has to be correct against toolchains on both sides of the change. It is, because
5bf967cb132b landed before fc6829a3:

A backend that enforces the module DataLayout has fc6829a3, hence also
5bf967cb132b — so it has p10–p15 and requires them.
A backend lacking p10–p15 predates 5bf967cb132b, hence fc6829a3 — so it
still overwrites the module DataLayout and tolerates the longer string.

Verified empirically in both directions, including that an older /opt/rocm clang accepts
the fixed IR (rc=0).

This is why the drift guard skips with a named reason rather than failing when it
meets a hipcc predating 5bf967cb132b. A bare failure there would push someone toward
"fixing" the constant by deleting the fields — which re-breaks every kernel on a current
toolchain. The skip arm is deliberately narrow: it triggers only when the sole difference
is exactly the p10–p15 block. Negative-tested — p8 drift, a missing m:e, A5→A4,
a partial p10–p13, and a wrong-width p10 all still fail.

Verification

Compile-time, all green:

Byte-identity gate at llvm20 / llvm22 / llvm23 — all families, each flavor.
C++ SSOT test (flavor_datalayout_ssot) green across all three flavors, holding the
ISA-layer and lowerer tables together by construction.
Python constant byte-exact against IR emitted by staging clang.
Drift guard: passes on the host; on a container whose hipcc predates
5bf967cb132b, 1 passed / 7 skipped with the named reason.
Lint: ruff check (no --fix, per the emitter-code rule) at baseline — 0 new
findings. pre-commit green (black, clang-format, bandit, whitespace/EOF).

End-to-end on the enforcing toolchain — a real MFMA attention forward
(kernels.common.fmha_mfma.build_fmha_fwd_mfma) lowered through rocKE at the llvm23
flavor, then codegen'd and linked by staging clang:

Case	Result
gfx942	builds + links
gfx942, causal	builds + links
gfx950	builds + links

Control — same kernel, same toolchain, layout reduced to the pre-fix form by deleting
the six fields: rc=1 with a backend crash reproducer, versus rc=0 with the fix. One
variable changed, so the check demonstrably has teeth.

Reviewer notes

The comments intentionally reference upstream commits rather than dates: the
safety argument is about which commit landed first, and commit ordering stays checkable
against the upstream repo.
If a future toolchain drifts further, regenerate via
test_datalayout_matches_hipcc_emitted_ir rather than hand-editing.
Any change to an emitted byte must move both engines together and re-run the
byte-identity gate.

Test plan

cd platform && export ROCKE=$PWD PYTHONPATH=$ROCKE/python

# #1 invariant, every flavor
python3 tools/check_byte_identity.py
ROCKE_LLVM_FLAVOR=llvm22 python3 tools/check_byte_identity.py
ROCKE_LLVM_FLAVOR=llvm23 python3 tools/check_byte_identity.py

# drift guard + datalayout unit tests
python3 -m pytest tests/instances/test_rocke_multiarch.py -k datalayout -q

# on a healthy gfx942/gfx950, the numerics gate that is still outstanding
PYTHONPATH=$ROCKE/python:$PWD/../library python3 \
  library/builders/common/fmha_fwd_verify_hip.py --arch gfx942

rocKE hardcodes a per-flavor AMDGPU target datalayout. The llvm23 constant was stale: upstream 5bf967cb132b added address spaces p10-p15, and rocKE never picked them up. This was latent until fc6829a3 ("clang: Do not overwrite a module's DataLayout in the backend"). Before it, the backend silently replaced the module's DataLayout with the target's, so the short string was corrected on the way through and nothing failed. After it, a module carrying the stale form fails codegen outright: Can't create a MachineFunction using a Module with a Target-incompatible DataLayout attached which breaks every attention kernel on a current amd-staging toolchain. Add the six fields to all three copies of the constant -- the Python lowerer, the C++ lowerer, and the ISA-layer accessor -- keeping them byte-identical per the #1 invariant. Verified gfx-invariant across every wired arch, so one value covers all of them. Carrying p10-p15 is safe on older toolchains too, because 5bf967cb132b landed before fc6829a3: a backend that enforces the module DataLayout also has p10-p15 and requires them, while a backend lacking p10-p15 predates the enforcement change and so still overwrites and tolerates the longer string. The drift guard gains a narrow skip arm for that second case, so a hipcc predating 5bf967cb132b reports a named skip instead of a failure that would tempt someone to "fix" the constant by deleting the fields again. Verification: byte-identity green at llvm20/llvm22/llvm23; the C++ SSOT test green across all three flavors; a real MFMA attention forward lowers, codegens, and links under staging clang for both gfx942 and gfx950, masked and unmasked, where the pre-fix layout fails on the same kernel and toolchain. On-silicon numerics are not yet covered -- the harness is ready but no healthy device was reachable. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

therock-pr-bot · 2026-09-28T22:31:26Z

✅ 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 · 2026-09-28T22:31:29Z

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

tenpercent

Reviewed d51c37b. The three updated constants are byte-identical, and the added fields match the cited LLVM change. I found no blocking defect in this diff; the two inline comments are non-blocking suggestions to make the compatibility explanation more precise.

Validation: the focused datalayout test passes on ROCm 7.1, with gfx1250 skipped as unsupported. Controlled inputs to the changed test confirm that it accepts the exact layout, skips the complete missing p10-p15 block, and rejects p8 drift, missing m:e, A4, a partial block, and an incorrect p10 width. Minimal kernel compilation accepted both layout forms on the two LLVM 23 builds and one LLVM 24 nightly I checked. I did not reproduce the reported failure on an enforcing staging compiler or rerun full engine parity or GPU numerics.

CongMa13

Reviewed d51c37bc743eabdd2374d8fbdf6a0286d934c1f7. Two host-test regressions need accompanying updates: the existing CPU datalayout assertion still expects only the m:e difference, and the committed llvm23 golden hashes still describe the previous IR header. Python/C++ parity alone does not cover either failure.

Static review only; no build, tests, or project-code execution performed.

Review follow-up on the p10-p15 datalayout fix. Scope the safety argument to what was actually observed. The previous text deduced that carrying p10-p15 is safe on every toolchain from the fact that 5bf967cb132b landed before fc6829a3. That does not follow: LLVM's DataLayout compatibility check predates fc6829a3 -- it is already in LLVM 22's MachineFunction.cpp -- and fc6829a3 only stops clang replacing an explicitly supplied layout, which exposes the pre-existing check. So the commit ordering does not establish that every enforcing backend has fc6829a3. Replace the deduction with the two clang builds actually exercised: staging clang rejects the short form and accepts the long one; an older /opt/rocm clang accepts the long form because it overwrites the supplied layout before codegen. Label that as an observation about those builds, not a proof about every toolchain. Distinguish rocKE's llvm23 flavor from a claim about all LLVM 23+ / ROCm 7.13+ builds, since compiler builds vary and older ones -- including some numbered LLVM 23 or later -- omit p10-p15. Carry the same distinction into the drift guard, whose skip comment and message now describe only what the comparison shows: this hipcc's layout differs from the constant by exactly the p10-p15 block, treated as compiler-build variation. It says nothing about whether that build enforces DataLayout compatibility, so the skip is no longer worded as if it did. Shorten both C++ copies to a summary deferring to the Python constant, so the rationale lives in one place. Drop the specific staging version string from the comment rather than name an unreleased build. Update the CPU-side contract test, which still asserted llvm23 == llvm22 plus m:e and so failed on every host, and rename it to match. It also now asserts the p10-p15 block on its own, so dropping the fields names the field instead of diffing two long strings. Re-bless tests/golden/rocke_representative_ir_sha256.json for the intended +60-byte llvm23 header. Verified the llvm20 and llvm22 sections are unchanged (0 of 146 cases) and all 150 llvm23 cases differ in sha256/bytes only, each by exactly +60. Verification: byte-identity GREEN at llvm20/llvm22/llvm23; C++ flavor_datalayout_ssot passes; test_rocke.py + test_rocke_ci_static.py 273 passed / 60 skipped; drift guard 1 passed / 7 skipped on this host; ruff check at baseline, 0 new findings. On-silicon numerics remain outstanding, as before. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

cgmillette

PR #12672 — add p10-p15 to LLVM23 datalayout · Verdict: 🟢 merge-worthy (fixes fatal attention-kernel build break; minimal, 3 copies byte-synced, well-tested)

#	Sev	Location	Net-new finding
1	🟡	PR body	Issue key `JIRA-2416` malformed → won't auto-link Jira; should be `AICK-2416`
2	🟢	both test files	`assertIn(p9+p10-p15+i64)` encodes strict p9↔i64 adjacency; a future field insert there fails the test even when p10-p15 correct
3	🟢	multiarch skip arm	`replace("-i64:64",…)` no-ops if anchor absent → benign p10-p15-only mismatch fails (not skips); fail-closed
V1	🟢	golden json	Verified: only 150 llvm23 entries, each +60 bytes; llvm20/22 untouched
V2	🟢	repo-wide	Verified: literal exists only in the 3 updated files — no missed copy

Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor-nit · 🟢 info/confirmed. Existing tenpercent + CongMa13 threads all
addressed in head — not re-filed.

…ut-p10-p15

Merging develop brought in #10999, which added four dense-attention cases: attention_dense/gfx950/bottom_right_bm128_sq512 attention_dense/gfx950/bottom_right_sinks_sq512 attention_dense/gfx950/bottom_right_sq512 attention_dense/gfx950/ragged_bottom_right_sq500 Their goldens were blessed before this branch added p10-p15 to the llvm23 datalayout. The four are new JSON keys, so the golden merges cleanly and no conflict flags the staleness -- but their llvm23 hashes still describe the short header, and rocke_golden_static fails on the merge result even though both sides pass on their own. Re-bless so they carry the fixed header. Verified the delta is exactly those four llvm23 entries, each +60 bytes, with llvm20 and llvm22 unchanged (0 of 150 cases) -- the same +60 every other llvm23 case took. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…ut-p10-p15 Conflict in the representative-IR golden: develop renamed five conv_wgrad_reduce / conv_wgrad_two_stage kernels (wsr8 suffix) and re-hashed them, while this branch re-blessed the llvm23 flavor for the p10/p15 datalayout change. Took develop's kernel names and re-blessed the llvm23 sub-doc with the parity harness; the five entries pick up the same +60-byte datalayout delta as every other llvm23 case. rocke_ir_parity_harness --check: IR parity OK [llvm20, llvm22, llvm23]

…ut-p10-p15 Conflict: tests/golden/rocke_representative_ir_sha256.json, llvm23 only, on the four gfx1250 block-scaled WMMA GEMM cases: gemm/gfx1250/mxfp8/wmma_scale/bf8e5m2 gemm/gfx1250/mxfp8/wmma_scale/fp8e4m3 gemm/gfx1250/mxfp8/wmma_scale16/bf8e5m2 gemm/gfx1250/mxfp8/wmma_scale16/fp8e4m3 Both sides re-blessed the same keys for different reasons, so neither side is correct on its own: - develop #11960 (packed FP4 scaled WMMA) rewrote the emission itself (89055 -> 35091 and 95649 -> 41653 bytes); - this branch added p10-p15 to the llvm23 datalayout (+60 bytes of header on every llvm23 case). Resolved by taking develop's emission and re-running the parity harness so the header fix lands on top of it, rather than hand-picking either side. Verified the regenerated golden differs from develop's only in llvm23, in all 154 cases, by exactly +60 bytes each, with no non-sha field moving and llvm20/llvm22 byte-identical. Gates: rocke_ir_parity_harness --check GREEN for llvm20/llvm22/llvm23; tests/test_rocke_ci_static.py 5 passed; all 154 dumped llvm23 modules parse clean under llvm-as from the llvm23 amd-staging build. check_byte_identity.py reports DRIFT on gfx942_attention_tiled_3d cfg[3] (69 GREEN / 1 DRIFT). That drift reproduces on pristine develop (812bcc6) in a clean worktree, so it is pre-existing upstream and not introduced by this merge. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…ck-2416-llvm23-datalayout-p10-p15

Adding address spaces p10-p15 to the llvm23 datalayout (d51c37b) shifted every byte of llvm23-flavored IR, but only the main byte-stability golden (tests/golden/rocke_representative_ir_sha256.json) was re-blessed. tests/core/gfx1250_scaled_wmma_llvm23.json is a second, standalone llvm23 sha pin consumed by test_gfx1250_fp6_scaled_wmma.py, and it was left on the pre-p10-p15 hashes -- so all six cases failed under the rocke_pytest CTest entry. Verified the delta is exactly the datalayout block and nothing else: for all six cases, deleting the p10-p15 substring from today's emitted IR reproduces the previously committed sha bit-for-bit. No kernel body moved. tests/core/test_gfx1250_fp6_scaled_wmma.py: 82 passed, 12 skipped. rocke_ir_parity_harness.py --check: IR parity OK [llvm20, llvm22, llvm23]. Co-Authored-By: Claude Code <noreply@anthropic.com>

…ck-2416-llvm23-datalayout-p10-p15 Conflict: tests/golden/rocke_representative_ir_sha256.json, edited on both sides in orthogonal ways: * ours -- llvm23 only, all 154 cases, exactly +60 bytes each (the length of the "-p10:32:32...-p15:32:32" block). Pure datalayout header shift. * develop -- 57 real kernel changes across all three flavors (varied byte deltas, 5 kernel_name moves per flavor), from the conv AOT switch. The golden is a derived artifact, so a textual merge of two sets of hashes is meaningless. Resolved by seeding from develop's side and regenerating with rocke_ir_parity_harness.py --write against the merged source tree. Verified the result is develop's golden plus our change and nothing else: * llvm20: 0/150 cases differ from develop * llvm22: 0/150 cases differ from develop * llvm23: 154/154 differ, every one by exactly +60 bytes, only the bytes/sha256 fields moved -- no kernel_name, no case added or removed, expected_failures and summary unchanged on all three flavors. Harness counts match the expected baseline (llvm20 150 ok/4 fail, llvm22 150/4, llvm23 154/0), confirming a sane env rather than a silently broken one. Gates: rocke_ir_parity_harness --check -> IR parity OK [llvm20, llvm22, llvm23]; test_rocke_ci_static.py 5 passed; platform pytest 1725 passed, 355 skipped; library gfx950 llvm23 goldens + layering 12 passed. Co-Authored-By: Claude Code <noreply@anthropic.com>

…layout change #12672 (merged from develop) added address spaces p10-p15 to the LLVM 23 datalayout, which changes every llvm23 module header (+60 bytes). It re-blessed the cases on develop, but the nontemporal copy_bf16x8 cases added by this PR still carried the pre-#12672 llvm23 hashes. llvm20/llvm22 are unchanged. Verified: check_golden reports 0 drift across llvm20/22/23, and the file equals a fresh build_golden() (identical to harness --write).

assistant-librarian · 2026-10-07T12:35:17Z

TheRock Submodule Bump Activity

Newest first

2026-10-07 — Included in TheRock via ROCm/TheRock#8777.

…decode (#13226) JIRA ID : AICK-2522 ## Motivation Three library IR goldens have llvm23 entries that never matched `develop`: | Golden | Stale llvm23 entries | Origin | |---|---|---| | `attention_dense_gfx942` | 20/20 | #12735, blessed on commits made before #12672 merged | | `kda_gfx950` | 20/20 | on `develop` when #12672 merged | | `gdn_decode_gfx950` | 63/63 | on `develop` when #12672 merged | #12672 added address spaces `p10`-`p15` to `_DATALAYOUT_LLVM23` and re-blessed only the platform goldens. `library/tests` isn't in CI, and the KDA and GDN decode tests compared only the host's LLVM flavor, so nothing failed outside an llvm23 host. `test_attention_dense_gfx942_golden` checks every flavor and has failed on `develop` since #12735 merged. ## Technical Details 1. **Re-bless.** Only the llvm23 sub-documents change. Each entry grows by exactly 60 bytes (`-p10:32:32` through `-p15:32:32`), and every previous hash is reproduced by putting the pre-#12672 datalayout back into today's IR. No emitted code changed. 2. **Check every flavor.** `test_kda_gfx950_golden.py`, `test_kda_gfx942_golden.py` and `test_gdn_decode_golden.py` now use `rocke.core.ir_golden.check_golden`, like the dense test after #12735, instead of hand-rolled flavor tuples. A datalayout or intrinsic change for any flavor in `LLVM_FLAVORS` now fails them from any host. `kda_gfx942` gains llvm23 entries, identical to llvm22 apart from the datalayout line. Closes the hand-rolled-flavor-list follow-up (N2) from the #12735 review: all three remaining library golden tests now go through `check_golden`, and `kda_gfx942` records llvm23. Wiring `library/tests` into CI (N1) remains a separate change. ## Test Plan - [x] The 4 library golden tests + `test_library_layering.py`, with the host flavor at llvm20, llvm22 and llvm23 - [x] With `develop`'s stale llvm23 hashes restored, the converted tests fail on an llvm20 host (before this change, `kda_gfx950` and `gdn_decode` passed) - [x] `platform/tests/core/test_ir_golden.py`, `test_rocke.py::TestLlvmFlavorEnumeration` - [x] `rocke_ir_parity_harness.py --check tests/golden/rocke_representative_ir_sha256.json`: OK at llvm20, llvm22, llvm23 No emitter change, so no GPU lane or byte-identity gate is needed. ## Test Result All of the above pass. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Determine the LLVM compiler version by querying the loaded COMGR binary. This makes compiler identification independent of ROCm release metadata and installation-directory conventions. Distribution layout describes how software is packaged and installed. The loaded library can answer which compiler version it contains or uses. This change makes that binary the source of version information for automatic LLVM flavor selection and runtime compatibility checks. ```mermaid flowchart TD A[Existing library discovery] --> B[Load COMGR] B --> C{LLVMGetVersion available?} C -->|Yes| D[Query LLVM version through the COMGR handle] C -->|No| E[Preprocess Clang version macros through COMGR] D --> F[Compiler version and provenance] E --> F F --> G[Map version to an emission flavor] F --> H[Check runtime IR compatibility] ``` The query uses `LLVMGetVersion` when COMGR or its loaded dependencies expose it. For builds that hide LLVM symbols, it preprocesses Clang's built-in version macros through COMGR. Neither method reads release files, invokes an external compiler executable, or requires a GPU. Python exposes the query's version and provenance for diagnostics. Compiler ownership and emission policy remain separate: - Standalone C++ AUTO discovers and queries a native candidate. The query remains internal to the implementation. A native caller that owns a separate compilation stage supplies the flavor for that compiler explicitly. - Explicit flavors support offline emission without loading COMGR. Runtime compatibility checks use compiler evidence independently of emission overrides. - Fallbacks remain caller-specific: core lowering keeps `llvm22`; replay/parity drivers keep `llvm20`. Existing environment/argument precedence and best-effort error handling are preserved. - The legacy `resolved_lib_rocm_version()` metadata API keeps its existing behavior; compiler decisions use the new binary query. AUTO now loads COMGR to obtain compiler evidence. Failed loads remain retryable; the first successfully loaded library and its query result are retained for the process lifetime. Configure library selection before automatic lowering or compilation; changing a retained compiler requires a new process. Parity and replay utilities also use the binary-derived version information. Detection selects from the existing flavor DataLayout constants. The LLVM23 `p10`–`p15` fields from develop are preserved in both engines. The detected version selects a baseline flavor; it does not establish every capability of a particular compiler build. Related to #12672 The separate Python/native compiler-selection bug is handled in #12680. Both PRs target `develop` and can land independently. Native detector fixtures compile their wrapper and fake COMGR sources from stdin and link the production archive. Linux installs carry that archive beside the pytest module so the filtered test artifact retains native coverage without private detector source or headers. Source pytest can build a fresh temporary archive or use a matching archive supplied explicitly. Raw archive consumers and CMake exports include the required loader/thread dependencies. Validation at `fa5df1f88b23926940dffa1073c1bb4899e64066`, after merging develop at `69d7ff3ed300dbdcd64c088b1e7f7224e649953a`: - [x] LLVM20/22/23 Python/C++ byte-identity gates: all 73 families and 854 configurations per flavor passed. No golden updates were needed. Both engines retain develop's LLVM23 DataLayout additions. - [x] Fresh native archive and extension build; 14 native CTests passed. Five focused source backend modes passed, including strict C++ and differential mode. Default and Python each reported 326 passed / 60 skipped; permissive C++, strict C++, and both each reported 356 passed / 14 skipped. - [x] Relocated installed artifact with matching non-editable platform/library wheels: 17 selected CTests passed. Installed platform pytest reported 1,751 passed, 338 skipped, and 1,632 subtests passed, including all six native detector cases. Import provenance was verified outside the source tree, without private C++ sources/headers or the optional native binding. GPU reference installation and library-specific CTests were excluded from this host packaging check. - [x] Relocated CMake consumer and raw-link tests: 7 passed. Earlier checks also covered BUILD_TESTING=OFF, in-tree and standalone prebuilt-archive binding imports, raw micro-parity, and online shared-library build/load. Missing detector archives and invalid overrides fail explicitly. - [x] Caller-policy regressions cover absent/unqueryable compilers, detected versions, override precedence, explicit-flavor bypass, query errors, recovery after failed loading, successful-load retention, and the legacy metadata fallback. - [x] Earlier live compiler checks used static LLVM20 COMGR and two shared LLVM23 COMGR builds: matching Python/native version evidence, byte-identical automatic IR, and successful buffer-resource compilation to a code object. - [x] Repository pre-commit hooks passed over the full PR range. - [ ] Fully green broad source suite: with the low-bit dependency installed, the run reported 1,878 passed, 335 skipped, 5,230 subtests passed, and three compiler timeout failures. Two portable-IR failures passed a separate 10-test rerun with a longer limit, including launch units and HSACO parity. Catalog regeneration still reaches the generator's documented nonterminating LLVM20 `raw.ptr.buffer.load.async.lds` probes; that lane was stopped and its compiler children cleaned up. - [ ] Relative-path guard: it flags a DVC diagnostic path confirmed identical on develop. The inherited diagnostic and guard are unchanged. Host validation used a temporary uv CPython 3.12.12 environment with GCC 13.3.0, CMake 3.28.3, and glibc 2.39. The final develop refresh changed only CI workflow/tool-pin files; the validated rocKE subtree is unchanged. GPU numerical validation, Windows execution, and pre-glibc-2.34 execution were not run. Linux link-command regressions fail on the old raw-link commands and pass after repair; simulated Windows checks verify Linux flags are omitted.

…nd guard i32 KV addressing at launch (#12446) JIRA ID : AICK-2144 ## Motivation The forced-2D parity lane in `parity_unified_attention.py` reported correctness failures for `gqa_nqk3_bf16_d128_b64` and `gqa_nqk5_bf16_d128_b64`, while the NQK=8 divisor case passed. That pointed at the padded block-q / `VALID_ROWS < BLOCK_M` path from #9017. After bisection, the failure turned out to be a harness bug. The kernel is correct. Review then showed that the launch guard added for it trusted the problem's i64 decision. Several kernels don't honour that decision, and the i32 limit was slightly below 2 GiB. A second round found that some i64 kernels still read part of K or V through i32 offsets. This PR now handles all three. A third round made the guard's error name what actually bounds the kernel, and pinned the remaining untested paths. ## Technical Details ### Root cause The direct-2D lane (`_run_rocke` with `path="2d"`) builds `UnifiedAttentionProblem` itself instead of going through `run_unified_attention_torch`, and it left out `num_kv_blocks`, so the value stayed at the default `0`. `_enable_i64_kv_addr` reads `0` as "cache size unknown, assume small" and keeps the i32 buffer-voffset path. The parity harness allocates `num_blocks=32768`. The cache size per scenario: | Scenario | `num_kv_heads` | K cache | Path without the fix | |---|---|---|---| | NQK=3, NQK=5 | 8 | 32768 × 64 × 8 × 128 × 2 B = 4 GiB | i32, far past the buffer range | | NQK=7, NQK=8 | 4 | 2 GiB exactly | i32, only the top 64 KiB out of range | The i32 path puts the whole byte offset in a buffer voffset. The tiled builders bound K/V buffer loads at `num_records = 0x7FFF0000`, and a load at or past that offset returns zero without faulting. On the 4 GiB caches every high physical block reads as zeros, so NQK=3 and NQK=5 failed. Only the top 64 KiB of the 2 GiB caches is out of range, and NQK=7/8 passed. The failure tracks cache size, not the NQK ratio. ### Changes 1. **Harness fix.** Every direct-launch harness that builds `UnifiedAttentionProblem` and launches without `run_unified_attention_torch` now passes `num_kv_blocks=key_cache.shape[0]`: - `builders/gfx950/attention/prefill/parity_unified_attention.py` - `builders/gfx942/attention/prefill/parity_unified_attention.py` - `builders/gfx950/attention/decode/fp8_decode_3d_verify.py` - `builders/gfx1250/attention/decode_3d_verify.py` - `benchmarks/gfx950/attention/decode/fp8_decode_nseg_sweep.py` - `benchmarks/gfx{942,950}/attention/prefill/benchmark_prefill2d_live.py` - `benchmarks/gfx950/attention/prefill/benchmark_prefill2d_traces.py` - `benchmarks/gfx950/attention/prefill/benchmark_prefill2d_fastkv_regp.py` The benchmarks that build specs by hand also set `use_i64_kv_addr` from `_enable_i64_kv_addr(problem)`. They now key their launcher cache on that flag, because `kernel_name()` doesn't encode it. 2. **Launch-time guard.** `_attn_values` takes a required `kv_addr_limit: _KvAddrLimit | None`. It holds the largest K cache, in bytes, that the compiled kernel reads correctly, and what sets that bound. It is `None` for i64 addressing with no i32-only loader. `_check_kv_addr_width` compares the real cache (`k.shape[0]` blocks) against that limit and raises `ValueError` instead of reading zeros or wrong data. Production 2D and scalar launches and the torch direct-launch harnesses pack kernargs through `_attn_values`. The 3D path and the gfx1250 `tiled_2d_verify.py` harness pack their own and are not checked. The limit is what the kernel does, not what the problem decided: | Kernel | Limit | |---|---| | gfx942/gfx950 tiled 2D, i64 | none | | gfx942 tiled 2D, i64 with a flat-load K/V loader that has no i64 branch (cfv, cfvst, vhbm) | 2³¹ bytes | | gfx942/gfx950 tiled 2D, i64 with a buffer-load K/V loader that has no i64 branch (gfx942 khbm; fp8-PV on both) | `0x7FFF0000` bytes | | gfx942/gfx950 tiled 2D, i32 | `0x7FFF0000` bytes (buffer `num_records`) | | scalar 2D, gfx942 4-warp GQA, gfx1250 tiled 2D (no i64 path, i32 element offsets) | 2³¹ elements: 4 GiB bf16/fp16, 2 GiB fp8 | `_spec_kv_addr_limit(problem, spec, builder_kind=...)` computes it from the compiled spec. `_kv_addr_limit(problem, use_i64_kv_addr)` gives only the limit implied by the address width, and does not see the i32-only loaders. Where each launch gets it: - **Production tiled 2D:** carried in `_Attention2DLaunchMeta`. It is read off the same spec `_get_2d_launcher` compiles and memoized per cache key, so a builder branch that drops `use_i64_kv_addr` fails safe. The gfx942 4-warp route and gfx1250 (whose spec has no `use_i64_kv_addr` field) get the element limit. - **Explicit tuning spec:** from the bound `kernel_spec` and `builder_kind`. - **Scalar fallback:** the element limit. Production fills `num_kv_blocks` before this point, so the old problem-based check passed here while the kernel wrapped. - **Hand-built harness specs:** `_spec_kv_addr_limit(problem, spec)`. The error hint follows the cause: - **i32 spec:** a missing or stale count points at the problem. A correct count says to rebuild a hand-built spec as i64. - **No i64 path:** fixing `num_kv_blocks` won't help, so use a smaller cache. - **i32-only loader:** the hint names the loader flag (e.g. `use_conflict_free_v_store`). It suggests a smaller cache, or turning the flag off on a hand-built spec. Overcounting is allowed, since an i64 kernel on a small cache is still correct. `run_unified_attention_torch` keeps `num_kv_blocks=0` for small caches, so the fast path never pays for a per-call `dataclasses.replace`. That's why `0` stays a valid value instead of becoming a required field. The 3D split-KV path packs its own kernargs and isn't covered by the guard. Its three direct callers pass right-sized `num_kv_blocks`. 3. **i64 threshold moved to `> 0x7FFF0000` bytes.** It was `> 2³¹`, which left a 2 GiB cache on i32 with its top 64 KiB reading zeros. One helper, `kv_cache_needs_i64_addr`, now owns the threshold. `_enable_i64_kv_addr`, `run_unified_attention_torch` and the dispatcher's tuning specs (`AttentionTuningSpec.with_num_kv_blocks`, `_semantic_fields`) all use it. The C++ selector copy in `attention_unified_selectors.cpp` is updated to match. Only caches within 64 KiB of 2 GiB change path. 4. **The gfx942 4-warp GQA builder has one i64 rule.** `with_num_kv_blocks` used to set the flag on a 4-warp tuning spec, and the builder silently emitted i32 element addressing anyway. `kv_addr_width_for_builder` now decides the width for every tuning entry point (`with_num_kv_blocks`, and `make_explicit_attention_2d_spec` through `_semantic_fields`). A builder with no i64 path stays i32 up to 2³¹ elements, so a 4-warp spec on a 2-4 GiB bf16 cache still works. It raises `NotImplementedError` only past that. The builder itself still refuses `use_i64_kv_addr=True` as a backstop. Production is unaffected: `_gfx942_4warp_eligible` already sends i64 problems to the default builder. 5. **i64 specs with i32-only K/V loaders keep the i32 limit.** Under `use_i64_kv_addr`, some loaders still form i32 offsets from the cache base: the gfx942 cfv, cfvst, vhbm and khbm loaders, and the fp8-PV V loaders on gfx942 and gfx950. `_I32_ONLY_KV_LOADERS` maps each flag to what its loader still reads: 2³¹ bytes for flat loads (cfv, cfvst, vhbm), `0x7FFF0000` for buffer loads (khbm, fp8-PV). A spec gets the tightest limit among the flags it sets. The check is conservative: if the builder declines a flagged loader for some geometry, the flag still caps the cache. On gfx942, a cfvst spec on a 2.125 GiB cache faults the GPU at `06be846`; it now raises before launch. Giving these loaders an i64 branch changes emitted IR and is left for a separate PR. 6. **Docs.** The paged-cache section of `builders/gfx950/attention/README.md` says direct-launch harnesses must pass `num_kv_blocks` themselves. It also describes the guard and each kernel family's limit. ### Behaviour change Launches that used to read zeros or wrong data, or fault the GPU, now raise `ValueError` at launch: - the scalar 2D fallback on bf16/fp16 caches over 4 GiB or fp8 caches over 2 GiB - gfx1250 tiled 2D on fp8 caches over 2 GiB - gfx942 specs that use cfv, cfvst or vhbm on caches over 2³¹ bytes - khbm or fp8-PV specs (fp8-PV on either arch) over `0x7FFF0000` bytes ### Known limitation cfvst is the default for gfx942 bf16 D64 prefill and fp16 D64/D128 prefill. It is also the only capped loader any shipped gfx942 i64 spec selects. Deployments on those configurations with a per-layer K cache over 2³¹ bytes now get a `ValueError` on every prefill instead of a GPU fault, and dispatch can't route around it yet. This is common for small or fp16 models with vLLM-style cache sizing. The next PR gives `_cfvst_load_v_regs` an `offset_i64` branch, following the fp8 sync loader, which lifts the cap for every shipped configuration. Two more follow-ups: - Exposing the limit to `supports_native_unified_attention*`, so dispatch can route around a capped kernel. - Replacing the hardcoded `0x7FFF0000` K/V `num_records` with one named limit. No emitted IR changes for any existing kernel, so no re-bless. ### Test coverage CPU-only, in `library/tests/test_selector_parity.py`: - `TestProductionRouteGuard` (gfx950) runs `run_unified_attention_torch` with mocked launchers and a shape-only K: - `num_kv_blocks=0` on a 4 GiB cache launches an i64 kernel with the refilled count. - A stale caller-supplied count raises before launch. - A tuning spec is re-bound to the real cache and built with `use_i64_kv_addr=True`. - The scalar fallback launches at exactly 2³¹ elements and raises one block later. - `TestGfx942RouteGuard` runs the same route on gfx942: - the shipped cfvst spec launches at exactly 2³¹ bytes and raises one block over, with the cfvst hint - a 4-warp tuning spec launches as i32 past the buffer range - the graph-replay miss path runs the guard - `TestKvAddrLimit` covers the limit and its cause for each kernel family: - The gfx942 4-warp route and gfx1250 meta get the element limit. - Each i32-only loader flag caps an i64 spec at its own limit. - The meta follows the built spec rather than the problem. - One memo serves i32, cfvst-capped and uncapped kernels correctly. - A second `total_q` on the same cache key reuses the built spec. - Every bool `use_*` field on the gfx942 and gfx950 tiled specs must be classified as an i32-only K/V loader, i64-safe, or not K/V addressing, so a new loader flag can't ship uncapped by accident. - `TestFourWarpTuningSpec` covers the 4-warp builder rule: i32 up to the element limit, refused past it. It covers `with_num_kv_blocks`, `_semantic_fields`, and a 4-warp config through `make_explicit_attention_2d_spec`. - `TestKvAddrWidthGuard` covers the guard itself: - the buffer-range and fp8 boundaries - the hint for each cause and count - that `_attn_values` requires `kv_addr_limit` - `TestI64KvAddr` cache-key cases are regression guards against dropping the i64 flag from `_tiled_cache_key` or `_tiled_3d_cache_key`. Their boundaries move to 65534/65535 blocks. A new case pins that 65536 blocks (exactly 2 GiB) now gets i64. Another checks, for gfx942 and gfx950, that the spec builders, launch meta and guard agree at every boundary count. `tests/dispatch/attention/test_tuning_space.py` boundaries are updated for the new threshold and the 4-warp element limit. Each guarded path was also mutation-checked. Each of these mutations fails at least one test: - dropping the tiled or scalar check - treating gfx1250 or the 4-warp route as i64-capable - reverting the threshold - passing no limit at the graph-replay site - ignoring the i32-only loader flags - inferring the meta's flag from the problem - dropping `builder_kind` from the tuning-spec meta or from `make_explicit_attention_2d_spec` - keying the limit memo on a constant or without the i64 flag ## Test Plan Current head (cause-aware guard hint and round-3 tests; host Python only, no emitter change, so the GPU and gate results under `9ac3b902` still apply): - [x] `pytest library/tests/test_selector_parity.py`: 83 passed, 10 skipped (the skips are the existing PR #9220 bias-guard cases) - [x] Attention set (`test_selector_parity`, `tests/dispatch/attention`, `test_per_engine_spec_fns`, `test_attention_builds`, `test_library_layering`, `test_gfx942_gqa_head_fold`, `test_arch_binding_guard`): 587 passed, 28 skipped - [x] Full `library/tests` (CPU), run locally since CI doesn't run the rocKE suite: 2428 passed, 775 skipped. The only failure is the unrelated `test_conv_dgrad_correctness::test_fp32_stride1`, where the torch MIOpen reference fails on the test host. - [x] Mutation check: a constant memo key, a memo key without the i64 flag, and `builder_kind` dropped from `make_explicit_attention_2d_spec` each fail a test Previous head (`9ac3b902`, second round of review fixes): - [x] `pytest library/tests/test_selector_parity.py`: 77 passed, 10 skipped - [x] Full `library/tests` (CPU): 1984 passed, 194 skipped. The only failures are the two unrelated ones below. - `test_conv_dgrad_correctness`, where the torch MIOpen reference fails on the test host. - `test_attention_dense_gfx942_golden`, whose llvm23 entries had been stale on `develop` since #12672 (re-blessed in #13226). - [x] Mutation check of each new guarded path (listed under Test coverage) - [x] gfx942 GPU tests `test_attn_bf16_d128_ring.py`, `test_gfx942_gqa_head_fold_numeric.py`: 40 passed - [x] gfx942 parity (default group), torch reference: 22 PASS / 6 SKIP / 0 FAIL - [x] gfx942 boundary check through `run_unified_attention_torch`. Setup: bf16 D64 GQA 32/8 causal prefill on cfvst, block table on the top blocks, fp32 reference. The i64 control is bf16 D64 MHA with no i32-only loader on a 2.125 GiB cache. - this head: 65534, 65535 and 65536 blocks (exactly 2³¹ bytes) correct, 65537 and 2.125 GiB raise `ValueError`, i64 control correct - `06be846`: 65534 and 65535 correct, 2.125 GiB faults the GPU, i64 control correct - [x] Byte-identity gate, all families, fresh engine build: GREEN at llvm20, llvm22, llvm23 - [x] gfx942 at 65535 blocks: covered by the boundary check above (the parity harness right-sizes its caches). NQK=3/5/7/8 shapes pass in the default group - [x] gfx950 forced-2D and auto parity, NQK=3/5/7 (4 GiB and 2 GiB caches), torch reference, on the merged branch: `ck-2d` and `ck-auto` at `max_abs=0.0156` on all three ``` python builders/gfx950/attention/prefill/parity_unified_attention.py --paths 2d,auto --skip-triton \ --scenario gqa_nqk3_bf16_d128_b64 --scenario gqa_nqk5_bf16_d128_b64 \ --scenario gqa_nqk7_bf16_d128_b64 ``` Earlier head (`06be846b`, first round of review fixes): - [x] `pytest library/tests/test_selector_parity.py library/tests/dispatch/attention library/tests/test_per_engine_spec_fns.py library/tests/test_attention_builds.py library/tests/test_library_layering.py library/tests/test_gfx942_gqa_head_fold.py library/tests/test_arch_binding_guard.py`: 535 passed, 25 skipped - [x] Full `library/tests` (CPU). The only failures are unrelated to this PR: - 8 × `test_conv_dgrad_correctness`, where the torch MIOpen reference raises `miopenStatusInternalError` on the test host. - `test_attention_dense_gfx942_golden`, whose llvm23 entries had been stale on `develop` since #12672 changed the llvm23 datalayout. - [x] `rocke_ir_parity_harness.py --check tests/golden/rocke_representative_ir_sha256.json`: OK at llvm20, llvm22, llvm23 - [x] Byte-identity gate, `gfx942_attention_tiled_3d`: GREEN (5/5 configs) at llvm20, llvm22, llvm23 - [x] gfx942 parity (default group), torch reference: 22 PASS / 6 SKIP / 0 FAIL Earlier commits (harness fix and first version of the guard): - [x] gfx950 forced-2D parity, NQK=3/5/7/8, before and after the harness fix: ``` python builders/gfx950/attention/prefill/parity_unified_attention.py --paths 2d \ --scenario gqa_nqk3_bf16_d128_b64 --scenario gqa_nqk5_bf16_d128_b64 \ --scenario gqa_nqk7_bf16_d128_b64 --scenario gqa_nqk8_bf16_d128_b64 ``` - [x] gfx950 parity (`--skip-triton --paths auto,2d`), torch reference: - NQK=3/5/7/8: `ck-auto` and `ck-2d` both at `max_abs=0.0156` - full default set (27 scenarios): every lane within tolerance and the guard never fired - [x] gfx942 parity (`--scenario all`), torch reference: 94 PASS / 13 SKIP (unsupported on gfx942) / 0 FAIL - [x] gfx942 `benchmark_prefill2d_live.py` vs live AITER Triton: bf16 and fp16, a small cache plus a 4 GiB K cache (`--cap-blocks 131072`). `prod` and `fallback` correct on all 4 shapes. ## Test Result Before the fix: | Scenario | `max_abs` | `mean_abs` | Bad positions | Zero positions | |---|---|---|---|---| | NQK=3 (4 GiB cache) | 3.516 | 0.0723 | 1536 / 1536 | — | | NQK=5 (4 GiB cache) | 3.250 | 0.0493 | 1407 / 1536 | 64 / 1536 | | NQK=8 (2 GiB cache, control) | 0.0156 | 3.4e-4 | 0 / 1536 | 0 / 1536 | After the fix: | Scenario | `max_abs` | `mean_abs` | Zero positions | |---|---|---|---| | NQK=3 | 0.0156 | 3.4e-4 | 0 / 1536 | | NQK=5 | 0.0156 | 3.3e-4 | 0 / 1536 | | NQK=7 | 0.0156 | 3.4e-4 | 0 / 1536 | | NQK=8 | 0.0156 | 3.4e-4 | 0 / 1536 | ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Claude <noreply@anthropic.com>

github-actions Bot added project: hip-kernel-provider rocke: platform Touches rocKE platform labels Sep 28, 2026

CongMa13 marked this pull request as ready for review September 28, 2026 22:36

CongMa13 requested review from a team as code owners September 28, 2026 22:36

assistant-librarian Bot added the organization: ROCm label Sep 28, 2026

tenpercent approved these changes Sep 28, 2026

View reviewed changes

Comment thread dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py Resolved

Comment thread dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py Resolved

tenpercent mentioned this pull request Sep 28, 2026

refactor(rocke): query LLVM version from loaded COMGR #12678

Merged

7 of 9 tasks

CongMa13 commented Sep 29, 2026

View reviewed changes

Comment thread dnn-providers/hip-kernel-provider/rocke/platform/tests/instances/test_rocke_multiarch.py Resolved

Comment thread dnn-providers/hip-kernel-provider/rocke/platform/python/rocke/core/lower_llvm.py Resolved

cgmillette approved these changes Sep 29, 2026

View reviewed changes

CongMa13 and others added 3 commits September 29, 2026 11:41

Merge branch 'develop' into users/congma/ck/aick-2416-llvm23-datalayo…

7418063

…ut-p10-p15

CongMa13 enabled auto-merge (squash) September 30, 2026 14:38

CongMa13 and others added 4 commits October 1, 2026 19:15

Merge remote-tracking branch 'origin/develop' into users/congma/ck/ai…

73673be

…ck-2416-llvm23-datalayout-p10-p15

CongMa13 merged commit 94b5013 into develop Oct 6, 2026
106 checks passed

CongMa13 deleted the users/congma/ck/aick-2416-llvm23-datalayout-p10-p15 branch October 6, 2026 01:10

This was referenced Oct 7, 2026

fix(rocke/attention): pass num_kv_blocks in direct-launch harnesses and guard i32 KV addressing at launch #12446

Merged

feat(rocke/attention): expand golden flavor coverage for KDA and GDN decode #13226

Merged

This was referenced Oct 7, 2026

feat(rocke): enable mixed scaled GEMM on shared operand storage #12279

Merged

feat(rocke): add gfx942 bf16 MLA prefill kernel and dispatch #12464

Merged

feat(rocke): Lightning Indexer for DSA #12548

Merged

cgmillette mentioned this pull request Oct 9, 2026

feat(rocke): prototype device-side printing #10841

Draft

9 tasks done

shumway mentioned this pull request Oct 13, 2026

rocke: make portable-IR local tests reject or avoid stale engine artifacts #13543

Open

8 tasks

Navigation Menu

Repository navigation

fix(rocke): add address spaces p10-p15 to the LLVM 23 datalayout - #12672

fix(rocke): add address spaces p10-p15 to the LLVM 23 datalayout#12672
CongMa13 merged 9 commits into
developfrom
users/congma/ck/aick-2416-llvm23-datalayout-p10-p15

CongMa13 commented Sep 28, 2026 •

edited by tenpercent

Loading

Uh oh!

therock-pr-bot Bot commented Sep 28, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 28, 2026 •

edited

Loading

Uh oh!

tenpercent left a comment

Uh oh!

Uh oh!

Uh oh!

CongMa13 left a comment

Uh oh!

Uh oh!

Uh oh!

cgmillette left a comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 7, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

CongMa13 commented Sep 28, 2026 • edited by tenpercent Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

fix(rocke): add address spaces p10-p15 to the LLVM 23 datalayout

Summary

Why it surfaced only now

What changed

Why the superset is safe on older toolchains

Verification

Reviewer notes

Test plan

Uh oh!

therock-pr-bot Bot commented Sep 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Sep 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

tenpercent left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

CongMa13 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

cgmillette left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 7, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

CongMa13 commented Sep 28, 2026 •

edited by tenpercent

Loading

therock-pr-bot Bot commented Sep 28, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 28, 2026 •

edited

Loading