Repository navigation
fix(hipblaslt): preserve FNUZ zero and NaN during negation - #13119
Merged
Merged
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #13119 +/- ##
===========================================
- Coverage 72.10% 72.10% -0.00%
===========================================
Files 2861 2861
Lines 472618 472620 +2
Branches 69810 69812 +2
===========================================
Hits 340755 340755
- Misses 108125 108127 +2
Partials 23738 23738
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Alex-Vasile
approved these changes
Oct 6, 2026
Alex-Vasile
left a comment
Alex-Vasile
left a comment
Contributor
There was a problem hiding this comment.
Looks good, thanks for the fix.
tony-davis
approved these changes
Oct 6, 2026
tony-davis
left a comment
tony-davis
left a comment
Contributor
There was a problem hiding this comment.
Thanks newling! I reviewed this against the hipBLASLt PR quality checklist, and I think it's ready, so I'm approving.
What I checked:
- The fix reaches every caller. The host fills (
hipblaslt_init.hpp,hipblaslt_vector.hpp) and the device fills (hipblaslt_init_device.cpp) all go through thisnegate(), so integer_exact B can no longer pick up a NaN from a negated zero. - The tests catch the bug. I built this branch for gfx942 (ROCm 10.2, TheRock 10.2.0a20260928) and ran your four tests on an MI300A, both as written and with the two fnuz
negate()specializations reverted to the plain sign flip. With the fix, 4 of 4 pass. Reverted, 4 of 4 fail: the scalar tests at 0x00 and 0x80, and the GEMMs at the norm assertion (testing_matmul.hpp:1363). - The norm check is the right guard.
unit.hpp:156-158passes a NaN when the CPU reference is NaN too, sounit_checkalone can't see this. - Nearby tests: 16,903 standard-tier fnuz tests passed with the fix, with none failing, before my 3-hour job ran out of time.
- In this PR's own CI, all four new tests passed in TheRock's gfx94X hipBLASLt shards (3, 4 and 5).
The CI failures don't look related to me. Shard 5's one failure is matmul_size_t_output_overflow_bf16_ROCM26455, which couldn't allocate 9 GB of host memory on the runner and then segfaulted. The ASAN quick job hung before any test started, after amd-smi static failed on the runner. codecov/patch reads 0% because headers under clients/common/include aren't instrumented (lcov lists hipblaslt_math.hpp with 0 measured lines), so nothing in this diff can count.
A few optional suggestions:
smoke_gtest.yamlholds the other 35 smoke cases. This one might fit better there than inmatmul_gtest.yaml, where it's the only smoke case.norm_check_assert: trueis already the default (hipblaslt_common.yaml:804), so it could go. I don't mind it staying for clarity, since the comment explains why.- In the description, the "Shared automated checks have not run yet" line is out of date now, and the required checklist box is unchecked.
Once this lands, I'll rebase our stack so #13013 drops its copy of the change.
newling
enabled auto-merge (squash)
October 6, 2026 16:53
Extract the two FNUZ negate guards from #13013 (4e3c347). Alternating-sign client initialization could otherwise turn zero into NaN. Add smoke regressions covering all 256 encodings of each FNUZ format, comparing finite nonzero values with float negation and preserving the zero and NaN encodings. Both tests fail on the original helper and pass with the fix; no GEMM is needed.
newling
force-pushed
the
users/newling/fix-fnuz-negation
branch
from
October 6, 2026 18:34
296581c to
b23ef99
Compare
tony-davis
added a commit
that referenced
this pull request
Oct 8, 2026
…GEMMs (#12876) JIRA ID : AIHPBLAS-4988 Stack, in merge order: **#12876 (this PR)**, #12960, #13005, #13013, #13015, #13021, #13022, #13025, #13031. Independent of the stack: #13027 (rocisa 64-bit add lowering) and #13060 (address-carry lint). ## Summary Recently we fixed some bugs where address, index or launch-size math was done in 32 bits when it needed 64. That kind of bug only shows up at big sizes or unlucky addresses, and it usually lives in a handful of solutions out of hundreds. Our tests have a hard time reaching it, because `unit_check` compares against an OpenBLAS GEMM on the host. That costs M·N·K and copies every padded buffer to the host, so the big-shape tests stay small and run only the heuristic's top pick, once. `integer_exact` already got rid of the tolerance, but not the cost. This PR adds fast_check, an opt-in check (`fast_check: 1`) for `integer_exact` tests. Instead of recomputing the GEMM, it compares a few random "fingerprints" of D with the same fingerprints computed from the inputs (Freivalds' algorithm, in exact modular arithmetic). That costs about as much as reading the matrices once, D never leaves the GPU, and a second set of fingerprints pins down which elements are wrong. Padding is filled with poison so stray reads show up, and D starts as a sentinel so missed writes show up. Every failure names the solution, its kernel and where each buffer sits in memory. Only test code changes. To make sure it catches what it's meant to, we reverted the ROCM-31016 fix (dropped stores past 2 GiB). fast_check flagged the affected solutions and pointed at the exact element where the stores stopped. It has also found real defects: on gfx950 the ROCM-31016 shape and the bf16 every-solution case fail today (ROCM-32152 and ROCM-32151), so those two cases are quarantined there, and running it next to `unit_check` found a bug in the CPU reference (ROCM-32062). **CI time added**, measured on one MI300X / one MI350X under ROCm 10.2 at the stack tip (the sum of gtest's per-test times, not counting test-data parsing at startup): - This PR: about 32 s / 15 s in pre_checkin (108 cases; the 51 self-tests take under a second), and about 3.5 min / 20 s in nightly (35 / 33 cases; gfx950 runs fewer because of the two quarantines above). - The whole stack: about 37 s / 58 s in pre_checkin, about 15 min / 9 min in nightly, and about 3.5 min / 1 min in the new `stress` tier (#13021), which needs large-memory runners. On MI300X the Stream-K cases skip unless `TENSILE_SOLUTION_SELECTION_METHOD=2` is set; with it they add about 5.5 min. ## Motivation hipBLASLt's gtest suite checks a matrix multiply (GEMM) result with `unit_check`, which compares the GPU output against a reference computed on the CPU. That reference costs time proportional to M·N·K, and the harness copies every input and output buffer, padding included, into pinned host memory. Tests for out-of-bounds indexing need large leading dimensions or large shapes, so this cost limits them. On gfx942, `matmul_size_t_index_overflow` with ldd = 1049100 spends 12 s of its 14.4 s allocating host memory. The existing overflow tests therefore keep K small and run only the heuristic's first solution, even though indexing defects such as ROCM-31016 and ROCM-31258 affect specific kernels. ## Technical Details This PR adds an opt-in YAML argument, `fast_check: 1`, for tests that use `initialization: integer_exact`. Under integer_exact, every input is a small integer, so the correct result is an exact integer as well. Only client test code and test data change. The library, TensileLite and kernel selection are untouched. - **Every element of D is checked without a CPU GEMM.** fast_check uses Freivalds' algorithm, which multiplies D by a random vector `r` and compares the result with `A·(B·r) + C·r + bias·sum(r)`. That costs time proportional to M·K + K·N + M·N instead of M·N·K. It does the same from the left with a second vector, to locate bad columns. The sums are taken modulo the prime 2^61 − 1, so the comparison is exact, and a wrong result passes one probe with probability below 2^−60. The input side is computed on the host once per test. D's side is computed on the GPU for each solution, so D is never copied to the host. When the probes disagree, the failure message lists the mismatched elements with their row, column and element offset, and marks offsets at or above 2^31 or 2^32. - **Failure context.** Each failure names the solution: its position in the test's list, its library index and its kernel. It also lists every operand's address range and flags any that crosses a 4 GiB boundary, since some address defects depend on placement rather than size. - **Out-of-bounds reads.** The padding of A, B and C (rows past the matrix inside the leading dimension, and gaps between batches) is filled with 40, a value outside the integer_exact range [−2, 2]. A kernel that reads padding therefore produces a wrong D. - **Out-of-bounds writes and missed stores.** Before each solution runs, D's whole allocation is filled with 0xFF bytes, which form a NaN in f16, bf16 and f32. Afterwards a device kernel checks that every element outside the M × N × batch region still holds that pattern, and that no element inside it does. A failed launch of any of these fills or scans fails the test; it is never read as clean. - **Results above the exact range.** bf16 holds integers exactly only up to 256. fast_check compares each larger result with its correctly rounded value (round to nearest, ties to even, pinned by a self-test), so f16 and bf16 tests can use K above the 256 limit that `unit_check` imposes on integer_exact. Those elements are recomputed exactly, in parallel, on the host. - **Supported configurations:** alpha, beta, bias, `scaleAlpha_vector`, strided batches, and every transpose, including conjugate transpose, which stores real data the same way as transpose. An output with zero rows, columns or batches needs no check, so it reads no operands and launches nothing. The test fails with a stated reason for timing runs, activations, scale factors, grouped GEMM and pointer-array batches; later PRs in the stack lift several of these. Output types run against real kernels here are f32, f16 and bf16; the stack adds int8 and fp8 cases (#13013, #13022). The code also accepts f64 and fp8 outputs, which only the self-tests exercise. New files are `clients/common/include/fast_check.hpp`, `clients/common/src/fast_check.cpp` and `clients/tests/src/fast_check_gtest.cpp`. `testing_matmul.hpp` calls them, and in fast_check-only runs it skips the CPU reference and the padded host copies. Running both checks on the same results found ROCM-32062, a separate defect in the CPU reference. When `scaleAlpha_vector` is combined with a leading dimension of A larger than its row count, `cast_mul` in `cblas_interface.cpp` scales the wrong rows of A, so `unit_check` fails a correct GPU result. No existing test combined the two. That case, `matmul_fast_check_epilogue_padded_lda_ROCM32062`, is quarantined in `known_bugs.yaml` until the fix PR removes the entry; #13015 adds a fast_check-only copy that keeps checking the GPU there. Related: ROCM-26455, ROCM-31016 (#12281), ROCM-31230 (#12332), ROCM-31258 (#12482), and the original overflow tests (#8231). ## Test Plan - `FastCheck` (host only) and `FastCheckDevice` (small device buffers, no hipBLASLt kernels) in `fast_check_gtest.cpp`. These build correct results, inject faults, and check that each fault is reported: a wrong element, a pair of errors that cancel in a plain row sum, a shifted row, a duplicated column, an unwritten element, an out-of-bounds write, wrong rounding of a large bf16 result (including a tie), an empty buffer, an empty output (zero M, N or batch count, with null unused operands), and an unsupported type. They also require the device and host passes to reach the same verdict. - pre_checkin, 126 cases: fast_check and `unit_check` on the same results across precisions and transposes, including all nine N/T/C pairs on two rectangular f32 shapes with padded leading dimensions and three batches; the size_t index-overflow shapes in f32, f16 and bf16; and the ROCM-26455 shape with bias and `scaleAlpha_vector`. - nightly, 35 cases: element offsets and a batch base offset above 2^32; shapes up to 16385 × 16383 × 4097 with partial tiles; and every solution at the output-overflow and ROCM-31016 shapes. - Regression replay: revert the ROCM-31016 threshold and run `matmul_fast_check_store_extent_ROCM31016`. - `standard` tier (smoke, quick, pre_checkin) to confirm existing tests are unchanged. ## Test Result **ROCm 10.2, at the stack tip (3364b79, which includes this PR).** Built and run in a ROCm 10.2.0 image (TheRock 10.2.0a20260928). Every fast_check case and self-test, with none failing: | GPU | Ran | Passed | Skipped | | --- | --- | --- | --- | | gfx90a (MI210) | 496 | 410 | 86 | | gfx942 (MI300X) | 539 | 473 | 66 | | gfx950 (MI350X) | 559 | 486 | 73 | **After newling's review.** With the empty-output fix, at stack tip 6296886715 (ROCm 10.2, before the rebase onto develop), the same filter ran 501 tests on gfx90a, 544 on gfx942 and 564 on gfx950, with 0 failures. newling's conjugate-transpose case failed 10 of its 18 cases before the fix (all those with a C operand) and passed all 18 after, on gfx1201 with ROCm 7.1. On the current tip (86d4ac3, with both fixes, rebased onto develop with #13119), the full fast_check filter on gfx942 (MI300X) ran 561 tests: 495 passed, 66 skipped for a stated reason, and none failed, including all 18 conjugate-transpose cases and 28 fp8 fnuz cases. The skips are quarantined known bugs, shapes the library declines, Stream-K cases that need `TENSILE_SOLUTION_SELECTION_METHOD=2`, and on the shared MI210, cases that need more free memory than it had. Earlier development runs, on gfx942 (MI300A): - Self-tests: each of six deliberate source mutations, four in the host code and two in the GPU kernels, made at least one self-test fail. - With the ROCM-31016 fix reverted, 170 of the 1,093 solutions failed. The first reported 1793 elements never written, starting at element offset 2147481599, which is the 2^32-byte boundary for bf16. With the fix in place, all 1,093 passed in 144 s. - With the fix reverted, each failure names the kernel, for example `Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x448x32_…` (library index 75677). The placement record shows C and D, each about 4 GiB, crossing a 4 GiB boundary. - Same shapes, `unit_check` before and fast_check after: ldd = 1049100 f32 went from 14.4 s to 0.7 s, and lda = 1049100 f32 from 8.3 s to 1.2 s. - `standard` tier (smoke, quick, pre_checkin), run with the stack's second PR: 71,379 run, one failure, `matmul_fallback_compute_fp16`, which passes at the stack tip under ROCm 10.2 on an MI300X (see #12960). ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5): test code and test data only, with no library or device behavior change, and existing tests take the same paths unless they set `fast_check`. It is rated 3 rather than 2 because it adds shared machinery to `testing_matmul.hpp` and new GPU kernels that every architecture runs, and about half a minute of pre_checkin. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: James Newling <james.newling@gmail.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
tony-davis
added a commit
that referenced
this pull request
Oct 9, 2026
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>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : AIHPBLAS-5016
Motivation
Users can reach this bug through
hipblaslt-bench --initialization integer_exactwith FNUZ FP8 inputs, and throughhipblaslt-testcases using the same initializer. While generating matrix B, the client negates selected values by flipping the sign bit. In both FNUZ formats, that turns zero (0x00) into NaN (0x80), introducing NaNs before the matrix reaches GEMM. It also turns NaN into zero.This is a bug in the benchmark/test input generator. Applications supplying their own matrices to
hipblasLtMatmuldo not call this helper.Technical Details
Preserve zero and NaN when negating FNUZ values, while continuing to flip the sign of finite nonzero values. This extracts the helper fix from #13013 and can land independently of that stack.
Test Plan
Add scalar smoke tests covering all 256 encodings of each format, plus two 32×32×32 GEMM smoke tests using the existing device initializer and FP32 accumulation/output on gfx942. The GEMM tests enable the existing norm assertion: ordinary unit comparisons accept matching CPU/GPU NaNs and otherwise miss this corruption.
Run with the original and fixed helpers:
hipblaslt-test --gtest_filter='Math_smoke.*:*matmul_integer_exact_fnuz_negate_zero*'Test Result
With the original helper, both scalar tests fail at zero/NaN and both GEMMs fail at the norm assertion. With the fix, all four tests pass with no skips on an MI300X. The filtered process takes approximately 5.6 seconds including startup.
Validation used a source-built client and host library with ROCm 7.2.1 and the official gfx942 device-kernel package
10.2.0a20261005. Shared automated checks have now run on296581c545f. TheRock multi-architecture CI and Math CI passed, including all gfx942 and gfx950 hipBLASLt shards in TheRock. The separate hipBLASLt ASAN, Codecov, and gfx1250 checks have not passed.Submission Checklist
Risk level
Low (2/5): changes client input generation for zero and NaN in two formats, with scalar and real-GEMM regression coverage. The change is confined to client helpers and tests.