AISPARSE-703: clamp bsr2csr / csr2bsr / csr2bsr_nnz block-row grids and add grid-stride loops - #12007
Open
kliegeois wants to merge 1 commit into
Open
AISPARSE-703: clamp bsr2csr / csr2bsr / csr2bsr_nnz block-row grids and add grid-stride loops#12007kliegeois wants to merge 1 commit into
kliegeois wants to merge 1 commit into
Conversation
…e loops The BSR conversion routines launched one block (or wavefront) per block row with the grid taken directly from the 64-bit template parameter J and no clamp. Reached via the generic API (rocsparse_sparse_to_sparse -> spmat_bsr2csr -> gbsr2csr, and rocsparse_gbsr2csr) with J=int64_t, the grid overflows the 32-bit x-dimension limit at ~2.1e9 block rows. Clamp every block-per-row / wavefront-per-row grid against handle->properties.maxGridSize[0] and add a grid-stride loop over block rows to the corresponding kernels so a clamped grid still covers all block rows. The block_dim==1 launches already use divisor 1024/256 with a grid-stride kernel and are left unchanged. Part of ROCM-28974 (AISPARSE-703). Co-authored-by: Cursor <cursoragent@cursor.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project check has failed because the head coverage (74.61%) 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 #12007 +/- ##
========================================
Coverage 70.55% 70.55%
========================================
Files 2810 2810
Lines 462487 462501 +14
Branches 68090 68093 +3
========================================
+ Hits 326286 326298 +12
Misses 112661 112661
- Partials 23540 23542 +2
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
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
1 participant
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-703
The BSR conversion routines launch one block, or one wavefront, per block row,
with the count taken from the 64-bit template parameter
Jand no clamp againstthe device grid limit. Ten launch sites across three files, none of them touched
by any other ticket in the epic.
How the 64-bit path is reached
The legacy
rocsparse_Xbsr2csrentry point takesrocsparse_intand so selectsJ = int32_t; it cannot reach the defect. The generic API can:rocsparse_sparse_to_sparseat stage compute callsspmat_bsr2csr, which callsgbsr2csrpassingsource->rowsas anint64_t, and that instantiatesbsr2csr_core<T, I, int64_t>.rocsparse_gbsr2csris also directly callable andtakes an
int64_t mb.What changed
Each affected grid is clamped against
handle->properties.maxGridSize[0]and thecorresponding kernel gains a grid-stride loop over block rows. Most of the
launches sit inside
#definemacro bodies expanded once per block-dimensionbranch, so clamping the macro covers several call sites at once.
The kernels that use global scratch (
csr2bsr_65_inf_kernel,csr2bsr_nnz_65_inf_kernel) needed one extra change: theirtemp1/temp2/temp3segments were indexed by the block row, which goes out of bounds as soon as the
grid is clamped below the block-row count. They are now partitioned by the
physical block index instead, which is bounded by the launched grid by
construction. Each iteration re-initialises its segment, so sequential reuse
across grid-stride iterations is safe.
The two launches that already divide by 256 and by 1024 are left alone — they
cannot overflow the grid limit and the ticket says to leave them.
Loop bounds are block-uniform in every kernel, so the
__syncthreads()callsstay convergent. The wavefront-per-row kernel has no block-wide barrier at all,
only
__threadfence_block().Reviewing
Wrapping the kernel bodies in the grid-stride loop re-indents them, which
inflates the raw diff to +175/-132. Ignoring whitespace it is +108/-65. Please
review with
?w=1appended to the Files changed URL.
Testing
Built and tested on gfx1201 (Radeon RX 9070 XT): 1,639 conversion tests, 0
failing.
Coverage caveat:
grid.xon this hardware allows about 2.1 billion blocks, whichat divisor 1 corresponds to roughly 17 GB of row pointers alone, so the clamp
cannot be driven past one grid-stride iteration on any real device. The loop is a
correctness guard for the clamp rather than a path the test suite can enter.