test(hipblaslt): run fast_check on edge, bookkeeping and per-kernel configurations by tony-davis · Pull Request #13022 · ROCm/rocm-libraries

tony-davis · 2026-10-02T23:33:41Z

JIRA ID : AIHPBLAS-5020

Stack, in merge order: #12876, #12960, #13005, #13013, #13015, #13021, #13022 (this PR), #13025, #13031. The diff here is only the edge-configuration layer; the base moves down the stack as those merge.

Summary

Some bugs depend on configuration, not size: a prefetch that reads past a very short K loop, Stream-K partials that reuse batch 0's accumulators, mixed fp8 formats that select their own kernels, one output type selecting a different kernel than another, and intermittent failures on ragged shapes. Each of those needs to be named in a test, and the intermittent ones need repeated, checked runs.

This PR adds targeted nightly cases for each, built on the earlier layers (every solution, repeated launches, exact long K). They include the exact shapes from the stress-test reports that started this work. Several of those reports came from gfx1250 (MI45X and MI455X), which I couldn't run here, so these cases check the same mechanisms on gfx942 and gfx950; they will run on gfx1250 once CI does. It also keeps a Stream-K case from passing when the heuristic returned no Stream-K kernel at all.

CI time added (one MI300X / one MI350X, ROCm 10.2, sum of per-test times): nothing in pre_checkin, and about 3 min / 2 min in nightly (64 / 60 cases). With TENSILE_SOLUTION_SELECTION_METHOD=2 on MI300X, the Stream-K batch cases add about 1 min more.

Motivation

Classes E, F and G on AIHPBLAS-4988 depend on the configuration rather than the size:

a prefetch that reads past the end of a short K loop (ROCM-31105, gfx1250 fp8);
Stream-K partials that reuse batch 0's accumulators (GitHub PR 11962);
mixed fp8 formats selecting their own kernels, which then fail (ROCM-31905, MI45X);
one output type selecting a failing kernel where another passes (ROCM-32072, MI455X bf16 Stream-K);
intermittent failures on ragged shapes (ROCM-31901, which needs other GPU work running at the same time; test(hipblaslt): add an on-demand fast_check SDC-hunt recipe #13025 adds that).

Each needs its configuration named explicitly, and the class G races also need repeated, checked launches.

Technical Details

Each case uses fast_check with repeated launches (fast_check_repeat from #13005), and the large-K or 16-bit cases use sparse_k (#13013) so every result stays exact in the output type:

matmul_fast_check_short_k: K = 1, 15, 16, 33, 64 and 129 against every solution, so the loop runs fewer iterations than common prefetch depths. Kernels with PrefetchGL2 on and off are both covered where the library has them. These run f32 and bf16; the fp8 kernels ROCM-31105 names are on gfx1250.
matmul_fast_check_ragged: partial tiles in every dimension, including driftcheck's 8191 × 8193 × 4097 (ROCM-31901) and single-row and single-column shapes. This checks the shapes, not the concurrency that triggers ROCM-31901.
matmul_fast_check_output_types: bf16 and f16 inputs, each with its own type and an f32 output (ROCM-32072).
matmul_fast_check_streamk_batches: every Stream-K solution with 3 batches on ragged shapes (GitHub PR 11962). It needs TENSILE_SOLUTION_SELECTION_METHOD=2, except on gfx950.
matmul_fast_check_mixed_fp8_fnuz (gfx942) and matmul_fast_check_mixed_fp8 (gfx950): E4M3 × E5M2 and the reverse, f32 and bf16 outputs, at the ROCM-31905 failing shapes (2560 × 16384 × 2560 and 5120 × 12288 × 5120), at 10240³ where it passed, and at two sizes in between.

All are nightly. A case marked requires_streamk skips, with the reason, unless at least one returned solution is a Stream-K kernel (TileProcessingStrategy StreamK, TPSSK in the kernel name); that check now lives in #13005, where requires_streamk first appears.

Test Plan

Every new case on gfx90a, gfx942 and gfx950; the Stream-K cases with and without the environment variable.

Test Result

ROCm 10.2, at the stack tip (3364b79, which includes this PR). Every fast_check case and self-test passed or skipped for a stated reason on gfx90a, gfx942 and gfx950, with none failing (counts in #12876). This PR's nightly cases: 64 ran on gfx942 and 60 on gfx950, all passing, with the quarantined ones skipped. With TENSILE_SOLUTION_SELECTION_METHOD=2, the Stream-K batch cases ran and passed on gfx942 and gfx950; on gfx90a all 40 Stream-K cases in the stack now skip with "none of the N solutions for this case is a Stream-K kernel".

Math CI found ROCM-32277. Precheckin on #13031 failed the f32 Stream-K batch case on gfx950: one Stream-K solution wrote wrong values in the last partial M tile on one launch of three. On a gfx950 node it failed 4 of 30 runs, through two f32 Stream-K solutions (MT128x64x32 and MT64x64x32), and hipblaslt-bench -v reproduces it on 4 of 80 single launches. It looks like ROCM-32151 but in f32 and intermittent, so it has its own ticket, and the f32 case at M = 4099 is quarantined on gfx950 under it. The f32 case at M = 1031 passed 60 of 60 runs on two gfx950 nodes, one of which fails the M = 4099 case 4 times in 30, so it stays active.

Earlier development runs:

Cluster runs built with ROCm 7.1 and run on ROCm 7.2: the 20 mixed fp8 cases passed on each GPU (fp8 fnuz on gfx942 after the test(hipblaslt): keep fast_check exact at large K #13013 negate fix, OCP fp8 on gfx950), as did ragged and output types. The bf16 short-K TN cases on gfx950 and the bf16 Stream-K batch cases on both GPUs found ROCM-32151 (about half the right value in the last partial M tile) and are quarantined under it; their f32 cases passed. The quarantine covers the whole configuration; there is no way yet to exclude single solutions.

Submission Checklist

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

Risk level

Low (2/5): nightly test data, plus a skip in the shared harness that only affects cases marked requires_streamk.

Copilot

Copilot review overview

🔵 Needs a closer look

The Stream-K regression coverage is not reliably exercising the intended paths.

Review effort: Lite
Findings: None

What changed in this PR

Adds nightly hipBLASLt fast_check coverage for configuration-sensitive edge cases.

Changes:

Adds short-K, ragged-shape, and output-type tests.
Adds Stream-K batching and repeated-launch checks.
Adds architecture-specific mixed-FP8 cases.

File	Summary
`projects/hipblaslt/clients/tests/data/matmul_gtest.yaml`	Adds the configuration-focused nightly test matrix. Review found issues with batch configuration, guaranteed internal Stream-K batching, and architecture guarding.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

therock-pr-bot · 2026-10-03T04:06:19Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	PR does not contain code files — Unit Test auto-passed
🔎 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?

codecov · 2026-10-03T06:18:34Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@                                  Coverage Diff                                   @@
##           users/todavis_amdeng/AIHPBLAS-5018-size-thresholds   #13022      +/-   ##
======================================================================================
- Coverage                                               75.90%   69.22%   -6.69%     
======================================================================================
  Files                                                     224      292      +68     
  Lines                                                   89717   112405   +22688     
  Branches                                                16481    19269    +2788     
======================================================================================
+ Hits                                                    68097    77803    +9706     
- Misses                                                  17910    30014   +12104     
- Partials                                                 3710     4588     +878

Flag	Coverage Δ
TensileLite-CPP	`49.33% <ø> (ø)`
TensileLite-Unit	`83.32% <ø> (ø)`
hipBLASLt	`42.78% <ø> (?)`
see 68 files with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

Copilot

Copilot review overview

🟡 Changes recommended

The bf16 short-K sweep lacks sparse exactness coverage, and the batch-stride claim does not match its configuration.

Review effort: Lite
Findings: 1

Open (1)

Test does not configure the claimed strided batches · New

therock-pr-bot · 2026-10-03T16:48:02Z

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

Copilot

Copilot review overview

🟢 Approval recommended

No unresolved review issues were identified.

Review effort: Lite
Findings: None

Resolved since last review (1)

Test does not configure the claimed strided batches

tony-davis · 2026-10-07T19:55:25Z

Thanks @newling! I checked the scope. The f32 M = 1031 case passed 60 of 60 runs on two gfx950 nodes, one of which fails M = 4099 four times in 30, so the quarantine now covers M = 4099 only.

…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>

…configurations Classes E, F and G on AIHPBLAS-4988 need specific configurations rather than big sizes. New nightly cases, each launched and checked several times: - K of 1 to 129 against every solution, so the loop runs fewer iterations than the prefetch depth (ROCM-31105, PrefetchGL2 on and off); - partial tiles in every dimension, including driftcheck's ragged 8191 x 8193 x 4097 (ROCM-31901); - each output type for bf16 and f16 inputs (ROCM-32072); - every Stream-K solution with several batches (GitHub PR 11962); - mixed fp8 formats in both orders with a shape sweep between the known-bad and known-good sizes of ROCM-31905, FNUZ on gfx942 and OCP on gfx950. Co-authored-by: Cursor <cursoragent@cursor.com>

…atch cases on ROCM-32151 The short-K case on MI350X and the Stream-K batch case on MI300X and MI350X find bf16 solutions writing about half the right value, in the last partial M tile (rows 512 to 516 at M = 517, rows 1024 to 1030 at M = 1031) or, for gfx950 solution 354888, in every row (ROCM-32151). Quarantine the bf16 cases on those platforms; the f32 cases keep running. Co-authored-by: Cursor <cursoragent@cursor.com>

…wo batches The ragged-shape case said it covered strided batches but ran one batch. It now runs two, with the default packed batch strides; the 12 cases take about 8 s on gfx90a. Co-authored-by: Cursor <cursoragent@cursor.com>

… on gfx950 Math CI precheckin on PR 13031 failed matmul_fast_check_streamk_batches f32 on gfx950: one Stream-K solution wrote wrong values in the last partial M tile on one launch of three. It reproduces on about 1 run in 8 of the case, and in hipblaslt-bench -v, through two f32 Stream-K solutions (MT128x64x32 and MT64x64x32). Filed as ROCM-32277; quarantined like the bf16 variant, which ROCM-32151 already covers. Co-authored-by: Cursor <cursoragent@cursor.com>

…= 4099 The f32 Stream-K batch case at M = 1031 passed 60 of 60 runs on two gfx950 nodes, one of which fails the M = 4099 case 4 times in 30, so only M = 4099 stays quarantined. Co-authored-by: Cursor <cursoragent@cursor.com>

…12960) JIRA ID : AIHPBLAS-4994 Stack, in merge order: #12876, **#12960 (this PR)**, #13005, #13013, #13015, #13021, #13022, #13025, #13031. The diff here is only the placement harness; the base moves to `develop` when #12876 merges. ## Summary Some address bugs don't depend on how big the problem is. They depend on where the buffer sits. If an operand straddles a 4 GiB boundary, a kernel that drops the carry reads or writes exactly 4 GiB away. With `hipMalloc` that happens rarely and at random, so a test can pass for months and then fail once with no way to reproduce it. fast_check's poison and sentinel can't see an access that far outside the buffer, and on some GPUs that stray access faults hard enough to need a reset. This PR adds controlled placement. A `placement` argument maps a chosen operand (A, B, C, D, bias, the scale vector or the workspace) across a 4 GiB boundary using HIP virtual memory. It also maps poison windows exactly 4 GiB above and below. A bad read gets poison, so fast_check reports a wrong answer instead of the GPU faulting. A bad write lands in the poison, where a scan finds it. So a rare, random crash becomes a deterministic, safe test. A self-test simulates a dropped carry and checks that it's caught. Two deliberate mutations (checking only one window, and never filling the poison) each make that self-test fail. The motivating defect, ROCM-32046, is in gfx1250 MXFP8 Stream-K kernels, which no GPU here can run; MX operands get placement in #13031, and the cases here check f32 and bf16 kernels on every GPU CI runs. **CI time added** (one MI300X / one MI350X, ROCm 10.2, sum of per-test times): about 1 s / under 1 s in pre_checkin (12 cases), and about 2 min / 40 s in nightly (14 every-solution cases). ## Motivation Some hipBLASLt address defects depend on where a buffer sits in memory rather than on the problem size. ROCM-32046 drops the carry out of the low 32 bits of an address in two MXFP8 Stream-K solutions on gfx1250, so it goes wrong only when an operand's base crosses a 4 GiB boundary. With `hipMalloc` that crossing is rare and random, so a test can pass for months and then fail once with no reproduction. The wrong access lands exactly 4 GiB away, outside every allocation. fast_check's padding poison and D sentinel cannot see it, and on gfx1250 the resulting page fault can leave the GPU unusable until it is reset (ROCM-32049). ## Technical Details Only client test code and test data change. - **Placement.** A new test argument, `placement`, names one operand: `a`, `b`, `c`, `d`, `bias`, `scale_alpha_vec` or `workspace`. `placement_offset` sets how many bytes of it lie below the boundary; 0, the default, means half. `PlacedRegion` in the new `hip_placement.hpp` uses HIP virtual memory management (`hipMemAddressReserve`, `hipMemCreate`, `hipMemMap`). It reserves address space and maps the operand so that it crosses a 4 GiB boundary. - **Poison windows.** Two more mappings of the same size sit exactly 4 GiB below and above the operand. Both are filled with the fast_check poison value, and so is the rest of the operand's own mapping past its last element. A read whose carry was dropped returns poison, so fast_check reports a wrong D instead of the GPU faulting. A write whose carry was dropped, or that runs past the end, changes the poison, and a device scan after every solution reports the element it was meant for. - **The workspace.** Solutions need different amounts of workspace, and any of them may use the whole size it is given. The placed workspace is twice the largest size, and each solution starts it so that its own reported range straddles the boundary with the whole workspace size still mapped after it. A self-test checks that for sizes from 300 bytes to the full workspace. - **No silent fallback.** A placement case fails if no buffer was placed, or if the placed range does not cross the boundary. - **`HipDeviceBuffer`.** A new constructor allocates a buffer this way. Every place that passes the operand to the library uses the placed pointer without further changes. Its move operations now clear the moved-from pointer. Before this, a moved-from buffer re-checked guard bytes in memory owned by the moved-to buffer; that was harmless with `hipMalloc`, but would fault on a placed buffer. - **Requirements.** Placement requires `fast_check`. If the GPU or platform lacks virtual memory management (Windows, for one), the test skips with that reason. An invalid request fails the test. ## Test Plan - `FastCheckDevice` self-tests: - a kernel that computes addresses with the carry dropped on purpose must read the poison value and have its write reported at the right element; - each solution's share of a placed workspace must straddle the boundary and stay mapped, and a write past the end of a placed buffer must be reported; - requests that cannot cross a boundary are rejected; - a placed buffer keeps its address when `std::vector` growth moves it. - pre_checkin cases: on a 1031 × 517 × 263 GEMM in f32 and bf16, place A, B, C and D in turn, and bias and the scaleAlpha vector with their epilogue, using the heuristic's first 16 solutions. - nightly cases: the same placements, plus the workspace, with every supported solution. - `standard` tier, because the `HipDeviceBuffer` move change affects every test. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Every fast_check case and self-test passed or skipped for a stated reason on gfx90a (MI210), gfx942 (MI300X) and gfx950 (MI350X), with none failing; the counts are in #12876. The placement cases all ran on all three, including the workspace placement with every solution (more than 2,000 solutions on gfx90a, each with its own straddling range). Earlier development runs: - Self-tests (gfx942, MI300A): two deliberate mutations, one checking only the window above and one never filling the poison, each made the carry-drop self-test fail. - The first version of the per-solution workspace placed each solution's range near the end of a workspace-sized mapping, and solution 1 of the every-solution case faulted the GPU: hipBLASLt may use the whole workspace size it is given, not just the size its heuristic reports. Doubling the mapping fixed it. - `standard` tier (gfx942, MI300A, before the per-solution workspace change): 71,379 run, 71,373 passed, 5 skipped, 1 failed, `matmul_fallback_compute_fp16`, a workspace request larger than the test allocates. At the stack tip under ROCm 10.2, all 16 of its cases pass on an MI300X, so that failure looks tied to the MI300A or the ROCm 7.1 environment rather than to this PR. - ROCM-32046 itself is on gfx1250 MXFP8 kernels, which none of these machines can run. ## 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, but the `HipDeviceBuffer` move change affects every client test that stores buffers in a `std::vector`, so the `standard` tier must pass before merge. 🤖 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>

…13005) JIRA ID : AIHPBLAS-5015 Stack, in merge order: #12876, #12960, **#13005 (this PR)**, #13013, #13015, #13021, #13022, #13025, #13031. The diff here is only the every-solution, every-iteration layer; the base moves down the stack as those merge. ## Summary Address-arithmetic bugs tend to live in a few solutions out of hundreds, and some only show up on some launches. A test that runs the heuristic's top pick once can miss both kinds. And when a failure does show up, nobody can reproduce it unless the report says which solution and which launch failed. This PR lets a test run every solution the library offers for a problem, and launch and check each one several times (`fast_check_repeat`). D is reset to the sentinel and the workspace to zero before every launch, so each check sees only that launch's work. One summary at the end lists every failing solution, its kernel and which launches failed. Stream-K needs a word of care. On gfx942 the library offers Stream-K solutions only when the process starts with `TENSILE_SOLUTION_SELECTION_METHOD=2`, which no CI lane sets today, so the Stream-K cases skip there and are for manual runs. On gfx950 Stream-K is the only library, so they always run, and they overlap the default cases. A self-test corrupts one launch on purpose, to prove the report points at the right one. Cases that hit open defects are quarantined in `known_bugs.yaml`; see Defects found below. **CI time added** (one MI300X / one MI350X, ROCm 10.2, sum of per-test times): about 1 s / 2 s in pre_checkin (14 / 26 cases), and about 6 min / 3 min in nightly (28 / 29 cases). With `TENSILE_SOLUTION_SELECTION_METHOD=2` on MI300X, this PR's Stream-K cases add about 4.5 min more. ## Motivation Each address-arithmetic defect on AIHPBLAS-4988 lives in a handful of the hundreds of shipped solutions: 2 of 244 in ROCM-32046, 39 of 545 in ROCM-31290, 4 of 679 in ROCM-31258. Some defects are also intermittent. ROCM-31901 and ROCM-31905 fail only on some launches, and ROCM-31105 needed repeated launches to show up. A test that runs only the heuristic's top pick, and checks only the first call, misses both kinds. When a failure does show up, the report also needs to say which solution and which launch failed, or nobody can reproduce it. ## Technical Details Only client test code and test data change. - **Every iteration.** A new argument, `fast_check_repeat`, launches and checks each solution that many times, one solution at a time. Before every launch D is refilled with the sentinel, and the workspace is returned to the zeros of a fresh allocation, so a repeat cannot pass on partial sums an earlier launch left there (a Stream-K fixup that reads too early, for example). Zero rather than poison, because kernels may rely on their flags starting at zero. A failure message names the solution, its library index and kernel, and the iteration ("iteration 2 of 4"). - **Summary per solution.** `FastCheckSolutionLog` records pass or fail for each solution and iteration. At the end of the test, one failure lists every failing solution with its library index, kernel name and failing iterations. With hundreds of solutions, that list is the triage starting point. - **Self-test.** `fast_check_inject` corrupts one element of D after the launch on the named iteration. The case passes only if that iteration fails for every solution, its report names the corrupted element, and every other iteration passes. The corrupted value is the sentinel, so this exercises the per-iteration sentinel check; the probe-sum path is covered by the host self-tests in #12876. - **Stream-K.** A new argument, `requires_streamk`, skips the case with a message unless the process started with `TENSILE_SOLUTION_SELECTION_METHOD=2`. On gfx950 the case always runs, because Stream-K is the only library there. TensileLite reads that variable once per process (`Debug::Instance`), so `tensile_solution_selection_method` in a YAML case only has an effect if no earlier test in the process has run a GEMM. A Stream-K case whose solutions include no Stream-K kernel (TileProcessingStrategy StreamK, `TPSSK` in the kernel name) also skips with the reason, rather than passing while testing none. - **Every solution.** These cases use `algo_method: all`, which already includes the hand-written custom kernels (ROCM-31258). A new xf32 case on gfx950 runs them all. New YAML cases: - pre_checkin: the self-test (f32 and bf16, 4 solutions × 4 iterations); the heuristic's first 16 solutions × 4 iterations in f32, f16 and bf16 for every transpose; and the same under Stream-K. - nightly: every supported solution × 3 iterations in f32, f16 and bf16 on two ragged shapes; xf32 with custom kernels on gfx950; and every Stream-K solution × 3 iterations. ## Defects found The cluster runs found two library defects. Each is filed and quarantined in `known_bugs.yaml` on the platforms it affects, with a time-box: - **ROCM-32151**: some bf16 Stream-K solutions on gfx942 and gfx950 write about half the right value in the last partial M tile (rows 1024 to 1030 at M = 1031), and gfx950 solution 354888 does it in every row. On gfx942 these solutions are offered only with `TENSILE_SOLUTION_SELECTION_METHOD=2`, so the bf16 every-solution case stays active there and fails only in that manual mode. `hipblaslt-bench -v` pinned to each solution index fails the CPU reference; the heuristic's default passes. The quarantine covers the whole bf16 configuration on gfx950; there is no way yet to exclude single solutions. - **ROCM-32153**: gfx950 f16 solutions 405799 and 407668 return wrong, sometimes non-integer, values at 4099 x 2053 x 129. `hipblaslt-bench -v` confirms 407668. ## Test Plan - Host self-test: the solution log lists only failing solutions, with their kernels and iterations, and is empty when everything passed. - Device self-test: the corruption helper changes exactly one element, and never to the value it held. - `matmul_fast_check_repeat_reports_injected_iteration`, plus a mutation check: corrupting iteration 3 instead of 2 must make the self-test fail on both counts. - All existing fast_check cases still pass with `fast_check_repeat` at its default of 1. - The Stream-K cases skip without the variable and run with it, and skip on a GPU whose library offers no Stream-K kernel. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Every fast_check case and self-test passed or skipped for a stated reason on gfx90a, gfx942 and gfx950, with none failing; the counts are in #12876. With `TENSILE_SOLUTION_SELECTION_METHOD=2`: - gfx942 (MI300X): 52 passed, 12 quarantined, and the 6 bf16 every-solution cases found ROCM-32151 again (all 32 failing solutions are Stream-K kernels). - gfx950 (MI350X): 43 passed and the other 28 are quarantined. - gfx90a (MI210), which has no Stream-K library: all 40 Stream-K cases now skip with "none of the N solutions for this case is a Stream-K kernel", where before they passed while testing none. Earlier development runs on gfx90a (MI210): - Mutation check: with the wrong iteration corrupted, the self-test failed for all 8 solutions with "the corrupted iteration 2 passed" and "iteration 3 was not corrupted but failed". - One run of the full filter crashed with a segmentation fault in a `matmul_fast_check_matches_unit_check` case that also runs the host reference; that case passed alone, its 64-case group passed, and later full runs passed, so I could not reproduce it. - gfx1250 was not run. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): test code and test data only, but it restructures the shared solution loop in `testing_matmul.hpp`. With `fast_check_repeat` at its default of 1 that loop runs each solution once, as before, and the workspace reset happens only for fast_check cases. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

JIRA ID : AIHPBLAS-5016 Stack, in merge order: #12876, #12960, #13005, **#13013 (this PR)**, #13015, #13021, #13022, #13025, #13031. The diff here is only the large-K exactness layer; the base moves down the stack as those merge. ## Summary fast_check is only exact while the numbers stay exact. Some of the shapes we need to cover have K of 28,672 or 32,768. At that length, `integer_exact` results can outgrow what a 16-bit output stores exactly (bf16 holds integers only up to 256), and nothing stopped a configuration from also outgrowing the accumulator. Split-K and Stream-K add partial sums in an order the test doesn't control. So "exact" could quietly stop meaning anything. This PR adds a safety check and two new data patterns. Before any kernel runs, fast_check works out the largest value any partial sum could reach, and refuses the test with a clear reason if it's too big. `ternary` uses only -1, 0 and 1. `sparse_k` keeps each row of A mostly zero, with a few values spread across K that move from row to row, so every K index still matters to some row and the kernel has to read the whole length correctly. The fnuz fp8 generator fix this PR first carried landed on its own in #13119. **CI time added** (one MI300X / one MI350X, ROCm 10.2, sum of per-test times): about 1 s / under 1 s in pre_checkin (32 cases), and about 25 s / 20 s in nightly (35 cases). ## Motivation fast_check is exact only while integer_exact data stays exact. Every partial sum has to fit what the accumulator holds exactly (2^24 for f32). Every result also has to round predictably into D, and bf16 holds integers exactly only up to 256. At K = 28672 or 32768, which are the shapes the AIHPBLAS-4988 defects need, the standard values break the second limit for 16-bit outputs. Each such result then costs a K-length dot product on the host, and nothing stopped a configuration that broke the first limit too. Split-K and Stream-K make that worse, because they add partial sums over subsets of K in an order the test does not control. ## Technical Details Only client test code and test data change. - **Up-front bound.** `fast_check_expected` now bounds, from the actual inputs, every partial sum over any subset of K and every result before rounding: the sum over K of |a| times the largest |b| in that row of B, scaled by alpha and the scaleAlpha vector, plus |beta·C| and |bias|. The expected sums are now computed before the first launch. A configuration whose bound reaches the compute type's exact range (2^24 f32, 2^11 f16, 2^31 int32) fails with that reason before any kernel runs. - **`integer_exact_pattern`.** This new argument selects the data: - `ternary`: A, B and C in {-1, 0, 1}. - `sparse_k`: each row of A is zero except at one K index in each sixteenth of K, plus the last K index, and those kept values are 1 or 2, never 0. The kept offset changes from stretch to stretch and from row to row, so with at least K/16 rows every K index is nonzero in some row, and a wrong read of B anywhere along K changes D. Every result stays at most 140 whatever K is. The large-K cases here have 1,031 or 2,053 rows at K = 32,768 (K/16 = 2,048), so about half or all of the K indices are covered; a K loop that skips a 64-deep iteration is caught either way. The patterns are applied in the device fill, so fast_check's host copies and unit_check's reference both see them. A scope guard restores the standard pattern after each test, and the selected pattern is per thread, so one test thread's guard cannot reset another's. For A stored transposed or conjugate-transposed (the same layout for real data), the pattern runs K along A's rows. `sparse_k` also lifts the existing K ≤ 256 skip for 16-bit unit_check. - **Clearer refusal.** fast_check's message now names the first non-integer element, for example `op(B)(0, 0) = nan`. That message is how the FNUZ fp8 negate bug was found; its fix is in develop now (#13119). - **Coverage notes.** `TESTING.md` now records which operands integer_exact covers today (A, B and C; bias, the scaleAlpha vector, E and the scales get the generic 1-to-10 fill, and amaxD starts at zero) and which pattern each output type needs at large K. New YAML cases: - pre_checkin: both patterns cross-checked against unit_check (f32 and bf16, every transpose). - nightly: f16 with `ternary` and `sparse_k`, and bf16 with `sparse_k`, at K = 32768 and 28672, every transpose; int8 → int32 with the standard pattern and int8 → int8 with `sparse_k` at K = 32768; and fp8 → f32, bf16 and fp8 with `sparse_k` at K = 32768 (FNUZ on gfx942, OCP on gfx950). ## Test Plan - Host self-test: the bound equals the hand-computed maximum, with alpha 1 and beta 0 and again with alpha 2, beta −2, a scaleAlpha vector and a bias; a K that pushes it past 2^11 under f16 compute is refused while a smaller K is accepted. Mutation check: disabling the refusal must fail this test. - Device self-test: `sparse_k` leaves each row of A nonzero only at its chosen K indices (15 to 17 per row, including the last), for both storage orders, and with 16 rows at K = 200 every K index is nonzero in some row, counted from the values written rather than the positions kept; `ternary` gives B values in {-1, 0, 1} only, all three present. - The new YAML cases, plus every existing pre_checkin fast_check case and the existing large and offset nightly cases, to show that none is refused. - Device self-test: a pattern selected on one thread survives another thread's pattern scope ending. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Every fast_check case and self-test passed or skipped for a stated reason on gfx90a, gfx942 and gfx950, with none failing (counts in #12876). Every large-K case passed on gfx942 and gfx950, including fp8 fnuz on gfx942 and OCP fp8 on gfx950, and the 32 pattern cross-checks passed on all three GPUs. With the old `sparse_k` (the same K indices in every row), the new coverage check in the self-test fails. **Review fixes, October 7.** With newling's two fixes for this PR, at stack tip 6296886715 (ROCm 10.2, before the rebase onto develop), the full fast_check filter ran 501 tests on gfx90a, 544 on gfx942 and 564 on gfx950, with 0 failures; the rest passed or skipped for a stated reason. newling's new self-tests each fail before their fix. The conjugate-transpose layout fix and the rebase onto develop came after that run; on the current tip (86d4ac3), the full fast_check filter on gfx942 ran 561 tests with none failing, including the 28 fp8 fnuz cases, which now use develop's negate fix. Earlier development runs on gfx90a (MI210): - With the refusal disabled, the bound test failed as expected. - A first version that used `ternary` with bf16 output took 45 s per case; at K = 32768 a typical ternary result is near 256, so most results rounded and went to the host. bf16 now uses only `sparse_k`. - One run crashed with SIGILL in the bf16 `ternary` unit_check case, and a #13005 run crashed with a segfault in another bf16 unit_check case. Both were in the host reference, on a login node shared with other GPU jobs (load average 44). The same case passed in 6 repeated runs, so I could not reproduce it. ## 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, but it changes the shared integer_exact fill, and the new refusal applies to every fast_check case. Every existing case is far below the limit. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

…as gradients) (#13015) JIRA ID : AIHPBLAS-5017 Stack, in merge order: #12876, #12960, #13005, #13013, **#13015 (this PR)**, #13021, #13022, #13025, #13031. The diff here is only the side-operand layer; the base moves down the stack as those merge. ## Summary fast_check started with A, B, C, alpha, beta, bias and the scaleAlpha vector. Everything else a real call reads or writes was outside the check: scaleA through scaleE, the E output, amaxD and the bias gradients. A wrong-address read of a scale is invisible when every scale is 1, so these need data that makes a wrong read visible. This PR folds the scale factors into the fingerprints, using integer scales from 1 to 10 so a wrong read changes D. It checks E (including relu and clamp) element by element, and checks amaxD and the bias gradients exactly. Every output fast_check checks (D, E, the bias gradient and amaxD) is reset before every launch, so a solution cannot inherit the previous one's correct values. On MI300X and MI350X many of these options are only offered for fp8 inputs, so the tests aim there. Each new case also runs `unit_check` on the same result, as a cross-check, and a scaleB mutation makes the scale self-test fail. Cases that hit open defects are quarantined in `known_bugs.yaml`; see Defects found below. **CI time added** (one MI300X / one MI350X, ROCm 10.2, sum of per-test times): about 2 s / 1 s in pre_checkin (31 cases). Nothing in nightly. ## Motivation fast_check covered only A, B, C, alpha, beta, bias and the scaleAlpha vector. Every other operand a GEMM reads or writes was outside the check: scaleA through scaleE, the E output, amaxD and bias gradients. The AIHPBLAS-4988 defects include wrong-address reads of exactly these (ROCM-32046 reads the wrong MX scales), and a read of the wrong scale is invisible when every scale is 1. ## Technical Details Only client test code and test data change. - **Scales.** D = scaleD · (alpha · diag(scaleAlpha .* scaleA) · A · B · diag(scaleB) + beta · scaleC · C + bias). That is still linear in the result, so the scales fold into the probe sums: scaleA (one value or one per row) joins the row factor, scaleB (one value or one per column) scales the row probe and the column sums, and scaleC and scaleD fold into beta, alpha and the bias. This matches the client CPU reference's order. The bound from #13013 includes them, and a product of factors that reaches 2^61 is refused rather than overflowed. The values come from the generic integer fill (1 to 10), so reading the wrong one changes D. Non-integer scales (`norm_check`, fp8 outputs) are refused with a message. I kept integers rather than the 2⁻² to 2² suggested on the ticket: fractions would need dyadic rather than integer modular arithmetic, and integer scales from 1 to 10 already make a wrong read visible. - **E.** E, the pre-activation result times scaleE, is checked as a second linear problem with scaleE in place of scaleD. It is filled with the sentinel before every launch and its padding is scanned. - **Activations.** relu and clamp have exact integer forms. D is then not linear, so it is checked element by element as scaleD · act(E / scaleE) from the verified E. GELU, SiLU and sigmoid are refused, and an activation without `use_e` is refused. A scaleE of 0 leaves nothing to recover the pre-activation from, and the check says so. - **amaxD** must equal the largest |activation| (or |D| / |scaleD| without one) over the verified result. It is filled with all ones (NaN) before every launch, so a skipped store fails. The kernels store amaxD outright from the last workgroup rather than taking an atomic max into it, so the starting value is free. If the activation check cannot read D or E back, the expected amax is NaN rather than a stale zero, so the caller sees it as unavailable. - **Bias gradients.** BGRADA must equal the exact row sums of op(A) over K, and BGRADB the column sums of op(B). Each gradient's partial sums are bounded on their own before they are added up, since the GEMM bound includes the other operand, which may be zero; a gradient that cannot be computed exactly in the compute type is refused. The gradient buffer gets the sentinel before every launch, as D and E do. Bias source d with no activation is refused: the harness keeps a plain BIAS epilogue for it, so there is no gradient to check. - **Bias padding.** When `bias_stride` exceeds M, the gap before the next batch's vector holds poison, so a kernel that reads past M produces a wrong D. New pre_checkin cases each run fast_check and unit_check on the same result: - fp8 (gfx950) and fp8 fnuz (gfx942) inputs with scaleA and scaleB as scalars (bf16 and f32 output) and as vectors (bf16 output, lda and ldb packed); - E with relu and clamp; - BGRADA and BGRADB in every transpose for f16 (excluded on gfx11 and gfx12 other than gfx125x, which offer no f16 bias gradient, as the existing gradient cases are), and BGRADB for bf16; - a strided bias with a poisoned gap; - fast_check-only copies (`unit_check: 0`) of the cases quarantined only because the CPU reference is wrong (ROCM-32062, ROCM-32139), so fast_check keeps checking the GPU there; the aux copy also pads E (lde 140), so a kernel that writes E with D's leading dimension shows. The cases are aimed at configurations the libraries actually offer. A `hipblaslt-bench` probe on MI300X and MI350X found scaleA and scaleB only for fp8 inputs, vector scales only with a 16-bit output, no mix of one scale value and a vector, scaleC and scaleD only with an fp8 output (which the integer scales saturate), amaxD only with an fp8 output on gfx942 and not at all on gfx950, and a bf16 bias gradient only from B with B transposed. gfx90a offers none of the scales or the f32/bf16 gradients. So the scaleC, scaleD and amaxD checks are exercised by device self-tests only; no supported configuration reaches them on gfx90a, gfx942 or gfx950. ## Defects found Three defects turned up while cross-checking with unit_check. Each is filed and quarantined in `known_bugs.yaml` with a time-box: - **ROCM-32138**: hipBLASLt accepts `HIPBLASLT_MATMUL_DESC_EPILOGUE_AUX_SCALE_POINTER` but never passes it to the kernel (it stops at `RocblasltContractionProblem::scaleE`), so E is written unscaled on every GPU. fast_check and unit_check both report E off by exactly the scale factor. Case: `matmul_fast_check_aux_scale_e_ROCM32138`. - **ROCM-32139**: the client CPU reference writes E with D's leading dimension, so unit_check fails a correct GPU result whenever lde differs from ldd; fast_check passes the same result. Case: `matmul_fast_check_aux_padded_ldd_ROCM32139`. - **ROCM-32062** (filed earlier, for the scaleAlpha vector): the same loop in the client CPU reference reads `scaleAVec[i / k]`, so with vector scales and a padded lda the reference returns NaN or Inf. `hipblaslt-bench -v` agrees, and passes with lda packed; fast_check accepts the GPU result either way. Cases: `matmul_fast_check_scales_fp8_fnuz_vec_lda_ROCM32062` and `matmul_fast_check_scales_fp8_vec_lda_ROCM32062`, quarantined under ROCM-32062. ## Test Plan - Host self-tests: a D built with every scale passes, and one built with a single wrong scaleB entry fails and is located in that column; BGRADA and BGRADB pass when exact, and one wrong entry each is reported by index; factors whose product reaches 2^61 are refused; a gradient row of {2^24, 1, -2^24}, whose f32 sum depends on the order of addition, is refused. - Device self-tests: relu and clamp D pass against E with scaleD and scaleE, amax equals the largest |activation|, amaxD uses |scaleD| and reports a zero scale, and one wrong D element is reported with its coordinates; amax is NaN when E cannot be copied back, and the clamp check uses the float bound the kernel receives. - Mutation check: building the row probe without scaleB must fail the scale self-test. - The new pre_checkin cross-check cases, and every existing pre_checkin fast_check case. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Every fast_check case and self-test passed or skipped for a stated reason on gfx90a, gfx942 and gfx950, with none failing (counts in #12876). All 31 of this PR's pre_checkin cases ran on gfx942 and on gfx950, including the fast_check-only copies; the 3 quarantined cases skipped. The f16 bias-gradient cases now skip on the gfx1201 runner, where Math CI's precheckin had failed them with "NO solution found". **After newling's review.** With the gradient bound and the amax reset, at stack tip 6296886715 (ROCm 10.2, before the rebase onto develop), the full fast_check filter ran 501 tests on gfx90a, 544 on gfx942 and 564 on gfx950, with 0 failures. newling's new self-tests each fail before their fix. Earlier development runs: - gfx90a (MI210): with the scaleB mutation, the scale self-test failed. - Cluster runs built with ROCm 7.1 and run on ROCm 7.2: the first versions of the scales, amaxD and 32-bit gradient cases found no solution on either GPU, which led to the probe described above. ## 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, but it adds new checking paths to the shared harness, and the amaxD, scaleC and scaleD paths are covered by self-tests rather than by real kernels. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

…hresholds (#13021) JIRA ID : AIHPBLAS-5018 Stack, in merge order: #12876, #12960, #13005, #13013, #13015, **#13021 (this PR)**, #13022, #13025, #13031. The diff here is only the size-threshold layer; the base moves down the stack as those merge. ## Summary Index, offset and launch-size bugs only appear past a line: 2^31 or 2^32 elements, 2 or 4 GiB, or 2^32 work items. A test that sits below the line, or only on one side of it, can't tell a fix from luck. And the existing big-shape tests had no tier: their category matched no ctest tier, so only TheRock's unfiltered runs picked them up, and on a machine without enough memory they crashed instead of skipping. This PR adds pairs of cases just below and just above each threshold, so a failure brackets the overflow. A new `stress` tier gives these large-memory tests a place to run, and a memory gate skips cleanly (with the numbers) when the device or host is too small. The threshold cases ran on MI300X and MI350X. Cases that hit open defects are quarantined on the platforms they affect; see Defects found below. **CI time added** (one MI300X / one MI350X, ROCm 10.2, sum of per-test times): nothing in pre_checkin or nightly. The new `stress` tier runs these 49 / 39 cases in about 3.5 min / 1 min, with 10 / 12 more skipping for memory or no solution. That tier also picks up 37 older `category: stress` tests that use the host reference, which I haven't timed. These fast_check cases skip in a run with no gtest filter (TheRock runs `hipblaslt-test` unfiltered in its standard lanes), so they run only where a tier or filter selects them. ## Motivation Classes A, B and D on AIHPBLAS-4988 only appear past a size threshold: - an element index or byte offset at 2^31 or 2^32 (ROCM-31230, ROCM-31016); - a 2 or 4 GiB buffer limit (ROCM-31258, ROCM-32050); - a launch of 2^32 work items (ROCM-31290). A test that sits well below the threshold, or only on one side of it, can't tell a fix from luck. These cases put one shape just below each threshold and one just above it, so a failure brackets the overflow. Two things stood in the way, both in GitHub issue 12337. `category: stress` matched no tier in `test_categories.yaml`, so no ctest tier ran the existing stress cases (TheRock's unfiltered runs did). And a host-memory shortfall crashed the harness instead of skipping. ## Technical Details Only client test code, test data and the test tier configuration change. - **Stress tier.** `test_categories.yaml` gains a `stress` tier that selects `*stress*`, with a 4 h timeout and `stress` / `weekly` / `large_memory` labels, for large-memory runners or a weekly run. No CI workflow selects those labels today; scheduling the tier is a separate decision. This also picks up the existing stress cases. - **No surprise in unfiltered runs.** TheRock's `test_hipblaslt.py` runs `hipblaslt-test` with no filter outside its quick lane, which would have run these cases there. fast_check cases in the `stress` category now skip when no filter selects anything (the positive part of the filter is empty or `*`); the ctest tier patterns and any explicit filter still run them. The older host-reference stress cases are unchanged. A self-test covers which filters count. - **Memory gate.** Before allocating, a fast_check case adds up the device memory (A, B, C, D, E, bias, the scale vectors and the workspace, times the rotating-buffer count) and host memory (fast_check's unpadded copies, the padded copies when a host reference also runs, and the scale buffers) it needs. It skips with the amounts when hipMemGetInfo's free memory, or the host's available memory, is short. On Linux that is `MemAvailable`, which counts reclaimable page cache where sysinfo's freeram does not (the second half of issue 12337); on Windows it is `GlobalMemoryStatusEx`. The estimate errs high (a C copy is counted even when beta is 0), and the older host-reference stress cases are not gated. - **`allow_no_solution`.** A sweep can reach shapes the library declines. With this set, the case skips instead of failing, because declining a shape is a correct answer. A regression that stops offering solutions for a supported shape won't show up in these cases; the pre_checkin and nightly cases still cover that. The cases (all `category: stress`, fast_check only): | Case | Threshold | | --- | --- | | `threshold_lda`, `threshold_ldb` | last element of A or B at 2^31 and 2^32, through the leading dimension with K = 32 (f32, bf16) | | `threshold_ldd` | last element of C and D at 2^31 and 2^32, every transpose (bf16) | | `threshold_batch_stride` | batch 1 of A or B ending just below 2^31 and 2^32, or starting just past 2^31 | | `threshold_batch_stride_2_32_ROCM32140` | batch 1 of A, B, or C and D starting just past 2^32 (quarantined, below) | | `threshold_d_bytes_*` | D at 2 GiB and 4 GiB for f32, bf16 and int8, N = 256, M not a multiple of 256 (the ROCM-31016 trigger); the bf16 2 GiB pair with every solution | | `threshold_custom_kernel_ROCM31258` | xf32 on gfx950, every solution including custom kernels, store extent at and past 2 GiB | | `threshold_cd_4gib_ROCM32050` | f32 32768 × {32767, 32768, 32769} × 4096, C and D around exactly 4 GiB (quarantined on gfx1250, below) | | `threshold_tiles`, `threshold_tiles_streamk` | D of 65536 × 65536 (exactly 2^32 elements) and one column of 16 more. The heuristic picks the tile, so these bracket D's element count rather than a fixed launch size; the Stream-K variant is the one that has found a launch-size defect (ROCM-31826) | | `threshold_batch_count` | 65535 and 65537 batches | ## Defects found **ROCM-32140**: `ContractionSolution.cpp` packs batch strides into the kernel arguments as `uint32_t`, and no predicate refuses a stride of 2^32 elements or more. A stride of 2^32 + 2 becomes 2, so the GEMM returns success with batch 1 wrong. On gfx90a bf16, all 16 of the heuristic's first solutions failed for stride_a and for stride_b; strides that end just below 2^32 pass. `hipblaslt-bench -v` confirms it independently: CPU-reference error 0 at stride_a = 4294963200 and 1.09 at 4294967298. Quarantined in `known_bugs.yaml` with a time-box. The C and D entry of that case was added after the quarantine and has not run, since a quarantined case always skips. The cluster runs found two more, both quarantined on the platforms they affect: - **ROCM-32152**: on gfx950, some direct-to-LDS kernels (`DTLA1_DTLB1`, DepthU 32) write wrong values into large outputs, a different set of elements on each run. It fails the d_bytes, tiles and custom-kernel cases, and `hipblaslt-bench -v` with solution 357903 at 4,194,303 x 256 x 48 failed 5 runs out of 5 on plain allocations, on two nodes. gfx942 passes the same cases. - **ROCM-31826** (already filed for gfx950): MT16x16 Stream-K kernels leave 256 elements of D unwritten at 65,536 x 65,552, starting just past element 2^32. The Stream-K tile case shows it on gfx942 (4 solutions) and gfx950 (51 solutions). **ROCM-32050** (filed earlier): an illegal memory access on gfx1250 at the C and D 4 GiB shape, which kills the test process. That case is quarantined on gfx1250 and gfx1250-strict; no gfx1250 GPU was available to run it here. ## Test Plan - Self-tests: the memory gate returns empty when everything fits, and names device or host memory when one is short; which gtest filters count as selecting the stress cases. - Every new stress case on gfx90a (MI210, 64 GiB), gfx942 and gfx950. - Every existing fast_check self-test. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Every fast_check case and self-test passed or skipped for a stated reason on gfx90a, gfx942 and gfx950, with none failing (counts in #12876). The stress cases with an explicit filter: 49 ran on gfx942 and 39 on gfx950, all passing; the rest skipped for memory, no solution, or a quarantine. A stress case selected by name runs; under an unfiltered run it skips with the reason. Earlier development runs on gfx90a (MI210): - Stress cases: 43 passed, none failed. 21 skipped: 3 quarantined ROCM-32140 cases, 14 shapes the library declines (large lda in bf16, large ldb), and 4 tile cases that needed 16 GiB while another user's job held most of this shared GPU, which shows the gate working as intended. - Before quarantine, the stride_a and stride_b entries past 2^32 failed on all 16 solutions. - The Windows gfx110X lane failed the memory-gate self-test before the gate learned to read Windows host memory. ## 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, test data and a new test tier, with cases that need up to about 33 GiB. The memory gate, `allow_no_solution` and the unfiltered-run skip apply only to fast_check cases, and the new tier only selects `category: stress`. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

JIRA ID : AIHPBLAS-5021 Stack, in merge order: #12876, #12960, #13005, #13013, #13015, #13021, #13022, **#13025 (this PR)**, #13031. This layer adds only the on-demand SDC-hunt recipe. ## Summary Intermittent corruption reports from our stress tools are hard to chase, because the open questions are about the environment. Does it need a second workload on the GPU? Does preemption matter? Does XNACK matter? Does a 4 GiB crossing matter? Without a repeatable way to vary those and write down what happened, each run is an anecdote. This PR adds a script that runs fast_check cases across those conditions: no load, a competing GEMM, a kernel that occupies some of the CUs, or any command you give it, under each `HSA_XNACK` setting. It records the full environment with every result (driver and ROCm versions, CWSR and noretry settings, the GPUs the test sees, the failing solutions and iterations, and any buffer crossing a 4 GiB boundary), so failures can be lined up against conditions. It also adds `sdc_hunt` cases at the stress-tool sizes and types. No tier selects the `sdc_hunt` cases, and an unfiltered run skips them, so this adds nothing to CI; the README explains why it's for dedicated hardware. On an MI210 it ran cleanly with and without a competing GEMM. The load made each run about 14 times slower, so it really does contend for the GPU. **CI time added**: none. One pass of the f32 and bf16 cases takes about 75 s on an idle MI210 and much longer under load; the fp8 cases take about 8 s on an MI300X or MI350X. ## Motivation ROCM-31901 and ROCM-31905 are intermittent silent data corruption (SDC) reports from Achilles and driftcheck. Several failures needed a second GPU workload running, and the open questions are environmental: does it depend on concurrency, CWSR preemption (ROCM-31285), the XNACK mode, or a buffer that crosses a 4 GiB boundary (a class C carry drop, like ROCM-32046)? fast_check is the cheapest way to answer them. It checks every element on every launch and names the bad elements, solution and iteration. What was missing was a repeatable way to run it across those axes and record the environment with each result. ## Technical Details A client script with its CPU tests and README, test data, and one harness skip change. - **`clients/scripts/sdc_hunt/fast_check_sdc_hunt.py`** runs hipblaslt-test fast_check cases under each combination of: - `--load`: none; a second GEMM (`hipblaslt-bench`); the CU-occupying cotenant kernel (`cotenant:<CUs>`, waiting for its READY); or any command, such as a copy loop or the EDP helper; - `--xnack`: `HSA_XNACK` unset, 0 or 1. It appends one JSON line per run with the GPUs `hipblaslt-test` sees (following `HIP_VISIBLE_DEVICES`) and their architecture, the kernel, amdgpu and ROCm versions, the amdgpu `cwsr_enable` and `noretry` parameters, `HSA_XNACK`, `TENSILE_SOLUTION_SELECTION_METHOD`, the load, the pass, fail and skip counts, the failing solutions and iterations, and every buffer a failure reports crossing a 4 GiB boundary. CWSR needs a driver reload to change, so it is recorded rather than toggled. A failing run's full output is saved next to the results. - **A run is clean only if it really ran.** A load that fails to start or exits early fails the run, as does a run in which no test ran; `--fail-on-skip` also fails a run with any skipped test. Stopping the script with Ctrl-C, `kill` or a dropped session stops its background load (which runs in its own session), any of the load's children that outlive it, and the test it was running. The test binary and load commands are checked before the first run, and an empty `command:` load is refused rather than recorded as an uncontended pass. - **New cases, `category: sdc_hunt`:** the Achilles GEMM sizes and types (f32, bf16 with bf16 and f32 output, and fp8 in all four format orders, using `sparse_k` so results stay exact), 20 checked launches on each of the heuristic's first 4 solutions. No tier selects `sdc_hunt`, and like the stress cases (#13021) these skip in a run with no gtest filter, so they run only when asked for. - **Logs.** Each invocation gets a unique ID, so two invocations appending to the same results file within a second keep separate failure logs. Log names use only the load's kind, not its full command, which stays in the JSON record. - A README describes the axes, how to run it, how to run the script's own tests, and why it is for dedicated hardware rather than CI (ROCM-32049). A failure that lines up with a 4 GiB crossing is described as a lead for controlled placement, not proof of a dropped carry. ## Test Plan - Run the recipe on gfx90a over two XNACK settings, with and without a concurrent GEMM. - The fp8 cases on gfx942 (FNUZ) and gfx950 (OCP). - The script's own behavior: skip counting and `--fail-on-skip`, a load that cannot start, `cotenant:0` rejected, and cleanup after SIGTERM. - `test_fast_check_sdc_hunt.py`: seven CPU-only tests using real subprocesses and a fake test binary, covering the failure parser, skip counting, an empty load command, a relative test binary, log separation across invocations, a long load command, and cleanup of a load child that ignores SIGTERM. No CI lane runs them yet; wiring them in is left for a follow-up, since it means a new workflow. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** The fp8 SDC-hunt cases, all four format orders, passed on gfx942 (FNUZ, 8 cases) and gfx950 (OCP, 8 cases). Script behavior, in the ROCm 10.2 image on gfx90a: - a filter with two quarantined cases records `tests_skipped: 2`, and `--fail-on-skip` fails that run; - a load command that does not exist is recorded as `load_failed` with the reason, instead of stopping the script; - `cotenant:0` is rejected at startup; - SIGTERM to the script during a run under a `sleep` load: the script, the load and `hipblaslt-test` were all gone afterwards. After newling's review: the seven runner tests pass (`python -m unittest discover -s clients/scripts/sdc_hunt -v`, from `projects/hipblaslt`), and each of newling's new tests failed before its fix. Load and XNACK runs on gfx90a (MI210, amdgpu 6.16.6, ROCm 10.2.0, cwsr_enable 1, noretry -1), with the script before this round's cleanup changes: | HSA_XNACK | Load | Result | Time | | --- | --- | --- | --- | | unset | none | 9 passed | 73 s | | unset | bf16 8192³ GEMM loop | 9 passed | 1032 s | | 0 | none | 9 passed | 76 s | | 0 | bf16 8192³ GEMM loop | 9 passed | 1035 s | The concurrent GEMM makes each run about 14 times slower, which shows the load really contends for the GPU. Each time was within 6 s of the same run under ROCm 7.1. On the synthetic report, the parser extracted the failing test, the failing solution with its iterations, and the crossing buffer. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): a standalone script and on-demand test data that no tier or CI job runs, plus the one-line category check in the harness's unfiltered-run skip. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com> Co-authored-by: randyh62 <42045079+randyh62@users.noreply.github.com>

JIRA ID : AIHPBLAS-5019 Stack, in merge order: #12876, #12960, #13005, #13013, #13015, #13021, #13022, #13025, **#13031 (this PR, the top)**. This layer adds only MX block-scale coverage. ## Summary MX formats store small values with a shared E8M0 scale for each block of 32. Some address bugs read the scale tensors from the wrong place. fast_check refused MX entirely, and with the usual MX test data every scale is 1.0. So a read from the wrong scale block often changes nothing, and even a full host check can miss it. This PR adds exact MX data under `integer_exact`: small-integer elements, and scales of 1, 2 or 4 chosen per block, so a misread scale almost always changes D (two blocks share a scale one time in three). The scale buffers can also be placed across 4 GiB boundaries using #12960, and failure reports include them. All 16 MX cases pass on MI350X. Every gfx950 MXFP8 solution is a Stream-K kernel, so these cases run the same kernel family as ROCM-32046; that defect itself is on gfx1250, which I couldn't run, and the cases are ready to run there. The PR also fixes two harness issues found while bringing it up: MX operands now keep their host buffer when fast_check runs without `unit_check`, and the edge shape now respects the MX kernels' multiple-of-16 requirement for M and N. **CI time added** (MI350X only; MI300X doesn't run MX; ROCm 10.2, sum of per-test times): about 38 s in pre_checkin (12 cases), and about 3 min in nightly (4 cases). The MX cases are slow for their count because their K is long and their shapes are big. ## Motivation ROCM-32046 drops the carry in the address of both the A-scale and B-scale tensors of MXFP8 kernels. fast_check refused MX block scales, so nothing in the fast_check suite read those tensors. And with the usual MX test data, where every E8M0 scale is 1.0, a read of the wrong scale block usually changes nothing, so even a full reference check can miss it. ## Technical Details Only client test code and test data change. - **Exact MX data.** With `integer_exact`, the MX generator replaces the generated fp8 elements with small integers (A in {0, 1, 2}, B in {-2, ..., 2}) and the E8M0 scales with 1, 2 or 4, chosen per block of 32. The values are replaced before the bytes are copied for the GPU and before the scale swizzle and alignment, and the reference floats are recomputed from those bytes, so the dequantized values the GPU sees are the ones fast_check checks against. Every element of D is an exact integer. - **fast_check on MX.** fast_check now accepts MX scales on both A and B with fp8 (E4M3 or E5M2) elements, UE8M0 scales, K a multiple of the scale block, one batch, and lda and ldb equal to the rows stored. It checks D against the dequantized A and B. MX operands keep their host copy, since the generator writes them on the host, and the host-memory preflight counts that copy and both float copies of each MX input. Every device case is TN; the other layouts take a different host path, which the self-test checks. - **Scale placement.** `placement: scale_a` and `scale_b` put the MX scale buffers across a 4 GiB boundary (AIHPBLAS-4994). The poison 4 GiB away is 0x28, which E8M0 reads as 2^-87, so a scale read whose carry was dropped collapses its block toward zero and D comes out wrong instead of faulting. - **New cases** (gfx950 and gfx1250): - `matmul_fast_check_mx_fp8` (pre_checkin): every fp8 format pair at 256 x 256 x 512 and 1040 x 528 x 1024, the heuristic's first 16 solutions, each launched and checked twice. The MX kernels need M and N to be multiples of 16, so 1040 x 528 is the edge-tile shape. - `matmul_fast_check_mx_fp8_placement` (pre_checkin): the scale buffers, and A and B, placed across 4 GiB. - `matmul_fast_check_mx_fp8_large_k` (nightly): K = 16,384 and 2,560 at 3072 x 3072 and 2560 x 16384, where the scale tensors span many blocks. - **Host self-test.** The generated elements and scales are the intended values and more than one scale is used. It also decodes every element and its block's scale from the data and scale buffers, independently of the generator, and requires the reference floats to match, in all four A/B transpose layouts, for both E4M3 and E5M2. With the reference reading the neighbouring block's scale, it fails at element 0. ## Test Plan - Host self-test `integer_exact_mx_values_are_exact_with_varied_scales`, plus that mutation. - Every new MX case on an MI350X (gfx950). - Every fast_check case on gfx942 and gfx950, to check that the stack as a whole still passes with this layer on top. ## Test Result **ROCm 10.2, at the stack tip (3364b79, this PR).** Every fast_check case and self-test passed or skipped for a stated reason on gfx90a, gfx942 and gfx950, with none failing (counts in #12876). All 16 MX cases passed on gfx950: `matmul_fast_check_mx_fp8` (8), `_placement` (4) and `_large_k` (4). After newling's review (the preflight count and the E5M2 decode), at stack tip 6296886715 (ROCm 10.2, before the rebase onto develop), the full fast_check filter on gfx950 ran 564 tests with 0 failures, and the MX self-test passed for both formats. Earlier development runs on gfx950: - The first run crashed in every MX case: with fast_check and no unit_check, `testing_matmul` gave A and B an empty host buffer, and the MX generator writes into it. MX operands now keep their host copy. - 1031 x 517 found no MX solution; `hipblaslt-bench` showed the MX kernels need M and N to be multiples of 16, so the edge shape is now 1040 x 528. K = 1056 also found none, so K stays at 1024 for that shape. - gfx1250 was not run. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): test code and test data only. The shared MX generator changes only when `integer_exact` is requested, which only these fast_check cases do; existing MX tests get the same data as before. --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>

assistant-librarian · 2026-10-12T14:18:05Z

TheRock Submodule Bump Activity

Newest first

2026-10-12 — Included in TheRock via ROCm/TheRock#8933.

…oLds prefetch (#13268) JIRA ID : ROCM-32152, ROCM-32277 Targets `develop`. It removes the quarantine that the fast_check stack (#13022, AIHPBLAS-5020, now merged) added for these cases. ## Motivation Two gfx950 defects that looked different turned out to be one race. ROCM-32152: some bf16 and MX solutions that load A and B directly to LDS (direct-to-LDS, `DTLA1_DTLB1`) wrote wrong values into large outputs, different ones each run. ROCM-32277: two f32 Stream-K solutions sometimes wrote small wrong values in the last, partial tile of rows, about one launch in 20. Both are persistent kernels, which loop over many output tiles in one launch. Right after storing one tile, such a kernel issues the first prefetch of the next tile. With direct-to-LDS loads, that prefetch writes LDS at offset 0, where the store had just staged the bias and scaleAlpha vectors. Kernels with three LDS buffers (`DtlPlusLdsBuf`) and PrefetchGlobalRead 2 or more had no wait or barrier there. So a fast wave's prefetch could overwrite those vectors while a slower wave was still reading them, and a few rows of D came out scaled by A or B data instead. ## Technical Details - `directToLdsM0Update` in `KernelWriterAssembly.py` now takes its existing `s_waitcnt lgkmcnt(0)` and `s_barrier` path before a persistent tile's first direct-to-LDS prefetch, when PrefetchGlobalRead is 2 or more and Prefetch Across Persistent is off (that mode issues the prefetch from the previous tile). PrefetchGlobalRead 1 and one-buffer kernels already sync there, and their emitted code doesn't change. - A codegen test, `tensilelite/Tests/unit/characterization/_codegen/test_streamk_dtl_prefetch_barrier.py`, emits a designed gfx950 Stream-K `DtlPlusLdsBuf` PGR2 kernel (`_designed/gfx950/streamk_dtl_pgr2.yaml`) and checks for the wait and barrier before its first direct-to-LDS load. - `known_bugs.yaml`: the four ROCM-32152 gfx950 size-threshold entries (`d_bytes_bf16`, `d_bytes_all_solutions`, `tiles`, `custom_kernel_ROCM31258`), the ROCM-32277 f32 Stream-K batch entry and the gfx950 `store_extent_ROCM31016` entry come out. That last one passed 10 of 10 runs on MI355X and 10 of 10 on MI350X with this fix. Entries that also list ROCM-31826, ROCM-32151 or ROCM-32153 stay until those land. **CI time added**: one codegen test in the TensileLite unit lane (about 30 s locally, most of it emitting one Stream-K kernel), and nothing in `pre_checkin`. On gfx950 the cases this un-quarantines run again: about 2 min more per nightly run (`store_extent_ROCM31016`, every solution, 100 to 120 s, and the f32 Stream-K batch case, about 15 s) and about 3.5 min more per stress run (the four size-threshold cases). gfx942 is unchanged, since these cases already ran there. ## Test Plan - The codegen test, with and without the fix, and the TensileLite unit suite. - On gfx950 (MI355X and MI350X): the formerly quarantined ROCM-32152 cases with 10 repeats, and the ROCM-32277 f32 Stream-K batch case with 60 repeats, with this branch's Tensile library and with the stack tip's. - A combined build with the ROCM-31826, ROCM-32151 and ROCM-32153 fixes and all five tickets' entries removed, on gfx942 and gfx950, to confirm which entries each PR can drop. - `hipblaslt-bench` timing of the affected kernels before and after. ## Test Result ROCm 10.2.0 (TheRock 10.2.0a20260928): - The codegen test fails without the fix and passes with it. TensileLite unit suite: 9,947 passed; no characterization snapshot changed. - ROCM-32152 cases (11 tests, 10 repeats): all pass on MI355X and MI350X. With the stack tip's library, 9 of 12 size-threshold cases failed on MI355X, and 13 of 40 solutions failed a 20-launch bench loop. - ROCM-32277 f32 Stream-K batch case: 0 of 120 runs fail across MI355X and MI350X with this branch, against 5 of 30 with the stack tip's library. The ticket's bench loop goes from 4 of 60 wrong to 0 of 60. - Combined build: every case these entries cover passes on gfx942, MI355X and MI350X, except the gfx950 bf16 TN every-solution and short-K cases, which fail through the ROCM-32343 and ROCM-32344 kernels. - Timing on MI355X: unchanged within noise at large K. At K = 48, where per-tile overhead dominates, one MT128x64x32 PGR4 kernel is about 5.7% slower and another is about 12% faster. - Math CI precheckin passed against develop. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Moderate (3/5): a codegen change that adds one wait and barrier per output tile to persistent direct-to-LDS kernels with PrefetchGlobalRead 2 or more (about 3,000 gfx950 logic entries). Results change only where the race made them wrong; the cost is small at large K and mixed at very small K. --------- Co-authored-by: Cursor <cursoragent@cursor.com>

… batches (#13267) JIRA ID : ROCM-32151 Targets `develop`. It removes the quarantine that the fast_check stack (#13022, AIHPBLAS-5020, now merged) added for these cases. ## Motivation Some bf16 GEMM solutions on gfx942 and gfx950 return success with about half the right value in parts of D, usually the last, partial tile of rows. The ticket first suspected the Stream-K partial-tile path. It's actually the store path, in kernels that split the summation across waves (LocalSplitU, or LSU, greater than 1). For beta not equal to 0, kernels with StoreSyncOpt or StorePriorityOpt and arch-VGPR accumulation apply alpha early, at the top of each store batch. With LSU > 1, each batch after the first copies its reduced sums into the registers alpha was meant to scale only after that point. So alpha scales stale registers, and the real values are stored unscaled: `alpha·AB` becomes `AB`, which is half with alpha = 2. Edge tiles store in smaller batches, so they're the ones that have a second batch. ## Technical Details - `KernelWriterAssembly.py` sets `alphaBeforeLoadC` only when `LocalSplitU == 1`. LSU > 1 kernels apply alpha after the copy, as the default path already does. This is a one-line condition plus a comment. - A codegen test, `tensilelite/Tests/unit/characterization/_codegen/test_lsu_alpha_before_loadc.py`, emits a real gfx942 LSU = 2 solution (one of the ticket's kernels, copied into `test_data/_logic/gfx942/`) and checks that no store batch scales `ValuC` before its LSU copy. - `known_bugs.yaml`: the ROCM-32151 entries come out on gfx942. On gfx950 these bf16 cases still fail through kernels from other defects, so their entries stay under those tickets: the custom-schedule MT256x160x64 kernels (ROCM-32343) and the custom MT256x256x64 kernels at K = 64 (ROCM-32344) in the TN cases. The bf16 Stream-K batch case on gfx950 also failed through persistent direct-to-LDS kernels (ROCM-32152 and ROCM-32277). #13268 has since fixed those, so that entry comes out, along with the gfx950 bf16 `size_t_output_overflow_all_solutions` entry #13268 left under ROCM-32151. Only kernels with LSU > 1, StoreSyncOpt or StorePriorityOpt and arch-VGPR accumulation change; the emitted code for neighbouring kernels is byte-identical. **CI time added**: one codegen test in the TensileLite unit lane (about 30 s locally), and nothing in `pre_checkin`. On gfx942 the two bf16 Stream-K cases this un-quarantines skip under CI's default solution selection, since the library offers Stream-K solutions there only with `TENSILE_SOLUTION_SELECTION_METHOD=2`. With it set they take about 1.7 min per nightly run on an MI325X. On gfx950 two nightly cases run again: the bf16 Stream-K batch case (about 25 s) and the bf16 `size_t_output_overflow_all_solutions` case (about 9 s). ## Test Plan - The codegen test, with and without the fix, and the TensileLite unit suite. - The ROCM-32151 bf16 fast_check cases on gfx942 (`TENSILE_SOLUTION_SELECTION_METHOD=2`) and gfx950, with this branch's Tensile library and with the stack tip's. - The ticket's four `hipblaslt-bench -v` repros on gfx942. ## Test Result ROCm 10.2.0 (TheRock 10.2.0a20260928): - The codegen test fails without the fix and passes with it. TensileLite unit suite: 9,947 passed; no characterization snapshot changed. - gfx942 (MI325X): 26 of 30 bf16 cases fail with the stack tip's library, all through LSU > 1 kernels; 30 of 30 pass with this branch. The ticket's four bench repros go from CPU-reference errors of 0.0106, 0.299, 0.028 and 0.0194 to 0. - gfx950 (MI355X): 16 failing tests with the stack tip's library, 5 with this branch, none of them through LSU > 1 kernels. The 5 are the ROCM-32343, ROCM-32344 and ROCM-32152 kernels above. - With #13268's fix combined, both newly un-quarantined gfx950 cases pass on MI355X: the bf16 Stream-K batch case 20 of 20 and `size_t_output_overflow_all_solutions` 10 of 10. - Math CI precheckin passed against develop. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Moderate (3/5): a codegen change in the store path, limited to LSU > 1 kernels with the store-optimization flags and beta not equal to 0, where it was wrong. Those kernels apply alpha a little later in the store, which may cost a few cycles; I haven't benchmarked it. --------- Co-authored-by: Cursor <cursoragent@cursor.com>

… by #14147 and #12676 (#14390) JIRA ID : ROCM-32343, ROCM-32344 ## Motivation The gfx950 custom-schedule `MT256x160x64` kernels used to write wrong values in the TN every-solution tests (ROCM-32343 for bf16, ROCM-32153 for f16). We quarantined the failing cases in `known_bugs.yaml` so the nightlies could stay green. #14147 fixed the cause: the schedule's prologue issued the A loads before the B loads, and the main loop's partial wait counts assumed B first. That fix is on develop, so the quarantines have done their job. This PR removes the four entries they left behind, so the cases run again as real regression gates. The ROCM-32344 entry can go too. It quarantined `matmul_fast_check_short_k` (bf16 TN) because the custom `MT256x256x64` kernels gave wrong values at K = 64. #12676 fixed those kernels on Oct 7, and the quarantine landed a day later in #13022. Davood's retest on 59fb065 found no issue. ## Technical Details - `known_bugs.yaml` loses the two bf16 entries (`matmul_fast_check_every_solution_repeat` and `matmul_fast_check_streamk_every_solution_repeat`, `a_type: bf16_r`) and the two f16 entries (the same two names, `a_type: f16_r`, `M: 4099`, TN) on `gfx950`, along with their comments. - `known_bugs.yaml` also loses the ROCM-32344 entry (`matmul_fast_check_short_k`, `a_type: bf16_r`, `transA: T`) on `gfx950`, along with its comment. - Test data only. No library, kernel, or client code changes. ## Test Plan - Compare the test records `hipblaslt_gentest.py` generates on develop and on this branch. Only the previously quarantined cases should change. - Build develop plus this change for gfx950 and run the formerly quarantined cases on MI350X and MI355X. ## Test Result - `hipblaslt_gentest.py` generates the same 143,631 records on both sides. 24 records change, and all of them just lose `gfx950` from `known_bug_platforms`: the two every-solution names at 9 records each (8 bf16 plus the f16 TN 4,099 case) and `matmul_fast_check_short_k` at 6 (bf16 TN, K = 1, 15, 16, 33, 64 and 129). Nothing else changes. - MI350X (Device 75a0): `*fast_check_every_solution_repeat*` and `*fast_check_streamk_every_solution_repeat*` (without the large-K and xf32 variants) ran 48 tests, 48 passed. - MI355X (Device 75a3): the same filter ran 48 tests, 48 passed. - For comparison, before the fix my A/B on #13276 showed the bf16 TN defaults wrong in 33 of 41 runs on MI350X on a develop that predated #14147, and 0 of 70 with the fix. - I did not run these on MI350P (Device 75a8). The cluster has none. The 75a8 tables carry copies of the same schedule (2 f16 and 19 bf16 solutions), which #14147 should fix the same way. The first check of those is the MI350P nightly after TheRock's rocm-libraries pin moves past #14147 (it is at 565c4c7 today). - Required checks passed before this update. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (1/5): test data only, and these cases are nightly-category, so PR checks are unaffected. If the MI350P nightly still shows failures on the 75a8 copies, we put the quarantine back for that device's cases and look at those kernels again. MI350X and MI355X already pass. ## Risk Acceptance / Waivers None needed. Related: ROCM-32153, #13269, #13793, #13940 (closed), #14147. --------- Co-authored-by: Cursor <cursoragent@cursor.com>

github-actions Bot added the project: hipblaslt label Oct 2, 2026

assistant-librarian Bot added the organization: ROCm label Oct 2, 2026

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 908e5ee to af51666 Compare October 3, 2026 01:39

tony-davis added this pull request to stack #12961 October 3, 2026 01:40

tony-davis mentioned this pull request Oct 3, 2026

test(hipblaslt): add an on-demand fast_check SDC-hunt recipe #13025

Merged

1 task done

tony-davis requested a lite review from Copilot October 3, 2026 02:55

Copilot started reviewing on behalf of tony-davis October 3, 2026 02:56

View session

Copilot AI reviewed Oct 3, 2026

View reviewed changes

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from af51666 to 15fb6bd Compare October 3, 2026 04:01

tony-davis mentioned this pull request Oct 3, 2026

test(hipblaslt): keep fast_check exact at large K #13013

Merged

1 task done

tony-davis requested a lite review from Copilot October 3, 2026 13:41

Copilot started reviewing on behalf of tony-davis October 3, 2026 13:42

View session

Copilot AI reviewed Oct 3, 2026

View reviewed changes

Comment thread projects/hipblaslt/clients/tests/data/matmul_gtest.yaml Resolved

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 15fb6bd to 6867bdf Compare October 3, 2026 14:51

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 6867bdf to 269e941 Compare October 3, 2026 16:19

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 269e941 to 49246ac Compare October 3, 2026 16:28

tony-davis requested a lite review from Copilot October 3, 2026 19:31

Copilot started reviewing on behalf of tony-davis October 3, 2026 19:32

View session

Copilot AI reviewed Oct 3, 2026

View reviewed changes

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 49246ac to 4ff7765 Compare October 3, 2026 19:44

This was referenced Oct 7, 2026

fix(hipblaslt): apply alpha after the LocalSplitU copy in later store batches #13267

Merged

fix(tensilelite): sync waves before a persistent tile's first DirectToLds prefetch #13268

Merged

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 97e1481 to c0560d9 Compare October 7, 2026 18:56

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from c0560d9 to 1778fa2 Compare October 7, 2026 20:03

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 1778fa2 to 147bf65 Compare October 7, 2026 21:05

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 147bf65 to 144726c Compare October 8, 2026 16:17

tony-davis and others added 5 commits October 8, 2026 12:03

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-5020-edge-configs branch from 144726c to d9daa77 Compare October 8, 2026 17:03

tony-davis merged commit eae510a into develop Oct 9, 2026
122 of 123 checks passed

tony-davis mentioned this pull request Oct 30, 2026

test(hipblaslt): remove the gfx950 TN custom-kernel quarantines fixed by #14147 and #12676 #14390

Merged

1 task done

Navigation Menu

Repository navigation

test(hipblaslt): run fast_check on edge, bookkeeping and per-kernel configurations - #13022

test(hipblaslt): run fast_check on edge, bookkeeping and per-kernel configurations#13022
tony-davis merged 5 commits into
users/todavis_amdeng/AIHPBLAS-5018-size-thresholdsfrom
users/todavis_amdeng/AIHPBLAS-5020-edge-configs

tony-davis commented Oct 2, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

therock-pr-bot Bot commented Oct 3, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Oct 3, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

therock-pr-bot Bot commented Oct 3, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

tony-davis commented Oct 7, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 12, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

tony-davis commented Oct 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Risk level

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🔵 Needs a closer look

Uh oh!

therock-pr-bot Bot commented Oct 3, 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!

codecov Bot commented Oct 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Uh oh!

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

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟢 Approval recommended

Uh oh!

tony-davis commented Oct 7, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 12, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

tony-davis commented Oct 2, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 3, 2026 •

edited

Loading

codecov Bot commented Oct 3, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 3, 2026 •

edited

Loading