Repository navigation
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ 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. |
This was referenced Sep 18, 2026
Contributor
Author
|
Suggestions passed on by this PR for general consideration so closing now |
tony-davis
added a commit
that referenced
this pull request
Sep 25, 2026
) JIRA ID : ROCM-31230 ## Motivation On gfx1201 an fp32 NN GEMM with a large leading dimension either page-faults or returns a wrong answer while reporting success. `hipblasLtMatmul` gives `HSA_STATUS_ERROR_MEMORY_FAULT` at `lda = 67,108,864` with `K = 32`, and at `lda = 134,217,727` it returns success with a norm error of 0.911 against a 1e-06 tolerance. The trigger is a byte count that outgrows a signed 32-bit register. `GlobalReadIncs<tc>+unrollIdx` is one SGPR holding `stride * DepthU * bpeGR`, the number of bytes one unroll iteration advances the global-read pointer. `calculateStagger` widens it to 64 bits before it reaches the global-read buffer descriptor and used `s_mul_i64_i32`, a signed widen. Once `stride * DepthU * bpeGR` reaches `2^31` the register's bit 31 is set, the signed widen sign-extends it to a negative 64-bit value, and `incrementSrd` moves the read base backwards by about 2 GB instead of forwards. Where that displaced base is unmapped the kernel faults; where it is mapped the kernel reads other memory and returns a wrong result with no error. ## What is actually being completed here Each stagger function widens a **pair** of quantities per tensor: the stagger byte offset, `StaggerUIter * GlobalReadIncs`, and the unroll-loop span, `LoopCounter * GlobalReadIncs`, which becomes `WrapU`. There are four such pairs in tensilelite, and two of them were already fixed. This PR does the remaining two. | | main increment | sparse-metadata increment | |---|---|---| | `removeStagger` | fixed Dec 2024, `9446f2356a3` | fixed Dec 2024, `9446f2356a3` | | `calculateStagger` | **this PR** | **this PR** | The history explains why this took so long, and it is not carelessness. Tensile [PR 1672](ROCm/Tensile#1672) (February 2023) made exactly this change to Tensile's `calculateStagger`, and fixed `removeStagger` at the same time. Tensile has **no sparse-metadata arm at all**, so there was never anything there to port for that half. tensilelite's metadata stagger is its own, added in October 2022 (`bdb5085ec21`) using the signed idiom then in use, so it was born with the defect. tensilelite later received the `removeStagger` half of the Tensile fix in December 2024, but not the `calculateStagger` half. Every fix since 2022 happened somewhere that could not see all four pairs at once. ## Technical Details All four sites multiply operands that are non-negative by construction: `StaggerUIter` is an iteration mask and `LoopCounter` an iteration count, while `GlobalReadIncs` and `GlobalReadIncsMetadata` are byte strides. They now use `s_mul_u64_u32`, which the file already provides. `WrapU` still becomes negative in the subtraction that follows, and the borrowing `s_sub_u32` / `s_subb_u32` pair there already handles that, so only the intermediate product changes sign treatment. The widen is selected per kernel rather than applied unconditionally. `GSU.graIncrements` negates the increment when the unroll dimension is mirrored, so that operand is genuinely signed there and keeps `s_mul_i64_i32`. For A/B the predicate covers every tensor char `calculateStagger` receives: `A`, `B`, `Metadata`, `MXSA` and `MXSB`, the last two because the function is also called for the MX scale tensors. The sparse arm needs its own predicate against `MirrorDimsMetadata` rather than reusing `widenIncs`, because the block reassigns `tc` to `"Metadata"` internally and the outer predicate would be evaluated for the wrong tensor. Emitted instruction count is unchanged. `s_mul_hi_i32` becomes `s_mul_hi_u32`; the low-word `s_mul_i32` is identical. There is no register-pressure or scheduling change. The defect is kernel-specific rather than shape-specific, which is why it survived in the field. Reaching the multiply needs a non-zero `StaggerUIter`, and two things independently zero it: the loop-count clamp in `calculateStagger` drops the stagger for larger `StaggerUStrideShift` at low iteration counts, and `StaggerUMapping` picks which workgroup dimension feeds the mask, so a problem with a single workgroup in that dimension has no stagger either. Of 14 gfx1201 solutions measured at the boundary, 3 were affected. ## Test Plan Three tests, at two levels. Results and full evidence are in a [separate comment](#12332 (comment)). **Codegen, CPU-only, `-m unit` lane.** `test_r3_stagger_incs_unsigned_gfx1201_char.py` covers the A/B pair and `test_r3_stagger_incs_unsigned_sparse_gfx942_char.py` covers the metadata pair. Both assert on emitted assembly rather than source, because the signedness is only observable in the instruction: `SMulInt64to32` selects `s_mul_hi_i32` or `s_mul_hi_u32` from its `sign` argument and both carry the same comment into the listing. The gfx1201 test adds a third assertion tying the registers the stagger multiply writes to the ones the SRD increment reads, scoped to a single `addr += (StaggerUIter) * GlobalReadIncs<tc>` block, so the other two cannot pass vacuously. The sparse case needs its own config, `sparse_stagger.yaml`, because the arm is gated on `Sparse >= 1` and the existing `sparse_dtvsm.yaml` sets `StaggerU: 0`, which compiles the stagger path out; gfx942 rather than gfx1201 because SMFMA sparse shapes are gfx9. **Hardware, gfx1201.** `matmul_size_t_stagger_incs_overflow_f32_ROCM31230` in `clients/tests/data/matmul_gtest.yaml`: fp32 NN, `M=2048, N=1, K=32, lda=67108864`, `integer_exact` with `unit_check`, gated `gpu_arch: '1201'`. That `lda` puts `lda * DepthU * bpe` on `2^31` exactly for the affected `DepthU=8` kernels, one representable step above the ticket's passing case. It runs every supported solution (`algo_method: [1]` with `requested_solution_num: -1`) rather than one heuristic pick, so a future retune cannot silently move selection off the affected kernel. Category is `nightly` because A is 8.6 GiB on host and device and cannot be made smaller, the A span being `(K-1)*lda + M` whatever `M` is. ## Test Result Green at every level, with the detail in the [evidence comment](#12332 (comment)). In summary: both codegen tests fail before their respective change and pass after, each confirmed by negative control; the hardware gate fails on an unpatched library and passes on a patched one, three runs each; and the patched generator's output ran on four architectures in math-ci with zero failures. ## Risk level High overall, but the two halves differ and should not be averaged. **The A/B pair is the high-risk part.** It changes generated code for every kernel that compiles the stagger path, on every architecture, not only the gfx1201 kernels where the defect was found. `KernelWriterAssembly.py` offers no isolation smaller than the whole library: no feature flag, no per-arch gate. In the single gfx1201 fp32 family rebuilt for verification, 44 instructions changed; across all families and architectures it is considerably more. **The sparse-metadata pair is provably inert.** All 40,302 `Sparse:` declarations across the 16 architecture logic directories are `Sparse: 0`, so no tuned solution reaches that arm on any architecture. Rebuilding the gfx1201 fp32 family with and without that hunk gives a byte-identical code object. It is included for correctness and symmetry, not because anything ships through it. What bounds the A/B risk: the change is confined to the signedness of a widening of a quantity that cannot be negative, and is selected per kernel so mirrored unroll dimensions keep their previous behaviour; emitted instruction count and instruction addresses are unchanged, so there is no code motion, register-pressure or scheduling change; no shipped solution declares mirrored dims (0 of 537,486), so the conditional resolves to unsigned everywhere in practice today; and runtime coverage is cross-architecture rather than only the reporting arch. The residual risk is performance rather than correctness, on kernels and architectures not represented in those suites. The instruction-for-instruction substitution argues against any, but nothing here measures throughput. ## Related - Tensile [PR 1672](ROCm/Tensile#1672), the same change in Tensile, plus the `removeStagger` half that tensilelite already has. - PR #12212 reports the same defect and takes a different approach; see the comparison on ROCM-31230. - Issue #12337 covers two pre-existing CI problems found while adding the hardware gate, neither caused by this PR. - Still open on ROCM-31230: the mirrored path, which needs a synthetic config since no shipped solution mirrors, and which stays wrong past `2^31` in either signedness because the magnitude is unchanged. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Cursor <cursoragent@cursor.com>
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.
Motivation
This pull request introduces improvements to host memory allocation checks for large matrices in the test suite, and enhances the handling of 64-bit shadow buffer limits in the Tensile assembly code generation. The changes help avoid unnecessary host memory usage during benchmarking, improve error handling for memory allocation failures, and ensure correct register allocation and address calculations when 64-bit addressing is enabled.
Host memory allocation and test improvements:
CHECK_HOST_ALLOCATIONmacro to check host buffer allocations and handle failures gracefully, with different behaviors for test and bench builds (hipblaslt_test.hpp,utility.hpp).testing_matmul_with_biasto only allocate large host buffers when required for verification or special features, reducing memory usage for benchmark runs.matmul_common.yamlto exercise boundary conditions for host/device memory allocation.Tensile assembly code generation and 64-bit addressing:
use64bShadowLimitStoreflag and logic to support 64-bit shadow buffer limits for store operations, and adjusted register allocation for global read offsets when 64-bit addressing is enabled (KernelWriter.py).AsmAddressCalculation.py).KernelWriterAssembly.py).JIRA ID : ROCM-31230
Can be used or not as you see fit