Repository navigation
fix(tensilelite): stop offering solutions whose output stores are dropped - #12281
Merged
tony-davis merged 5 commits intoSep 25, 2026
Merged
tony-davis merged 5 commits into
tony-davis merged 5 commits into
Conversation
…pped A kernel that writes its output matrix D with buffer stores addresses it through a buffer resource descriptor. That descriptor's num_records field is 32 bits wide and bounds every store: the hardware discards a store at or past it and raises no fault. allocPostLoopSrd in KernelWriterAssembly.py sets that field to the BufferOOB sentinel, 0xfffff000. BufferStoreOffsetLimitCheck is the host-side check that keeps solution selection away from shapes a kernel cannot write in full, but it compared the required extent against 2^32 instead. Any shape whose extent landed between 0xfffff000 and 2^32 was reported as supported and then only partly written, with hipblasLtMatmul returning success. Measured on an MI350X against the ROCm 7.13 library at M = 8,388,607, N = 256, K = 48 in bfloat16: 22 of the 64 algorithms the heuristic offered returned success while leaving part of D unwritten. Every one of them had MacroTile1 = 256, which for this shape makes a workgroup's store extent span the whole tensor. The unwritten elements were the trailing partial tile of each solution's own MacroTile0, matching the pattern reported in ROCm/hipBLASLt#2299. Compare against the sentinel the generator emits, and name the constant so the two can be seen to agree. With the same build and only this threshold changed, the reproducing shape goes from 22 broken of 64 to 0, while a control one step below the sentinel keeps all 1629 of its candidates. The hand-written kernels under Tensile/CustomKernels still set BufferOOB to 0x80000000, so this single global threshold remains too permissive for them. That is recorded in the comment and tracked as ROCM-31258. debugEval printed strides()[1] * value, skipping the min() against N that operator() applies, so it could report a larger extent than the one that decided the result. It now prints the value actually compared. ROCM-31016 Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Policy Checks Passed
📖 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. |
tony-davis
requested review from
nakajee and
talumbau
and
a lite review from Copilot
September 17, 2026 23:42
tony-davis
marked this pull request as ready for review
September 17, 2026 23:42
tony-davis
marked this pull request as draft
September 17, 2026 23:42
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
The updated extent computation should be made overflow-safe and the debug string should avoid hard-coded threshold literals to prevent future divergence.
Get a fresh assessment by requesting another Copilot review.
Pull request overview
This PR fixes a host-side predicate in TensileLite (hipBLASLt) so solution selection rejects GEMM problems whose output stores would be silently dropped by the hardware due to the SRD num_records limit (BufferOOB), and adds C++ unit regression coverage for the reported failure mode.
Changes:
- Update
BufferStoreOffsetLimitCheckto compare the computed per-workgroup D-store extent against the generator’s BufferOOB sentinel (0xfffff000) instead of2^32, and aligndebugEval()output with the compared quantity. - Add predicate unit tests that cover “below / at / above sentinel” boundaries plus a narrow-output case where
min(MacroTile1, N)matters.
File summaries
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/tests/Predicates_test.cpp | Adds regression unit tests for BufferOOB-based rejection/acceptance boundaries and the narrow-output min() behavior. |
| projects/hipblaslt/tensilelite/include/Tensile/ContractionProblemPredicates.hpp | Fixes BufferStoreOffsetLimitCheck threshold constant and makes debugEval() report the exact compared extent. |
Review details
- Files reviewed: 2/2 changed files
- Comments generated: 2
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
debugEval printed the buffer-store limit as a literal 0xfffff000 while operator() compared against BufferOOBBytes. The two would disagree the moment the constant changed, and the debug output is what someone reads when they are working out why a solution was rejected. Format the constant into the string instead. ROCM-31016 Co-authored-by: Cursor <cursoragent@cursor.com>
tony-davis
marked this pull request as ready for review
September 18, 2026 00:49
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (48.80%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #12281 +/- ##
===========================================
+ Coverage 69.94% 71.12% +1.17%
===========================================
Files 2933 2846 -87
Lines 482179 467710 -14469
Branches 69511 69129 -382
===========================================
- Hits 337260 332635 -4625
+ Misses 120713 111123 -9590
+ Partials 24206 23952 -254
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
Contributor
Author
Measured: the predicate excludes exactly the solutions that drop outputThe risk when tightening a selection predicate is over-exclusion, since discarding solutions that would have run correctly costs performance for no benefit. I measured this directly on an MI355X (gfx950) against ROCm 10.2.0a20260918 with the stock tuned library. At the shape this PR is about, M = 8,388,607, N = 256, K = 48, bf16, ldd = M:
The clean count and the untouched count are the same number. Every solution that ran clean is one the predicate still admits, and every solution with That is not a coincidence of this shape. The extent is Method: fill D with NaN, run every algorithm the heuristic offers, count what is still NaN afterwards. A and B are finite and beta = 0, so any NaN left behind is an element the kernel never stored. No reference result or tolerance is involved. Blast radiusThe threshold moves from The ExactlyAtSentinel test is right, and now measured
That reasoning checks out on hardware. M = 8,388,600 gives an extent of exactly So the conservatism costs exactly one One forward-looking noteThis threshold is global across all kernels. That is correct for generated kernels, and too permissive for the 21 hand-written kernels under It would become too restrictive if #12212 lands. That PR introduces |
…ore-sentinel-mismatch
…ore-sentinel-mismatch
Contributor
There was a problem hiding this comment.
Copilot review overview
🟢 Approval recommended
The fix is narrowly scoped to host-side predicate logic, aligns the threshold with the generator’s sentinel, and adds targeted regression tests covering the defect and boundary conditions.
Review effort: Lite
Findings: None
1 task
…ore-sentinel-mismatch
newling
approved these changes
Sep 25, 2026
newling
left a comment
newling
left a comment
Contributor
There was a problem hiding this comment.
This looks good to me. My only optional request is for a numerical test. An AI agent review is here: https://github.com/newling/agent-public-reviews/blob/main/rocm-libraries/review_pr12281.md and it has linked a numerical test it created (runs on MI300X, fails without this fix and passes with it). It is quite a heavy looking test, and I don't know what the memory/runtime footprints are. Perhaps they're unacceptably high, and that's why there isn't such a test in this PR in the first place.
Approving as is, because this simple change fixes a real customer bug.
Contributor
|
Clarification from AI on the end-to-end test it added:
|
nakajee
approved these changes
Sep 25, 2026
nakajee
left a comment
nakajee
left a comment
Contributor
There was a problem hiding this comment.
The current change looks reasonable.
Contributor
Author
|
Thanks @newling! I appreciate the review, and the tests your agent put together are really useful. I'd like to merge this one as is (it fixes a real customer bug) and add tests in a follow-up. The thing is, I think this is a good moment to rethink how we test large outputs in general. I tried an The two CPU predicate tests (the wide output and the |
tony-davis
deleted the
users/todavis_amdeng/rocm-31016-buffer-store-sentinel-mismatch
branch
September 25, 2026 22:12
Contributor
TheRock Submodule Bump ActivityNewest first
|
tony-davis
added a commit
that referenced
this pull request
Sep 28, 2026
JIRA ID : ROCM-31258 ## Motivation hipBLASLt chooses a GPU kernel for each matrix multiplication from a tuned list, and some entries are hand-written kernels stored under `tensilelite/Tensile/CustomKernels`. A kernel writes its output matrix `D` through a buffer resource descriptor, a small structure that bounds the memory its store instructions may touch. The hardware silently discards any store at or past the descriptor's length field (`num_records`), and it raises no fault. Code-generated kernels set that length to `0xfffff000` (commit `ae0ca64a9c8`), and the host admits shapes up to that bound. 120 hand-written kernels still used `0x80000000`, or 2 GiB, so any admitted shape whose store range passed 2 GiB returned success with part of `D` never written. On gfx950 this is reachable through an ordinary fp32 matmul on the xf32 fast path: at `M` = 3,000,000, `N` = 256, `K` = 32, the kernel the heuristic ranks first leaves 14,976 elements of `D` unwritten. ## Technical Details Each affected kernel uses the old value in two roles. It is the length of the `D`, `C` and workspace descriptors, and it is also the store offset given to out-of-bounds lanes so the hardware drops their writes. Both roles must change together: raising only the length would bring those lanes' offsets back in bounds, so they would write past the end of `D`. - Readable kernels (22 files; 12 gfx950, 10 gfx942): `.set BufferOOB, 0x80000000` becomes `0xfffff000`. Nine of the gfx942 files also wrote `0x80000000` into the descriptor directly; those lines now use `BufferOOB`. - Raw-encoded kernels (98 gfx942 files, stored as `.long` instruction words): only the 32-bit literal of each descriptor-length `s_mov_b32` and each out-of-bounds `v_mov_b32` changes to `0xFFFFF000`. - Reassembling all 120 files before and after shows the same instruction count and addresses. The only differences are 277 `s_mov_b32` and 3,383 `v_mov_b32` literals going from `0x80000000` to `0xfffff000`. - The new value leaves at least 4 KiB of headroom below 2^32. Every out-of-bounds store in these kernels uses a scalar offset of 0 and an instruction offset of at most 512 bytes, so a discarded lane's address cannot wrap back into bounds. Workspace accesses never receive the out-of-bounds offset. - No host code or tuned logic changes. Two new tests are added (see Test Plan). A build note for reviewers: the `.s` files are not tracked as build dependencies, so an incremental build does not reassemble them. Delete `device-library/tensilelite-device-libraries.stamp` in the build tree to force it. ## Test Plan - `Tensile/Tests/unit/test_CustomKernels_buffer_oob.py`, in the TensileLite `-m unit` lane, with no GPU needed. It reads the value `KernelWriterAssembly.py` emits and checks every bundled custom kernel against it. For readable kernels it checks the `BufferOOB` definition and the `D`, `C` and workspace descriptor lengths. For raw-encoded kernels it decodes descriptor-length and out-of-bounds moves from the `.long` words. Its decoder counts matched an independent classification from disassembly exactly. - `matmul_integer_exact_custom_kernel_store_extent_xf32_ROCM31258` in `clients/tests/data/matmul_gtest.yaml`, category `pre_checkin`, gfx950 only. It runs fp32 NN on the xf32 path at `M` = 2,097,152, 2,097,153 and 3,000,000 with `N` = 256 and `K` = 32. These are the sizes just at, one past, and well past 2 GiB of checked store range. It uses `integer_exact` inputs, so the check is bit-exact, and `c_equal_d`, so an unwritten element keeps its `C` value and fails the check. - Hardware evidence on an MI350X (gfx950), not automated: the reproducer attached to the ticket fills `D` with NaN (not-a-number), runs each custom kernel the heuristic offers, and counts elements still NaN afterward. ## Test Result Unpatched and patched builds of the same `develop` commit (`fdec6c677df`), standalone hipBLASLt for gfx950, ROCm 10.2.0a20260922, MI350X: - Unit test: unpatched, 120 failed, which is every affected kernel, and 17 passed. Patched, 137 passed. - gtest: unpatched, 2,097,152 passed, and 2,097,153 and 3,000,000 failed. Patched, all 3 passed, in about 38 seconds total. - Full TensileLite unit suite, patched: 8,506 passed, 47 skipped, 2 xfailed, 1 xpassed, 0 failed; 765 saved expected results unchanged. - The `hipblaslt-bench` command from the ticket: `norm_error` 8.9e+32, failed, becomes 3.75e-06, passed, with the same custom kernel selected. Unwritten elements of `D` per custom kernel offered, fp32 xf32 NN, `K` = 32: | `M` × `N` | Unpatched | Patched | |---|---|---| | 2,097,152 × 256 | 0 | 0 | | 2,097,153 × 256 | 1 | 0 | | 3,000,000 × 256 | 14,976 | 0 | | 6,000,000 × 128 | 231,129,088 | 0 | | 12,000,000 × 64 | 231,129,088 | 0 | | 4,194,303 × 256 | 32,640 | 255 | The last row's store range falls between `0xfffff000` and 2^32. The host still admits it because its threshold on `develop` is 2^32, and #12281 lowers that threshold to `0xfffff000`. The two changes are independent and together close the gap. Performance and output were compared on gfx950 for the six custom kernels that the harness could select directly (fp32 xf32 NN, NT and TN; bf16 TN, two sizes; fp16 TN), each at a size from the tuned table that selects it. Each ran alternately against the unpatched and patched kernel libraries, three runs apiece. Output `D` was bit-identical, and the median time changed by between -2.1% and +0.5%. The six fp8 custom kernels were not compared, because the harness hits a host-side error in the library's algorithm query for that type combination, identical on both libraries. In the rebuilt gfx950 kernel library, 21 of 315 code objects changed, and the other 294 were byte-identical. gfx942 was verified on an MI300A; see [this comment](#12482 (comment)). In short, there's no regression on hardware, and on gfx942 the fix is defensive rather than reachable. Automated GitHub checks are pending. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5). This changes shipping assembly in 120 kernels on the default gfx942 and gfx950 selection paths. The reassembly comparison bounds the change to the two constants described above. Outside the 2 GiB to `0xfffff000` store range, output was bit-identical in every comparison run. Inside it, stores that were dropped are now written. gfx942 has been run on hardware, but the six gfx942 kernels whose stores could be affected can't currently be exercised past 2 GiB (see ROCM-31692), which is why this is not rated lower. ## Related - ROCM-31016 and #12281, the matching host-threshold fix for generated kernels. The two PRs are independent and should land close together; either order is safe. - ROCM-31207, removing the 4 GiB ceiling - #11385, draft build-time validation of custom kernel metadata; a per-kernel host threshold would build on it - AIHPBLAS-4650, custom kernels drifting from the metadata that selects them - `ae0ca64a9c8` (#6664), the generator change this aligns with --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: tony-davis <7938412+tony-davis@users.noreply.github.com>
1 task done
tony-davis
added a commit
that referenced
this pull request
Oct 8, 2026
…GEMMs (#12876) JIRA ID : AIHPBLAS-4988 Stack, in merge order: **#12876 (this PR)**, #12960, #13005, #13013, #13015, #13021, #13022, #13025, #13031. Independent of the stack: #13027 (rocisa 64-bit add lowering) and #13060 (address-carry lint). ## Summary Recently we fixed some bugs where address, index or launch-size math was done in 32 bits when it needed 64. That kind of bug only shows up at big sizes or unlucky addresses, and it usually lives in a handful of solutions out of hundreds. Our tests have a hard time reaching it, because `unit_check` compares against an OpenBLAS GEMM on the host. That costs M·N·K and copies every padded buffer to the host, so the big-shape tests stay small and run only the heuristic's top pick, once. `integer_exact` already got rid of the tolerance, but not the cost. This PR adds fast_check, an opt-in check (`fast_check: 1`) for `integer_exact` tests. Instead of recomputing the GEMM, it compares a few random "fingerprints" of D with the same fingerprints computed from the inputs (Freivalds' algorithm, in exact modular arithmetic). That costs about as much as reading the matrices once, D never leaves the GPU, and a second set of fingerprints pins down which elements are wrong. Padding is filled with poison so stray reads show up, and D starts as a sentinel so missed writes show up. Every failure names the solution, its kernel and where each buffer sits in memory. Only test code changes. To make sure it catches what it's meant to, we reverted the ROCM-31016 fix (dropped stores past 2 GiB). fast_check flagged the affected solutions and pointed at the exact element where the stores stopped. It has also found real defects: on gfx950 the ROCM-31016 shape and the bf16 every-solution case fail today (ROCM-32152 and ROCM-32151), so those two cases are quarantined there, and running it next to `unit_check` found a bug in the CPU reference (ROCM-32062). **CI time added**, measured on one MI300X / one MI350X under ROCm 10.2 at the stack tip (the sum of gtest's per-test times, not counting test-data parsing at startup): - This PR: about 32 s / 15 s in pre_checkin (108 cases; the 51 self-tests take under a second), and about 3.5 min / 20 s in nightly (35 / 33 cases; gfx950 runs fewer because of the two quarantines above). - The whole stack: about 37 s / 58 s in pre_checkin, about 15 min / 9 min in nightly, and about 3.5 min / 1 min in the new `stress` tier (#13021), which needs large-memory runners. On MI300X the Stream-K cases skip unless `TENSILE_SOLUTION_SELECTION_METHOD=2` is set; with it they add about 5.5 min. ## Motivation hipBLASLt's gtest suite checks a matrix multiply (GEMM) result with `unit_check`, which compares the GPU output against a reference computed on the CPU. That reference costs time proportional to M·N·K, and the harness copies every input and output buffer, padding included, into pinned host memory. Tests for out-of-bounds indexing need large leading dimensions or large shapes, so this cost limits them. On gfx942, `matmul_size_t_index_overflow` with ldd = 1049100 spends 12 s of its 14.4 s allocating host memory. The existing overflow tests therefore keep K small and run only the heuristic's first solution, even though indexing defects such as ROCM-31016 and ROCM-31258 affect specific kernels. ## Technical Details This PR adds an opt-in YAML argument, `fast_check: 1`, for tests that use `initialization: integer_exact`. Under integer_exact, every input is a small integer, so the correct result is an exact integer as well. Only client test code and test data change. The library, TensileLite and kernel selection are untouched. - **Every element of D is checked without a CPU GEMM.** fast_check uses Freivalds' algorithm, which multiplies D by a random vector `r` and compares the result with `A·(B·r) + C·r + bias·sum(r)`. That costs time proportional to M·K + K·N + M·N instead of M·N·K. It does the same from the left with a second vector, to locate bad columns. The sums are taken modulo the prime 2^61 − 1, so the comparison is exact, and a wrong result passes one probe with probability below 2^−60. The input side is computed on the host once per test. D's side is computed on the GPU for each solution, so D is never copied to the host. When the probes disagree, the failure message lists the mismatched elements with their row, column and element offset, and marks offsets at or above 2^31 or 2^32. - **Failure context.** Each failure names the solution: its position in the test's list, its library index and its kernel. It also lists every operand's address range and flags any that crosses a 4 GiB boundary, since some address defects depend on placement rather than size. - **Out-of-bounds reads.** The padding of A, B and C (rows past the matrix inside the leading dimension, and gaps between batches) is filled with 40, a value outside the integer_exact range [−2, 2]. A kernel that reads padding therefore produces a wrong D. - **Out-of-bounds writes and missed stores.** Before each solution runs, D's whole allocation is filled with 0xFF bytes, which form a NaN in f16, bf16 and f32. Afterwards a device kernel checks that every element outside the M × N × batch region still holds that pattern, and that no element inside it does. A failed launch of any of these fills or scans fails the test; it is never read as clean. - **Results above the exact range.** bf16 holds integers exactly only up to 256. fast_check compares each larger result with its correctly rounded value (round to nearest, ties to even, pinned by a self-test), so f16 and bf16 tests can use K above the 256 limit that `unit_check` imposes on integer_exact. Those elements are recomputed exactly, in parallel, on the host. - **Supported configurations:** alpha, beta, bias, `scaleAlpha_vector`, strided batches, and every transpose, including conjugate transpose, which stores real data the same way as transpose. An output with zero rows, columns or batches needs no check, so it reads no operands and launches nothing. The test fails with a stated reason for timing runs, activations, scale factors, grouped GEMM and pointer-array batches; later PRs in the stack lift several of these. Output types run against real kernels here are f32, f16 and bf16; the stack adds int8 and fp8 cases (#13013, #13022). The code also accepts f64 and fp8 outputs, which only the self-tests exercise. New files are `clients/common/include/fast_check.hpp`, `clients/common/src/fast_check.cpp` and `clients/tests/src/fast_check_gtest.cpp`. `testing_matmul.hpp` calls them, and in fast_check-only runs it skips the CPU reference and the padded host copies. Running both checks on the same results found ROCM-32062, a separate defect in the CPU reference. When `scaleAlpha_vector` is combined with a leading dimension of A larger than its row count, `cast_mul` in `cblas_interface.cpp` scales the wrong rows of A, so `unit_check` fails a correct GPU result. No existing test combined the two. That case, `matmul_fast_check_epilogue_padded_lda_ROCM32062`, is quarantined in `known_bugs.yaml` until the fix PR removes the entry; #13015 adds a fast_check-only copy that keeps checking the GPU there. Related: ROCM-26455, ROCM-31016 (#12281), ROCM-31230 (#12332), ROCM-31258 (#12482), and the original overflow tests (#8231). ## Test Plan - `FastCheck` (host only) and `FastCheckDevice` (small device buffers, no hipBLASLt kernels) in `fast_check_gtest.cpp`. These build correct results, inject faults, and check that each fault is reported: a wrong element, a pair of errors that cancel in a plain row sum, a shifted row, a duplicated column, an unwritten element, an out-of-bounds write, wrong rounding of a large bf16 result (including a tie), an empty buffer, an empty output (zero M, N or batch count, with null unused operands), and an unsupported type. They also require the device and host passes to reach the same verdict. - pre_checkin, 126 cases: fast_check and `unit_check` on the same results across precisions and transposes, including all nine N/T/C pairs on two rectangular f32 shapes with padded leading dimensions and three batches; the size_t index-overflow shapes in f32, f16 and bf16; and the ROCM-26455 shape with bias and `scaleAlpha_vector`. - nightly, 35 cases: element offsets and a batch base offset above 2^32; shapes up to 16385 × 16383 × 4097 with partial tiles; and every solution at the output-overflow and ROCM-31016 shapes. - Regression replay: revert the ROCM-31016 threshold and run `matmul_fast_check_store_extent_ROCM31016`. - `standard` tier (smoke, quick, pre_checkin) to confirm existing tests are unchanged. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Built and run in a ROCm 10.2.0 image (TheRock 10.2.0a20260928). Every fast_check case and self-test, with none failing: | GPU | Ran | Passed | Skipped | | --- | --- | --- | --- | | gfx90a (MI210) | 496 | 410 | 86 | | gfx942 (MI300X) | 539 | 473 | 66 | | gfx950 (MI350X) | 559 | 486 | 73 | **After newling's review.** With the empty-output fix, at stack tip 6296886715 (ROCm 10.2, before the rebase onto develop), the same filter ran 501 tests on gfx90a, 544 on gfx942 and 564 on gfx950, with 0 failures. newling's conjugate-transpose case failed 10 of its 18 cases before the fix (all those with a C operand) and passed all 18 after, on gfx1201 with ROCm 7.1. On the current tip (86d4ac3, with both fixes, rebased onto develop with #13119), the full fast_check filter on gfx942 (MI300X) ran 561 tests: 495 passed, 66 skipped for a stated reason, and none failed, including all 18 conjugate-transpose cases and 28 fp8 fnuz cases. The skips are quarantined known bugs, shapes the library declines, Stream-K cases that need `TENSILE_SOLUTION_SELECTION_METHOD=2`, and on the shared MI210, cases that need more free memory than it had. Earlier development runs, on gfx942 (MI300A): - Self-tests: each of six deliberate source mutations, four in the host code and two in the GPU kernels, made at least one self-test fail. - With the ROCM-31016 fix reverted, 170 of the 1,093 solutions failed. The first reported 1793 elements never written, starting at element offset 2147481599, which is the 2^32-byte boundary for bf16. With the fix in place, all 1,093 passed in 144 s. - With the fix reverted, each failure names the kernel, for example `Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x448x32_…` (library index 75677). The placement record shows C and D, each about 4 GiB, crossing a 4 GiB boundary. - Same shapes, `unit_check` before and fast_check after: ldd = 1049100 f32 went from 14.4 s to 0.7 s, and lda = 1049100 f32 from 8.3 s to 1.2 s. - `standard` tier (smoke, quick, pre_checkin), run with the stack's second PR: 71,379 run, one failure, `matmul_fallback_compute_fp16`, which passes at the stack tip under ROCm 10.2 on an MI300X (see #12960). ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5): test code and test data only, with no library or device behavior change, and existing tests take the same paths unless they set `fast_check`. It is rated 3 rather than 2 because it adds shared machinery to `testing_matmul.hpp` and new GPU kernels that every architecture runs, and about half a minute of pre_checkin. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.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.
JIRA ID : ROCM-31016
Motivation
hipBLASLt picks a GPU kernel for each matrix multiplication from a tuned list. Before a kernel is offered, host-side checks called predicates remove kernels that cannot run the requested problem.
A kernel writes its output
Dthrough a buffer resource descriptor. Its 32-bitnum_recordsfield bounds every store, and the hardware silently discards any store at or past it.allocPostLoopSrdinKernelWriterAssembly.pysets that field to theBufferOOBsentinel,0xfffff000.BufferStoreOffsetLimitCheckis the predicate meant to keep selection away from shapes a kernel cannot write in full, but it compared against2^32instead of the sentinel. A shape whose store extent fell between the two was reported as supported and then only partly written, whilehipblasLtMatmulreturned success. This is the failure reported in ROCm/hipBLASLt#2299, where a convolution in a video model returned a partly unwritten result with no error.Technical Details
BufferStoreOffsetLimitChecknow compares against0xfffff000, held in a named constant with a comment pointing at where the generator emits it.The extent formula is unchanged:
ldd * min(MacroTile1, N) * elementBytes. Each workgroup advances its descriptor base alongNbefore storing (computeStoreSrdStart), so only oneMacroTile1of columns has to fit, not all ofD. That is why outputs larger than 4 GiB keep their solutions. Bounding the whole tensor instead would leave a 6.04 GiB bf16 matmul (M= 12,667,846,N= 256) with no solution at all, for a size that runs correctly today.The comparison stays strict, so an extent landing exactly on the sentinel is rejected even though its last store is still in range. That costs one
lddvalue per tile geometry and avoids reasoning about how wide a solution's final store is. A test records the choice.debugEvalprintedstrides()[1] * value, skipping themin()againstN, so for outputs narrower thanMacroTile1it reported a larger extent than the one that decided the result. It now prints the compared value and formats the threshold from the constant.This changes host-side selection only. No kernel code changes, and the threshold is evaluated at run time rather than stored in the tuned library files, so a host-library rebuild applies it to libraries that have already shipped.
Out of scope: the 21 hand-written kernels under
Tensile/CustomKernelssetBufferOOBto0x80000000, so this global threshold is still too permissive for them. That is tracked in ROCM-31258 and #12482.Test Plan
tensilelite-tests, C++ predicate tests intensilelite/tests/Predicates_test.cpp. Five cases: the measured broken shape is rejected, the shape one step below the sentinel is accepted, an extent exactly on the sentinel is rejected (the strict comparison), the shape from the original report stays selectable, and a tall output narrower thanMacroTile1stays selectable. These run in the shared TensileLite lane without a GPU.Dwith NaN, run each algorithm the heuristic offers, and count elements never written. A and B are finite and beta = 0, so any NaN left behind is a store that never happened.Test Result
Predicate tests, local build on an MI350X host:
tensilelite-testssuite: 499 ran, 496 passed, 3 skipped. The skips are unrelated:RangeLibraryTest.SpecificSizes/0andMXDataGenFP4Test.ZeroFrequencyWithinBounds/3also skip ondevelop, andDataInitializationReferenceRecompute.Gfx1250ReusesPreswizzledMXAcrossSolutionsis gfx1250-only.2^32, 2 of the 5 fail: past-sentinel and exactly-at-sentinel. Those are the shapes whose extent falls between the sentinel and2^32.Hardware, MI350X, two host builds differing only in this threshold, against the tuned library from TheRock nightly
10.1.0a20260822. The first 64 algorithms the heuristic offered were run for each shape:N= 256,K= 48,ldd=M)Dpartly unwrittenM= 8,388,6072^32M= 8,388,6070xfffff000M= 8,388,5992^32M= 8,388,5990xfffff000All 22 had
MacroTile1= 256, including the heuristic's first choice. The unwritten elements were the trailing partial tile of each solution'sMacroTile0in the last column. The shape from the issue (M= 5,326,464) and the 6.04 GiB shape above were clean before and after and keep every candidate.A later full sweep on an MI355X against ROCm
10.2.0a20260918ran all 2,207 offered solutions atM= 8,388,607. 605 leftDpartly unwritten and 1,601 ran clean. The change excludes the 606 withMacroTile1>= 256 and still admits exactly the 1,601 that ran clean. The one-solution gap between 606 and 605 is a counting difference between the two tools, not a working solution being excluded.Submission Checklist
Risk level
Low (2/5). Host-side selection only, with no kernel code change and no change to the extent formula. The predicate becomes stricter only for store extents in the 4,096 bytes between
0xfffff000and2^32. Shapes outside that window get the same answer as before. Inside it, the measured solutions it removes were the ones leavingDpartly unwritten, plus the one conservative exactly-at-sentinel extent, which would have run correctly.Related