Repository navigation
fix(rocsparse): clamp extra-nightly launches to the 32-bit dispatch limit - #12546
Merged
kliegeois merged 13 commits intoOct 8, 2026
Merged
Conversation
…imit CDNA rejects grid.x * blockDim.x above UINT32_MAX, and axpby_extra used nnz > size. Clamp axpyi, scatter, ELL-N, and COO AoS atomic launches and make the extra test descriptors valid. Fixes ROCm#12545 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❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #12546 +/- ##
===========================================
- Coverage 71.93% 70.60% -1.33%
===========================================
Files 2863 2939 +76
Lines 472710 484516 +11806
Branches 69829 69781 -48
===========================================
+ Hits 340010 342058 +2048
- Misses 109117 118582 +9465
- Partials 23583 23876 +293
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
ntrost57
reviewed
Sep 29, 2026
ntrost57
reviewed
Sep 29, 2026
ntrost57
reviewed
Sep 29, 2026
# Conflicts: # projects/rocsparse/library/src/include/rocsparse_common.hpp
axpby scales y before axpyi. For size > 2^32 - 256 the scale_kernel<256> grid exceeded the 32-bit dispatch limit, and scale_device formed its index with a 32-bit multiply. Clamp the grid with get_grid_size_x and grid-stride in 64-bit. axpby_extra now uses beta = 3 so the scale kernel is launched, and checks the last element of y as well as the probe.
The LOOPS=2 fallback launched ceil(nnz / 512) blocks without a clamp, which is illegal again past nnz = 2^33 - 512, and the transposed coomvt_aos launch was never clamped. The atomic kernel now strides over LOOPS * BLOCKSIZE tiles and coomvt_aos grid-strides over nnz, so both launches use a get_grid_size_x-clamped grid and the LOOPS=2 variant is removed.
ellmvt launched ceil(m / 1024) blocks with no clamp, which exceeds the 32-bit dispatch limit for m > 2^32 - 1024. Clamp the grid with get_grid_size_x and grid-stride over rows in 64-bit, as ellmvn does.
roti_template clamped grid.x only to maxGridSize[0]. For a 64-bit index and nnz just past 2^32 that is 8388609 blocks of 512 threads, which exceeds the 2^32 - 1 work-item dispatch limit, so the launch is rejected with hipErrorInvalidConfiguration. Clamp with rocsparse::get_grid_size_x and select the kernel variant on the host: a straight-line kernel with a 32-bit element index when the grid is not clamped, and a grid-stride loop with a 64-bit index only when it is. The 64-bit index also removes the signed overflow that the I-typed grid-stride step could hit for 32-bit indices and nnz > 2^30.
rot_extra created a sparse vector of size 2 holding 2^32 + 512 entries. A sparse-vector descriptor requires nnz <= size, so rocsparse_rot returned rocsparse_status_invalid_size before reaching the kernel. Use size = nnz as axpby_extra does, and raise device_memory_gb to 70 for the 64 GiB of device data this now needs.
On gfx1151 (Windows, Strix Halo APU) hipMemGetInfo reports about 91 GiB free because it counts shared system memory, so the memory guard lets the 64-80 GiB extras run. hipMalloc then fails between 48 and 64 GiB: scatter_extra, spmv_ell_extra, spmv_coo_extra and spmv_coo_aos_extra threw in the nightly. Skip the six extras there, as bsrgemm and spgemm_bsr already do, rather than change the guard for every APU. Also correct the axpby_extra comment: 16 * nnz bytes is 64 GiB.
A failed hipMalloc leaves hipErrorOutOfMemory as the last error. The client turned it into std::bad_alloc without consuming it, so the next test failed in rocsparse_create_handle, whose hipGetLastError check "prior to hipLaunchKernelGGL" picked up the stale error. On gfx1151 the first quick test after each throwing extra failed this way. Consume the error when reporting the failure. Also drop the %' printf flag, which MSVC does not support and which garbled the size in the message.
…array when the grid is not clamped The grid-stride loops added for the clamped-grid fixes cost up to 12% on MI300A even when the grid is not clamped: the float scale_array loop is unrolled with 64-bit trip-count logic, and the COO AoS atomic kernel pays for the tile loop and its extra barrier. scale_kernel, coomvn_aos_atomic_loops and ellmvt_kernel take a GRID_STRIDE template parameter. The host launches the straight-line instantiation, one element, tile or row per thread or block as on develop, whenever get_grid_size_x returns the full block count, and the looping instantiation only when it clamps.
ntrost57
approved these changes
Oct 2, 2026
jsandham
approved these changes
Oct 5, 2026
kliegeois
enabled auto-merge (squash)
October 5, 2026 22:09
newling
approved these changes
Oct 7, 2026
kliegeois
added a commit
to kliegeois/rocm-libraries
that referenced
this pull request
Oct 8, 2026
…E-699/700) ROCm#12546 rewrites scale_device, scale_kernel and the scale_array launch with a clamped grid.x and a compile-time GRID_STRIDE fast path. This PR rewrote the same code with an always-on stride loop, so the two conflicted in rocsparse_common.cpp. Restore those three to develop so ROCm#12546 owns them. valset_2d, scale_2d, axpby and axpby_batched keep their 64-bit counts, clamped grids and stride loops; none of them used the removed scale code. The "why length is widened" comment moves from scale_kernel to axpby_kernel. The gemmi m * n fix is unchanged and still passes a 64-bit length to scale_array. scale_array_grid_stride stays: it passes against develop's unclamped launch and covers ROCm#12546's GRID_STRIDE path once both are merged. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
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.
Fixes #12545
Summary
grid.xwithrocsparse::get_grid_size_x, i.e. tomin(maxGridSize[0], (2^32-1)/blockDim.x)(the dispatch packet's grid size is a 32-bit work-item count; HIP rejects larger launches withhipErrorInvalidConfiguration), for axpyi / axpby (scale_array), scatter, ELL SpMV (non-transposed and transposed), COO AoS SpMV (atomic non-transposed, and both transposed launches), and roti /rocsparse_rot. Atnnz = 2^32 + 512the roti grid was 8,388,609 blocks of 512, just over the limit.scale_kernel,coomvn_aos_atomic_loops,ellmvt_kerneland the roti kernel take abool GRID_STRIDEtemplate parameter. The host launches the develop-shaped straight-line kernel wheneverget_grid_size_xreturns the full block count, and the looping kernel only when it actually clamps. A plain grid-stride conversion was measurably slower on MI300A even when the clamp did not engage (floatscale_array-12.4% at 1e8+, COO AoS atomic up to -6.7%); with the fast paths every measured config is back within noise of develop (see Performance).axpby_extraandrot_extrausesize = nnzso the sparse-vector descriptor is valid (nnz <= size).rot_extrahadsize = 2; itsdevice_memory_gbrises from 55 to 70 for the 64 GiB it needs.axpby_extranow usesbeta = 3soscale_arrayactually runs, and checks an element past 2^32. COO extram/nsizing is already on develop via fix(rocsparse): fix sizes of two SPMV COO tests #12345.hipMemGetInforeports ~91 GiB free because it counts shared system memory, so the memory guard let the 64-80 GiB extras run andhipMallocfailed between 48 and 64 GiB. The six 2^32-element extras now useskip_hardware: gfx1151, as bsrgemm and spgemm_bsr already do. The guard itself is unchanged (capping it at host memory would also change filtering on MI300A).hipErrorOutOfMemoryas the last HIP error, so the next test failed inrocsparse_create_handle(itshipGetLastErrorcheck). The allocator now consumes the error, and drops the%'printf flag that MSVC garbled.invalid_size(4) orhipErrorInvalidConfiguration/internal_error(6); gfx1151 withinvalid_size,thrown_exception(13) and the follow-on quick-testhipErrorOutOfMemoryfailures.Performance (MI300A, alola job 68007075)
10 interleaved repeats with rotated order, medians of GB/s, develop vs this PR:
scale_array(beta = 3): geomean 0.999 over 8 configs (range 0.994-1.006); float at 1e8 / 3e8 1.006 (was 0.876 / 0.878).Occupancy and LDS match develop for every instantiation on gfx942 and gfx1201; the fast-path COO AoS kernels have the same instruction count as develop.
Test plan
rocsparse-test --gtest_filter='nightly/axpby.level1/axpby_extra:nightly/scatter.level1/scatter_extra:nightly/spmv_ell.level2/spmv_ell_extra:nightly/spmv_coo.level2/spmv_coo_extra:nightly/spmv_coo_aos.level2/spmv_coo_aos_extra'— 5/5 PASSED.nightly/*extra*— all six 2^32-element extras filtered by the memory guard (rot at 70 GB); 2513 other nightly.extratests PASSED.*axpby*:*coomv*:*ellmv*:*spmv_coo*:*spmv_ell*— 2879 PASSED, also with every changed launch capped at 3 blocks to force the grid-stride path; quick + pre_checkinrot/roti/axpby/sctr/scatter— 1094 PASSED;rocsparse-unit-test-device— 504 PASSED, including the newdevice_level1_roti.grid_stride.--yaml: forcedrot_extrafails its 32 GiB allocation and the next test in the suite passes. Withskip_hardware: gfx1201the entry is not registered.rot_extraat full size on gfx942 (the nightly will exercise it).Made with Cursor