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
On gfx1201 an fp32 NN GEMM with a large leading dimension either page-faults or returns a wrong answer while reporting success. hipblasLtMatmul gives HSA_STATUS_ERROR_MEMORY_FAULT at lda = 67,108,864 with K = 32, and at lda = 134,217,727 it returns success with a norm error of 0.911 against a 1e-06 tolerance.
The trigger is a byte count that outgrows a signed 32-bit register. GlobalReadIncs<tc>+unrollIdx is one SGPR holding stride * DepthU * bpeGR, the number of bytes one unroll iteration advances the global-read pointer. calculateStagger widens it to 64 bits before it reaches the global-read buffer descriptor, at two sites, and both used s_mul_i64_i32, a signed widen. Once stride * DepthU * bpeGR reaches 2^31 the register's bit 31 is set, the signed widen sign-extends it to a negative 64-bit value, and incrementSrd moves the read base backwards by about 2 GB instead of forwards. Where that displaced base is unmapped the kernel faults; where it is mapped the kernel reads other memory and returns a wrong result with no error.
Technical Details
Both sites multiply operands that are non-negative by construction: StaggerUIter is an iteration mask and LoopCounter an iteration count, and GlobalReadIncs is a byte stride. They now use s_mul_u64_u32, which the file already provides. WrapU still becomes negative in the subtraction that follows, and the borrowing s_sub_u32 / s_subb_u32 pair there already handles that, so only the intermediate product changes sign treatment.
Emitted instruction count is unchanged. s_mul_hi_i32 becomes s_mul_hi_u32; the low-word s_mul_i32 is identical. There is no register-pressure or scheduling change.
The defect is kernel-specific rather than shape-specific, which is why it survived. Reaching the multiply needs a non-zero StaggerUIter, and two things independently zero it: the loop-count clamp in calculateStagger drops the stagger for larger StaggerUStrideShift at low iteration counts, and StaggerUMapping picks which workgroup dimension feeds the mask, so a problem with a single workgroup in that dimension has no stagger either. Of 14 gfx1201 solutions measured at the boundary, 3 were affected.
Test Plan
Tensile/Tests/unit/characterization/_codegen/test_r3_stagger_incs_unsigned_gfx1201_char.py, CPU-only, in the -m unit lane. It emits gfx1201 kernel assembly and asserts that neither widening site uses a signed high multiply, plus a third assertion pinning that the widened stagger offset still reaches the global-read SRD, so the first two cannot pass vacuously if the offset stops being used.
The signedness is only observable in the emitted instruction, not in the Python source: SMulInt64to32 selects s_mul_hi_i32 when its sign argument is true and s_mul_hi_u32 when false, and both carry the same comment into the listing. The test therefore asserts on generated assembly rather than on source text.
Test Result
The new test fails on both sites before this change and passes after.
Full TensileLite unit lane (pytest Tensile/Tests/unit -m unit): 7746 passed, 32 skipped, 2 xfailed, 1 xpassed, 0 failed, 765 snapshots passed. The _codegen and CodegenResidue characterization suites specifically: 513 passed, 67 snapshots passed, 0 failed. No .ambr golden captures the affected instructions, so no snapshot was re-recorded and no ADR is required.
Hardware confirmation of the mechanism came from gfx1201 before the fix. Holding lda, K, m and the kernel fixed with lda * DepthU * elementBytes on the 2^31 boundary, solution 140189 is correct at n = 64, where dimension 1 has a single workgroup so its WorkGroup1 stagger input is always zero, and faults at n = 128, where a second workgroup makes an odd WorkGroup1 reachable. The same switch reproduces on the reported kernel 140192, which takes its stagger input from WorkGroup0: correct at m = 256, faults at m = 512.
Risk level
Medium. KernelWriterAssembly.py is a high-coupling file and this changes generated code for every kernel that compiles the stagger path. The change is confined to the signedness of two 64-bit widenings of a quantity that cannot be negative, the instruction count is unchanged, and the full unit lane is green. Post-merge verification on gfx1201 hardware with the shapes above is still wanted, because the unit test proves the emitted instruction rather than the runtime result.
GlobalReadIncs<tc>+unrollIdx is one SGPR holding stride*DepthU*bpeGR, a byte
count that is non-negative by construction. calculateStagger widened it to 64
bits with s_mul_i64_i32 at two sites: the stagger byte offset
(StaggerUIter * GlobalReadIncs), and the unroll-loop span
(LoopCounter * GlobalReadIncs, which becomes WrapU).
Once that product reaches 2^31 the register's bit 31 is set, the signed widen
sign-extends it to a negative 64-bit value, and incrementSrd moves the
global-read base backwards by about 2 GB. Measured on gfx1201: an fp32 NN GEMM
at lda = 67108864 with DepthU = 8 faults with HSA_STATUS_ERROR_MEMORY_FAULT, and
at lda = 134217727 the displaced base lands in mapped memory and hipblasLtMatmul
returns success with a norm error of 0.911.
Both operands are non-negative at both sites, so widen with s_mul_u64_u32. WrapU
still goes negative in the subtraction below it, which the borrowing
s_sub_u32 / s_subb_u32 pair already handles. Emitted instruction count is
unchanged; s_mul_hi_i32 becomes s_mul_hi_u32.
JIRA ID: ROCM-31230
Co-authored-by: Cursor <cursoragent@cursor.com>
Fixes gfx1201 large-leading-dimension GEMM failures by using unsigned widening for StaggerU global-read offsets and adding assembly regression coverage.
Changes:
Switches stagger and WrapU products to unsigned 64-bit multiplication.
The previous commit widened GlobalReadIncs unsigned at both stagger sites.
GSU.graIncrements negates that increment when the unroll dimension is in
MirrorDimsA/B, so for a mirrored unroll dimension the operand is genuinely
signed and an unsigned widen would move the global-read descriptor forward
instead of backward.
Select the widen per kernel: signed when the unroll summation index is
mirrored, unsigned otherwise. The mirrored case keeps its existing behavior;
only the non-mirrored case changes, which is where ROCM-31230 lives.
No shipped solution is affected either way: across gfx942, gfx950, gfx1200 and
gfx1201, 0 of 537,486 solutions declare a non-empty MirrorDimsA/B.
JIRA ID: ROCM-31230
Co-authored-by: Cursor <cursoragent@cursor.com>
Ran this on a gfx1201 R9700 (32 GB) at both commits. Every shape that failed before the fix is now correct, and the shapes that already worked are bit-for-bit unchanged.
Method
Only the Tensile device library needed regenerating, and only for the single logic family holding the affected solutions (gfx1201_Cijk_Ailk_Bljk_S_B_Bias_HA_S_SAV_UserArgs, 24 solutions). Client and host library came from TheRock nightly 35289346844-linux. I ported the hunk onto that nightly's own tensilelite rather than building the branch's tree, so the fix is the only thing differing between the baseline and patched runs; the ported hunk is byte-identical to the branch's. The unpatched baseline library was generated the same way first and reproduces every failure, including norm_error 0.911159 to the digit, so it's a valid control. Every non-faulting run is validated against the CPU reference, because above the boundary this defect sometimes returns a wrong answer instead of faulting.
Solution indices below are local to a single-family build, matched by libraryLogicIndex: 5 is 140192, 8 is 140195, 2 is 140189, 9 is 140196.
Shapes that failed before the fix
solution
shape
lda*DepthU*4
before
after
5 (140192)
m=lda=67,108,864, k=32, n=1
2^31
FAULT
CORRECT, norm_error 1.0544e-07
5 (140192)
m=lda=100,000,000, k=32, n=1
3.2e9
FAULT
CORRECT, norm_error 9.92358e-08
5 (140192)
m=lda=134,217,727, k=32, n=1
2^32 - 32
WRONG, norm_error 0.911159
CORRECT, norm_error 9.85013e-08
8 (140195)
m=lda=22,369,622, k=96, n=1
2^31 + 64
FAULT
CORRECT, norm_error 1.68249e-07
2 (140189)
m=262,144, lda=16,777,216, k=128, n=128
2^31
FAULT
CORRECT, norm_error 1.96929e-07
The third row is the one I care most about. It was returning a wrong answer while reporting success, and it now agrees with the CPU reference to 1e-07.
Regression guards
solution
shape
before
after
5 (140192)
m=lda=67,108,863, k=32, n=1
CORRECT 9.93493e-08
CORRECT 9.93493e-08
9 (140196)
m=lda=134,217,727, k=32, n=1
CORRECT 0
CORRECT 0
8 (140195)
m=lda=22,369,621, k=96, n=1
CORRECT 1.69081e-07
CORRECT 1.69081e-07
2 (140189)
m=262,144, lda=16,777,216, k=128, n=64
CORRECT 0
CORRECT 0
Identical norm_error on all four, so nothing that already worked moved.
The two commits are equivalent for shipped kernels
58a790730b claims no shipped solution declares mirrored dims. For this family that's confirmed empirically and not just argued: the code objects generated from ff7d1154f0 (unconditional unsigned) and from 58a790730b (conditional) are byte-identical, both sha256 a66968deae6bbb52, while the baseline is 5435fcf1d21a1370. The conditional always resolves to unsigned here. Both commits also went through the full table above and produced identical results down to the norm_error digits.
Instruction-level check
Unbundled the gfx1201 code object from both libraries and disassembled:
library
s_mul_hi_i32
s_mul_hi_u32
disassembly lines
baseline
44
275
163,929
patched
0
319
163,929
Diffing the two disassemblies gives 180 lines of difference: the 44 instruction substitutions plus the objdump filename header, and nothing else. Instruction addresses are unchanged on both sides (the substituted instruction sits at 0x14348 either way), so there's no code motion, no register-pressure change and no scheduling change. That's direct support for the "emitted instruction count is unchanged" claim in the description.
Scope
One logic family on one architecture, covering the three solutions known to be affected plus four controls, so this isn't a library-wide statement. All shapes are fp32 NN with n = 1 except the two 140189 rows, which need n >= 128 before the stagger fires at all. I didn't construct a shape that isolates a WrapU-specific failure, so that site is exercised here only because it shares widenIncs. One thing still open and not for this PR: the sparse-metadata stagger multiply (" stagger byte offset of metadata") is still s_mul_i64_i32 and has the same exposure, though it isn't reachable from these shapes.
… check
calculateStagger is called for the MX scale tensors as well as A and B
(KernelWriter.py, tensorParametersA["MX"] and tensorParametersB["MX"]), and
MirrorDimsMXSA/MXSB are copied from MirrorDimsA/B in Solution.py. The previous
commit's predicate listed only A, B and Metadata, so a mirrored unroll dimension
on an MX scale tensor would have taken the unsigned widen and moved that
descriptor forward instead of backward.
Add MXSA and MXSB to the check. All five MirrorDims keys exist and default to
empty, so the predicate is safe for every tensor char calculateStagger sees.
JIRA ID: ROCM-31230
Co-authored-by: Cursor <cursoragent@cursor.com>
Same gfx1201 rig as above. The hardware table is unchanged: all five previously failing shapes correct, all four guards bit-for-bit identical to the commit-2 run. For the gfx1201 fp32 family the code object is byte-identical across all three commits (sha256 a66968deae6bbb52), so this commit is a no-op there, which is expected since no MX tensors are involved.
Worth adding, because it's the part this commit is actually about: the MXSA/MXSB path is real and reachable with shipped logic, so commit 2's tuple omitted a live path rather than a hypothetical one. I instrumented a throwaway copy of the generator to log the tensor char reaching the widen site. For gfx1250/Equality/gfx1250_Cijk_Alik_Bljk_F8BS_MXAE8B32_MXBE8B32_BH_UserArgs, calculateStagger is reached with tc = A, B, MXSA and MXSB. For the gfx950 S_MX families it is reached with A and B only, because the MX calls are gated on MXBlockA/MXBlockB (KernelWriter.py:3113-3116) and those solutions don't set them.
It changes no emitted code today. On that MX family all four tensor chars report mirrored=False, and the disassembly is identical between commit 2 and commit 3: 62 s_mul_hi_u32 and 0 s_mul_hi_i32 in both. So this is correctness for a future mirrored MX kernel, not a behaviour change now.
One methodology note in case anyone repeats this: the gfx1250 MX build is not byte-reproducible. Two builds from the same unmodified generator gave different code object hashes, so I compared disassembly rather than sha256 there. The gfx1201 build is reproducible, which is what makes the sha comparison above meaningful.
This also confirms the safety argument in the commit message. All five MirrorDims keys are declared with empty defaults (SolutionStructs/Problem.py:485-489) and MirrorDimsMXSA/MXSB are copied from A/B (SolutionStructs/Solution.py:2830 and 2845). Generation is clean on both the MX and non-MX families.
Confirms the expectation. Commit 3 is a no-op on gfx1201 because all 24 solutions in that family have mxBlockA, mxBlockB and sparse false, so calculateStagger only ever sees tc = A or B there.
The MX instrumentation is the valuable part: MXSA/MXSB are reached by shipped gfx1250 logic, so commit 2's tuple omitted a live path, not a hypothetical one.
Worth recording beyond this PR: gfx1250 MX builds are not byte-reproducible, so sha256 is not a valid before/after technique on that family. Diff disassembly instead.
Still uncovered, both on ROCM-31230: the mirrored path itself, which needs a synthetic config since no shipped solution mirrors, and the deferred sparse-metadata multiply. The gfx1201 coverage Copilot asked for is now satisfied in substance but not in form, since none of it runs in a shared CI lane.
This is the coverage asked for in review. matmul_size_t_stagger_incs_overflow_f32_ROCM31230 in clients/tests/data/matmul_gtest.yaml, with its shape anchor in matmul_common.yaml. fp32 NN, M=2048, N=1, K=32, lda=67108864, integer_exact plus unit_check, gated to gpu_arch: '1201'.
lda * DepthU * bpe is 2^26 * 8 * 4 = 2^31 exactly for the MT256x128x8 kernel the heuristic picks, which is the first failing value. lda = 2^26 - 1 is the ticket's passing case and I confirmed it stays correct at 1.02e-07, so the gate sits one representable step off a verified control.
Before and after on gfx1201
library
exit
result
baseline, unpatched
1
FAILED, memory fault at testing_matmul.hpp:1057 plus d_vector guard trips
patched
0
[ PASSED ] 1 test, 4.3 s
Three runs each, fully deterministic. Both libraries were generated from the nightly's own tensilelite (shipped copy is byte-identical to 4606f831) differing only by this PR's hunk, which I checked is byte-identical to the branch's. The .dat.zlib metadata came out bit-identical between baseline and patched, so selection and solution indices don't move and only the code objects differ. Disassembly matches the table already in this PR: baseline 44 s_mul_hi_i32 / 275 s_mul_hi_u32, patched 0 / 319. As a regression guard, 506 fp32 NN cases pass identically on both.
I also regenerated hipblaslt_gtest.data from the unmodified branch yaml first and got a byte-identical file to the one shipped in the nightly, which is what makes the harness comparison trustworthy.
Two choices worth flagging
Selection is left to the heuristic rather than pinned with solution_index. Indices aren't stable across library retunes, so pinning one buys a gate that silently rots. The trade is the mirror image: a future retune could move selection off the affected kernel and weaken this gate instead. I think that's the better failure mode, but it's a real trade.
Category is nightly, not pre_checkin. A is 8.6 GiB on both the host and the device and can't be made smaller, since the A span is (K-1)*lda + M whatever M is. The silent-corruption face of this defect (wrong answer, no fault) needs lda near 2^32 and 17 GiB for A, so it won't fit on a 16 GiB part and isn't what's gated here. unit_check with integer_exact is bit-exact, so it would catch that variant too on a larger card.
One behaviour to know about: when this test fails, the process tears down on the poisoned HIP context and never prints gtest's final summary. CI still gets exit code 1 and the per-assertion failures, just not the [ FAILED ] block.
Out of scope, filed separately
Getting this landed turned up two pre-existing CI problems, now tracked in #12337. Short version: category: stress matches no TheRock tier, so the three existing matmul_size_t_* stress cases never run, and an unknown category is silently accepted because match_test_category() is bypassed. Separately, a host-memory shortfall segfaults rather than skipping, because get_available_host_memory() uses freeram and there's no host-side equivalent of CHECK_DEVICE_ALLOCATION. I hit that once on this box during verification. Neither is caused by this PR, but the second one is a flake risk for any large-shape case.
…rflow
The unit test in this PR asserts on emitted assembly, which proves the
instruction changed but not that the kernel is correct at runtime. Add the
hardware reproducer to close that gap.
matmul_size_t_stagger_incs_overflow_f32_ROCM31230 is fp32 NN with
M=2048, N=1, K=32, lda=67108864. For the MT256x128x8 kernel the heuristic
selects, lda * DepthU * bpe is 2^26 * 8 * 4 = 2^31 exactly, the first value
whose bit 31 is set and so the first one the signed widen sign-extends.
lda = 2^26 - 1 is the ticket's passing case and stays correct at 1.02e-07,
so the gate sits one representable step off a verified control.
M only has to be large enough for a second workgroup in dimension 0, since
the selected kernel takes its stagger input from WorkGroup0 and a single
workgroup zeroes it. N = 1 keeps B, C and D negligible. A cannot be made
small: its span is (K-1)*lda + M elements whatever M is, so A is 8.6 GiB on
both the host and the device. That is why the case is nightly rather than
pre_checkin.
Selection is left to the heuristic rather than pinned with solution_index,
because indices are not stable across library retunes.
Verified on a gfx1201 RX 9070 XT against two Tensile libraries generated
from the same generator differing only by this PR's hunk: FAIL pre-fix
(exit 1, memory fault), PASS post-fix (exit 0), three runs each. The
.dat.zlib metadata is bit-identical between the two, so selection does not
move; only the code objects differ, 44 s_mul_hi_i32 becoming 0. 506 fp32 NN
cases pass identically on both.
JIRA ID: ROCM-31230
Co-authored-by: Cursor <cursoragent@cursor.com>
This regex does not prove that the newly computed stagger offset reaches the SRD: removeStagger() emits the same s_addc_u32 ... gra SRD += inc(upper) pattern later (KernelWriterAssembly.py:6832-6860). Deleting the calculateStagger increment at line 6638 would still satisfy this assertion. Match the add within the SRDs += (StaggerUIter) ... block so this non-vacuity check covers the production path it is intended to guard.
@nakajee Thank you, this sent me down the right path. Tensile PR 1672 is the precedent I should have found myself, and your calculateStagger hunk there is the same change this PR makes. Reassuring to have converged on what you already concluded three years ago.
Tail loop: you were right to ask, and it's fixed
My gate uses K=32 with DepthU=8, which divides evenly, so there was no tail loop in it at all. Measured on gfx1201 with solution 140192 forced and lda = 2^26, against two libraries differing only by this PR's hunk:
K
tail iterations
baseline
patched
32
0
FAULT
PASS, norm 1.67e-07
33
1
FAULT
PASS, norm 1.68e-07
36
4
FAULT
PASS, norm 1.58e-07
Controls hold: K=36 just below the boundary and at a small lda pass on either library.
On removeStagger
I think it may already be covered, and the helper name is what makes it look otherwise. Both of its multiplies call s_mul_i64_i32_u32, which reads as signed but computes |src0| * src1 unsigned and negates if src0 was negative, so GlobalReadIncs (src1) is treated as unsigned. The signed operand is (3 or 2) - StaggerUIter, which really can go negative.
It arrived by a different route than yours: commit 9446f2356a3 (December 2024, "memory access fault in tail loop with extreme larger size") rather than your algebraic expansion into two unsigned multiplies. Same property, and it's why the tail rows above pass without a second fix. So my reading is that tensilelite got the removeStagger half of 1672 back then and never got the calculateStagger half, which this PR supplies. Please push back if I'm misreading that helper, since it's the load-bearing claim and you know this code better than I do.
Still open, and I'd value your view: the sparse-metadata pair in calculateStagger (lines 6657 and 6662) is still s_mul_i64_i32. I deferred it because no sparse shape was measured here. After this PR it's the last of the four sites still sign-extending. Tracked on ROCM-31230.
The new regression case ran and passed on the gfx1201 R9700 node in math-ci precheckin (4803 ms), so it's exercised on real hardware in a shared lane.
…fy the SRD assertion
Two review findings on the regression coverage.
The gfx1201 hardware gate selected its kernel through the heuristic, so a future
library retune could pick a solution with StaggerU disabled and report success
without executing either widening. Pinning solution_index is not an option
because indices are not stable across retunes. Use algo_method 1 with
requested_solution_num -1 instead, the pattern matmul_heuristic_all_solutions
already uses, which runs every supported solution. At this shape that is 17
solutions rather than 1, for about 3 extra seconds.
The third assertion in the codegen test searched the whole kernel for
"s_addc_u32 s[sgprSrd{A,B}+1] ... gra SRD += inc(upper)". incrementSrd emits that
from nine call sites, and the ordinary global-read increment among them fires in
every kernel, so the assertion was close to unconditionally true: one kernel
emits 44 such lines, 8 of which match the regex, none from calculateStagger.
Scope it to a single "addr += (StaggerUIter) * GlobalReadIncs<tc>" block and tie
the registers the stagger multiply writes to the ones the SRD increment reads.
Verified by negative control. With calculateStagger's incrementSrd call replaced
by pass, the old assertion still passed while the new one fails on "the A stagger
offset in s16 is no longer added into sgprSrdA+0". Generator restored afterwards;
the _codegen suite is 302 passed, 1 xfailed, 1 xpassed, 67 snapshots, and the
hardware gate still FAILs pre-fix / PASSes post-fix on gfx1201.
JIRA ID: ROCM-31230
Co-authored-by: Cursor <cursoragent@cursor.com>
Both actionable findings are addressed in the latest commit. Thanks Copilot, the third one in particular was a good catch.
Kernel selection. Fair point, and pinning solution_index isn't an option since indices move across library retunes. Used algo_method: [1] with requested_solution_num: -1 instead, which is what matmul_heuristic_all_solutions already does: it runs every supported solution rather than the heuristic's pick. At this shape that's 17 solutions instead of 1, for about 3 extra seconds. Still FAILs pre-fix and PASSes post-fix on gfx1201. A retune can't weaken the gate now, because nothing is being selected.
The SRD assertion. Correct, and worse than described. incrementSrd emits that instruction and comment from nine call sites, and the ordinary global-read increment among them fires in every kernel. The kernel this test emits contains 44 gra SRD += inc(upper) lines, 8 of which match the old regex, and none of those 8 come from calculateStagger.
I proved it with a negative control rather than taking it on argument. Replacing calculateStagger's incrementSrd call with pass and re-emitting: the old assertion still passed, matching s_addc_u32 s[sgprSrdA+1], s[sgprSrdA+1], s17 from /* global read inc A loopL */; the new one fails with "the A stagger offset in s16 is no longer added into sgprSrdA+0". So the assertion really did prove nothing, and the PR description's claim that it stopped the first two passing vacuously was wrong. The repair scopes everything to one addr += (StaggerUIter) * GlobalReadIncs<tc> block and ties the registers the stagger multiply writes to the ones the SRD increment reads. It also skips the sparse arm, which reuses the same phrase and has its own increment. Generator restored; _codegen is 302 passed, 1 xfailed, 1 xpassed, 67 snapshots, 0 failed.
Host memory. Agreed, and thanks for the precise site, hA[i] at testing_matmul.hpp:2648-2651 with no host-side counterpart to CHECK_DEVICE_ALLOCATION. This is pre-existing and harness-wide, so it's tracked separately in #12337 rather than fixed here. get_available_host_memory() reads sysinfofreeram rather than available memory, so it also trips spuriously when reclaimable page cache is warm; I hit exactly that during verification on a 125 GB host with 111 GB in cache. Keeping this case in nightly limits the blast radius but doesn't remove it, which is why the issue exists.
Thanks for your comment.
I did not realize that removeStagger was already fixed.
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
tony-davis
changed the title
JIRA ID : ROCM-31230
Motivation
On gfx1201 an fp32 NN GEMM with a large leading dimension either page-faults or returns a wrong answer while reporting success.
hipblasLtMatmulgivesHSA_STATUS_ERROR_MEMORY_FAULTatlda = 67,108,864withK = 32, and atlda = 134,217,727it returns success with a norm error of 0.911 against a 1e-06 tolerance.The trigger is a byte count that outgrows a signed 32-bit register.
GlobalReadIncs<tc>+unrollIdxis one SGPR holdingstride * DepthU * bpeGR, the number of bytes one unroll iteration advances the global-read pointer.calculateStaggerwidens it to 64 bits before it reaches the global-read buffer descriptor, at two sites, and both useds_mul_i64_i32, a signed widen. Oncestride * DepthU * bpeGRreaches2^31the register's bit 31 is set, the signed widen sign-extends it to a negative 64-bit value, andincrementSrdmoves the read base backwards by about 2 GB instead of forwards. Where that displaced base is unmapped the kernel faults; where it is mapped the kernel reads other memory and returns a wrong result with no error.Technical Details
Both sites multiply operands that are non-negative by construction:
StaggerUIteris an iteration mask andLoopCounteran iteration count, andGlobalReadIncsis a byte stride. They now uses_mul_u64_u32, which the file already provides.WrapUstill becomes negative in the subtraction that follows, and the borrowings_sub_u32/s_subb_u32pair there already handles that, so only the intermediate product changes sign treatment.Emitted instruction count is unchanged.
s_mul_hi_i32becomess_mul_hi_u32; the low-words_mul_i32is identical. There is no register-pressure or scheduling change.The defect is kernel-specific rather than shape-specific, which is why it survived. Reaching the multiply needs a non-zero
StaggerUIter, and two things independently zero it: the loop-count clamp incalculateStaggerdrops the stagger for largerStaggerUStrideShiftat low iteration counts, andStaggerUMappingpicks which workgroup dimension feeds the mask, so a problem with a single workgroup in that dimension has no stagger either. Of 14 gfx1201 solutions measured at the boundary, 3 were affected.Test Plan
Tensile/Tests/unit/characterization/_codegen/test_r3_stagger_incs_unsigned_gfx1201_char.py, CPU-only, in the-m unitlane. It emits gfx1201 kernel assembly and asserts that neither widening site uses a signed high multiply, plus a third assertion pinning that the widened stagger offset still reaches the global-read SRD, so the first two cannot pass vacuously if the offset stops being used.The signedness is only observable in the emitted instruction, not in the Python source:
SMulInt64to32selectss_mul_hi_i32when itssignargument is true ands_mul_hi_u32when false, and both carry the same comment into the listing. The test therefore asserts on generated assembly rather than on source text.Test Result
The new test fails on both sites before this change and passes after.
Full TensileLite unit lane (
pytest Tensile/Tests/unit -m unit): 7746 passed, 32 skipped, 2 xfailed, 1 xpassed, 0 failed, 765 snapshots passed. The_codegenandCodegenResiduecharacterization suites specifically: 513 passed, 67 snapshots passed, 0 failed. No.ambrgolden captures the affected instructions, so no snapshot was re-recorded and no ADR is required.Hardware confirmation of the mechanism came from gfx1201 before the fix. Holding
lda,K,mand the kernel fixed withlda * DepthU * elementByteson the2^31boundary, solution 140189 is correct atn = 64, where dimension 1 has a single workgroup so itsWorkGroup1stagger input is always zero, and faults atn = 128, where a second workgroup makes an oddWorkGroup1reachable. The same switch reproduces on the reported kernel 140192, which takes its stagger input fromWorkGroup0: correct atm = 256, faults atm = 512.Risk level
Medium.
KernelWriterAssembly.pyis a high-coupling file and this changes generated code for every kernel that compiles the stagger path. The change is confined to the signedness of two 64-bit widenings of a quantity that cannot be negative, the instruction count is unchanged, and the full unit lane is green. Post-merge verification on gfx1201 hardware with the shapes above is still wanted, because the unit test proves the emitted instruction rather than the runtime result.Related
Submission Checklist