Skip to content

AISPARSE-677: clamp bsrgemm_calc and csrgemm_calc bucket grids and add grid-stride loops #11824

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
kliegeois merged 17 commits into from
Oct 6, 2026
Merged

AISPARSE-677: clamp bsrgemm_calc and csrgemm_calc bucket grids and add grid-stride loops #11824

Show file tree
Hide file tree
Changes from 10 commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
248e7c0
AISPARSE-677: clamp bsrgemm_calc bucket grids and add grid-stride loops
kliegeois Sep 8, 2026
cf3fdbe
AISPARSE-677: clamp csrgemm_calc/nnz/symbolic/numeric bucket grids an…
kliegeois Sep 8, 2026
cc4a35d
Clamp and grid-stride csrgemm_intermediate_products
kliegeois Sep 10, 2026
4217bc4
fix(rocsparse): widen the bsrgemm/csrgemm calc grid-stride induction …
kliegeois Sep 16, 2026
9b55e9d
style(rocsparse): clang-format the csrgemm calc grid-stride loops (AI…
kliegeois Sep 22, 2026
b5b2ec7
perf(rocsparse): GRID_STRIDE fast path for csrgemm_calc fill kernels …
kliegeois Sep 22, 2026
6570ab5
perf(rocsparse): GRID_STRIDE fast path for bsrgemm/csrgemm numeric+sy…
kliegeois Sep 23, 2026
baf32b4
fix(rocsparse): barrier between grid-stride rows in bsrgemm/csrgemm h…
kliegeois Sep 25, 2026
d26b7d3
Merge branch 'develop' into users/kliegeois/aisparse-677-bsrgemm-calc…
kliegeois Sep 29, 2026
fe8bcea
refactor(rocsparse): use rocsparse_grid.hpp helpers (AISPARSE-677)
kliegeois Sep 29, 2026
8430a11
fix(rocsparse): 64-bit grid-stride counters in csrgemm_intermediate_p…
kliegeois Sep 29, 2026
35a1e8f
docs(rocsparse): name get_grid_size_x as the gemm calc clamp (AISPARS…
kliegeois Sep 29, 2026
bfd4cca
refactor(rocsparse): dispatch GRID_STRIDE through one launch per site…
kliegeois Sep 29, 2026
deb36c4
style(rocsparse): use static_cast instead of C-style casts (AISPARSE-…
kliegeois Sep 29, 2026
76dfa7a
test(rocsparse): forced-clamp unit tests for csrgemm/bsrgemm calc (AI…
kliegeois Sep 29, 2026
b1c1d26
fix(rocsparse): 64-bit induction variable in the block-per-row grid-s…
kliegeois Oct 2, 2026
451f29c
perf(rocsparse): GRID_STRIDE fast path for the csrgemm nnz calc kerne…
kliegeois Oct 2, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
48 changes: 30 additions & 18 deletions projects/rocsparse/library/src/extra/bsrgemm_device.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -230,7 +230,8 @@ namespace rocsparse
typename J,
typename T>
ROCSPARSE_DEVICE_ILF void
bsrgemm_fill_wf_per_row_2x2_device(rocsparse_direction dir,
bsrgemm_fill_wf_per_row_2x2_device(J block_offset,
rocsparse_direction dir,
J mb,
J nkb,
const J* __restrict__ offset,
Expand Down Expand Up @@ -266,8 +267,9 @@ namespace rocsparse
// Wavefront id
int wid = hipThreadIdx_x / WFSIZE;

// Each (sub)wavefront processes a row
J row = hipBlockIdx_x * (BLOCKSIZE / WFSIZE) + wid;
// Each (sub)wavefront processes a row (block_offset supplied by the grid-stride
// loop in the kernel wrapper so a grid clamped to maxGridSize[0] still covers all rows)
J row = block_offset + wid;

// Hash table in shared memory
__shared__ J stable[(BLOCKSIZE / WFSIZE) * HASHSIZE];
Expand Down Expand Up @@ -503,7 +505,8 @@ namespace rocsparse
typename J,
typename T>
ROCSPARSE_DEVICE_ILF void
bsrgemm_fill_block_per_row_2x2_device(rocsparse_direction dir,
bsrgemm_fill_block_per_row_2x2_device(J block_id,
rocsparse_direction dir,
J mb,
J nkb,
const J* __restrict__ offset,
Expand Down Expand Up @@ -557,8 +560,9 @@ namespace rocsparse

__syncthreads();

// Each block processes a row (apply permutation)
J row = perm ? perm[hipBlockIdx_x + *offset] : hipBlockIdx_x;
// Each block processes a row (block_id supplied by the grid-stride loop in the
// kernel wrapper so a grid clamped to maxGridSize[0] still covers all rows)
J row = perm ? perm[block_id + *offset] : block_id;

// alpha * A * B part
if(mul == true)
Expand Down Expand Up @@ -762,7 +766,8 @@ namespace rocsparse
typename I,
typename J,
typename T>
ROCSPARSE_DEVICE_ILF void bsrgemm_fill_wf_per_row_device(rocsparse_direction dir,
ROCSPARSE_DEVICE_ILF void bsrgemm_fill_wf_per_row_device(J block_offset,
rocsparse_direction dir,
J mb,
J nkb,
J block_dim,
Expand Down Expand Up @@ -804,8 +809,9 @@ namespace rocsparse
int c = slid & (BLOCKDIM - 1);
int r = slid / BLOCKDIM;

// Each (sub)wavefront processes a row
J row = hipBlockIdx_x * (BLOCKSIZE / WFSIZE) + wid;
// Each (sub)wavefront processes a row (block_offset supplied by the grid-stride
// loop in the kernel wrapper so a grid clamped to maxGridSize[0] still covers all rows)
J row = block_offset + wid;

// Hash table in shared memory
__shared__ J stable[(BLOCKSIZE / WFSIZE) * HASHSIZE];
Expand Down Expand Up @@ -993,7 +999,8 @@ namespace rocsparse
typename I,
typename J,
typename T>
ROCSPARSE_DEVICE_ILF void bsrgemm_fill_block_per_row_device(rocsparse_direction dir,
ROCSPARSE_DEVICE_ILF void bsrgemm_fill_block_per_row_device(J block_id,
rocsparse_direction dir,
J mb,
J nkb,
J block_dim,
Expand Down Expand Up @@ -1030,8 +1037,9 @@ namespace rocsparse
int c = lid & (BLOCKDIM - 1);
int r = lid / BLOCKDIM;

// Each block processes a row (apply permutation)
J row = perm ? perm[hipBlockIdx_x + *offset] : hipBlockIdx_x;
// Each block processes a row (block_id supplied by the grid-stride loop in the
// kernel wrapper so a grid clamped to maxGridSize[0] still covers all rows)
J row = perm ? perm[block_id + *offset] : block_id;

// Hash table in shared memory
__shared__ J table[HASHSIZE];
Expand Down Expand Up @@ -1207,7 +1215,8 @@ namespace rocsparse
typename J,
typename T>
ROCSPARSE_DEVICE_ILF void
bsrgemm_block_per_row_atomic_multipass_device(rocsparse_direction dir,
bsrgemm_block_per_row_atomic_multipass_device(J block_id,
rocsparse_direction dir,
J nb,
J block_dim,
const J* __restrict__ offset_,
Expand Down Expand Up @@ -1244,8 +1253,9 @@ namespace rocsparse
int c = lid & (BLOCKDIM - 1);
int r = lid / BLOCKDIM;

// Each block processes a row (apply permutation)
J row = perm ? perm[hipBlockIdx_x + *offset_] : hipBlockIdx_x;
// Each block processes a row (block_id supplied by the grid-stride loop in the
// kernel wrapper so a grid clamped to maxGridSize[0] still covers all rows)
J row = perm ? perm[block_id + *offset_] : block_id;

// Row entry marker and value accumulator
__shared__ int table[CHUNKSIZE];
Expand Down Expand Up @@ -1538,7 +1548,8 @@ namespace rocsparse
typename J,
typename T>
ROCSPARSE_DEVICE_ILF void
bsrgemm_block_per_row_multipass_device(rocsparse_direction dir,
bsrgemm_block_per_row_multipass_device(J block_id,
rocsparse_direction dir,
J nb,
J block_dim,
const J* __restrict__ offset_,
Expand Down Expand Up @@ -1572,8 +1583,9 @@ namespace rocsparse
// Wavefront id
int wid = hipThreadIdx_x / (BLOCKSIZE / BLOCKDIM);

// Each block processes a row (apply permutation)
J row = perm ? perm[hipBlockIdx_x + *offset_] : hipBlockIdx_x;
// Each block processes a row (block_id supplied by the grid-stride loop in the
// kernel wrapper so a grid clamped to maxGridSize[0] still covers all rows)
J row = perm ? perm[block_id + *offset_] : block_id;

// Row entry marker and value accumulator
__shared__ bool table[CHUNKSIZE];
Expand Down
Loading
Loading