test(hipblaslt): place fast_check operands across 4 GiB boundaries by tony-davis · Pull Request #12960 · ROCm/rocm-libraries

tony-davis · 2026-10-02T17:02:13Z

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

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

codecov · 2026-10-02T18:29:42Z

Codecov Report

❌ Patch coverage is 89.06250% with 7 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...ojects/hipblaslt/clients/common/src/fast_check.cpp	88.33%	3 Missing and 4 partials ⚠️

❌ Your project check has failed because the head coverage (49.33%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage.

Additional details and impacted files

@@                              Coverage Diff                               @@
##           users/todavis_amdeng/AIHPBLAS-4988-fast-check   #12960   +/-   ##
==============================================================================
  Coverage                                          71.97%   71.97%           
==============================================================================
  Files                                               2864     2864           
  Lines                                             473750   473780   +30     
  Branches                                           70045    70046    +1     
==============================================================================
+ Hits                                              340962   340991   +29     
- Misses                                            109156   109157    +1     
  Partials                                           23632    23632

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`49.33% <ø> (ø)`
TensileLite-Unit	`83.29% <ø> (ø)`
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from f13a6f5
hipBLASLt	`41.45% <89.06%> (+0.07%)`	⬆️
hipCUB	`82.68% <ø> (ø)`		Carriedforward from f13a6f5
hipDNN	`86.64% <ø> (ø)`		Carriedforward from f13a6f5
hipFFT	`43.43% <ø> (ø)`		Carriedforward from f13a6f5
hipRAND	`76.12% <ø> (ø)`		Carriedforward from f13a6f5
hipSOLVER	`68.96% <ø> (ø)`		Carriedforward from f13a6f5
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from f13a6f5
rocBLAS	`48.25% <ø> (ø)`		Carriedforward from f13a6f5
rocFFT	`47.62% <ø> (ø)`		Carriedforward from f13a6f5
rocRAND	`57.45% <ø> (ø)`		Carriedforward from f13a6f5
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from f13a6f5
rocSPARSE	`74.64% <ø> (ø)`		Carriedforward from f13a6f5
rocThrust	`93.15% <ø> (ø)`		Carriedforward from f13a6f5

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
...slt/clients/common/include/hipblaslt_arguments.hpp	`30.00% <100.00%> (+0.38%)`	⬆️
...pblaslt/clients/common/src/hipblaslt_arguments.cpp	`89.68% <100.00%> (+0.13%)`	⬆️
...ojects/hipblaslt/clients/common/src/fast_check.cpp	`91.24% <88.33%> (+0.12%)`	⬆️

🚀 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.

therock-pr-bot · 2026-10-02T20:42:10Z

❌ PR Check — Action Required

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	❌ Fail	Error: Check concluded with `cancelled`.
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

⚠️ 1 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

❌ pre-commit

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-10-02T20:42:13Z

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

Copilot

Copilot review overview

🟡 Changes recommended

Moderate unresolved issues affect poison detection, workspace coverage and allocation, diagnostics, and test correctness.

Review effort: Lite
Findings: 1

Open (1)

Unused workspace allocation can cause placement-run OOM · New

What changed in this PR

Adds opt-in HIP virtual-memory placement across 4 GiB boundaries to hipBLASLt fast_check tests.

Changes:

Adds placed buffers, poison windows, and scan helpers.
Integrates placement arguments and test cases for operands and workspace.
Updates buffer moves, self-tests, and documentation.

File	Reviewed change
`projects/hipblaslt/TESTING.md`	Documents placement coverage.
`projects/hipblaslt/clients/tests/src/matmul_gtest.cpp`	Adds placement to test names.
`projects/hipblaslt/clients/tests/src/fast_check_gtest.cpp`	Adds placement self-tests.
`projects/hipblaslt/clients/tests/data/matmul_gtest.yaml`	Adds placement test cases.
`projects/hipblaslt/clients/tests/data/hipblaslt_common.yaml`	Defines placement arguments.
`projects/hipblaslt/clients/common/src/hipblaslt_arguments.cpp`	Initializes placement settings.
`projects/hipblaslt/clients/common/src/fast_check.cpp`	Implements poison scanning helpers.
`projects/hipblaslt/clients/common/include/testing_matmul.hpp`	Integrates placed operands and workspace.
`projects/hipblaslt/clients/common/include/hipBuffer.hpp`	Adds placed-buffer ownership and safe moves.
`projects/hipblaslt/clients/common/include/hipblaslt_arguments.hpp`	Adds placement fields.
`projects/hipblaslt/clients/common/include/hip_placement.hpp`	Implements virtual-memory placement and verification.
`projects/hipblaslt/clients/common/include/fast_check.hpp`	Declares scan helpers.

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

Copilot

Copilot review overview

🟡 Changes recommended

Unresolved findings remain in validation, failure handling, workspace allocation, self-test correctness, and MXFP8 coverage.

Review effort: Lite
Findings: 2

Open (2)

Avoid undefined float-to-integer type punning · New
Unused workspace allocation can cause placement-run OOM

Copilot

Copilot review overview

🔵 Needs a closer look

Unresolved validation, portability, and invalid-request handling issues remain.

Review effort: Lite
Findings: None

Resolved since last review (2)

Avoid undefined float-to-integer type punning
Unused workspace allocation can cause placement-run OOM

newling

Just the one agent's finding about a situation where RAII can save a memory leak: https://github.com/newling/agent-public-reviews/blob/main/rocm-libraries/review_pr12960.md

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

Good catch @newling, thanks! I took the unique_ptr fix as is, with you as author.

… boundaries JIRA ID : AIHPBLAS-4994 Some address defects depend on where a buffer sits rather than on the problem size. ROCM-32046 drops the carry out of the low 32 bits of an address, so it goes wrong only when an operand crosses a 4 GiB boundary, and the wrong access lands 4 GiB away, outside every allocation. fast_check's padding poison and D sentinel cannot see such an access. On gfx1250 the page fault it causes can leave the GPU unusable (ROCM-32049). - A new placement test argument names one operand (a, b, c, d, bias, scale_alpha_vec or workspace), plus an offset. PlacedRegion, in hip_placement.hpp, uses HIP virtual memory management to map that operand across a 4 GiB boundary, with poison windows 4 GiB below and above it. A dropped-carry read returns the poison value, so fast_check reports a wrong D instead of the GPU faulting. A dropped-carry write changes the poison, which is checked after every solution. - HipDeviceBuffer gains a placement constructor. Its move operations now clear the moved-from pointer, so a moved-from buffer no longer re-checks guard bytes in memory it no longer owns. - Self-tests use a kernel that drops the carry on purpose, plus a placed buffer moved by vector growth. - New YAML cases place each operand in turn on a small shape. The pre_checkin cases run the heuristic's first 16 solutions and the nightly cases run every solution. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

…kspace is placed A placement run that places the workspace still allocated the full normal workspace, which nothing used: placement requires fast_check, fast_check refuses timing runs, and the checked solutions use the placed workspace. For a large workspace that extra allocation could run out of memory before the placed one. The normal workspace is now empty in that case. Co-authored-by: Cursor <cursoragent@cursor.com>

fast_check_poison_bits and fast_check_count_changed_device had no self-test, and Codecov reported 69% patch coverage for this layer. A new self-test checks the poison bits and the changed-element count, with its first index, in every supported type. Co-authored-by: Cursor <cursoragent@cursor.com>

…ion uses a workspace On gfx12 every solution of the all-solutions workspace-placement cases needs no workspace, and the case failed with "placing the workspace requires a nonzero workspace size" in the Math CI precheckin run. Whether a workspace is needed depends on the library and GPU, so that case now skips. Co-authored-by: Cursor <cursoragent@cursor.com>

…e placement self-test The self-test read a float's bits through a uint32_t pointer, which breaks strict aliasing and can be miscompiled. It now copies them with memcpy. Co-authored-by: Cursor <cursoragent@cursor.com>

…orted type run_scan launched nothing for an element size it does not handle and returned success with zero counters, so the padding scan and the changed-poison count reported a clean result for a type they never looked at. It now returns hipErrorInvalidValue, and a self-test checks both callers fail. Co-authored-by: Cursor <cursoragent@cursor.com>

…e the boundary A placed workspace put the 4 GiB boundary at half the largest solution's workspace, so any solution needing less than half of it never reached the boundary and still passed as placed. The workspace region is now twice the largest size, and each solution starts so that its own reported range straddles the boundary while the whole workspace size, which any solution may use, stays mapped after it. The harness now fails a placement case whose buffer was not placed or whose range does not cross the boundary, instead of passing while testing nothing. The mapped tail of a placed buffer, after its last element, is poisoned and scanned like the windows 4 GiB away, and the poison fill reports launch errors. Co-authored-by: Cursor <cursoragent@cursor.com>

(cherry picked from commit 81db146)

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

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

…onfigurations (#13022) 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; #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 #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 - [x] 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`. --------- Co-authored-by: Cursor <cursoragent@cursor.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-09T18:03:52Z

TheRock Submodule Bump Activity

Newest first

2026-10-09 — Included in TheRock via ROCm/TheRock#8881.

tony-davis added this pull request to stack #12961 October 2, 2026 17:02

github-actions Bot added documentation project: hipblaslt labels Oct 2, 2026

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

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from de4efac to 9164eca Compare October 2, 2026 20:34

tony-davis changed the title ~~test(hipblaslt): AIHPBLAS-4994 place fast_check operands across 4 GiB boundaries~~ test(hipblaslt): place fast_check operands across 4 GiB boundaries Oct 2, 2026

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from 9164eca to f160088 Compare October 3, 2026 01:02

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

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

View session

Copilot AI reviewed Oct 3, 2026

View reviewed changes

Comment thread projects/hipblaslt/clients/common/include/testing_matmul.hpp Resolved

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from f160088 to 1de16cc Compare October 3, 2026 04:01

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

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/src/fast_check_gtest.cpp Resolved

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from 1de16cc to b5eb839 Compare October 3, 2026 14:51

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

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

View session

Copilot AI reviewed Oct 3, 2026

View reviewed changes

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from 39bfa67 to c641e13 Compare October 3, 2026 19:44

tony-davis requested a review from newling October 6, 2026 15:08

newling reviewed Oct 6, 2026

View reviewed changes

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from 8351692 to d0482ee Compare October 7, 2026 18:56

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from d0482ee to f0c3e90 Compare October 7, 2026 20:03

tony-davis and others added 8 commits October 7, 2026 16:55

fix(hipblaslt): release workspace on placement exits

4e3dcc9

(cherry picked from commit 81db146)

tony-davis force-pushed the users/todavis_amdeng/AIHPBLAS-4994-placement branch from f0c3e90 to 4e3dcc9 Compare October 7, 2026 21:05

tony-davis requested a review from newling October 7, 2026 21:20

newling approved these changes Oct 7, 2026

View reviewed changes

tony-davis merged commit 4fb52ae into develop Oct 8, 2026
131 of 149 checks passed

tony-davis deleted the users/todavis_amdeng/AIHPBLAS-4994-placement branch October 8, 2026 17:03

tony-davis mentioned this pull request Oct 9, 2026

test(tensilelite): add an address-carry lint for generated kernels #13060

Closed

1 task done

tony-davis mentioned this pull request Oct 12, 2026

fix(hipblaslt): sync the matmul test stream before destroying it #13517

Merged

1 task

Navigation Menu

Repository navigation

test(hipblaslt): place fast_check operands across 4 GiB boundaries - #12960

test(hipblaslt): place fast_check operands across 4 GiB boundaries#12960
tony-davis merged 8 commits into
users/todavis_amdeng/AIHPBLAS-4988-fast-checkfrom
users/todavis_amdeng/AIHPBLAS-4994-placement

tony-davis commented Oct 2, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Oct 2, 2026 •

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

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

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Copilot AI left a comment

Uh oh!

newling left a comment

Uh oh!

tony-davis commented Oct 7, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 9, 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!

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

Uh oh!

Codecov Report

Uh oh!

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

Uh oh!

❌ PR Check — Action Required

Uh oh!

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

🟡 Changes recommended

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🟡 Changes recommended

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Copilot review overview

🔵 Needs a closer look

Uh oh!

newling left a comment

Choose a reason for hiding this comment

Uh oh!

tony-davis commented Oct 7, 2026

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 9, 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

codecov Bot commented Oct 2, 2026 •

edited

Loading

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

edited

Loading

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

edited

Loading