Repository navigation
fix(rocke): wave_all is never true on wave32 - #13896
Merged
Merged
Conversation
ballot.i64 sets only the low wave_size bits, so comparing it against -1 can only succeed on wave64. Compare against the full-wave mask (0xffffffff on wave32, -1 on wave64) in both the Python and C LLVM lowerings; wave64 output is unchanged. On gfx11 this re-enables the lazy-rescale skip in the transposed-QK attention path (its 16 parity anchors plus 2 replays are re-blessed). Replaces the IR-text test with an on-GPU wave_all/wave_any test that runs a full wave on the detected device.
github-actions
Bot
added
documentation
project: hip-kernel-provider
rocke: platform
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Add a host-only lowering test and a C++ ctest case that pin the wave_all full-wave mask (-1 on wave64, 4294967295 on wave32), and correct the wave_all/wave_ballot docstrings and helpers README.
dpeabody-amd
enabled auto-merge (squash)
October 18, 2026 22:25
dpeabody-amd
disabled auto-merge
October 18, 2026 22:26
cgmillette
reviewed
Oct 19, 2026
cgmillette
left a comment
cgmillette
left a comment
Contributor
There was a problem hiding this comment.
PR #13896 — wave_all is never true on wave32 (dpeabody-amd, AICK-1847) — 🟢 ACCEPT: correct, CI-verified wave32 fix
| Prior item | Status |
|---|---|
| Reviewers | 🟢 None yet: no reviews or threads, only the policy bot |
| Bots / CI | 🟢 Policy bot pass. Multi-Arch CI at head: hipkernelprovider green on gfx94X (wave64) and on gfx110X / gfx120X (wave32). On Linux gfx120X the 10 Python wave-vote GPU cases PASSED; the cpp cases skip because CI has no binding. The gfx120X summary is red only because of other components (rocprofiler-sdk, hipthreads, miopen) |
| Our checks | 🟢 Clean merge with develop (+25): same 32 platform failures as develop. Differential sweep at llvm20 over 333 configs in 3 wave32 families: only the 18 re-blessed gfx1151_wmma_fmha_fwd configs change, and all match. Develop's new fmha_fwd_v2 / decode families are unchanged because they vote with wave_any. Representative goldens ×3 = develop's. IR guards exist in both engines, and the native ctest is in the tier YAML. Binary caches key on the LLVM text. Rules 1–3 clean (ran) |
| # | Sev | file:line | Finding |
|---|---|---|---|
| N1 | 🟡 | lower_llvm.py:4405, wave_and_cross_lane.md:55 |
The new comment says HIP's __all makes the same full-EXEC assumption, but __all counts only active lanes; rocKE's own HIP lowerer comment says so too. Under partial EXEC (a divergent region or a sub-wave workgroup), LLVM wave_all returns 0 while HIP returns 1 → state the full-wave precondition in the user docs and the ir.py docstring, or vote ballot(p == 0) == 0 on wave32, which needs neither the full-wave mask nor full EXEC |
Utility: 🟢 Positive — real wave32 bug fixed in both engines, minimal, wave64 IR unchanged.
Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good
Contributor
Author
|
Thanks, N1 addressed in 1d602ff (docs/comments only, no IR change). You're right that HIP's I kept the |
cgmillette
approved these changes
Oct 20, 2026
cgmillette
left a comment
cgmillette
left a comment
Contributor
There was a problem hiding this comment.
I re-reviewed #13896, and it's still an ACCEPT. The nit from our 10-19 review is fixed, and the merge with today's develop is clean.
PR #13896 — wave_all is never true on wave32 (dpeabody-amd, AICK-1847) — 🟢 ACCEPT: nit fixed, merges clean
| Prior item | Status |
|---|---|
Ours (10-19) N1: the comment claims HIP's __all makes the same full-EXEC assumption |
🟢 FIXED 1d602ff4. The lowering comment (lower_llvm.py:4405), the IRBuilder.wave_all docstring (ir.py:2975) and wave_and_cross_lane.md:62 now say __all counts only active lanes, and that wave_all needs full EXEC and returns 0 otherwise. The C++ comment (crosslane.cpp:168) never made the claim (read) |
Author's question (10-19): switch to active-lane ballot(p == 0) == 0? |
🟢 Answered: no. Keeping the full-wave mask leaves wave64 IR and goldens untouched, and the precondition is now documented where callers will see it |
| Bots / CI | 🟢 Policy bot and pre-commit pass. Multi-Arch at head: hipkernelprovider is green on gfx94X (wave64) and Windows gfx110X (wave32), and both hipDNN superbuilds are green. No one has approved yet; our 10-19 review was a COMMENT |
| Our checks | 🟢 Clean merge with develop 061b772a (+72). A sweep of all 11 wave32 families at llvm20 (500 configs) changes only the 18 re-blessed gfx1151_wmma_fmha_fwd anchors, and all of them match. Develop's 30 wmma_gemm drifts are identical on both trees (ran: simulated merge). The wave-vote numerics passed 20/20 on our gfx1151 with both engines, including the 10 C++-engine cases CI never runs (ran). The native rocke_future_intrinsic_lowering test, which pins 4294967295 on wave32, passes (ran). Rules 1–3 are clean (read) |
Net-new: none.
Utility: 🟢 Positive — real wave32 bug fixed in both engines; C++ path now hardware-verified.
Legend: 🔴 blocking · 🟠 should-fix · 🟡 minor nit · 🟢 info/good
cgmillette
added a commit
that referenced
this pull request
Oct 20, 2026
The differential golden had gone stale: 627 of its 1240 LLVM IR entries and 537 of its 1206 canonical IR entries no longer matched the tree, because nothing ran the check. Recorded afresh in the new format, with the Python reference emitters, from a tree whose two-engine gate is green for all 82 families: - ll at llvm20, llvm22 and llvm23: 82 families, 1518 configs per flavor; 1615 at llvm23, where gfx1250_lowering has configs that need LLVM 23 and are refused below it. - ir_canonical: 81 families, 1598 configs (one family's emitter has no IR mode). Canonical IR was checked to be identical at llvm20 and llvm23. What the re-bless absorbs: every mismatch that could be traced is an intended code generation change that did not re-bless this file -- the shared LDS pool (#8844, about 425 configs), the ds.swizzle declaration (#9583, about 177), renamed values in the attention transpose reads (#9133, 57), the smaller LDS pool (#10112, about 39), the flat work group size of one-wave wave32 kernels (#12710, about 39), #9220 (6) and the conv changes of #13109 (6). 441 entries were already stale at the start of the history available here and cannot be attributed; no unexplained change was found after it. The two repeated keys (conv_direct_grouped, conv_wgrad_workspace_reduce), which hid their newer copies, are gone. It also takes in the code generation changes merged since then, each traced by recording the file at the commit and at its parent: - the magic division sum emitted without signed-overflow poison (#13802): 215 configs in each of the four sections, in 19 families -- attention_unified 6, attention_unified_stress 32, conv_direct_grouped 16, conv_implicit_gemm 26, conv_implicit_gemm_dgrad 31, conv_implicit_gemm_wgrad 42, deep_fused_conv_pool 2, gfx1201_deep_fused_conv_pool 5, gfx950_deep_fused_conv_pool 5, dynamic_helpers 2, fmha_splitkv_decode 6, img2col 5, moe_fused_mega 5, moe_gemm_fused 6, moe_sorting 6, pooling 4, sage_attention 4, sparse_attention 6, transpose 6; - attention decoding on gfx942 and gfx950 (#13492): 23 changed and 25 new configs per section in attention_strided_kv, gfx942_attention_tiled_3d and gfx950_attention_tiled_3d; - wave_all on wave32 (#13896): 18 gfx1151_wmma_fmha_fwd configs per LLVM flavor, canonical IR unchanged; - new configs: the barrier_fence family (#13575, 10) and one conv_implicit_gemm_dgrad config (#13782). Co-Authored-By: Claude <noreply@anthropic.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
ISSUE ID : AICK-1847
Summary
IRBuilder.wave_alllowers tollvm.amdgcn.ballot.i64and compares the result against-1. On wave32 the ballot only sets the low 32 bits, so the comparison can never succeed andwave_allis always false on gfx11/gfx12. The fix compares against the full-wave mask (0xffffffffon wave32,-1on wave64) in both the Python (rocke/core/lower_llvm.py) and C++ (cpp/core/lower_llvm/crosslane.cpp) lowerings, and documents the semantics indsl_docs/primitives/wave_and_cross_lane.md.This was found during the RDNA SDPA work (AICK-1847).
Impact
wave_allnow take their fast path. In the current catalog this is the lazy-rescale skip in the gfx11 transposed-QKwmma_fmha_fwdattention path.Golden change
18
llanchors ofgfx1151_wmma_fmha_fwdintests/instances/differential/golden/llvm_gfx_all.jsonare re-blessed (16 transposed-QK parity configs plus 2 replays). These are exactly the configs whose emitted.lldiffers from develop; no other anchor changes, and their --canonicalanchors are untouched.Tests
The old IR-text assertion in
test_rocke.pyis replaced bytests/numeric/test_wave_vote_numeric.py, which runswave_all/wave_anyover one full wave on the detected GPU (Python and C++ engine backends; all-true, non-zero-true, first-lane-false, last-lane-false, all-false).python tools/check_byte_identity.pywithROCKE_LLVM_FLAVOR=llvm20andllvm22: GREEN (C++ engine.llbyte-identical to Python).python tests/instances/differential/run_diff.py --mode ll --check-goldenand--mode ir --canonical --check-golden(llvm20): unblessed counts are identical to develop's baseline; the only emitted differences vs develop are the 18 re-blessed anchors, and each matches the new golden.rocke_arch_target_identity+pytest tests/core/test_arch_target_identity.py: pass.pytest tests/numeric/test_wave_vote_numeric.pywith therocke_enginebinding built (-DROCKE_BUILD_PYBIND=ON): 20 passed on gfx1100 (wave32) and 20 passed on gfx942 (wave64).pytest tests/test_rocke.py: same results as develop (no new failures).pre-commit run --files <changed files>: pass.