AISPARSE-699/700: 64-bit element counts and block indexing in the common launches by kliegeois · Pull Request #12315 · ROCm/rocm-libraries

kliegeois · 2026-09-18T17:27:25Z

AISPARSE-699, AISPARSE-700

Two commits, one per ticket. They are paired because both touch
common/rocsparse_common.cpp and the launches around it.

AISPARSE-699 — 32-bit dense element counts (P2, reachable)

Both forms trigger at a 46,341 x 46,341 dense matrix: 8.6 GB in single precision,
which fits on a single MI300X.

Grid right, kernel guard wrong. rocsparse_common.cpp:231 and :324 size their
grids correctly as dim3((int64_t(m) * n - 1) / 256 + 1) — that promotion is
deliberate. But the kernels recompute the same product in the 32-bit template type
for their bounds check (gid >= m * n at :40 in valset_2d_device, :132-135 in
scale_2d_device). The launch covers the whole matrix and the kernel then refuses to
process most of it. Half the problem had been solved, host-side only.

Caller overflows before the helper sees it. rocsparse_gemmi.cpp:82 and :93
pass m * n computed as rocsparse_int, so scale_array receives an
already-overflowed length. The helper is fine; the call site is not.

The element count is now formed once as int64_t and the same expression feeds the
grid, the kernel bound and both gemmi call sites.

gid is widened too, and that is the load-bearing half.
hipBlockIdx_x * BLOCKSIZE is unsigned-int arithmetic that wraps at 2^32 regardless
of the template index type, so a widened guard over a 32-bit index would have changed
nothing.

axpby_batched_kernel (:285) and scale_kernel (:259) take a single extent rather
than a product, so neither was ever exposed. They are widened and given the same
stride loop anyway so all five kernels in the file share one idiom; axpby_kernel
likewise. This is stated in the code and the commit message.

AISPARSE-700 — 64-bit lengths vs 32-bit kernel indices (P3, unreachable)

valset drops its index template parameter so the length is structurally 64-bit.
singularity's device-pointer path gets a clamp plus stride, with the grid.x limit
threaded in as a parameter since the inner overload takes a stream, not a handle. The
transpose pair keeps its existing >INT32_MAX dispatch and adds a clamped grid.x with
a row-panel stride loop in the __global__ wrapper.

The existing chunking path was not reused for singularity. It exists to stage
results through the handle buffer, which a device-pointer destination does not need;
reusing it would add a device-to-device copy per chunk and serialise every call on a
shared buffer. The host path is untouched.

markers2position indexing is touched — the stride loop and a tid widened from
unsigned int to int64_t. For any batch_count below 2^32, i.e. everything
reachable, each block runs exactly one iteration with base == hipBlockIdx_x * BLOCKSIZE,
so every thread handles bit-for-bit the element it handled before. This cannot affect
the pre-existing pre_checkin/csrsv fault in that kernel (issue is present on
unmodified develop); test runs were filtered away from it.

Ticket inaccuracy. The ticket says a 32-bit instantiation sees a truncated bound
in valset. It does not: launch_valset always passed an int64_t length, so the
kernel's index template parameter was only ever deduced as int64_t. The wrap was
purely in the unsigned block-index arithmetic. Same fix, corrected reasoning.

Block uniformity

All ten grid-stride loops verified two ways. By inspection, every term in every bound
is hipBlockIdx_x, hipGridDim_x, a kernel argument (nelm/len/length/
batch_count/m) or a compile-time constant, and no bound mentions hipThreadIdx_x.
Empirically, the transposes call __syncthreads() and their tests run 6 stride sweeps
over a clamped grid and pass, which divergent trip counts would not survive. The
sdata tile is safe to reuse across outer iterations because the column loop opens
with a __syncthreads() separating the previous iteration's last reads from the next
one's writes.

Testing

17 cases in rocsparse-unit-test-device — 8 for 699, 9 new for 700. No device-memory
guards, no size-based skips; largest allocation ~740 kB. Regressions: 6451
*gemmi*:*dense*:*coo2dense*:*ell2dense*, 491 device unit tests, 14 host unit tests,
all passing.

UBSan on both 699 forms. The pre-fix guard does not merely mis-size — it returns
true for gid == 0:

== caller_after  (fixed)  : 2147488281
== guard_after   (fixed)  : 0
ubsan_probe.cpp:14:14: runtime error: signed integer overflow: 46341 * 46341 cannot be represented in type 'rocsparse_int'
== caller_before (pre-fix): -2147479015
ubsan_probe.cpp:27:21: runtime error: signed integer overflow: 46341 * 46341 cannot be represented in type 'int'
== guard_before  (pre-fix): 1

Break-validation: 6 rounds, each rebuilt with CCACHE_DISABLE=1, restored and
re-confirmed after. Every one of the 17 fails in at least one round.

Round	Break	Result
1	699 count + index back to 32-bit	5 failed
2	699 stride loops only, bound kept 64-bit	5 failed, exact complement of round 1
3	gemmi length short by one	1 failed
4	700 stride loops in all four files	8 failed
5	valset inner bound off by one	1 failed
6	both valset bounds off by one	2 failed

Round 1 separates the two halves of the defect:

valset_2d left elements of a 46341 x 46341 matrix unwritten: the kernel bound is not the 64-bit element count 2147488281
valset_2d wrote below the matrix: the flattened element index wrapped through int32_t past INT32_MAX

Notes for review

The clamp helper is duplicated between rocsparse_common.cpp and
rocsparse_valset.cpp as identical anonymous-namespace templates, and
rocsparse_singularity.cpp spells its clamp with std::min<int64_t> while the
others use rocsparse::min (that file does not include rocsparse_common.hpp and
already used std::min). Both follow from keeping rocsparse_common.hpp untouched
while AISPARSE-696/697/668: shared grid helpers, batch grid.y clamp, and csrmm/coomm large dense column counts #11512 is in flight; all four sites carry the substitution note.
New test header clients/unittests/unit_test_grid_clamp.hpp for the
ScopedMaxGridSizeX RAII helper, rather than duplicating it across two test files
or touching shared unit_test_utils.hpp. Commit 1 owns a header commit 2 also uses.
Pre-existing bug found and deliberately not fixed:
dense_transpose_strided_batched_kernel in device pointer mode computes the
per-batch scalar as alpha + i * alpha_stride after dereferencing the pointer,
adding the batch offset to the scalar value instead of indexing the array. Unrelated
to block indexing and out of scope, but it forced the batched transpose test into
host pointer mode; a device-pointer test with batch_count > 1 would be asserting
on the bug. There is a comment on the test.
Also adjacent and out of scope: the transpose launches pass batch_count straight
into dim3(grid_x, batch_count) without clamping grid.y to 65535.

A dense element count is the product of two index-typed extents, so it overflows int32_t at m = n = 46341 -- 8.6 GB in single precision, which fits on a single MI300X. The defect had two forms, and the library had already solved the host half of one of them. Form one, in library/src/common/rocsparse_common.cpp. valset_2d and scale_2d_array sized grid.x from the 64-bit product, deliberately, but valset_2d_device and scale_2d_device recomputed the same product in the 32-bit template index type for their bounds check (gid >= m * n). The launch covered the whole matrix and the kernel then refused to process any of it: at 46341 x 46341 the product wraps negative, so gid >= m * n is true for every thread. Confirmed under -fsanitize=undefined, which reports the signed overflow and shows the pre-fix guard returning true for gid == 0. Widening the guard alone accomplishes nothing, so the flattened index is widened too: hipBlockIdx_x * BLOCKSIZE is unsigned-int arithmetic that wraps at 2^32 whatever the index type is, and narrowing the result into a 32-bit I wraps sooner still. valset_2d is reached from coo2dense, coo2dense_aos, ell2dense and csx2dense_impl; scale_2d_array from csrmm, coomm, bsrmm, bellmm and gebsrmm. Form two, in library/src/level3/rocsparse_gemmi.cpp. The k == 0 and alpha == 0 shortcuts computed the element count of C as rocsparse_int in the caller and handed the already-truncated value to scale_array, which then sized a correspondingly wrong grid. The count is now formed in int64_t once, before the call. scale_kernel, axpby_kernel and axpby_batched_kernel take a single extent rather than a product, so they were never exposed to the overflow. They are widened and given the same stride loop anyway, so that every kernel in the file shares one indexing idiom. Every launch now clamps grid.x against handle->properties.maxGridSize[0] and grid-strides over what the clamp drops. The stride loops live in the __global__ wrappers and pass the element index down, leaving each __device__ function a single-element body (the AISPARSE-666 idiom). Every loop bound depends only on hipBlockIdx_x, hipGridDim_x, BLOCKSIZE and the element count -- all block uniform -- so all threads of a block run the same number of iterations. The clamp is computed by a file-local helper rather than a shared one: AISPARSE-696 (PR ROCm#11512) adds rocsparse::get_grid_size and rocsparse::ceil_div for exactly this, but it has not merged and rocsparse_common.h/.hpp are a live conflict zone with AISPARSE-677/678. The helper body is a one-line substitution once ROCm#11512 lands, and says so. Tests: clients/unittests/unit_test_internal_common_dense_grids.cpp, eight cases in rocsparse-unit-test-device. The 46341 x 46341 threshold is made affordable by folding the matrix onto m + n - 1 real elements with ld = 1 -- both operations under test are idempotent stores, so the aliasing cannot change the expected result -- and by allocating a shadow region in front of the matrix, which a 32-bit flattened index writes into and a 64-bit one does not. That is what makes the two halves of the defect separately visible. The grid-stride paths are covered by shrinking handle->properties.maxGridSize[0] (helper in the new clients/unittests/unit_test_grid_clamp.hpp). There is no device-memory guard and no size-based skip anywhere in the file: the largest allocation is about 740 kB, so every case runs on every GPU. Co-authored-by: Cursor <cursoragent@cursor.com>

…RSE-700) valset, singularity and the dense transpose pair each sized grid.x from an int64_t count, correctly, and then formed the element or row index in the kernel as hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x. hipBlockIdx_x and BLOCKSIZE are both unsigned int, so that product is evaluated in 32-bit unsigned arithmetic and wraps at 2^32 regardless of the type it is assigned to. None of the three had a grid-stride loop either, so a grid.x clamped against the device limit would have left a tail unprocessed. valset (library/src/common/rocsparse_valset.cpp) is the most widely reached of the three: 15+ conversion routines, csrgemm/csrgeam/bsrgeam/ bsrgemm and csrsv_analysis. One correction to the ticket here -- launch_valset always passed an int64_t length, so the kernel's index template parameter was only ever deduced as int64_t and the bound was never truncated. The wrap was purely in the index arithmetic. The template parameter is dropped so that the length is structurally 64-bit, the index is int64_t, and grid.x is clamped with a stride loop behind it. singularity (library/src/common/rocsparse_singularity.cpp) computed the batch index as const auto, which deduces plain unsigned int. Only the device-pointer path was exposed: the host-pointer path stages results through the handle buffer and only ever launches one buffer-sized chunk at a time. That chunking loop is deliberately NOT reused for the device path. It exists to stage results through the handle buffer, which a device-pointer destination does not need, and reusing it would add a device-to-device copy per chunk and serialize every call on a shared buffer. The device path gets the clamp and the stride loop instead; the host path is untouched. The grid.x limit is threaded in as a parameter because the inner overload takes a stream rather than a handle. dense_transpose / dense_transpose_back (library/src/common/rocsparse_dense_transpose*.cpp) already dispatch to a 64-bit index kernel above INT32_MAX; that is kept. It did not go far enough, because the 64-bit kernel still indexed blocks in 32 bits. grid.x is clamped and the row panel base is passed into the device function from a stride loop in the __global__ wrapper, which also leaves each device function a single-panel body (the AISPARSE-666 idiom). All five stride loops have block-uniform bounds: every term is hipBlockIdx_x, hipGridDim_x, a kernel argument or a compile-time constant, and no bound mentions hipThreadIdx_x, so all threads of a block run the same number of iterations. That matters for the transposes, whose device functions call __syncthreads(). Reusing the sdata tile across outer iterations is safe because the column loop opens with a __syncthreads() that separates the previous iteration's last reads of the tile from the next iteration's writes. markers2position, which a pre-existing gfx1201 fault in pre_checkin/csrsv is implicated in, is touched here. For any batch_count below 2^32 -- that is, everything reachable -- each block runs exactly one iteration with base == hipBlockIdx_x * BLOCKSIZE, so every thread handles bit for bit the element it handled before. This change cannot affect that fault. As in AISPARSE-699 the clamps are computed locally pending AISPARSE-696 (PR ROCm#11512), and each site says so. Tests: clients/unittests/unit_test_internal_common_large_grids.cpp, nine cases in rocsparse-unit-test-device. The 2^32 thresholds are not reachable -- 17.2 GB for an int32_t valset array, four billion rows or batch entries for the others -- so each case shrinks handle->properties.maxGridSize[0] and drives the real kernels through the clamp and the stride loop, which is what the fix actually consists of. Two further valset cases pin the ceiling division at the block boundary and at zero length. There is no device-memory guard and no size-based skip: the largest allocation is 5000 int64_t, so every case runs on every GPU. Co-authored-by: Cursor <cursoragent@cursor.com>

therock-pr-bot · 2026-09-18T17:31:56Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/rocsparse/clients/unittests/unit_test_internal_common_dense_grids.cpp`, `projects/rocsparse/clients/unittests/unit_test_internal_common_large_grids.cpp`, `projects/rocsparse/library/src/common/rocsparse_common.cpp`, `projects/rocsparse/library/src/common/rocsparse_dense_transpose.cpp`, `projects/rocsparse/library/src/common/rocsparse_dense_transpose_back.cpp` (+3 more); no test file found
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

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

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

🙋 Wish to Override Policy?

therock-pr-bot · 2026-09-18T17:31:59Z

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

CI runs static-analysis with clang-format 13 while the pre-commit hook pins 18.1.4, and the two disagree on two constructs this branch introduced. The sources were clean under 18 and red under 13, so the fix is to write the lines in a form that is a fixed point under both rather than to pick a side. AlignConsecutiveAssignments: 18 aligns the '=' across a run of adjacent '= delete;' declarations and pads the shorter one; 13 does not treat a deleted function as an assignment and wants a single space. Separating the deleted copy constructor from the deleted copy assignment with a blank line ends the alignment run, after which both versions emit the same single-space form. Macro multiply parsed as a cast: in '(bi) * block_dim', 13 reads the parenthesised token as a cast and emits '(bi)*block_dim', while 18 reads a multiply and emits '(bi) * block_dim'. Swapping the operands to 'block_dim * (bi)' removes the ')' immediately preceding '*', so neither version can see a cast. Multiplication is commutative, so the index arithmetic is unchanged. Verified by formatting every touched file with both 13.0.1 and 18.1.4 against projects/rocsparse/.clang-format and confirming byte-identical output.

codecov · 2026-09-22T05:24:54Z

Codecov Report

❌ Patch coverage is 89.13043% with 5 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...s/rocsparse/library/src/level3/rocsparse_gemmi.cpp	33.33%	2 Missing ⚠️
.../rocsparse/library/src/common/rocsparse_common.cpp	94.44%	0 Missing and 1 partial ⚠️
...parse/library/src/common/rocsparse_singularity.cpp	90.91%	1 Missing ⚠️
.../rocsparse/library/src/common/rocsparse_valset.cpp	88.89%	0 Missing and 1 partial ⚠️

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #12315      +/-   ##
===========================================
- Coverage    72.12%   71.86%   -0.27%     
===========================================
  Files         2868     2869       +1     
  Lines       473269   473844     +575     
  Branches     69803    69896      +93     
===========================================
- Hits        341336   340483     -853     
- Misses      108234   109756    +1522     
+ Partials     23699    23605      -94

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`49.33% <ø> (ø)`		Carriedforward from 3cd5218
TensileLite-Unit	`83.29% <ø> (-0.03%)`	⬇️	Carriedforward from 3cd5218
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 3cd5218
hipBLASLt	`38.91% <ø> (-<0.01%)`	⬇️	Carriedforward from 3cd5218
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 3cd5218
hipDNN	`86.66% <ø> (ø)`		Carriedforward from 3cd5218
hipFFT	`43.43% <ø> (+0.91%)`	⬆️	Carriedforward from 3cd5218
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 3cd5218
hipSOLVER	`68.96% <ø> (ø)`		Carriedforward from 3cd5218
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 3cd5218
rocBLAS	`48.24% <ø> (+0.07%)`	⬆️	Carriedforward from 3cd5218
rocFFT	`47.62% <ø> (-3.84%)`	⬇️	Carriedforward from 3cd5218
rocRAND	`57.45% <ø> (ø)`		Carriedforward from 3cd5218
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from 3cd5218
rocSPARSE	`74.10% <89.13%> (+0.03%)`	⬆️
rocThrust	`93.15% <ø> (ø)`		Carriedforward from 3cd5218

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

Files with missing lines	Coverage Δ
...e/library/src/common/rocsparse_dense_transpose.cpp	`86.36% <100.00%> (+0.10%)`	⬆️
...rary/src/common/rocsparse_dense_transpose_back.cpp	`78.38% <100.00%> (+0.30%)`	⬆️
.../rocsparse/library/src/common/rocsparse_common.cpp	`88.57% <94.44%> (+0.57%)`	⬆️
...parse/library/src/common/rocsparse_singularity.cpp	`59.58% <90.91%> (+1.39%)`	⬆️
.../rocsparse/library/src/common/rocsparse_valset.cpp	`56.25% <88.89%> (+1.41%)`	⬆️
...s/rocsparse/library/src/level3/rocsparse_gemmi.cpp	`79.91% <33.33%> (+0.09%)`	⬆️

... and 76 files with indirect coverage changes

🚀 New features to boost your workflow:

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

…-grid PRs (AISPARSE-699/700) Comment-only change. The header comment now describes the clamp as rocsparse::get_grid_size_x applies it, min(maxGridSize[0], (2^32 - 1) / blockDim.x) blocks, instead of the older "1e11" wording. This makes the file byte-identical to the copy carried by ROCm#12306, ROCm#12308, ROCm#11815, ROCm#11816 and the other large-grid PRs (blob 0ce9d85), so those PRs no longer conflict with this one on it. ScopedMaxGridSizeX is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

…E-699/700) ROCm#12546 rewrites scale_device, scale_kernel and the scale_array launch with a clamped grid.x and a compile-time GRID_STRIDE fast path. This PR rewrote the same code with an always-on stride loop, so the two conflicted in rocsparse_common.cpp. Restore those three to develop so ROCm#12546 owns them. valset_2d, scale_2d, axpby and axpby_batched keep their 64-bit counts, clamped grids and stride loops; none of them used the removed scale code. The "why length is widened" comment moves from scale_kernel to axpby_kernel. The gemmi m * n fix is unchanged and still passes a 64-bit length to scale_array. scale_array_grid_stride stays: it passes against develop's unclamped launch and covers ROCm#12546's GRID_STRIDE path once both are merged. Co-authored-by: Cursor <cursoragent@cursor.com>

# Conflicts: # projects/rocsparse/clients/unittests/CMakeLists.txt

…SE-699) valset_2d_kernel, scale_kernel and axpby_batched_kernel now hold their element bodies directly instead of calling a __device__ function used only there. scale_2d_kernel keeps scale_2d_device, as it loops over batches. axpby_kernel and axpby_device were never launched and are removed.

…rid_size_x (AISPARSE-699)

…in clients/unittests/CMakeLists.txt Every large-grid PR appended its entries after unit_test_internal_conversion_grid_clamp.cpp, so merging any one of them conflicted every sibling there. Move the entries next to related tests so sibling PRs touch different lines. No source or build behaviour change.

assistant-librarian · 2026-10-12T14:17:17Z

TheRock Submodule Bump Activity

Newest first

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

kliegeois and others added 2 commits September 18, 2026 00:47

kliegeois requested a review from a team as a code owner September 18, 2026 17:27

github-actions Bot added the project: rocsparse label Sep 18, 2026

assistant-librarian Bot added the organization: ROCm label Sep 18, 2026

kliegeois force-pushed the aisparse-699-700-common-helpers branch from 0cf9a53 to b8bb55b Compare September 18, 2026 19:43

Merge branch 'develop' into aisparse-699-700-common-helpers

e62583a

kliegeois and others added 4 commits September 29, 2026 13:31

Merge branch 'develop' into aisparse-699-700-common-helpers

addbee6

refactor(rocsparse): use rocsparse_grid.hpp helpers (AISPARSE-699-700)

6c8987a

kliegeois mentioned this pull request Oct 7, 2026

AISPARSE-687: clamp grid.x in the dense conversion paths and add grid-stride loops #11830

Merged

kliegeois added 3 commits October 8, 2026 03:59

Merge branch 'develop' into aisparse-699-700-common-helpers

4706504

# Conflicts: # projects/rocsparse/clients/unittests/CMakeLists.txt

Merge branch 'develop' into aisparse-699-700-common-helpers

2360a5c

# Conflicts: # projects/rocsparse/clients/unittests/CMakeLists.txt

Merge branch 'develop' into aisparse-699-700-common-helpers

7361e92

# Conflicts: # projects/rocsparse/clients/unittests/CMakeLists.txt