You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
Independent of the fast_check stack; based on develop. The address-carry lint that started out in this PR now has its own, #13060.
Summary
While looking at how rocisa lowers 64-bit adds on targets without a native one (gfx90a, gfx942, gfx950), we noticed that splitSrcs copies an immediate into both 32-bit halves. That's right for the packed-F32 operations that share it, but for a 64-bit ptr += imm it would also add imm × 2^32 and move the pointer by hundreds of GiB. The vector form also put the immediate where gfx9 can't encode a literal, so most immediates wouldn't even assemble.
This PR fixes that lowering with a new splitSrcs64, puts a vector immediate where the hardware can encode it, and adds tests that check, and assemble, the emitted carry chain on gfx90a, gfx942 and gfx950. Nothing on develop lowers an immediate through this path yet (the immediate callers are gfx1250 TDM code, and gfx1250 has a native 64-bit add), so this closes the door before anyone walks through it. The lowering tests failed 45 of the original 57 cases before the fix, and every case passes after it.
CI time added: about 7 s in the rocisa tests (tox -e rocisa), most of it assembling the emitted pairs.
Motivation
Classes C and C2 on AIHPBLAS-4988 are mistakes in emitted 64-bit address arithmetic, so some can be caught before a kernel ever runs. There have been at least three 64-bit address bugs in a few weeks: ROCM-31230; ROCM-32046, where the carry out of the low dword is dropped; and the launch-size narrowing in ROCM-31290. A lowering that adds imm × 2^32 to a pointer would be another, so it's worth fixing while nothing uses it.
Technical Details
Splitting the immediate. On targets without s_add_u64 or v_add_nc_u64, SAddU64 and VAddNCU64 lower to a 32-bit carry chain. A new splitSrcs64 splits each source for the two 64-bit adds:
a register pair splits into its two registers;
a 32-bit immediate gets its sign extension (0 or -1) in the high half;
a larger one arrives from Python as a double, and splits as a 64-bit integer up to 2^53 in magnitude, where that double is still exact; anything larger is rejected, since it may already have been rounded;
a symbol it cannot split is rejected.
The packed operations keep splitSrcs. An immediate above 2^31 used to reach the assembly as a double (s_add_u32 s4, s4, 4294967299.0).
Where a vector immediate goes. VOP2 takes a literal only in src0, and VOP3 on gfx9 takes none, so v_add_co_u32 v4, vcc, v4, <literal> does not assemble. The lowered VAddNCU64 now puts the immediate in src0 (the add is commutative). v_addc_co_u32 reads vcc as well, which leaves no room for a literal at all, so an immediate whose high half is not an inline constant (−16 to 64) is rejected with a message; every int immediate has a high half of 0 or −1. SAddU64 (SOP2) takes a literal in either position, so it is unchanged.
Tests.rocisa/test/test_lowering64.py checks the emitted pair on gfx90a, gfx942 and gfx950. For register operands, the low add must set the carry and the high add must consume it, from the next registers. For immediates (positive, negative, at ±2^31, above 2^32 and at ±(2^53 − 1)), the high half must be the sign extension or the true high 32 bits. Symbolic immediates, immediates of 2^53 or more in magnitude, and vector immediates with a non-inline high half must be rejected. Every emitted scalar and vector pair is also assembled with amdclang++ for each target, since a string check can't see an operand the hardware cannot encode.
Hand-written kernel limits are already checked on develop by Tensile/Tests/unit/test_CustomKernels_buffer_oob.py, from the ROCM-31258 fix, so this PR doesn't duplicate that.
Race detection (class G).rocjitsu race check (gfx94X-dcgpu) already runs on rocm-libraries PRs (therock-rocjitsu-race-check-linux.yml), and RocJITsu is what found the GitHub PR 12605 race. Two gaps: the job is gfx94X only, while the ROCM-31905 single-LDS-buffer kernels are on gfx1250; and on recent PRs it fails during setup (requirements-test.txt not found) before checking anything. I'd fix the setup and add gfx950 and gfx1250 coverage for those kernel families rather than add a second race detector. I haven't changed the workflow here.
Test Plan
test_lowering64.py before and after the fix, and with the vector immediate left in src1.
The full rocisa and TensileLite unit and characterization suites.
Test Result
In a ROCm 10.2.0 image (TheRock 10.2.0a20260928), on a gfx90a host, at 2ac9fca:
test_lowering64.py: 117 passed on gfx90a, gfx942 and gfx950, including the assembled pairs. With the vector immediate left in src1, 33 cases fail (assembler: "literal operands are not supported"). With the old splitSrcs, 45 of the original 57 cases failed (every immediate case, plus the symbol case).
Full TensileLite unit and rocisa suites: 10,174 passed, 380 skipped, 2 xfailed, 1 xpassed, none failed. With this host's ROCm 7.1 toolchain the same suites had 157 failures and 14 errors on develop, almost all gfx1250, TDM and MX code it cannot assemble.
Low (2/5): the change only affects 64-bit add lowering of immediates on targets without a native 64-bit add, which nothing on develop emits. rocisa is code generation, so per the hipBLASLt overlay this needs coverage on the affected architectures in CI and two code-owner approvals.
…igh halves
On targets without s_add_u64 or v_add_nc_u64 (gfx90a, gfx942, gfx950),
SAddU64 and VAddNCU64 lower to a 32-bit carry chain. splitSrcs copied an
immediate into both halves, which is right for packed operations but makes a
64-bit add of imm also add imm * 2^32, moving a pointer by hundreds of GiB
(AMD-ROCm-Internal PR 2557). An immediate above 2^31 also reached the
assembly as a double. splitSrcs64 gives the high half the sign extension of a
32-bit immediate, splits a larger one as a 64-bit integer, and rejects a
symbol it cannot split. Nothing on develop lowers an immediate this way
today; the immediate callers are gfx1250 TDM code with a native 64-bit add.
test_lowering64.py checks the emitted pair for register and immediate
operands (positive, negative, and above 2^31) on gfx90a, gfx942 and gfx950.
Tensile.Utilities.address_carry_lint reports a 64-bit address whose low
dword is updated without a carry into the high dword, the ROCM-32046 shape.
test_address_carry_lint.py checks it on synthetic cases and on kernels
generated from designed configs for gfx90a, gfx942, gfx950 and gfx1250.
Co-authored-by: Cursor <cursoragent@cursor.com>
…and constant adds
On gfx950 the lint reported every kernel with a long branch: s_getpc_b64,
s_add_i32 sN, <offset>, 4, a carried add into the PC and s_setpc_b64. The
linear scan walked past the jump and met an unrelated address use of the same
register in the next block. Following a written value now stops at an
unconditional jump (s_branch, s_setpc_b64, s_endpgm), and an add of two
constants is not an address update. Two synthetic tests cover both, and an
address update followed by a jump is still reported.
Co-authored-by: Cursor <cursoragent@cursor.com>
✅ All modified and coverable lines are covered by tests.
❌ Your project check has failed because the head coverage (49.02%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage.
…int's generated kernels
Develop now rejects every solution in gfx1250/mx_f4_nn_umlds0.yaml ("M-major
MX-scale local read for A requires VectorWidthA >= MatrixInstK // MXBlockA"),
so the generated-kernel lint case emitted no kernels and failed the
TensileLite test, coverage and Jenkins codecov checks. The gfx950 MX configs
still cover MX scale addressing.
Co-authored-by: Cursor <cursoragent@cursor.com>
…es, and split kernels in linear time
The gfx950 audit found 105 false findings in fp8 kernels where the scheduler
placed s_addc_u32 about 50 instructions after its s_add_u32, past the lint's
16-instruction window, with nothing in between writing SCC. A scalar carry is
now followed until the matching carry-in, a jump, or the first instruction
that writes SCC. Vector carries keep the short window.
The audit also timed out on 33 large code objects: kernels() kept line numbers
by padding each kernel's text with every newline before it, which is quadratic
in the file size. Kernels now carry their first line number instead.
New synthetic tests cover a far carry-in, a carry lost to an SCC write, and
line numbers in a later kernel.
Co-authored-by: Cursor <cursoragent@cursor.com>
splitSrcs64 rejected any immediate whose magnitude reached 2^63, which also
rejected -2^63, a valid signed 64-bit value. The bound is now [-2^63, 2^63).
The lowering test splits -2^63 and checks that 2^63 is still rejected.
Co-authored-by: Cursor <cursoragent@cursor.com>
…vector carry, skip unsupported targets
- The parser split the mnemonic from its operands on a space only, so
tab-separated hand-written kernels were skipped entirely and reported clean.
Tabs are now treated as spaces.
- A vector carry chain was accepted if any carry-in to the high half followed
within the window, even after another instruction overwrote the carry
register. The carry-in must now read the register the low add wrote, before
anything else writes it.
- The generated-kernel cases now skip a target the toolchain does not support
(the SupportedISA capability the other gfx1250 codegen tests use) instead of
failing because no kernels were emitted.
Co-authored-by: Cursor <cursoragent@cursor.com>
The no-carry set is incomplete for opcodes emitted by rocisa: v_add_nc_i32, v_add_i32, v_sub_nc_i32, and v_sub_i32 are defined in rocisa/include/instruction/common.hpp:3304-3308,4193-4201, and hand-written kernels also use v_subrev_u32. An address low dword updated by any of these instructions is skipped completely, so the lint can miss the exact bug it is intended to catch.
Detect TH_ATOMIC_RETURN when selecting atomic operands
The returning atomic syntax used by this tree is marked with th:TH_ATOMIC_RETURN (for example flat_atomic_dec_u32 and global_atomic_inc_u32), not a _rtn mnemonic suffix. Consequently a returning atomic takes the data = {1} branch here, excludes operand 1 (the address), and misses a pair address entirely; a dropped carry in that address will never be linted. Detect the return modifier when selecting the destination/data operands.
A Python integer above 2^31 reaches the 64-bit add lowering as a double,
which rounds integers at or beyond 2^53 in magnitude: 2^63 - 1 arrives as
2^63 and 2^53 + 1 as 2^53. Splitting such a value could add a different
offset than the one asked for, so the lowering now rejects any immediate
of magnitude 2^53 or more. The lowering test splits +/-(2^53 - 1) and
rejects 2^53, 2^53 + 1, -2^53, 2^63 - 1, -2^63 and 2^63.
Co-authored-by: Cursor <cursoragent@cursor.com>
NO_CARRY omits the valid v_sub_i32 and v_subrev_u32 instructions. Both update one VGPR dword without producing a carry, so an address decrement using either mnemonic is silently skipped instead of being checked before its next address use. Add these opcodes and regression coverage; they are supported by the gfx90a/gfx942/gfx950 ISA definitions.
The operand-index rule is wrong for the return form emitted by this tree. global_atomic_inc_u32 (and the flat atomic return form) uses dst, addr, data, saddr with th:TH_ATOMIC_RETURN, not a _rtn mnemonic; this branch therefore selects {1}, excludes the actual address, and can treat a data pair as an address. A dropped carry while updating vaddr is consequently missed. Recognize the return modifier (and add a regression using the existing emitted syntax).
…y returning-atomic form
The lint missed v_sub_i32 and v_subrev_u32 (and the gfx10+ v_subrev_nc_u32,
v_add_nc_i32 and v_sub_nc_i32), so a low-dword decrement with any of them
went unchecked. It also recognized a returning global or flat atomic only
by a _rtn suffix or glc; gfx94x and gfx950 mark it with sc0 and gfx12 with
th:TH_ATOMIC_RETURN. For those it took the address (operand 1) for data
and could take the data pair for an address. Returning atomics are now
recognized by any of the four, and stores never count as writing a
register. Tests cover each opcode and each return form.
Co-authored-by: Cursor <cursoragent@cursor.com>
Both of the previously-missed findings were real, so they're fixed in 685081e.
No-carry vector subtracts.NO_CARRY now has v_sub_i32 and v_subrev_u32, plus the gfx10+ v_subrev_nc_u32, v_add_nc_i32 and v_sub_nc_i32. A test checks each one on an address low dword.
Returning atomics. The lint now recognizes a returning global or flat atomic by _rtn, glc (gfx90a), sc0 (gfx94x and gfx950) or th:TH_ATOMIC_RETURN (gfx12 and later). For those the address is operand 1, and the data pair is no longer taken for an address. Tests cover each return form and the non-returning one.
The lint tests pass in a ROCm 10.2 image (123, including the generated gfx1250 kernels). I'm rerunning the full-library audit with the new rules and will update the description when it finishes.
v_add_co_ci_u32 and v_sub_co_ci_u32 are listed as carry consumers but not as carry producers. If either form updates the low dword of an address, the loop never enters the CARRY_OUT branch, so a missing high-dword carry is silently accepted; the current implementation returns no finding for such a case. Add these carry-producing forms to CARRY_OUT and cover a low-dword regression case (the add form is already emitted by Tensile code).
The carry-register invalidation compares raw operand strings, so it misses partial/aliased writes to a ranged carry register. Tensile emits vector carry chains with SGPR pairs such as s[20:21]; an intervening s_mov_b32 s20, ... (or another partial write) changes that carry, but this scan still accepts the later v_addc_co_u32 because s20 is not textually equal to s[20:21]. Parse the carry operand and stop when an intervening write overlaps any of its registers, with a regression test for a partial SGPR write.
…nd partial carry writes
A carry-in add (v_add_co_ci_u32, v_sub_co_ci_u32, v_addc_co_u32 and the
rest) also sets a carry out, but the lint only followed carries from the
plain carry-out adds, so one updating an address low dword without
carrying into the high dword was accepted. It also compared the carry
register as text, so a write to part of it, such as s_mov_b32 s20 while
the carry sits in s[20:21], or vcc_lo while it sits in vcc, did not
count as changing it. Carry registers are now compared by the registers
(or vcc halves) they occupy. Tests cover three carry-in forms and three
partial writes.
Co-authored-by: Cursor <cursoragent@cursor.com>
Both of the newest previously-missed findings were real too, so they're fixed in ab76f47.
Carry-in adds set a carry as well. The carry-in forms (v_add_co_ci_u32, v_sub_co_ci_u32, v_addc_co_u32 and the rest) now count as carry setters, so one that updates an address low dword has to carry into the high dword too. Tests cover three of them.
Partial writes to the carry register. Carry registers are now compared by the registers they occupy (and the halves of vcc), so s_mov_b32 s20 with the carry in s[20:21], or a write to vcc_lo with the carry in vcc, counts as changing it. Tests cover three partial writes.
The lint tests pass in a ROCm 10.2 image (129, including the generated gfx1250 kernels). The full-library audit is rerunning with these rules, and I'll update the description when it finishes.
…gfx9 can encode it
The lowered VAddNCU64 put an immediate in src1 of v_add_co_u32, where
VOP2 takes only a VGPR and VOP3 on gfx9 takes no literal, so any
immediate outside the inline-constant range failed to assemble (as it did
before the split fix). The immediate now goes in src0, which takes a
literal. v_addc_co_u32 reads vcc as well, which leaves no room for a
literal at all, so an immediate whose high half is not an inline constant
(-16 to 64) is rejected; every int immediate has a high half of 0 or -1.
The lowering test now also assembles every emitted pair for gfx90a,
gfx942 and gfx950. With the immediate left in src1, 33 of its cases fail.
Co-authored-by: Cursor <cursoragent@cursor.com>
…ts own PR
The lint and its tests move to #13060 for their own review, so this PR
is only the rocisa lowering fix and its test.
Co-authored-by: Cursor <cursoragent@cursor.com>
…GEMMs (#12876)
JIRA ID : AIHPBLAS-4988
Stack, in merge order: **#12876 (this PR)**, #12960, #13005, #13013,
#13015, #13021, #13022, #13025, #13031. Independent of the stack: #13027
(rocisa 64-bit add lowering) and #13060 (address-carry lint).
## Summary
Recently we fixed some bugs where address, index or launch-size math was
done in 32 bits when it needed 64. That kind of bug only shows up at big
sizes or unlucky addresses, and it usually lives in a handful of
solutions out of hundreds. Our tests have a hard time reaching it,
because `unit_check` compares against an OpenBLAS GEMM on the host. That
costs M·N·K and copies every padded buffer to the host, so the big-shape
tests stay small and run only the heuristic's top pick, once.
`integer_exact` already got rid of the tolerance, but not the cost.
This PR adds fast_check, an opt-in check (`fast_check: 1`) for
`integer_exact` tests. Instead of recomputing the GEMM, it compares a
few random "fingerprints" of D with the same fingerprints computed from
the inputs (Freivalds' algorithm, in exact modular arithmetic). That
costs about as much as reading the matrices once, D never leaves the
GPU, and a second set of fingerprints pins down which elements are
wrong. Padding is filled with poison so stray reads show up, and D
starts as a sentinel so missed writes show up. Every failure names the
solution, its kernel and where each buffer sits in memory. Only test
code changes.
To make sure it catches what it's meant to, we reverted the ROCM-31016
fix (dropped stores past 2 GiB). fast_check flagged the affected
solutions and pointed at the exact element where the stores stopped. It
has also found real defects: on gfx950 the ROCM-31016 shape and the bf16
every-solution case fail today (ROCM-32152 and ROCM-32151), so those two
cases are quarantined there, and running it next to `unit_check` found a
bug in the CPU reference (ROCM-32062).
**CI time added**, measured on one MI300X / one MI350X under ROCm 10.2
at the stack tip (the sum of gtest's per-test times, not counting
test-data parsing at startup):
- This PR: about 32 s / 15 s in pre_checkin (108 cases; the 51
self-tests take under a second), and about 3.5 min / 20 s in nightly (35
/ 33 cases; gfx950 runs fewer because of the two quarantines above).
- The whole stack: about 37 s / 58 s in pre_checkin, about 15 min / 9
min in nightly, and about 3.5 min / 1 min in the new `stress` tier
(#13021), which needs large-memory runners. On MI300X the Stream-K cases
skip unless `TENSILE_SOLUTION_SELECTION_METHOD=2` is set; with it they
add about 5.5 min.
## Motivation
hipBLASLt's gtest suite checks a matrix multiply (GEMM) result with
`unit_check`, which compares the GPU output against a reference computed
on the CPU. That reference costs time proportional to M·N·K, and the
harness copies every input and output buffer, padding included, into
pinned host memory. Tests for out-of-bounds indexing need large leading
dimensions or large shapes, so this cost limits them. On gfx942,
`matmul_size_t_index_overflow` with ldd = 1049100 spends 12 s of its
14.4 s allocating host memory. The existing overflow tests therefore
keep K small and run only the heuristic's first solution, even though
indexing defects such as ROCM-31016 and ROCM-31258 affect specific
kernels.
## Technical Details
This PR adds an opt-in YAML argument, `fast_check: 1`, for tests that
use `initialization: integer_exact`. Under integer_exact, every input is
a small integer, so the correct result is an exact integer as well. Only
client test code and test data change. The library, TensileLite and
kernel selection are untouched.
- **Every element of D is checked without a CPU GEMM.** fast_check uses
Freivalds' algorithm, which multiplies D by a random vector `r` and
compares the result with `A·(B·r) + C·r + bias·sum(r)`. That costs time
proportional to M·K + K·N + M·N instead of M·N·K. It does the same from
the left with a second vector, to locate bad columns. The sums are taken
modulo the prime 2^61 − 1, so the comparison is exact, and a wrong
result passes one probe with probability below 2^−60. The input side is
computed on the host once per test. D's side is computed on the GPU for
each solution, so D is never copied to the host. When the probes
disagree, the failure message lists the mismatched elements with their
row, column and element offset, and marks offsets at or above 2^31 or
2^32.
- **Failure context.** Each failure names the solution: its position in
the test's list, its library index and its kernel. It also lists every
operand's address range and flags any that crosses a 4 GiB boundary,
since some address defects depend on placement rather than size.
- **Out-of-bounds reads.** The padding of A, B and C (rows past the
matrix inside the leading dimension, and gaps between batches) is filled
with 40, a value outside the integer_exact range [−2, 2]. A kernel that
reads padding therefore produces a wrong D.
- **Out-of-bounds writes and missed stores.** Before each solution runs,
D's whole allocation is filled with 0xFF bytes, which form a NaN in f16,
bf16 and f32. Afterwards a device kernel checks that every element
outside the M × N × batch region still holds that pattern, and that no
element inside it does. A failed launch of any of these fills or scans
fails the test; it is never read as clean.
- **Results above the exact range.** bf16 holds integers exactly only up
to 256. fast_check compares each larger result with its correctly
rounded value (round to nearest, ties to even, pinned by a self-test),
so f16 and bf16 tests can use K above the 256 limit that `unit_check`
imposes on integer_exact. Those elements are recomputed exactly, in
parallel, on the host.
- **Supported configurations:** alpha, beta, bias, `scaleAlpha_vector`,
strided batches, and every transpose, including conjugate transpose,
which stores real data the same way as transpose. An output with zero
rows, columns or batches needs no check, so it reads no operands and
launches nothing. The test fails with a stated reason for timing runs,
activations, scale factors, grouped GEMM and pointer-array batches;
later PRs in the stack lift several of these. Output types run against
real kernels here are f32, f16 and bf16; the stack adds int8 and fp8
cases (#13013, #13022). The code also accepts f64 and fp8 outputs, which
only the self-tests exercise.
New files are `clients/common/include/fast_check.hpp`,
`clients/common/src/fast_check.cpp` and
`clients/tests/src/fast_check_gtest.cpp`. `testing_matmul.hpp` calls
them, and in fast_check-only runs it skips the CPU reference and the
padded host copies.
Running both checks on the same results found ROCM-32062, a separate
defect in the CPU reference. When `scaleAlpha_vector` is combined with a
leading dimension of A larger than its row count, `cast_mul` in
`cblas_interface.cpp` scales the wrong rows of A, so `unit_check` fails
a correct GPU result. No existing test combined the two. That case,
`matmul_fast_check_epilogue_padded_lda_ROCM32062`, is quarantined in
`known_bugs.yaml` until the fix PR removes the entry; #13015 adds a
fast_check-only copy that keeps checking the GPU there.
Related: ROCM-26455, ROCM-31016 (#12281), ROCM-31230 (#12332),
ROCM-31258 (#12482), and the original overflow tests (#8231).
## Test Plan
- `FastCheck` (host only) and `FastCheckDevice` (small device buffers,
no hipBLASLt kernels) in `fast_check_gtest.cpp`. These build correct
results, inject faults, and check that each fault is reported: a wrong
element, a pair of errors that cancel in a plain row sum, a shifted row,
a duplicated column, an unwritten element, an out-of-bounds write, wrong
rounding of a large bf16 result (including a tie), an empty buffer, an
empty output (zero M, N or batch count, with null unused operands), and
an unsupported type. They also require the device and host passes to
reach the same verdict.
- pre_checkin, 126 cases: fast_check and `unit_check` on the same
results across precisions and transposes, including all nine N/T/C pairs
on two rectangular f32 shapes with padded leading dimensions and three
batches; the size_t index-overflow shapes in f32, f16 and bf16; and the
ROCM-26455 shape with bias and `scaleAlpha_vector`.
- nightly, 35 cases: element offsets and a batch base offset above 2^32;
shapes up to 16385 × 16383 × 4097 with partial tiles; and every solution
at the output-overflow and ROCM-31016 shapes.
- Regression replay: revert the ROCM-31016 threshold and run
`matmul_fast_check_store_extent_ROCM31016`.
- `standard` tier (smoke, quick, pre_checkin) to confirm existing tests
are unchanged.
## Test Result
**ROCm 10.2, at the stack tip (3364b79, which includes this PR).**
Built and run in a ROCm 10.2.0 image (TheRock 10.2.0a20260928). Every
fast_check case and self-test, with none failing:
| GPU | Ran | Passed | Skipped |
| --- | --- | --- | --- |
| gfx90a (MI210) | 496 | 410 | 86 |
| gfx942 (MI300X) | 539 | 473 | 66 |
| gfx950 (MI350X) | 559 | 486 | 73 |
**After newling's review.** With the empty-output fix, at stack tip
6296886715 (ROCm 10.2, before the rebase onto develop), the same filter
ran 501 tests on gfx90a, 544 on gfx942 and 564 on gfx950, with 0
failures. newling's conjugate-transpose case failed 10 of its 18 cases
before the fix (all those with a C operand) and passed all 18 after, on
gfx1201 with ROCm 7.1. On the current tip (86d4ac3, with both fixes,
rebased onto develop with #13119), the full fast_check filter on gfx942
(MI300X) ran 561 tests: 495 passed, 66 skipped for a stated reason, and
none failed, including all 18 conjugate-transpose cases and 28 fp8 fnuz
cases.
The skips are quarantined known bugs, shapes the library declines,
Stream-K cases that need `TENSILE_SOLUTION_SELECTION_METHOD=2`, and on
the shared MI210, cases that need more free memory than it had.
Earlier development runs, on gfx942 (MI300A):
- Self-tests: each of six deliberate source mutations, four in the host
code and two in the GPU kernels, made at least one self-test fail.
- With the ROCM-31016 fix reverted, 170 of the 1,093 solutions failed.
The first reported 1793 elements never written, starting at element
offset 2147481599, which is the 2^32-byte boundary for bf16. With the
fix in place, all 1,093 passed in 144 s.
- With the fix reverted, each failure names the kernel, for example
`Cijk_Ailk_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x448x32_…` (library
index 75677). The placement record shows C and D, each about 4 GiB,
crossing a 4 GiB boundary.
- Same shapes, `unit_check` before and fast_check after: ldd = 1049100
f32 went from 14.4 s to 0.7 s, and lda = 1049100 f32 from 8.3 s to 1.2
s.
- `standard` tier (smoke, quick, pre_checkin), run with the stack's
second PR: 71,379 run, one failure, `matmul_fallback_compute_fp16`,
which passes at the stack tip under ROCm 10.2 on an MI300X (see #12960).
## Submission Checklist
- [x] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
## Risk level
Medium (3/5): test code and test data only, with no library or device
behavior change, and existing tests take the same paths unless they set
`fast_check`. It is rated 3 rather than 2 because it adds shared
machinery to `testing_matmul.hpp` and new GPU kernels that every
architecture runs, and about half a minute of pre_checkin.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
---------
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: James Newling <james.newling@gmail.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
newling
left a comment
JIRA ID : AIHPBLAS-5022
Independent of the fast_check stack; based on
develop. The address-carry lint that started out in this PR now has its own, #13060.Summary
While looking at how rocisa lowers 64-bit adds on targets without a native one (gfx90a, gfx942, gfx950), we noticed that
splitSrcscopies an immediate into both 32-bit halves. That's right for the packed-F32 operations that share it, but for a 64-bitptr += immit would also addimm × 2^32and move the pointer by hundreds of GiB. The vector form also put the immediate where gfx9 can't encode a literal, so most immediates wouldn't even assemble.This PR fixes that lowering with a new
splitSrcs64, puts a vector immediate where the hardware can encode it, and adds tests that check, and assemble, the emitted carry chain on gfx90a, gfx942 and gfx950. Nothing on develop lowers an immediate through this path yet (the immediate callers are gfx1250 TDM code, and gfx1250 has a native 64-bit add), so this closes the door before anyone walks through it. The lowering tests failed 45 of the original 57 cases before the fix, and every case passes after it.CI time added: about 7 s in the rocisa tests (
tox -e rocisa), most of it assembling the emitted pairs.Motivation
Classes C and C2 on AIHPBLAS-4988 are mistakes in emitted 64-bit address arithmetic, so some can be caught before a kernel ever runs. There have been at least three 64-bit address bugs in a few weeks: ROCM-31230; ROCM-32046, where the carry out of the low dword is dropped; and the launch-size narrowing in ROCM-31290. A lowering that adds
imm × 2^32to a pointer would be another, so it's worth fixing while nothing uses it.Technical Details
Splitting the immediate. On targets without
s_add_u64orv_add_nc_u64,SAddU64andVAddNCU64lower to a 32-bit carry chain. A newsplitSrcs64splits each source for the two 64-bit adds:The packed operations keep
splitSrcs. An immediate above 2^31 used to reach the assembly as a double (s_add_u32 s4, s4, 4294967299.0).Where a vector immediate goes. VOP2 takes a literal only in src0, and VOP3 on gfx9 takes none, so
v_add_co_u32 v4, vcc, v4, <literal>does not assemble. The loweredVAddNCU64now puts the immediate in src0 (the add is commutative).v_addc_co_u32reads vcc as well, which leaves no room for a literal at all, so an immediate whose high half is not an inline constant (−16 to 64) is rejected with a message; every int immediate has a high half of 0 or −1.SAddU64(SOP2) takes a literal in either position, so it is unchanged.Tests.
rocisa/test/test_lowering64.pychecks the emitted pair on gfx90a, gfx942 and gfx950. For register operands, the low add must set the carry and the high add must consume it, from the next registers. For immediates (positive, negative, at ±2^31, above 2^32 and at ±(2^53 − 1)), the high half must be the sign extension or the true high 32 bits. Symbolic immediates, immediates of 2^53 or more in magnitude, and vector immediates with a non-inline high half must be rejected. Every emitted scalar and vector pair is also assembled withamdclang++for each target, since a string check can't see an operand the hardware cannot encode.Hand-written kernel limits are already checked on develop by
Tensile/Tests/unit/test_CustomKernels_buffer_oob.py, from the ROCM-31258 fix, so this PR doesn't duplicate that.Race detection (class G).
rocjitsu race check (gfx94X-dcgpu)already runs on rocm-libraries PRs (therock-rocjitsu-race-check-linux.yml), and RocJITsu is what found the GitHub PR 12605 race. Two gaps: the job is gfx94X only, while the ROCM-31905 single-LDS-buffer kernels are on gfx1250; and on recent PRs it fails during setup (requirements-test.txtnot found) before checking anything. I'd fix the setup and add gfx950 and gfx1250 coverage for those kernel families rather than add a second race detector. I haven't changed the workflow here.Test Plan
test_lowering64.pybefore and after the fix, and with the vector immediate left in src1.Test Result
In a ROCm 10.2.0 image (TheRock 10.2.0a20260928), on a gfx90a host, at 2ac9fca:
test_lowering64.py: 117 passed on gfx90a, gfx942 and gfx950, including the assembled pairs. With the vector immediate left in src1, 33 cases fail (assembler: "literal operands are not supported"). With the oldsplitSrcs, 45 of the original 57 cases failed (every immediate case, plus the symbol case).Submission Checklist
Risk level
Low (2/5): the change only affects 64-bit add lowering of immediates on targets without a native 64-bit add, which nothing on develop emits. rocisa is code generation, so per the hipBLASLt overlay this needs coverage on the affected architectures in CI and two code-owner approvals.