AISPARSE-692/693/694: gtsv offsets, zero-batch guard and BSR row-map bounds checks by kliegeois · Pull Request #12316 · ROCm/rocm-libraries

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

JIRA ID: AISPARSE-692, AISPARSE-693, AISPARSE-694

Three commits, one per ticket, in that order. Grouped because they are the remaining
three tickets of the preconditioner epic and they overlap each other's files.

AISPARSE-692 — gtsv interleaved batch offsets (P2)

gtsv_interleaved_batch_device.h computed batch_count * k + gid in rocsparse_int.
The interleaved layout makes the arrays batch_count * m elements, so the product
walks the full extent of the data: at batch_count 50,000 with m 50,000 it reaches
2.5e9 and overflows. Reachable — four diagonals plus the RHS at 2.5e9 doubles is
roughly 100 GB, which fits one MI300X. The offsets are now int64_t.

The grid itself was already correct: rocsparse_gtsv_interleaved_batch.cpp puts
batch_count on grid.x with a ceiling divide and an in-kernel bounds check. The
defect was purely the index arithmetic.

Scope is the interleaved kernel only. The two non-interleaved sites the ticket
names — gtsv_device.h and gtsv_nopivot_strided_batch_medium_device.h — are
genuinely fixed by #12136, which widened bidy to int64_t across both files. They
remain 32-bit on develop. They are not redone here, and git merge-tree against
#12136 reports zero conflicts.

Also confirmed: gtsv_transpose_and_pad_array_kernel is dead code; only the
_shared_ variant has launch sites.

AISPARSE-693 — zero-batch guard and grid rounding (P3)

Four preprocess dispatches in rocsparse_csritilu0_async_inplace.cpp sized the grid
as (target_size_ + BLOCKSIZE - 1) / BLOCKSIZE with BLOCKSIZE 1024. target_size_ is
rocsparse_int, so the addition overflows signed 32-bit at 2,147,482,625 — undefined
behaviour, before the division can bring it back in range. Switched to the
(x - 1) / BLOCKSIZE + 1 form the rest of the directory uses, which cannot overflow
for any representable positive x. This is a real arithmetic fix.

The zero-batch guard is added to csric0 and csrildlt0, making all five entry
points agree. It is dead code, and the commit message says so.
rocsparse_create_csr_descr initialises batch_count = 1, and all seven
Xxx_set_strided_batch entry points reject batch_count <= 0 with
rocsparse_status_invalid_value before storing it. The suggested test — a descriptor
with rows and batch_count = 0 — cannot be written against the public API, so none
was written rather than ship something misleading. The quick return stays a success
rather than introducing a second, different answer for a condition the
argument-checking layer already rejects.

The csrilu0_kernel_binsearch zero-batch-stride gap is left out: it is a batch-axis
issue and belongs with #12137.

AISPARSE-694 — grid.x clamp and missing bounds checks (P2)

All eight missing BSR bounds guards are added. The ticket is exactly right here:
eight kernels index the row map with the raw block index and never compare it to mb,
and bsrilu0_kernel_general was the only one already checking. Adding
if (idx >= mb) return turns a potential out-of-bounds read into a clean early exit
and stands on its own merits.

All nine BSR launches are converted. Every BSR ic0/ilu0 launch clamps grid.x with get_grid_size_x and every block walks the contiguous chunk of rows grid_x_chunk gives it. The seven CSR launches the ticket also lists are left for a follow-up, to keep the conflict surface against #12137 small.

The recovery loop uses contiguous chunks, not a grid-stride loop, and that is
load-bearing. The map is a topological order, so contiguous chunks preserve the
invariant that a block only ever waits on itself or a lower-numbered block. A
grid-stride loop breaks it: block 0's second item waits on block grid_size-1's
first, which cannot start until a block retires, and none retires while spinning.
That is a deadlock.

The seven block-dim specific kernels (bsric0_kernel_2_8, _2_8_unrolled, _9_16, _17_32 and bsrilu0_kernel_2_8, _9_32, _33_64) hold __shared__ arrays, but they are only launched on wave64 devices, where a 64-thread block is a single wavefront, so reusing the arrays for the next row needs no barrier. The launchers send every wave32 device to the general kernels. Their grid.x limit is (2^32 - 1) / blockDim.x, which they reach well before 2^31 block rows (about 5.4e8 rows for the 8x8 kernels, 2.7e8 and 1.3e8 for bsrilu0_9_32 and 6.7e7 for bsrilu0_33_64).

The seven CSR launches are genuinely safe to convert (their shared memory is
partitioned per wavefront, so no barrier is needed); they are left for a follow-up to
keep the conflict surface against #12137 small.

Sequencing

None of the three referenced PRs is merged. #12136, #12137 and #11512 are all sibling
branches based on an older develop; this branch is based on current develop, so their
fixes are not in this tree and the affected sites are still broken on develop.

#12137 and 694 conflict. The kernel files from this PR conflict with #12137 on the dim3(...) expressions and the kernel wrapper (9 files, all in precond/bsric0 and precond/bsrilu0); the resolution is to nest the row chunk loop inside the batch loop. CHANGELOG.md auto-merges. #12137 should merge first.

The csritilu0 rounding fix duplicates #11512 (AISPARSE-696), which fixes the same
four lines using a new ceil_div helper. Four one-line expressions, trivially
resolved either way; kept here deliberately and confirmed empirically. Note the ticket
points at the wrong commit for this overlap — that commit is AISPARSE-668
(csrmm/coomm) and does not touch the file.

Hardware queried rather than assumed: maxGridSize = [2147483647, 65535, 65535].

Testing

Builds clean, no new warnings. 12,866 preconditioner tests pass, zero failures
across *gtsv*, *ic0*, *ilu0*, *ildlt0*, *itilu0*, excluding known_bug. 30
"Insufficient memory" skips, which is the suite's guard working.

Seven new host-side cases in rocsparse-unit-test (CPU-only binary, no memory
requirement declared, so nothing can drop them at instantiation). All 12 new gtsv
interleaved cases ran and passed.

Break-validation of the clamp — broken to a bare cast, rebuilt with
CCACHE_DISABLE=1:

grid_size
  Which is: 705032704
2147483647u

705032704 is 5,000,000,000 wrapped mod 2^32. Restored: 7/7 pass, tree clean.

No BSR test here is value-checking, and none was written. near_check_values only
compares when no pivot was reported, and the random BSR generator reports a pivot on
every batch — reducing bsric0_kernel_general to a no-op still passes all 131 quick
BSR spic0 tests. The BSR suite is used here only to confirm the eight new guards did
not change behaviour for valid input. The clamped-grid test (unit_test_internal_precond_grid_x_chunk.cpp) covers the general kernels on every device and the seven block-dim specific kernels on wave64 devices; on wave32 those cases skip because the launcher never selects those kernels. On an MI300A (wave64) its 36 cases pass and are not skipped.

Notes for review

694 clamps all nine BSR launches; the seven CSR launches are left for a follow-up. The eight bounds guards stand independently.
The duplicated csritilu0 fix against AISPARSE-696/697/668: shared grid helpers, batch grid.y clamp, and csrmm/coomm large dense column counts #11512 is a judgement call — drop it if
AISPARSE-696/697/668: shared grid helpers, batch grid.y clamp, and csrmm/coomm large dense column counts #11512 lands first.
clang-format 18.1.4 reformats pre-existing code in
rocsparse_csritilu0_async_inplace.cpp (the buffer_layout_inplace_t(){} and
light_run call). That churn was reverted to keep the commit clean, so the file is
not 18.1.4-clean as a whole. Worth confirming which version CI uses.

…t (AISPARSE-692) The interleaved batch layout stores each of the four diagonals and the right-hand side as batch_stride * m elements with the batch as the fastest varying axis, so the row offsets batch_count * k and batch_stride * k inside the three solver kernels walk the full extent of the data. Both operands were rocsparse_int, so the products overflowed signed 32-bit arithmetic at batch_count * m = 2^31 while each operand on its own was still far inside range. At batch_count = 50000 with m = 50000 the offsets reach 2.5e9; four diagonals plus the right-hand side at that size is about 100 GB in double precision, which fits on a single MI300X, so this is reachable rather than theoretical. The grid is not at fault: all three launches in rocsparse_gtsv_interleaved_batch.cpp put batch_count on grid.x with a ceiling divide and the kernels bounds check gid against batch_count. batch_count and batch_stride become int64_t in the thomas, lu and qr kernel signatures. Widening the operands rather than casting at each use promotes every one of the roughly ninety offsets in the file, including the ones written inline as array subscripts, so no site can be missed. The offsets those products feed (index, minus, ind_k, ind_k_1, ind_k_2) become int64_t as well, since a 64-bit product assigned back into a rocsparse_int would truncate and gain nothing. gid is widened for the same reason. It was computed as hipThreadIdx_x + BLOCKSIZE * hipBlockIdx_x in unsigned 32-bit arithmetic and then stored in a rocsparse_int, so for a batch_count within BLOCKSIZE of 2^31 the threads of the last block produced a negative gid that passed the gid >= batch_count guard and indexed out of bounds. It is now computed and held in 64 bit. The row counters k, i and s stay rocsparse_int: they are bounded by m, which the public interface already limits to rocsparse_int, and the pivot values in p are row indices rather than offsets. Only the interleaved batch family is left to fix. The other two sites named by the ticket are already covered by PR ROCm#12136. Its AISPARSE-689 commit rewrote seven of the eight gtsv_device.h kernels around an int64_t bidy grid-stride loop, which makes m_pad * bidy and gidx + bidy * stride 64 bit; the eighth, gtsv_transpose_and_pad_array_kernel, keeps a rocsparse_int bidy but has no launch site anywhere in the library and is dead code. Its AISPARSE-691 commit does the same for gtsv_nopivot_strided_batch_medium_device.h, where stride * bidy, m * bidy, ldb * bidy and num_spikes * bidy are all 64 bit once bidy is. The small paths of the no-pivot families were never affected, because their batch_stride and ldb parameters are already int64_t; the interleaved batch family is the outlier in taking batch_stride as rocsparse_int. The new test is a coverage case for the edited expressions, not a reproducer. Reaching 2^31 needs about 100 GB of device memory, so the overflow itself cannot be exercised on the 15 GB card available here: reverting this commit entirely leaves all twelve new cases passing. What the case does catch is a widened offset that picks up the wrong multiplier or truncates on the way back into a 32-bit local, because batch_stride is set well above batch_count; swapping batch_count for batch_stride in the thomas forward sweep faults the kernel on the batch_count-sized workspace. M = 513 is past the m > 2 and m > 3 special cases, and all three algorithms are covered, where the existing case at that shape covers thomas only. The case declares host_memory_gb and device_memory_gb of 1 rather than inheriting the suite default of 4. The default is compared against sysinfo freeram, which excludes page cache, so on a loaded machine the whole file is dropped at gtest instantiation with no failure and no test names. The real requirement here is about 70 MB. Co-authored-by: Cursor <cursoragent@cursor.com>

…w (AISPARSE-693) Two unrelated defects in the incomplete factorization preconditioners, both found by inspection rather than by a failing test. rocsparse::csric0 and rocsparse::csrildlt0 quick returned on A->rows == 0 but not on A->batch_count == 0, while rocsparse::bsric0, rocsparse::bsrilu0 and rocsparse::csrilu0 all test both. Every launch in the two odd ones out puts the batch count on grid.y, and zero is not a legal grid extent, so a descriptor carrying rows with a zero batch count would have reached an invalid launch configuration. Both now test both terms, which makes all five agree. That path is not reachable through the public API today, and this commit claims no user visible fix for it. rocsparse_create_csr_descr and its siblings initialise batch_count to 1, and all seven Xxx_set_strided_batch entry points reject batch_count <= 0 with rocsparse_status_invalid_value before it is stored, so no caller can present a descriptor with a zero batch count. The guard is defensive hardening against a future internal caller that constructs a descriptor directly, and the reason it is worth adding is consistency: the five entry points should not disagree about what a degenerate descriptor means. For the same reason the quick return is a success rather than a new invalid_value rejection. The library already rejects a zero batch count, and it does so in the argument checking layer at the setter; turning the unreachable internal case into an error would put a second, different answer for the same condition in a second place. It also means the suggested test, a descriptor with rows and batch_count = 0 asserting a clean return, cannot be written against the public API, so none is added here. The second defect is real arithmetic. Four preprocessing dispatches in rocsparse_csritilu0_async_inplace.cpp sized the grid as (target_size_ + BLOCKSIZE - 1) / BLOCKSIZE with BLOCKSIZE 1024 and target_size_ a rocsparse_int, so the addition overflowed signed 32 bit at 2147482625. That is undefined behaviour rather than a defined wrap, so the resulting grid extent is whatever the optimiser decides it is. The rest of the itilu0 directory writes the same rounding as (x - 1) / BLOCKSIZE + 1, which cannot overflow for any representable positive x because the subtraction moves away from the maximum. Switching these four to that form is the whole fix; no grid-stride loop is needed, since the kernels already bounds check against target_size_. That file is also edited by PR ROCm#11512 (AISPARSE-696), which fixes the same four lines with a new rocsparse::ceil_div helper in rocsparse_common.hpp. The fix here is semantically identical and textually different, so whichever lands second takes a four line conflict that is resolved by keeping either side. It is duplicated deliberately: ROCm#11512 is unmerged and based on an older develop, the overlap is four single-line expressions rather than a kernel rewrite, and leaving the only arithmetic defect in this ticket unfixed on the chance that another branch lands first is the worse trade. The ticket named 73c21ae as the overlapping commit; that commit is AISPARSE-668 and does not touch this file at all. The zero-batch-stride gap in rocsparse_csrilu0_kernel_binsearch.cpp is deliberately not included. It is an AISPARSE-690 finding about the done array being sized from the guarded count while the kernel walks it with an unguarded batch_stride. That is the batch axis, and it belongs with the rest of the batch axis work in PR ROCm#12137 rather than with either half of this ticket. Co-authored-by: Cursor <cursoragent@cursor.com>

…SE-694) The incomplete factorizations take the block row count from the sparse matrix descriptor, where it is an int64_t, and hand it to dim3, which takes a uint32_t. gfx1201 reports maxGridSize[0] == 2147483647, so a count above that is narrowed with no clamp and no diagnostic: between 2^31 and 2^32 the narrowing produces a value dim3 accepts but the driver rejects, and at 2^32 it wraps, so 2^32 block rows launch an empty grid and 2^32 + 7 launch seven blocks. The products themselves are safe, since int64_t times int promotes; the defect is purely the narrowing plus the absence of any way to recover the work above the clamp. Reachable only through rocsparse_spic0, rocsparse_spilu0 and rocsparse_spildlt0, which pass the descriptor's row count through untouched. The legacy entry points take a rocsparse_int row count that saturates below the cap, so this is confined to the generic API and needs no BUILD_ROCSPARSE_ILP64. Two changes here, of different weight. The first is a plain robustness fix and it is complete. Eight of the nine BSR kernels indexed the row map with the raw block index and never compared it against mb, so a grid rounded up past the row count read outside the map allocation. bsrilu0_kernel_general already checked, and so does every CSR kernel in the family, which left these eight as the only unguarded readers of a level-scheduled map. All eight now check, spelled the same way as the kernel that already did. The second is the clamp, and it is deliberately partial. rocsparse_grid_x.hpp adds get_grid_size_x, which clamps an extent to the grid.x maximum, and grid_x_chunk, which hands a block the half-open range of items it owns when the items are split into contiguous chunks. Both are applied to bsric0_kernel_general and bsrilu0_kernel_general and to nothing else. The chunks are contiguous rather than a grid-stride loop, and that is the whole design. These kernels have no barrier anywhere; they synchronise by spinning on per-row done flags with __threadfence, and the row map is a topological order, so an item only ever waits on items earlier in that order. Contiguous chunks keep the property the unclamped launch relies on today, that a block only ever waits on itself or on a lower numbered block, which the dispatcher started first. A grid-stride loop destroys it: block 0 on its second item waits on the first item of block grid_size - 1, which cannot start until a block retires, and no block retires while it is spinning. That is a deadlock, not a slowdown. Only these two kernels are converted because only these two hold no shared memory, so for them the recovery loop needs no barrier between iterations and the row axis argument is finished. Every scratch location either kernel touches (the done array, the zero-pivot position) is addressed by the logical row read from the map and never by the block index, so a block may be reused across iterations safely. The other seven BSR kernels are not converted, and the reason is a hazard rather than a shortage of effort. bsric0_kernel_2_8, _2_8_unrolled, _9_16 and _17_32 and bsrilu0_kernel_2_8, _9_32 and _33_64 each hold between one and seven __shared__ arrays that are written and read across the whole block, and they launch 64-thread blocks, which is two wavefronts on gfx1201 at wavefront 32. Wrapping those in a row loop introduces a write-after-read race on the shared arrays that does not exist in the current one-row-per-block code, and closing it needs a __syncthreads() in a family that deliberately has none. The barrier is probably safe, since the loop bound and the spin conditions are all block uniform, but "probably" is the wrong standard for a deadlock on a path that cannot be executed here: reaching the clamp needs more than 2^31 block rows, which is upwards of 17 GB of row pointers alone against a 15 GB card. Shipping an untestable barrier into a barrier-free spin-wait family is the worse risk, so these seven keep the narrowing and the ticket stays open for them. The seven CSR launches are also unconverted. They are safe to convert, and for a different reason: their shared memory is partitioned per wavefront, so each wavefront owns its own slice and no barrier is needed. They are left out only to keep this commit to what it can justify, and because all sixteen dim3 expressions are already being edited on the grid.y axis by PR ROCm#12137, so every site converted here is a conflict there. Testing is host-side and it is honest about its reach. The new unit tests live in rocsparse-unit-test, the CPU-only binary, and declare no memory requirement, so nothing can drop them at instantiation. They check the properties the kernels actually depend on: the clamp never returns an extent the hardware rejects and never truncates one it accepts, the chunks partition [0, count) exactly, the chunks are monotone in the block id, and at any unclamped extent the partition degenerates to one item per block so the recovery loop is a no-op for every problem size that fits on a device. There is no device-side test for the clamp and there cannot be one here. Nor is there a new BSR value test, because such a test would prove nothing: spmat_descr::near_check_values only compares values when no pivot was reported and the random BSR generator reports a pivot on every batch, which is why reducing bsric0_kernel_general to a no-op still passes all 131 existing quick BSR spic0 tests. The existing BSR suite is a launch assertion, and it is used here only to confirm the eight new bounds checks did not change behaviour for valid input. Co-authored-by: Cursor <cursoragent@cursor.com>

therock-pr-bot · 2026-09-18T17:32:47Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 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:32:50Z

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

codecov · 2026-09-22T05:34:10Z

Codecov Report

❌ Patch coverage is 68.83117% with 24 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...src/precond/bsric0/rocsparse_bsric0_kernel_2_8.cpp	0.00%	13 Missing ⚠️
...s/rocsparse/library/src/include/rocsparse_grid.hpp	0.00%	6 Missing ⚠️
...c/precond/bsric0/rocsparse_bsric0_kernel_17_32.cpp	95.00%	0 Missing and 1 partial ⚠️
...nd/bsric0/rocsparse_bsric0_kernel_2_8_unrolled.cpp	94.12%	0 Missing and 1 partial ⚠️
...rc/precond/bsric0/rocsparse_bsric0_kernel_9_16.cpp	92.31%	0 Missing and 1 partial ⚠️
...se/library/src/precond/csric0/rocsparse_csric0.cpp	0.00%	0 Missing and 1 partial ⚠️
...rary/src/precond/csrildlt0/rocsparse_csrildlt0.cpp	0.00%	0 Missing and 1 partial ⚠️

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

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #12316      +/-   ##
===========================================
- Coverage    71.96%   71.89%   -0.07%     
===========================================
  Files         2874     2871       -3     
  Lines       474907   474945      +38     
  Branches     70151    70151              
===========================================
- Hits        341766   341458     -308     
- Misses      109384   109770     +386     
+ Partials     23757    23717      -40

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`49.33% <ø> (ø)`		Carriedforward from 22e6dff
TensileLite-Unit	`83.29% <ø> (-0.04%)`	⬇️	Carriedforward from 22e6dff
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 22e6dff
hipBLASLt	`38.91% <ø> (-3.06%)`	⬇️	Carriedforward from 22e6dff
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 22e6dff
hipDNN	`86.65% <ø> (-0.01%)`	⬇️	Carriedforward from 22e6dff
hipFFT	`43.43% <ø> (+0.91%)`	⬆️	Carriedforward from 22e6dff
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 22e6dff
hipSOLVER	`68.96% <ø> (ø)`		Carriedforward from 22e6dff
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 22e6dff
rocBLAS	`48.24% <ø> (+0.07%)`	⬆️	Carriedforward from 22e6dff
rocFFT	`47.62% <ø> (-1.04%)`	⬇️	Carriedforward from 22e6dff
rocRAND	`57.45% <ø> (ø)`		Carriedforward from 22e6dff
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from 22e6dff
rocSPARSE	`74.30% <68.83%> (+0.23%)`	⬆️
rocThrust	`93.15% <ø> (ø)`		Carriedforward from 22e6dff

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

Files with missing lines	Coverage Δ
...precond/bsric0/rocsparse_bsric0_kernel_general.cpp	`44.62% <100.00%> (ø)`
...c/precond/bsrilu0/rocsparse_bsrilu0_kernel_2_8.cpp	`68.25% <ø> (ø)`
...precond/bsrilu0/rocsparse_bsrilu0_kernel_33_64.cpp	`50.82% <100.00%> (ø)`
.../precond/bsrilu0/rocsparse_bsrilu0_kernel_9_32.cpp	`60.14% <100.00%> (ø)`
...econd/bsrilu0/rocsparse_bsrilu0_kernel_general.cpp	`50.64% <100.00%> (+0.64%)`	⬆️
...c/precond/bsric0/rocsparse_bsric0_kernel_17_32.cpp	`48.74% <95.00%> (+0.43%)`	⬆️
...nd/bsric0/rocsparse_bsric0_kernel_2_8_unrolled.cpp	`54.61% <94.12%> (+0.32%)`	⬆️
...rc/precond/bsric0/rocsparse_bsric0_kernel_9_16.cpp	`48.74% <92.31%> (+0.43%)`	⬆️
...se/library/src/precond/csric0/rocsparse_csric0.cpp	`84.62% <0.00%> (-2.56%)`	⬇️
...rary/src/precond/csrildlt0/rocsparse_csrildlt0.cpp	`43.90% <0.00%> (ø)`
... and 2 more

... and 132 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.

kliegeois

Self-review after the #12617 rebase: I went through this PR for the issues reviewers raised on the other large-grid PRs. 4 items below that I will address.

State the binding grid.x limit as (2^32 - 1) / blockDim.x, and describe the csritilu0 extent change as a consistency change: the unsigned sum could not overflow.

…RSE-692-694) rocsparse_grid_x.hpp existed only so the host-only rocsparse-unit-test binary could include grid_x_chunk without HIP. Keep one grid API instead: grid_x_chunk now sits next to get_grid_size_x in rocsparse_grid.hpp, spelled with __device__ __host__ __forceinline__ like the other host and device helpers, and rocsparse_grid_x.hpp and its ROCSPARSE_GRID_X_ILF macro are gone. unit_test_internal_grid_x.cpp moves from rocsparse-unit-test to rocsparse-unit-test-device, since rocsparse_grid.hpp pulls in the handle. The tests themselves are unchanged. The device code of bsrilu0_kernel_general and bsric0_kernel_general is unchanged: the gfx1201 assembly of both translation units is identical before and after.

…ds (AISPARSE-692-694) Nothing launched bsrilu0_kernel_general or bsric0_kernel_general with grid.x below the row count, so the chunk loop never ran more than one item per block. Add device unit tests that shrink handle->properties.maxGridSize[0], which get_grid_size_x reads, to 1, 2 and 3 blocks for the numeric phase and compare against the full-grid factorization bit for bit. The matrix is double complex with block_dim 33, so both routines take the general kernel on wave32 and wave64 devices. Its 67 block rows each depend on rows 1, 2, 7, 23 and 41 before them, so the spin waits cross chunk boundaries for every grid tested and the last chunk is ragged. Both directions are covered.

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

…d issues list After merging develop, 5.1.0 is a released section; the entries belong in the unreleased 5.2.0 list, at a position that does not collide with the other open PRs' entries.

…ernels (AISPARSE-694) bsrilu0_{2_8,9_32,33_64} and bsric0_{2_8,2_8_unrolled,9_16,17_32} launched one block per block row with grid.x = rows. The dispatch limit is (2^32 - 1) / blockDim.x blocks, so these launches fail well before 2^31 rows: about 5.4e8 for the 8x8 kernels, 2.7e8 / 1.3e8 for bsrilu0_9_32 and 6.7e7 for bsrilu0_33_64. Clamp grid.x with get_grid_size_x and have each block walk the contiguous chunk of rows grid_x_chunk gives it, as bsrilu0_kernel_general and bsric0_kernel_general already do. These kernels are only launched on wave64 devices, where a block is a single wavefront, so reusing the shared arrays for the next row needs no barrier.

…on clamped grids (AISPARSE-694) Parametrize the bsrilu0/bsric0 clamped-grid test over block dimension and precision and add a case per block-dim specific kernel. Those kernels are only launched on wave64 devices, so the cases skip on wave32.

Resolve the bsric0/bsrilu0 kernel conflicts by nesting the grid_x_chunk row loop from #12316 inside the grid-y batch loop, and launching with both the clamped grid.x and the clamped grid.y.

assistant-librarian · 2026-10-16T13:04:29Z

TheRock Submodule Bump Activity

Newest first

2026-10-16 — Included in TheRock via ROCm/TheRock#9032.

kliegeois and others added 3 commits September 18, 2026 00:09

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

github-actions Bot added documentation project: rocsparse labels Sep 18, 2026

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

Merge branch 'develop' into aisparse-692-693-694-precond

b623b87

kliegeois added 2 commits September 29, 2026 13:33

Merge branch 'develop' into aisparse-692-693-694-precond

cf72834

refactor(rocsparse): use rocsparse_grid.hpp helpers (AISPARSE-692-694)

9ea4e71

kliegeois commented Sep 29, 2026

View reviewed changes

kliegeois added 6 commits September 29, 2026 18:27

docs(rocsparse): correct spic0/spilu0 and csritilu0 CHANGELOG entries

dd8cc6b

State the binding grid.x limit as (2^32 - 1) / blockDim.x, and describe the csritilu0 extent change as a consistency change: the unsigned sum could not overflow.

Merge branch 'develop' into aisparse-692-693-694-precond

ba88239

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

Merge remote-tracking branch 'upstream/develop' into HEAD

d2f439e

rocsparse: move the CHANGELOG entries to the unreleased 5.2.0 Resolve…

22e6dff

…d issues list After merging develop, 5.1.0 is a released section; the entries belong in the unreleased 5.2.0 list, at a position that does not collide with the other open PRs' entries.

ntrost57 reviewed Oct 9, 2026

View reviewed changes

Comment thread projects/rocsparse/library/src/precond/bsrilu0/rocsparse_bsrilu0_kernel_2_8.cpp Resolved

kliegeois added 2 commits October 9, 2026 18:03

kliegeois requested a review from ntrost57 October 12, 2026 19:25

ntrost57 approved these changes Oct 13, 2026

View reviewed changes

jsandham approved these changes Oct 13, 2026

View reviewed changes

kliegeois merged commit 39ed988 into ROCm:develop Oct 13, 2026
132 of 134 checks passed

jiagaoxiang mentioned this pull request Oct 15, 2026

testintel: fix severe issue in sellmvt_device (created by kliegeoi (@kliegeois)) #13339

Merged

1 task

Navigation Menu

Repository navigation

AISPARSE-692/693/694: gtsv offsets, zero-batch guard and BSR row-map bounds checks - #12316

AISPARSE-692/693/694: gtsv offsets, zero-batch guard and BSR row-map bounds checks#12316
kliegeois merged 14 commits into
ROCm:developfrom
kliegeois:aisparse-692-693-694-precond

kliegeois commented Sep 18, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 18, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Sep 18, 2026

Uh oh!

codecov Bot commented Sep 22, 2026 •

edited

Loading

Uh oh!

kliegeois left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 16, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

kliegeois commented Sep 18, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

AISPARSE-692 — gtsv interleaved batch offsets (P2)

AISPARSE-693 — zero-batch guard and grid rounding (P3)

AISPARSE-694 — grid.x clamp and missing bounds checks (P2)

Sequencing

Testing

Notes for review

Uh oh!

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

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Sep 18, 2026

Uh oh!

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

Uh oh!

Codecov Report

Uh oh!

kliegeois left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 16, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

kliegeois commented Sep 18, 2026 •

edited

Loading

therock-pr-bot Bot commented Sep 18, 2026 •

edited

Loading

codecov Bot commented Sep 22, 2026 •

edited

Loading