Repository navigation
AISPARSE-677: clamp bsrgemm_calc and csrgemm_calc bucket grids and add grid-stride loops - #11824
Merged
kliegeois merged 17 commits intoOct 6, 2026
Merged
Conversation
The bsrgemm_calc launch sites narrowed each nnz-per-row bucket size into a one-block(or wavefront-block)-per-row grid with no clamp against the device maximum grid dimension, and the block-per-row / wavefront device functions read their row straight from hipBlockIdx_x with no grid-stride. A bucket with more than maxGridSize[0] block rows (block-per-row truncates at 2,147,483,647; wavefront at ~8.6e9) would silently drop the tail of the bucket and produce a wrong result. This is only reachable with 64-bit index types (rocsparse_spgemm + rocsparse_indextype_i64), i.e. a P3 defect beyond current hardware. Add a small rocsparse::bsrgemm_clamp_grid_size helper (clamps against handle->properties.maxGridSize[0]) and use it for both the block-per-row and wavefront launches. Move the block index into a kernel-wrapper grid-stride loop and pass it to the device functions so a clamped grid still covers every row, mirroring the fixed-grid + grid-stride idiom already used by the bsrgemm setup kernels. The wavefront stride arithmetic is widened to the index type to avoid a 32-bit overflow at the clamped grid size. The two multipass kernel wrappers gain a bucket-size parameter for the loop bound. Part of epic AISPARSE-674 (extra routines large grid sizes). Co-authored-by: Cursor <cursoragent@cursor.com>
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
…d add grid-stride loops
The csrgemm nnz/symbolic/numeric-value launch sites narrowed each nnz-per-row
bucket size into a one-block(or wavefront-block)-per-row grid with no clamp
against the device maximum grid dimension, and the block-per-row / wavefront /
multipass kernels read their row straight from hipBlockIdx_x (or perm[hipBlockIdx_x
+ *offset]) with no grid-stride. A bucket with more than maxGridSize[0] rows would
silently drop the tail of the bucket and produce a wrong result. The block-per-row
path is the more reachable one: one block per row means the grid is the row count,
so it truncates at 2,147,483,647 rows; the wavefront path divides by
CSRGEMM_DIM / CSRGEMM_SUB first. As with bsrgemm this is only reachable with 64-bit
index types (rocsparse_spgemm + rocsparse_indextype_i64), i.e. a P3 defect beyond
current hardware.
Add a file-local rocsparse::csrgemm_clamp_grid_size helper to each of the four
extra/rocsparse_csrgemm_*_calc.cpp translation units (clamps against
handle->properties.maxGridSize[0]; deliberately not added to rocsparse_common.hpp
which AISPARSE-678 / AISPARSE-696 already touch) and use it at every bucket launch:
- rocsparse_csrgemm_calc.cpp: 4 sites (2 wavefront, block macro, multipass macro)
- rocsparse_csrgemm_nnz_calc.cpp: 11 sites (2 wavefront, 8 block-per-row, 1 multipass)
- rocsparse_csrgemm_symbolic_calc.cpp: 4 sites (2 wavefront, block macro, multipass macro)
- rocsparse_csrgemm_numeric_calc.cpp: 8 sites (2 wavefront, block macro, 2 launcher block/
multipass, 1 launcher multipass path, 2 group-10 multipass)
Move the block index into a grid-stride loop so a clamped grid still covers every
row. For the kernel-wrapper split (csrgemm_calc / symbolic / numeric) the loop lives
in the wrapper and passes the block index into the *_device function, mirroring
bsrgemm. For the block-per-row/multipass wrappers the loop strides the block index
(not the row) because the row comes from perm[block_id + *offset]; those wrappers and
their launches gain a bucket-size argument for the loop bound. The direct nnz kernels
in csrgemm_device.h (csrgemm_nnz_wf_per_row / _block_per_row / _block_per_row_multipass)
have no wrapper, so the grid-stride loop is added inside the kernel body (the
block-per-row ones gain a size parameter); the shared hash table is re-initialised
each iteration with an extra trailing __syncthreads so it can be reused across rows.
The wavefront stride arithmetic is widened to the index type to avoid a 32-bit
overflow at the clamped grid size.
The grid-stride loop for csrgemm_fill_wf_per_row_device rewrites the row-derivation
line at csrgemm_device.h:941 that AISPARSE-675 (#11817) widens to int64_t. The loop
init/stride are cast to the index type J, so the 64-bit loop subsumes that widening;
whoever merges #11824 and #11817 second should keep the loop form.
Part of epic AISPARSE-674 (extra routines large grid sizes).
Co-authored-by: Cursor <cursoragent@cursor.com>
kliegeois
changed the title
Codecov Report❌ Patch coverage is ❌ Your project check has failed because the head coverage (74.64%) 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 #11824 +/- ##
===========================================
- Coverage 70.61% 69.43% -1.18%
===========================================
Files 2939 2830 -109
Lines 483889 470781 -13108
Branches 69753 69446 -307
===========================================
- Hits 341695 326863 -14832
- Misses 118325 120183 +1858
+ Partials 23869 23735 -134
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
The csrgemm_intermediate_products launch in csrgemm_nnz_calc was the only remaining csrgemm grid whose extent scales with the row count but was not clamped to maxGridSize[0]; every other unclamped launch in the csrgemm calc files uses a constant grid (dim3(1) or dim3(CSRGEMM_DIM)) and needs nothing. Clamp the grid and wrap the kernel in a grid-stride loop over the (sub)wavefront rows, matching the form already used by the other wavefront-per-row kernels in this file. Note for merge ordering: AISPARSE-675 (PR #11817) widens the row derivation in this same kernel to int64_t. The grid-stride loop here subsumes that cast, so on conflict keep the loop form. JIRA ID: AISPARSE-677 Co-authored-by: Cursor <cursoragent@cursor.com>
…variable (AISPARSE-677)
The column/row-panel grid-stride loops added by this PR declared their
induction variable as J:
for(J block_offset = static_cast<J>(hipBlockIdx_x) * (BLOCKSIZE / WFSIZE);
block_offset < m;
block_offset += static_cast<J>(hipGridDim_x) * (BLOCKSIZE / WFSIZE))
With J = int32_t and a grid clamped to maxGridSize[0], the increment
hipGridDim_x * (BLOCKSIZE / WFSIZE) overflows, the induction variable wraps
negative and the loop re-enters, indexing out of bounds. That is the same
defect AISPARSE-675 fixes in the neighbouring kernels, so with both changes in
place the bug would have survived.
The loop variable and both products are now int64_t; the call site narrows back
to J explicitly, which is safe because the value is bounded by m/mb.
Five sites: bsrgemm_calc (x2), csrgemm_calc, csrgemm_numeric_calc,
csrgemm_symbolic_calc.
Co-authored-by: Cursor <cursoragent@cursor.com>
…mbolic calc kernels (AISPARSE-675/676/677)
jsandham
reviewed
Sep 25, 2026
jsandham
reviewed
Sep 25, 2026
jsandham
reviewed
Sep 25, 2026
Contributor
|
Can you make sure that we do not need any syncthronization when looping multiple times with the multi-pass kernels. I notice some of them write to shared memory |
…ash-table kernels The GRID_STRIDE variants of the block-per-row fill kernels reuse the same shared hash table for the next row. The device bodies end by reading that table (sorted-position scan into C) after their last __syncthreads(), so a fast wavefront could re-initialise the table for row i+1 while others were still reading row i. Add a __syncthreads() at the end of each grid-stride iteration, under if constexpr(GRID_STRIDE) so the single-pass instantiation is unchanged: - bsrgemm_fill_block_per_row_2x2 - bsrgemm_fill_block_per_row - csrgemm_fill_block_per_row - csrgemm_numeric_fill_block_per_row_kernel - csrgemm_symbolic_fill_block_per_row bsrgemm_fill_wf_per_row is launched with WF_SIZE=64, so on wave32 each per-(sub)wavefront table spans two hardware wavefronts and has the same hazard. Its barrier is skipped on the last row so it is never reached by some wavefronts after out-of-range ones have returned from the body. The multipass kernels already end every chunk iteration with a barrier after their last shared read, and the other wf-per-row kernels use WFSIZE <= 16 (table private to lanes of one hardware wavefront), so they need no change.
ntrost57
reviewed
Sep 29, 2026
ntrost57
reviewed
Sep 29, 2026
ntrost57
reviewed
Sep 29, 2026
kliegeois
commented
Sep 29, 2026
…roducts and csrgemm_nnz_wf_per_row (AISPARSE-677) Both loops used a 32-bit J row counter. Once the grid is clamped, the last row += gridDim.x * (BLOCKSIZE / WFSIZE) overflows int32 for m above ~1.07e9 (nnz_wf_per_row group 0) or ~1.61e9 (intermediate_products), and the negative row passes row < m and indexes out of bounds. Use an int64_t counter and stride.
…E-677) The grid-stride comments in the bsrgemm/csrgemm calc kernels said the grid is clamped to maxGridSize[0] (2^31 - 1), which reads as a bound only 64-bit indices can exceed. get_grid_size_x also caps grid.x at (2^32 - 1) / blockDim, 16,777,215 blocks at 256 threads, so 32-bit indices reach the strided path too. Also restore the truncated "covers all rows" endings.
… (AISPARSE-677) Each bsrgemm and csrgemm_numeric calc launch was written out twice, once per GRID_STRIDE value. Add rocsparse::dispatch_grid_stride_x, which computes the clamped grid once and calls a generic lambda with std::true_type or std::false_type, so each launch is written once. The csrgemm and csrgemm_symbolic sites already share one macro body. Both instantiations are kept. Kernel code, descriptors and resource usage are unchanged on gfx942 and gfx1201 (4560 instantiations compared).
…SPARSE-677) Shrink handle->properties.maxGridSize[0] to 1, 3, 7 and 64 blocks so the csrgemm nnz / compute / symbolic / numeric bucket kernels and every bsrgemm fill bucket take the clamped GRID_STRIDE path on small problems, and compare C exactly against a host reference and an unclamped run. Operands are built so every bucket is populated, with float/double and 32/64-bit indices, and bsrgemm block_dim 2, 4, 8 and 16 in both directions.
kliegeois
added a commit
to kliegeois/rocm-libraries
that referenced
this pull request
Sep 29, 2026
…673) The ELL dense-sample launch sized grid.x as (ell_width - 1) / 16 + 1 with 512 or 1024 threads and no clamp, and sddmm_ell_sample_kernel handled one group of 16 ell columns per block with no stride. Move the launch into sddmm_ell_sample_launch, which clamps grid.x with rocsparse::get_grid_size_x and, only when the clamp binds, runs a GRID_STRIDE instantiation that walks the column groups with an int64_t base. The straight-line kernel is unchanged; its resource usage is identical to before for all 28 instantiations on gfx942 and gfx1201. Add SddmmGridClamp.ell_dense_sample_wf32/wf64, which drive sddmm_ell_sample_launch with maxGridSize[0] at the device limit, 3 and 1 (the dense algorithm itself needs rocBLAS). Replace the test-local scoped_max_grid_x with ScopedMaxGridSizeX from the shared unit_test_grid_clamp.hpp, identical to the copy in ROCm#11824, ROCm#12307 and ROCm#12007.
kliegeois
added a commit
to kliegeois/rocm-libraries
that referenced
this pull request
Oct 1, 2026
…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.
jsandham
reviewed
Oct 1, 2026
amontoison
reviewed
Oct 2, 2026
amontoison
left a comment
amontoison
left a comment
Contributor
There was a problem hiding this comment.
The 64-bit induction-variable fix only landed in two of the fourteen grid-stride loops this PR adds. Still using J (12):
csrgemm_device.h:605 csrgemm_nnz_block_per_row
csrgemm_device.h:738 csrgemm_nnz_block_per_row_multipass
rocsparse_csrgemm_calc.cpp:162 csrgemm_fill_block_per_row
rocsparse_csrgemm_calc.cpp:242 csrgemm_fill_block_per_row_multipass
rocsparse_csrgemm_symbolic_calc.cpp:69 csrgemm_symbolic_fill_block_per_row_multipass
rocsparse_csrgemm_symbolic_calc.cpp:203 csrgemm_symbolic_fill_block_per_row
rocsparse_csrgemm_numeric_calc.cpp:160 csrgemm_numeric_fill_block_per_row_kernel
rocsparse_csrgemm_numeric_calc.cpp:244 csrgemm_numeric_fill_block_per_row_multipass_kernel
rocsparse_bsrgemm_calc.cpp:168 bsrgemm_fill_block_per_row_2x2
rocsparse_bsrgemm_calc.cpp:343 bsrgemm_fill_block_per_row
rocsparse_bsrgemm_calc.cpp:426 bsrgemm_block_per_row_atomic_multipass
rocsparse_bsrgemm_calc.cpp:505 bsrgemm_block_per_row_multipass
Already correct (7): csrgemm_device.h:96 and :496, plus rocsparse_csrgemm_calc.cpp:78, rocsparse_csrgemm_symbolic_calc.cpp:135, rocsparse_csrgemm_numeric_calc.cpp:77, rocsparse_bsrgemm_calc.cpp:84 and :252.
All twelve have the identical shape:
for(J block_id = hipBlockIdx_x; block_id < size; block_id += hipGridDim_x)
hipGridDim_x is unsigned int.
With J = int32_t the two operands have equal rank, so block_id converts to unsigned, the addition happens mod 2³², and the result converts back to a negative int32_t.
That negative value still satisfies block_id < size, which gives two failures at once: the loop never terminates, and every iteration from then on evaluates perm[block_id + *offset] roughly 8.6 GB below the array.
I did a small reproducer for myself to better understand it.
It needs no GPU.
// g++ -O2 -o wrap wrap.cpp && ./wrap
#include <cstdint>
#include <cstdio>
using J = int32_t; // rocsparse_int
int main()
{
// csrgemm_nnz_block_per_row, group 2 (CSRGEMM_DIM = 128).
// get_grid_size_x -> min(maxGridSize[0] = 2^31-1, (2^32-1)/128) = 33,554,431
const unsigned int hipGridDim_x = 33554431u;
const J size = 2113929218; // rows in the bucket
// The one block whose stride walk lands exactly on the last row, size - 1.
const unsigned int hipBlockIdx_x = (size - 1) % hipGridDim_x;
printf("size = %d, gridDim.x = %u, blockIdx.x = %u\n\n", size, hipGridDim_x, hipBlockIdx_x);
long long iters = 0;
for(J block_id = hipBlockIdx_x; block_id < size; block_id += hipGridDim_x)
{
++iters;
if(block_id > size - 2 * (J)hipGridDim_x || block_id < 0)
{
printf(" iter %lld: block_id = %12d -> perm[block_id] at byte offset %+lld\n",
iters, block_id, 4LL * block_id);
if(block_id < 0)
{
printf("\n (%d < %d) is still true -> the loop never exits,\n"
" and perm[] is read %.1f GB below the array.\n",
block_id, size, -4.0 * block_id / 1e9);
return 1;
}
}
}
printf("\n exited normally after %lld iterations\n", iters);
}
amontois@MKMAMONTOIS01:~/rocm-libraries$ g++ -O2 -o wrap wrap.cpp && ./wrap
size = 2113929218, gridDim.x = 33554431, blockIdx.x = 64
iter 63: block_id = 2080374786 -> perm[block_id] at byte offset +8321499144
iter 64: block_id = 2113929217 -> perm[block_id] at byte offset +8455716868
iter 65: block_id = -2147483648 -> perm[block_id] at byte offset -8589934592
(-2147483648 < 2113929218) is still true -> the loop never exits,
and perm[] is read 8.6 GB below the array.
Iteration 64 processes the last real row at block_id = size - 1; the next increment lands exactly on INT32_MAX + 1.
Switching the declaration to int64_t makes it exit cleanly after 64 iterations.
The following fix should do the job:
for(int64_t block_id = hipBlockIdx_x; block_id < size; block_id += hipGridDim_x)
{
..._device<...>(static_cast<J>(block_id), ...);
}
ntrost57
reviewed
Oct 2, 2026
…tride loops (AISPARSE-677) hipGridDim_x is unsigned int, so with a 32-bit J the stride add wraps mod 2^32 and yields a negative block_id that still satisfies block_id < size when a clamped grid walks a bucket near INT32_MAX rows: the loop never exits and perm[] is read far below the array. Widen the remaining twelve block-per-row loops in csrgemm nnz/fill/symbolic/numeric and bsrgemm calc to int64_t and narrow to J only at the device call.
…ls (AISPARSE-677) csrgemm_intermediate_products, csrgemm_nnz_wf_per_row, csrgemm_nnz_block_per_row and csrgemm_nnz_block_per_row_multipass always ran the grid-stride loop, and the two block kernels paid a barrier at the end of every row even when every row has its own block. Give them the same GRID_STRIDE flag as the fill kernels and launch them through rocsparse::dispatch_grid_stride_x, so the straight-line variant runs unless get_grid_size_x clamps. hipFuncSetAttribute is now set on the variant that is launched. The wf-per-row kernels start from the 32-bit row index again; it cannot overflow because the dispatch caps grid.x * BLOCKSIZE below 2^32. The barrier at the end of each multipass row is dropped: thread 0 alone reads and resets nnz, and the table and next_chunk accesses of a row are already fenced by the barriers that close its last chunk. On gfx942 and gfx1201 the straight-line variants match develop VGPR count and occupancy for every instantiation.
Contributor
Author
Fixed in b1c1d26, thanks! |
Contributor
Author
jsandham
approved these changes
Oct 5, 2026
amontoison
approved these changes
Oct 6, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
kliegeois
added a commit
that referenced
this pull request
Oct 8, 2026
…11817) JIRA ID : AISPARSE-675 ## Summary The csrgemm row-compute kernels in `projects/rocsparse/library/src/extra` derive the output row from the block index as `hipBlockIdx_x * BLOCKSIZE (+ hipThreadIdx_x) [/ WFSIZE + wid]`, and advance grid-stride loops with `row += hipGridDim_x * BLOCKSIZE`. Those products were evaluated in **32-bit unsigned** arithmetic and wrapped at `4,294,967,296` *regardless of how wide the index type `J` is*, because the multiply happens before the result is assigned to `J row`. The bucketed launchers size `grid.x` directly from the per-bucket row count (`dim3(group_size)` / `dim3((h_group_size[i]-1)/(DIM/SUB)+1)`), so both the block-index initialization and the grid-stride increment scale with the row count. Once a single bucket reaches a few hundred million rows (~2.1 GB row-pointer array, reachable on current large-memory hardware), the product overflows and the kernel silently misreads the row — reading/writing the wrong location and producing **wrong results with no error returned**. ## Fix (commit `933b4be`) Cast the leading `hipBlockIdx_x` / `hipGridDim_x` term to `int64_t` before the `* BLOCKSIZE` product in every row-index derivation and grid-stride increment in the csrgemm compute / symbolic / numeric kernels (`csrgemm_device.h`, `csrgemm_symbolic_device.h`, `csrgemm_numeric_device.h`). The widened arithmetic is off the inner loop and is a no-op for small values, so correct-sized problems are bit-for-bit unchanged. The element-copy kernels (`I idx = ...`) are the separate scaling copy paths handled in AISPARSE-679. ## Depends on #11824 (merge after it) At M = 552,870,912 the bucket-0 wavefront-per-row launches in csrgemm need 17,277,216 blocks × 256 threads = 4.42e9 work-items, which is above the per-axis dispatch limit (grid × block < 2^32), so HIP rejects the launch and `testing_csrgemm_extra_675` never reaches the row it probes. #11824 clamps those grids with `get_grid_size_x` and grid-strides the kernels with 64-bit counters, which makes rows past 536,870,912 reachable; this PR's index widening is still needed for those rows to be computed correctly. The two PRs conflict on the row-index line in `csrgemm_device.h`, `csrgemm_numeric_device.h` and `csrgemm_symbolic_device.h`; resolve by keeping #11824's `block_offset + wid` / grid-stride form. The merged pair was checked on gfx1201 with `maxGridSize[0]` forced down to 1, 3, 7 and 64: results were bit-identical to the unclamped run and matched the host reference, and `*csrgemm*` passed (1657 tests). ## Regression test Adds `testing_csrgemm_extra_675` (registered as `csrgemm_extra` in `test_csrgemm.yaml`). It runs `C = A * B` with `A = identity` and `B = 2 * identity`, sized so that a single bucket computes **more than `2^32 / WFSIZE = 536,870,912` rows** through the `fill_wf_per_row` kernel (BLOCKSIZE = 256, WFSIZE = 8). It then probes a row past the wrap boundary and asserts the diagonal result there has exactly one non-zero, with the expected column index and value `alpha * A * B = 2`. Before the fix that row is silently misread (wrong nnz / column / value); after the fix it is correct. The problem needs tens of GB of host + device memory, so the case is **nightly** and **memory-guarded** (`host_memory_gb: 14`, `device_memory_gb: 28`: each diagonal float/int32 CSR is 12·M bytes = 6.18 GiB, so the host holds A + B = 12.36 GiB and the device holds A + B + C plus the 8.27 GiB csrgemm buffer = 26.81 GiB). It auto-skips with `SKIPPING TEST: Insufficient memory` on cards that cannot host it (e.g. the 15 GB gfx1201), which is the established pattern for large-size regression tests in this suite. ## Testing - `librocsparse` + clients build cleanly (gfx1201, EXIT 0). - New test compiles and links (`testing_csrgemm_extra_675` present in `rocsparse-test`). - On the 15 GB gfx1201 dev card the case auto-skips via the memory guard, as expected. Made with [Cursor](https://cursor.com) --------- Co-authored-by: Cursor <cursoragent@cursor.com>
kliegeois
added a commit
that referenced
this pull request
Oct 8, 2026
…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>
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
JIRA ID: AISPARSE-677
Summary
The bsr/csr spgemm "extra" routines group output rows into nnz-per-row buckets and launch one block (or one wavefront-block) per row for each bucket. Every launch sized its grid directly from the bucket size, and the block-per-row / wavefront / multipass kernels read their row straight from
hipBlockIdx_x(orperm[hipBlockIdx_x + *offset]) with no grid-stride loop.The binding limit is the dispatch packet's 32-bit work-item count: HIP rejects any launch with
gridDim.x * blockDim.x >= 2^32(hipErrorInvalidConfiguration). That is 16,777,215 blocks at 256 threads, so a bucket with more than about 16.8M rows is enough, even with 32-bit indices. rocSPARSE discards the launch status in default builds, so the caller sees success while the rows of that bucket are never computed. (maxGridSize[0]= 2^31 - 1 is not the real bound.)Part of epic AISPARSE-674 (extra routines large grid sizes), ROCm SuperEpic ROCM-28956. This PR covers both halves: bsrgemm, and the csrgemm twin (nnz counting, symbolic and numeric phases).
Fix
rocsparse::get_grid_size_x(handle, count, block_size)fromrocsparse_grid.hpp(feat(rocsparse): per-axis grid-size helpers in their own header (AISPARSE-696) #12617), which clamps tomin(maxGridSize[0], (2^32 - 1) / block_size).nnzkernels), so a clamped grid still covers every row. For block-per-row / multipass paths the loop strides the block index, not the row, because the row comes fromperm[block_id + *offset]; those wrappers take the bucket size as the loop bound. All 19 grid-stride loops use a 64-bit counter and narrow toJonly at the device call, so the last increment cannot wrap for 32-bit index types.GRID_STRIDEinstantiation only; the loop bound is uniform across the block).csrgemm_nnz_block_per_row_multipassneeds none: its shared reads all finish before the barriers that close each chunk.bool GRID_STRIDEtemplate parameter (if constexpr). The newrocsparse::dispatch_grid_stride_x(handle, count, block_size, launch)inrocsparse_grid.hppcomputes the clamped grid once and calls the launch lambda withstd::true_typeorstd::false_type, so each launch site is written once and the looping instantiation only runs when the grid is actually clamped. This covers the 26 bsrgemm sites, the 12csrgemm_nnz_calcsites (csrgemm_intermediate_products,csrgemm_nnz_wf_per_row,csrgemm_nnz_block_per_rowand the multipass kernel) and the 7 csrgemm_numeric sites; csrgemm and csrgemm_symbolic already share one macro body per launch.rocsparse_bsrgemm_calc.cpprocsparse_csrgemm_calc.cpprocsparse_csrgemm_nnz_calc.cpprocsparse_csrgemm_symbolic_calc.cpprocsparse_csrgemm_numeric_calc.cppEvery remaining launch in these files uses a constant grid (
dim3(1)ordim3(CSRGEMM_DIM)).Performance
Grid-stride-only kernels cost measurable occupancy and time, which is why the fast path exists. Over 1,101
GRID_STRIDEtrue/false pairs per arch (-Rpass-analysis=kernel-resource-usage), the looping version has lower occupancy in 440 pairs on gfx942 (40%; e.g.bsrgemm_fill_block_per_row_2x28.0 -> 6.9 waves/SIMD) and in 4 pairs on gfx1201. On MI300X, grid-stride-only kernels made csrgemm 1.6% slower at the median (worst 6.4%) and bsrgemm 3.2% slower; with the fast path, csrgemm is at 1.019 and bsrgemm at 1.017 against the pre-fix head, and at parity with develop (csrgemm 1.0041, bsrgemm 0.9988).The cost of keeping both instantiations is device code size: the
trueversions are about 49% of the device code in these four files (4.07 MB of 8.26 MB on gfx942, 5.25 MB of 10.84 MB on gfx1201, uncompressed). Thedispatch_grid_stride_xrefactor left all 4,560 kernel instantiations (gfx942 + gfx1201) byte-identical and adds about 0.8 MB of host code.Merge order with AISPARSE-675 (#11817)
#11817's nightly test (
testing_csrgemm_extra_675, M = 552,870,912) needs this PR: its bucket-0 wavefront-per-row launches ask for about 4.42e9 work-items, which is over the dispatch limit without the clamp. Merge this PR first. The two PRs conflict on the row-index line incsrgemm_device.h,csrgemm_numeric_device.handcsrgemm_symbolic_device.h; resolve by keeping this PR'sblock_offset + wid/ grid-stride form.Testing (gfx1201)
rocsparse-test, quick and pre_checkin for*csrgemm*,*bsrgemm*and*spgemm*: 2,613 passed, 0 failed (about 47 oversized configs skipped by the suite's "Insufficient memory" guard on the 15 GB card).rocsparse-unit-test-devicegrid and gemm suites (GemmCalcGrids,internal_grid*,ExtraCsrgemm,ExtraBsrgemm,GenericSpgemm): 51/51 passed.clients/unittests/unit_test_internal_extra_gemm_calc_grids.cpp(suiteGemmCalcGrids, 21 tests, inrocsparse-unit-test-device) shrinkshandle->properties.maxGridSize[0]to 1, 3, 7 and 64 so every csrgemm nnz / compute / symbolic / numeric bucket (including the multipass fallback) and every bsrgemm fill bucket (block_dim 2/4/8/16, row and column direction) runs its clampedGRID_STRIDE=truevariant. C is compared exactly against a host reference and an unclamped run, for float/double and 32/64-bit indices; 210/210 passed over 10 repeats. HIP launch logs confirm every clamped launch uses grid.x of exactly 1, 3, 7 or 64, and breaking a grid-stride loop makes the tests fail.Co-authored-by: Cursor cursoragent@cursor.com