csrmv LRB: fix unsigned 32-bit overflow in grid_size (AISPARSE-662) by kliegeois · Pull Request #11215 · ROCm/rocm-libraries

kliegeois · 2026-08-24T17:30:12Z

Summary

Part of the level-2 large-grid / 64-bit indexing epic (AISPARSE-654). Fixes an unsigned 32-bit overflow in the CSR SpMV LRB (long-row-block) launch and hardens the affected kernels.

rocsparse_csrmv_template_lrb.cpp: widen the grid_size product to int64_t before assignment and clamp the launch grids to handle->properties.maxGridSize[0].
csrmv_device.h: add grid-stride loops to the affected LRB kernels so oversized grids are handled correctly.

Test

Added unit_test_internal_level2_csrmv_lrb_grids.cpp to rocsparse-unit-test-device. It shrinks maxGridSize[0] to 1, 2, 3 and 7 blocks so that all five LRB kernels and adaptive partial_scale_y run their grid-stride loops on a small matrix covering LRB bins 0–14, then checks y per bin against a host reference (f32/f64, i32/i64, beta = −3 and 0), with unclamped controls. HIP launch logs confirm the shrunk grids, and long_rows stays a whole multiple of blocks-per-row. The tests fail if any kernel's grid-stride loop is removed.

Test plan

rocsparse-test --gtest_filter=*csrmv* on gfx1201

JIRA ID : AISPARSE-662

Made with Cursor

The LRB csrmv variant computed its launch grid as a uint32_t product of a bin row count and a per-row workgroup count (nRowsBins[j] * num_wgs_per_row). This product wraps at 2^32, which both undersizes the kernel launch and, via max_required_grid, undersizes the wg_flags synchronization buffer. - Widen the grid products to int64_t before assignment in both the analysis buffer sizing and the long-rows launch; size max_required_grid (and thus wg_flags) from the unwrapped 64-bit value. - Clamp every LRB launch grid (short rows, short rows 2, medium warp-reduce, medium one-block, long rows) to handle->properties.maxGridSize[0]. The long rows cap is rounded down to a whole number of per-row workgroup groups so all cooperating workgroups of a row stay within one grid-stride wave, preserving the spin-loop hand-off. - Add grid-stride loops to the affected LRB device kernels so the clamped grid still covers the full logical work range; pass the total work count to the medium one-block and long-rows kernels. - Also size the adaptive partial_scale_y launch for the scale operation itself, clamp it to maxGridSize, and grid-stride its kernel. Contributes to AC3 of ROCM-28951 (AISPARSE-662). Co-authored-by: Cursor <cursoragent@cursor.com>

Co-authored-by: Cursor <cursoragent@cursor.com>

…2-csrmv-lrb-overflow

therock-pr-bot · 2026-08-24T17:42:50Z

✅ 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_grid_dispatch.cpp`, `projects/rocsparse/clients/unittests/unit_test_internal_level2_csrmv_lrb_grids.cpp`, `projects/rocsparse/library/src/level2/csrmv_device.h`, `projects/rocsparse/library/src/level2/rocsparse_csrmv_template_adaptive.cpp`, `projects/rocsparse/library/src/level2/rocsparse_csrmv_template_lrb.cpp`; 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-08-24T17:42:52Z

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

codecov · 2026-08-24T20:21:35Z

Codecov Report

❌ Patch coverage is 18.06084% with 431 lines in your changes missing coverage. Please review.

Files with missing lines	Patch %	Lines
...ojects/rocsparse/library/src/level2/csrmv_device.h	0.00%	284 Missing ⚠️
...ibrary/src/level2/rocsparse_csrmv_template_lrb.cpp	39.58%	142 Missing and 3 partials ⚠️
...y/src/level2/rocsparse_csrmv_template_adaptive.cpp	0.00%	2 Missing ⚠️

❌ Your project check has failed because the head coverage (74.27%) 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   #11215      +/-   ##
===========================================
- Coverage    71.93%   71.87%   -0.06%     
===========================================
  Files         2863     2864       +1     
  Lines       472710   473127     +417     
  Branches     69829    69833       +4     
===========================================
+ Hits        340010   340031      +21     
- Misses      109117   109513     +396     
  Partials     23583    23583

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`49.33% <ø> (ø)`		Carriedforward from 3c1d85f
TensileLite-Unit	`83.29% <ø> (ø)`		Carriedforward from 3c1d85f
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 3c1d85f
hipBLASLt	`38.91% <ø> (ø)`		Carriedforward from 3c1d85f
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 3c1d85f
hipDNN	`86.64% <ø> (-<0.01%)`	⬇️	Carriedforward from 3c1d85f
hipFFT	`43.43% <ø> (ø)`		Carriedforward from 3c1d85f
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 3c1d85f
hipSOLVER	`68.96% <ø> (ø)`		Carriedforward from 3c1d85f
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 3c1d85f
rocBLAS	`48.25% <ø> (ø)`		Carriedforward from 3c1d85f
rocFFT	`47.62% <ø> (ø)`		Carriedforward from 3c1d85f
rocRAND	`57.45% <ø> (ø)`		Carriedforward from 3c1d85f
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from 3c1d85f
rocSPARSE	`74.27% <18.06%> (-0.36%)`	⬇️
rocThrust	`93.15% <ø> (ø)`		Carriedforward from 3c1d85f

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

Files with missing lines	Coverage Δ
...y/src/level2/rocsparse_csrmv_template_adaptive.cpp	`77.17% <0.00%> (-0.16%)`	⬇️
...ibrary/src/level2/rocsparse_csrmv_template_lrb.cpp	`51.52% <39.58%> (-3.53%)`	⬇️
...ojects/rocsparse/library/src/level2/csrmv_device.h	`0.00% <0.00%> (ø)`

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

…short-rows grids (AISPARSE-662) The two LRB short-rows launches computed their workgroup count with a float-based ceiling (rocsparse::ceil((float)nRowsBins / block)), which loses precision for bins with more than 2^24 rows. Switch to the overflow-safe integer idiom (count - 1)/block + 1 to match the rest of the LRB overflow fix and the sibling-PR review conventions. Review notes for AISPARSE-662: - All ticket sites (analysis buffer sizing, long-rows launch, short-rows, medium warp-reduce, medium one-wg-per-row, and the adaptive scale_y launch) widen the workgroup product to int64 before multiplying, clamp the launch grid via rocsparse::min against maxGridSize[0], and the kernels grid-stride over the full logical range. - max_required_grid (wg_flags sizing) is accumulated in int64 from the unclamped product, so the sync buffer is never undersized. - The long-rows launch rounds the clamped grid down to a whole number of per-row workgroup groups, keeping every cooperating workgroup of a row in the same grid-stride wave and preserving the spin-loop hand-off. - The csrmv_device.h diff is large but is entirely grid-stride wrapping (bodies re-indented into for-loops); ignoring whitespace it is only ~74/23 lines and touches no untouched code, so there is no pure-format churn to restore. Co-authored-by: Cursor <cursoragent@cursor.com>

kliegeois

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

…SPARSE-662) partial_scale_y_device strided gid in J. Once the number of rows to scale exceeds 2^30, gid += stride overflows int32, wraps negative, takes the gid < first_row branch and writes y out of bounds. Use int64_t for gid and stride.

…SPARSE-662) The grid-stride loops iterate in J, but wg_row_start, wg_row_end and row_idx were narrowed to uint32_t, so with J = int64_t a short-row bin of 2^32 or more rows wrapped onto low rows. Keep the bin offset in J and the in-workgroup offset (bounded by BLOCKSIZE / rows_per_wg) in uint32_t.

…ISPARSE-662) n_rows_bins[bin_id] does not depend on the grid-stride index but was re-read every iteration. Load it once before the loop, as the other LRB kernels already do, instead of relying on the compiler to hoist it.

…oops (AISPARSE-662) blockreduce_sum already ends with a barrier; after it only lid 0 reads partialSums[0], and only lid 0 writes that slot in the next iteration, so the extra barrier guarded no hazard. In the long-rows kernel it also held the other waves while lid 0 spin-waited. The short-row kernels keep theirs (needed).

…SE-662)

YvanMokwinski · 2026-09-29T19:14:32Z

I was going through the source code modification, and wondering why there is so many C-style cast?

…662)

kliegeois · 2026-09-29T20:08:02Z

I was going through the source code modification, and wondering why there is so many C-style cast?

@YvanMokwinski the commit a6022d6 replaces them with static casts.

…artial_scale_y (AISPARSE-662) Shrink handle->properties.maxGridSize[0] to 1, 2, 3 and 7 blocks for the compute stage so every LRB kernel (short_rows, short_rows_2, medium_rows_warp_reduce, medium_rows, long_rows) and the adaptive partial_scale_y kernel run their grid-stride loops on a small matrix, and compare y against a host reference. float/double, int32/int64 indices, alpha = 2, beta = -3 (and beta = 0 for partial_scale_y), plus unclamped controls.

…lamped (AISPARSE-662) The grid-stride conversion made every csrmv LRB launch pay for the large-grid case: a 64-bit loop over hipGridDim_x, 64-bit row offsets, and in the short-row kernels an extra block barrier per iteration, even though the loop runs once whenever the grid is not clamped. On gfx942 this raised VGPR and SGPR counts in all five kernels and slowed normal SpMV. Give each LRB kernel a GRID_STRIDE template parameter. The host launches the straight-line variant whenever get_grid_size_x returns the full block count, and the grid-stride variant only when the clamp binds. The straight-line variant is develop's kernel, with 32-bit row offsets; they cannot wrap there because blocks * blockDim.x < 2^32. The grid-stride variant keeps the 64-bit loop, the barrier between short-row iterations, and the long-rows grid rounded to a multiple of num_wgs_per_row. Both variants share one per-block body. Launches go through rocsparse::dispatch_grid_stride_x, added here unchanged from the AISPARSE-677 branch; the long-rows site dispatches on its rounded grid directly. The forced-clamp unit tests run the grid-stride variants and the unclamped control runs the straight-line ones.

…ARSE-673) The grid-stride conversion made every default-algorithm sddmm launch pay for the large-grid case: sddmm_coox_kernel, sddmm_csx_kernel and sddmm_ell_kernel all walked a 64-bit grid-stride loop, although it runs once unless get_grid_size_x clamps grid.x, which needs more than (2^32 - 1) / blockDim.x blocks. On gfx942 this added 12 to 24 SGPRs to each of these kernels. Occupancy stayed at 8 waves per SIMD, but the extra loop and 64-bit index arithmetic slowed small-k SDDMM on MI300X by about 3% geomean, and by up to 11% (CSC float, k = 2). Give the three kernels a GRID_STRIDE template parameter. The host launches them through rocsparse::dispatch_grid_stride_x, which runs the straight-line variant whenever get_grid_size_x returns the full block count and the grid-stride variant only when the clamp binds. Both variants share one device body. The device functions take the block index as a deduced type and form the nonzero, row or column index with develop's expression: the straight-line variant passes hipBlockIdx_x as a uint32_t, and the grid-stride variant passes an int64_t block index. For COO and ELL the nonzero index is kept as wide as the block index, so the last grid-stride block cannot wrap a 32-bit I when nnz is near INT32_MAX. On gfx942 all 432 straight-line instantiations have the same SGPR, VGPR, scratch and occupancy counts as develop, and the same instruction stream once labels are normalized. rocsparse_grid.hpp is taken unchanged from the AISPARSE-662 branch (blob bd92dbc), so it merges cleanly with ROCm#11215 and ROCm#11824. The SddmmGridClamp tests with maxGridSize[0] at 3 and 1 run the grid-stride variants; the device-limit control runs the straight-line ones.

…v LRB kernels (AISPARSE-662) Hoist the 64-bit start index of the grid-stride loops in partial_scale_y_device and the GRID_STRIDE path of csrmvn_lrb_medium_rows_warp_reduce into a named const, next to the stride, so each loop header reads `for(idx = gid; idx < n; idx += stride)`. No functional change; gfx942 ISA is identical.

ntrost57 · 2026-10-02T07:44:16Z

all kernel instantiations doubled now, for performance reasons?

… LRB grid tests (AISPARSE-662)

…lection (AISPARSE-662)

kliegeois · 2026-10-05T18:42:49Z

all kernel instantiations doubled now, for performance reasons?

@ntrost57 Yes. Same reason as in #12307 (comment) .

YvanMokwinski · 2026-10-05T20:28:36Z

all kernel instantiations doubled now, for performance reasons?

@ntrost57 Yes. Same reason as in #12307 (comment) .

@kliegeois May I ask if you have some measurments of the performance drop it induces?

kliegeois · 2026-10-05T20:31:39Z

all kernel instantiations doubled now, for performance reasons?

@ntrost57 Yes. Same reason as in #12307 (comment) .

@kliegeois May I ask if you have some measurments of the performance drop it induces?

@YvanMokwinski yes, I have some measurements that were flagged as performance regression.

YvanMokwinski · 2026-10-05T20:42:11Z

all kernel instantiations doubled now, for performance reasons?

@ntrost57 Yes. Same reason as in #12307 (comment) .

@kliegeois May I ask if you have some measurments of the performance drop it induces?

@YvanMokwinski yes, I have some measurements that were flagged as performance regression.

@kliegeois What is roughly the order of magnitude? 10%?

kliegeois · 2026-10-05T20:54:38Z

all kernel instantiations doubled now, for performance reasons?

@ntrost57 Yes. Same reason as in #12307 (comment) .

@kliegeois May I ask if you have some measurments of the performance drop it induces?

@YvanMokwinski yes, I have some measurements that were flagged as performance regression.

@kliegeois What is roughly the order of magnitude? 10%?

@YvanMokwinski I am not allowed to comment on that here but the last comment on the JIRA ticket has the results.

…lock index (#12307) Resolves AISPARSE-671 and AISPARSE-672. Both csrmm paths compute a 64-bit block count and assign it straight into `dim3`, which is `unsigned int`, then read the index back into a plain `int` inside the kernel — so the value was truncated twice over. Neither narrowing is a launch error, so nothing reports it. On the nnz split path the clamp first binds at about 4.3e9 non-zeros (one 256-thread block per 256 non-zeros, so beyond 16,777,215 blocks), and in `csrmmnn_general_block_reduce` at 4,194,304 dense columns (1024-thread blocks). The merge path cannot reach its clamp: it already rejects m or nnz above `UINT32_MAX`, which keeps m + nnz below the roughly 1.7e10 (wave64) at which its grid would clamp. ### AISPARSE-671, merge path Four grid.x extents in `rocsparse_csrmm_template_merge.cpp` clamped with `rocsparse::get_grid_size_x`, i.e. to min(`maxGridSize[0]`, (2^32 - 1) / blockDim.x), the four `__global__` wrappers grid-stride over the full merge-block count, and `bid` is widened to match the type of `block_count`. ### AISPARSE-672, nnz split path Six grid.x extents clamped, four wrappers grid-strided, and the six `__device__` functions take the block index as an `int64_t` parameter instead of reading `hipBlockIdx_x`. Widening is half the fix: a clamped grid with a 32-bit index still truncates. **The clamp was not safe on its own, and this is the part worth reviewing closely.** On develop the reduction producer writes `val_block_red[hipGridDim_x * (colB + i) + bid]` while the consumer reads `val_block_red[idx + nblocks * col]`, and the host allocation is sized from `nblocks`. There is even a comment on develop recording the assumption: "hipGridDim_x equals n". Producer and consumer agree only while the grid is unclamped. Clamping grid.x without reconciling them would have silently turned the reduction into a wrong-answer bug. Both sides now derive the stride from `nblocks`, so they agree by construction. Also in this commit, and not mentioned by the ticket: `csrmmnn_general_block_reduce` launches one block per dense column and takes its `val_block_red` batch stride from `hipGridDim_x` for the same reason, so it is clamped and grid-strided over the dense columns too. The `alpha == 0 && beta == 1` shortcut now marks every nnz block the wrapper owns rather than only `row_block_red[hipBlockIdx_x]`. The pre-existing partial mitigation in `csrmmnn_general_block_reduce_device` (an inner loop over `nblocks`) was extended rather than duplicated. ### Where the helpers live The launches use `rocsparse::get_grid_size_x` / `get_grid_size_y` and `rocsparse::dispatch_grid_stride_x` from `rocsparse_grid.hpp` (the same helpers as #11215 and #11824), so there are no file-local helpers. When grid.x is not clamped, `dispatch_grid_stride_x` launches the original straight-line kernels (`GRID_STRIDE = false`); the grid-stride variants run only when the clamp binds. The wrappers drive the loops and pass the block identity down, which keeps the pre-instantiated kernel signatures and their 21 instantiation translation units untouched. ### Verification - New tests in `rocsparse-unit-test-device`, permanently enabled, no device-memory guard. The clamp is forced by shrinking `handle->properties.maxGridSize[0]` through a scoped RAII guard rather than by allocating an unreachable matrix. - `*csrmm*:*spmm_csr*` regression suite passes. - Break-validation with `CCACHE_DISABLE=1`: breaking the stride loops and breaking the reduction strides fail different, disjoint sets of cases, and the full-grid controls pass throughout. ### Notes for review - **Review with `git diff -w`.** Roughly 295 of the library insertions are macro-continuation reflow: a longer first macro line pushes every trailing backslash right, so two-line changes appear as eighty. - No interaction with the batching work; both paths are non-batched. - Despite the similar filename, this does not touch `level2/rocsparse_csrmv_template_nnzsplit.cpp` and has no overlap with the csrmv nnzsplit perf PR. - Reaching the clamp for real needs about 4.3e9 non-zeros on the nnz split path (tens of GB of matrix), so it is exercised through an artificially shrunk grid limit rather than end to end. --------- Co-authored-by: Cursor <cursoragent@cursor.com>

assistant-librarian · 2026-10-09T01:19:51Z

TheRock Submodule Bump Activity

Newest first

2026-10-09 — Included in TheRock via ROCm/TheRock#8860.

kliegeois and others added 3 commits August 19, 2026 16:53

style(rocsparse): apply clang-format-14 to large-grid changes

f543f2d

Co-authored-by: Cursor <cursoragent@cursor.com>

Merge remote-tracking branch 'upstream/develop' into task/aisparse-66…

db34ad7

…2-csrmv-lrb-overflow

kliegeois requested a review from a team as a code owner August 24, 2026 17:30

github-actions Bot added the project: rocsparse label Aug 24, 2026

therock-pr-bot Bot added the Not ready to Review label Aug 24, 2026

assistant-librarian Bot added the organization: ROCm label Aug 24, 2026

therock-pr-bot Bot removed the Not ready to Review label Sep 3, 2026

Merge branch 'develop' into aisparse-662-csrmv-lrb-overflow

4a78b6b

jsandham reviewed Sep 25, 2026

View reviewed changes

Comment thread projects/rocsparse/library/src/level2/rocsparse_csrmv_template_adaptive.cpp Resolved

jsandham reviewed Sep 25, 2026

View reviewed changes

Comment thread projects/rocsparse/library/src/level2/rocsparse_csrmv_template_lrb.cpp Resolved