From cd89e920a670f41778b4eadaaded2ea8632a8a51 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Wed, 19 Aug 2026 16:53:12 -0600 Subject: [PATCH 1/6] Fix sctr large-nnz index overflow (AISPARSE-652) The sctr kernel computed the element index as `hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x`, which is evaluated in 32-bit unsigned arithmetic and wraps at 2^32 even when the index type I is int64_t. Because the kernel mapped one thread to one element with no grid-stride loop, elements past the wrap point were silently never scattered. Cast the block index to I before the multiply and add a grid-stride loop over nnz so the kernel covers the full range regardless of grid size. Clamp the launch grid.x against handle->properties.maxGridSize[0] so an overflowing nnz cannot exceed the device grid limit. Fixes AC3/AC4 of ROCM-28948 for rocsparse_Xsctr and rocsparse_scatter. Co-authored-by: Cursor --- .../library/src/level1/rocsparse_sctr_impl.cpp | 9 ++++++++- projects/rocsparse/library/src/level1/sctr_device.h | 11 ++++++----- 2 files changed, 14 insertions(+), 6 deletions(-) diff --git a/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp index 0b1b24442bab..f30b82ddbff1 100644 --- a/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp +++ b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp @@ -68,7 +68,14 @@ rocsparse_status rocsparse::sctr_template(rocsparse_handle handle, hipStream_t stream = handle->stream; #define SCTR_DIM 512 - dim3 sctr_blocks((nnz - 1) / SCTR_DIM + 1); + // Clamp the grid dimension against the device limit; the kernel uses a + // grid-stride loop so any grid size covers the full range of nnz. + int64_t num_blocks = (nnz - 1) / SCTR_DIM + 1; + if(num_blocks > handle->properties.maxGridSize[0]) + { + num_blocks = handle->properties.maxGridSize[0]; + } + dim3 sctr_blocks(num_blocks); dim3 sctr_threads(SCTR_DIM); RETURN_IF_HIPLAUNCHKERNELGGL_ERROR((rocsparse::sctr_kernel), diff --git a/projects/rocsparse/library/src/level1/sctr_device.h b/projects/rocsparse/library/src/level1/sctr_device.h index f8fc8a93dea7..31309f530116 100644 --- a/projects/rocsparse/library/src/level1/sctr_device.h +++ b/projects/rocsparse/library/src/level1/sctr_device.h @@ -32,13 +32,14 @@ namespace rocsparse ROCSPARSE_KERNEL(BLOCKSIZE) void sctr_kernel(I nnz, const T* x_val, const I* x_ind, T* y, rocsparse_index_base idx_base) { - I idx = hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x; + // Cast to the (possibly 64-bit) index type I before the multiply so the + // thread index does not wrap in 32-bit arithmetic for large nnz. + const I stride = static_cast(hipGridDim_x) * BLOCKSIZE; - if(idx >= nnz) + for(I idx = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; idx < nnz; + idx += stride) { - return; + y[x_ind[idx] - idx_base] = x_val[idx]; } - - y[x_ind[idx] - idx_base] = x_val[idx]; } } From 05ec529defa4f34fa5e0b2fef2181b4b434cd881 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Fri, 21 Aug 2026 13:59:06 -0600 Subject: [PATCH 2/6] test(rocsparse): regression + changelog for sctr large-nnz overflow (AISPARSE-652) Add the sctr_extra regression test driving the 64-bit-index path of rocsparse_scatter with nnz = 2^32 + 512 and probing an element past the 2^32 boundary; before the fix the wrapped block index leaves it unscattered. Guarded to the nightly category with a device_memory_gb budget so it is skipped on smaller GPUs. Document the fix in CHANGELOG. Co-authored-by: Cursor --- projects/rocsparse/CHANGELOG.md | 1 + .../clients/testings/testing_sctr.cpp | 64 ++++++++++++++++++- .../rocsparse/clients/tests/test_sctr.yaml | 6 ++ 3 files changed, 70 insertions(+), 1 deletion(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index f61e7eeb2dc0..95f1f9c4aa70 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -15,6 +15,7 @@ Documentation for rocSPARSE is available at * Optimized architecture-aware launch configurations for RDNA (wave32) and CDNA (wave64) GPUs, improving performance and performance portability for several sparse level 2 and level 3 routines without algorithmic or numerical changes. Affected routines include `rocsparse_spmv` for the CSR adaptive, nnz-split, and LRB algorithms, the COO (SoA and AoS) formats, and the ELL format (`rocsparse_Xellmv`); `rocsparse_Xbsrmv`; `rocsparse_Xbsrxmv`; `rocsparse_Xgemvi`; `rocsparse_Xgemmi`; and `rocsparse_spmm` with the blocked-ELL format. ### Resolved issues +* Fixed `rocsparse_Xsctr` and the generic `rocsparse_scatter` routine (which delegates to it), whose device kernel computed the element index with 32-bit arithmetic. For sparse vectors with `nnz` of `2^32` (~4.3 billion) or more entries, reachable with 64-bit indices, the block-index multiply overflowed, leaving elements past the overflow point unscattered. The index is now computed in the 64-bit index type, and the kernel uses a grid-stride loop with the launch grid clamped to the device limit. * Fixed an integer overflow in `rocsparse_prune_dense2csr_by_percentage` and `rocsparse_prune_csr2csr_by_percentage`, which computed the matrix element count in 32-bit arithmetic. For matrices with more than `INT32_MAX` (~2.1 billion) elements the count overflowed to a negative value, resulting in out-of-bounds pointer construction and an invalid kernel launch grid. The element count is now computed in 64-bit arithmetic. * Fixed `rocsparse_spmm` with the segmented COO, atomic COO, segmented-atomic COO, and row-split CSR algorithms, which failed with `hipErrorInvalidConfiguration` for batch counts exceeding 65535 because the batch dimension of the kernel launch grid exceeded the maximum supported grid dimension. * Fixed an issue with `rocsparse_spmm` when using the nnz-split algorithm with the CSR or CSC format. The operation produced incorrect results because the segmented-block-reduction helper had shared-memory pointer parameters marked `__restrict__`, while threads in the block must read values written by other threads. The `__restrict__` attribute has now been removed. diff --git a/projects/rocsparse/clients/testings/testing_sctr.cpp b/projects/rocsparse/clients/testings/testing_sctr.cpp index d7a821d38de4..245245f000f2 100644 --- a/projects/rocsparse/clients/testings/testing_sctr.cpp +++ b/projects/rocsparse/clients/testings/testing_sctr.cpp @@ -128,4 +128,66 @@ INSTANTIATE(float); INSTANTIATE(double); INSTANTIATE(rocsparse_float_complex); INSTANTIATE(rocsparse_double_complex); -void testing_sctr_extra(const Arguments& arg) {} +void testing_sctr_extra(const Arguments& arg) +{ + // Regression test for AISPARSE-652. + // + // Before the fix, sctr_device computed the element index as + // hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x + // entirely in 32-bit unsigned arithmetic. Once nnz reaches 2^32 the + // block-index multiply wraps around, so the tail of the sparse vector past + // the wrap point is never scattered into y. The fix casts the block index + // to the (64-bit) index type before the multiply and iterates with a + // grid-stride loop. + // + // This drives the 64-bit-index path of rocsparse_scatter (which dispatches + // to sctr_template) with nnz just past the 2^32 boundary and checks that an + // element beyond that boundary is actually scattered. To stay within a + // single device allocation everything is initialized on the device and a + // single element is probed. + using I = int64_t; + using T = float; + + static constexpr int64_t two_pow_32 = static_cast(1) << 32; + + // nnz just beyond 2^32 so at least one block has a block index whose + // (blockIdx * BLOCKSIZE) product overflows 32-bit arithmetic. + const I nnz = two_pow_32 + 512; + const I size = 2; + + const rocsparse_index_base base = rocsparse_index_base_zero; + + rocsparse_local_handle handle(arg); + + device_vector dx_ind(nnz); + device_vector dx_val(nnz); + device_vector dy(size); + + // Filler elements all reference dense entry 0 and scatter value 0; y is 0. + CHECK_HIP_ERROR(hipMemset(dx_ind, 0, sizeof(I) * nnz)); + CHECK_HIP_ERROR(hipMemset(dx_val, 0, sizeof(T) * nnz)); + CHECK_HIP_ERROR(hipMemset(dy, 0, sizeof(T) * size)); + + // The probe lives past the 2^32 boundary. It scatters the only non-zero + // value into dense entry 1, which no filler element writes, so its result + // is isolated from the filler scatters to dense entry 0. + const I probe_idx = two_pow_32 + 5; + const I probe_ind = 1; + const T x_in = static_cast(3); + CHECK_HIP_ERROR(hipMemcpy( + static_cast(dx_ind) + probe_idx, &probe_ind, sizeof(I), hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(static_cast(dx_val) + probe_idx, &x_in, sizeof(T), hipMemcpyHostToDevice)); + + rocsparse_local_spvec x(size, nnz, dx_ind, dx_val, get_indextype(), base, get_datatype()); + rocsparse_local_dnvec y(size, dy, get_datatype()); + + CHECK_ROCSPARSE_ERROR(testing::rocsparse_scatter(handle, x, y)); + + // y[1] = x_in = 3. Before the fix the wrapped block index leaves the probe + // element unscattered, so y[1] stays 0. + T y_out = static_cast(0); + CHECK_HIP_ERROR(hipMemcpy(&y_out, static_cast(dy) + 1, sizeof(T), hipMemcpyDeviceToHost)); + + unit_check_scalar(x_in, y_out); +} diff --git a/projects/rocsparse/clients/tests/test_sctr.yaml b/projects/rocsparse/clients/tests/test_sctr.yaml index c16914414517..2637c0673480 100644 --- a/projects/rocsparse/clients/tests/test_sctr.yaml +++ b/projects/rocsparse/clients/tests/test_sctr.yaml @@ -63,3 +63,9 @@ Tests: nnz: [5, 10, 500] baseA: [rocsparse_index_base_zero, rocsparse_index_base_one] graph_test: true + +- name: sctr_extra + category: nightly + function: sctr_extra + host_memory_gb: 1 + device_memory_gb: 55 From ed7e6cb1e42343e9c565a7eab1758167f56f00af Mon Sep 17 00:00:00 2001 From: kliegeois Date: Wed, 2 Sep 2026 09:38:29 -0600 Subject: [PATCH 3/6] refactor(rocsparse): address review on sctr large-nnz fix (AISPARSE-652) - use rocsparse::min to clamp the grid size instead of an if-branch - name the grid-stride start index (gid) for readability - move the large-nnz regression test from sctr_extra to scatter_extra (it exercises the generic rocsparse_scatter API) Co-authored-by: Cursor --- .../clients/testings/testing_scatter.cpp | 64 ++++++++++++++++++- .../clients/testings/testing_sctr.cpp | 64 +------------------ .../rocsparse/clients/tests/test_scatter.yaml | 6 ++ .../rocsparse/clients/tests/test_sctr.yaml | 6 -- .../src/level1/rocsparse_sctr_impl.cpp | 10 ++- .../library/src/level1/sctr_device.h | 4 +- 6 files changed, 76 insertions(+), 78 deletions(-) diff --git a/projects/rocsparse/clients/testings/testing_scatter.cpp b/projects/rocsparse/clients/testings/testing_scatter.cpp index 08c2d5f7659a..37c441761a62 100644 --- a/projects/rocsparse/clients/testings/testing_scatter.cpp +++ b/projects/rocsparse/clients/testings/testing_scatter.cpp @@ -130,4 +130,66 @@ INSTANTIATE(int64_t, double); INSTANTIATE(int64_t, rocsparse_float_complex); INSTANTIATE(int64_t, rocsparse_double_complex); -void testing_scatter_extra(const Arguments& arg) {} +void testing_scatter_extra(const Arguments& arg) +{ + // Regression test for AISPARSE-652. + // + // Before the fix, sctr_device computed the element index as + // hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x + // entirely in 32-bit unsigned arithmetic. Once nnz reaches 2^32 the + // block-index multiply wraps around, so the tail of the sparse vector past + // the wrap point is never scattered into y. The fix casts the block index + // to the (64-bit) index type before the multiply and iterates with a + // grid-stride loop. + // + // This drives the 64-bit-index path of rocsparse_scatter (which dispatches + // to sctr_template) with nnz just past the 2^32 boundary and checks that an + // element beyond that boundary is actually scattered. To stay within a + // single device allocation everything is initialized on the device and a + // single element is probed. + using I = int64_t; + using T = float; + + static constexpr int64_t two_pow_32 = static_cast(1) << 32; + + // nnz just beyond 2^32 so at least one block has a block index whose + // (blockIdx * BLOCKSIZE) product overflows 32-bit arithmetic. + const I nnz = two_pow_32 + 512; + const I size = 2; + + const rocsparse_index_base base = rocsparse_index_base_zero; + + rocsparse_local_handle handle(arg); + + device_vector dx_ind(nnz); + device_vector dx_val(nnz); + device_vector dy(size); + + // Filler elements all reference dense entry 0 and scatter value 0; y is 0. + CHECK_HIP_ERROR(hipMemset(dx_ind, 0, sizeof(I) * nnz)); + CHECK_HIP_ERROR(hipMemset(dx_val, 0, sizeof(T) * nnz)); + CHECK_HIP_ERROR(hipMemset(dy, 0, sizeof(T) * size)); + + // The probe lives past the 2^32 boundary. It scatters the only non-zero + // value into dense entry 1, which no filler element writes, so its result + // is isolated from the filler scatters to dense entry 0. + const I probe_idx = two_pow_32 + 5; + const I probe_ind = 1; + const T x_in = static_cast(3); + CHECK_HIP_ERROR(hipMemcpy( + static_cast(dx_ind) + probe_idx, &probe_ind, sizeof(I), hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(static_cast(dx_val) + probe_idx, &x_in, sizeof(T), hipMemcpyHostToDevice)); + + rocsparse_local_spvec x(size, nnz, dx_ind, dx_val, get_indextype(), base, get_datatype()); + rocsparse_local_dnvec y(size, dy, get_datatype()); + + CHECK_ROCSPARSE_ERROR(testing::rocsparse_scatter(handle, x, y)); + + // y[1] = x_in = 3. Before the fix the wrapped block index leaves the probe + // element unscattered, so y[1] stays 0. + T y_out = static_cast(0); + CHECK_HIP_ERROR(hipMemcpy(&y_out, static_cast(dy) + 1, sizeof(T), hipMemcpyDeviceToHost)); + + unit_check_scalar(x_in, y_out); +} diff --git a/projects/rocsparse/clients/testings/testing_sctr.cpp b/projects/rocsparse/clients/testings/testing_sctr.cpp index 245245f000f2..d7a821d38de4 100644 --- a/projects/rocsparse/clients/testings/testing_sctr.cpp +++ b/projects/rocsparse/clients/testings/testing_sctr.cpp @@ -128,66 +128,4 @@ INSTANTIATE(float); INSTANTIATE(double); INSTANTIATE(rocsparse_float_complex); INSTANTIATE(rocsparse_double_complex); -void testing_sctr_extra(const Arguments& arg) -{ - // Regression test for AISPARSE-652. - // - // Before the fix, sctr_device computed the element index as - // hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x - // entirely in 32-bit unsigned arithmetic. Once nnz reaches 2^32 the - // block-index multiply wraps around, so the tail of the sparse vector past - // the wrap point is never scattered into y. The fix casts the block index - // to the (64-bit) index type before the multiply and iterates with a - // grid-stride loop. - // - // This drives the 64-bit-index path of rocsparse_scatter (which dispatches - // to sctr_template) with nnz just past the 2^32 boundary and checks that an - // element beyond that boundary is actually scattered. To stay within a - // single device allocation everything is initialized on the device and a - // single element is probed. - using I = int64_t; - using T = float; - - static constexpr int64_t two_pow_32 = static_cast(1) << 32; - - // nnz just beyond 2^32 so at least one block has a block index whose - // (blockIdx * BLOCKSIZE) product overflows 32-bit arithmetic. - const I nnz = two_pow_32 + 512; - const I size = 2; - - const rocsparse_index_base base = rocsparse_index_base_zero; - - rocsparse_local_handle handle(arg); - - device_vector dx_ind(nnz); - device_vector dx_val(nnz); - device_vector dy(size); - - // Filler elements all reference dense entry 0 and scatter value 0; y is 0. - CHECK_HIP_ERROR(hipMemset(dx_ind, 0, sizeof(I) * nnz)); - CHECK_HIP_ERROR(hipMemset(dx_val, 0, sizeof(T) * nnz)); - CHECK_HIP_ERROR(hipMemset(dy, 0, sizeof(T) * size)); - - // The probe lives past the 2^32 boundary. It scatters the only non-zero - // value into dense entry 1, which no filler element writes, so its result - // is isolated from the filler scatters to dense entry 0. - const I probe_idx = two_pow_32 + 5; - const I probe_ind = 1; - const T x_in = static_cast(3); - CHECK_HIP_ERROR(hipMemcpy( - static_cast(dx_ind) + probe_idx, &probe_ind, sizeof(I), hipMemcpyHostToDevice)); - CHECK_HIP_ERROR( - hipMemcpy(static_cast(dx_val) + probe_idx, &x_in, sizeof(T), hipMemcpyHostToDevice)); - - rocsparse_local_spvec x(size, nnz, dx_ind, dx_val, get_indextype(), base, get_datatype()); - rocsparse_local_dnvec y(size, dy, get_datatype()); - - CHECK_ROCSPARSE_ERROR(testing::rocsparse_scatter(handle, x, y)); - - // y[1] = x_in = 3. Before the fix the wrapped block index leaves the probe - // element unscattered, so y[1] stays 0. - T y_out = static_cast(0); - CHECK_HIP_ERROR(hipMemcpy(&y_out, static_cast(dy) + 1, sizeof(T), hipMemcpyDeviceToHost)); - - unit_check_scalar(x_in, y_out); -} +void testing_sctr_extra(const Arguments& arg) {} diff --git a/projects/rocsparse/clients/tests/test_scatter.yaml b/projects/rocsparse/clients/tests/test_scatter.yaml index 4b2dbc9abf5f..d5dd8701de81 100644 --- a/projects/rocsparse/clients/tests/test_scatter.yaml +++ b/projects/rocsparse/clients/tests/test_scatter.yaml @@ -149,3 +149,9 @@ Tests: nnz: [5, 10, 500] baseA: [rocsparse_index_base_zero, rocsparse_index_base_one] graph_test: true + +- name: scatter_extra + category: nightly + function: scatter_extra + host_memory_gb: 1 + device_memory_gb: 55 diff --git a/projects/rocsparse/clients/tests/test_sctr.yaml b/projects/rocsparse/clients/tests/test_sctr.yaml index 2637c0673480..c16914414517 100644 --- a/projects/rocsparse/clients/tests/test_sctr.yaml +++ b/projects/rocsparse/clients/tests/test_sctr.yaml @@ -63,9 +63,3 @@ Tests: nnz: [5, 10, 500] baseA: [rocsparse_index_base_zero, rocsparse_index_base_one] graph_test: true - -- name: sctr_extra - category: nightly - function: sctr_extra - host_memory_gb: 1 - device_memory_gb: 55 diff --git a/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp index f30b82ddbff1..dce0d3d76fa4 100644 --- a/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp +++ b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp @@ -70,12 +70,10 @@ rocsparse_status rocsparse::sctr_template(rocsparse_handle handle, #define SCTR_DIM 512 // Clamp the grid dimension against the device limit; the kernel uses a // grid-stride loop so any grid size covers the full range of nnz. - int64_t num_blocks = (nnz - 1) / SCTR_DIM + 1; - if(num_blocks > handle->properties.maxGridSize[0]) - { - num_blocks = handle->properties.maxGridSize[0]; - } - dim3 sctr_blocks(num_blocks); + const int64_t num_blocks + = rocsparse::min((static_cast(nnz) - 1) / SCTR_DIM + 1, + static_cast(handle->properties.maxGridSize[0])); + dim3 sctr_blocks(static_cast(num_blocks)); dim3 sctr_threads(SCTR_DIM); RETURN_IF_HIPLAUNCHKERNELGGL_ERROR((rocsparse::sctr_kernel), diff --git a/projects/rocsparse/library/src/level1/sctr_device.h b/projects/rocsparse/library/src/level1/sctr_device.h index 31309f530116..f5cf3e95653c 100644 --- a/projects/rocsparse/library/src/level1/sctr_device.h +++ b/projects/rocsparse/library/src/level1/sctr_device.h @@ -35,9 +35,9 @@ namespace rocsparse // Cast to the (possibly 64-bit) index type I before the multiply so the // thread index does not wrap in 32-bit arithmetic for large nnz. const I stride = static_cast(hipGridDim_x) * BLOCKSIZE; + const I gid = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; - for(I idx = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; idx < nnz; - idx += stride) + for(I idx = gid; idx < nnz; idx += stride) { y[x_ind[idx] - idx_base] = x_val[idx]; } From 7765e1db26ec8083fd95691be5dbb5ba15e66891 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 15 Sep 2026 11:20:34 -0600 Subject: [PATCH 4/6] docs(rocsparse): apply reviewer CHANGELOG wording for sctr overflow fix (AISPARSE-652) Applies the dso-amd review suggestion on projects/rocsparse/CHANGELOG.md verbatim: reword the 'Resolved issues' entry for rocsparse_Xsctr / rocsparse_scatter to the shorter, standardized overflow description. Co-authored-by: Cursor --- projects/rocsparse/CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index 95f1f9c4aa70..b8a307b04bd4 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -15,7 +15,7 @@ Documentation for rocSPARSE is available at * Optimized architecture-aware launch configurations for RDNA (wave32) and CDNA (wave64) GPUs, improving performance and performance portability for several sparse level 2 and level 3 routines without algorithmic or numerical changes. Affected routines include `rocsparse_spmv` for the CSR adaptive, nnz-split, and LRB algorithms, the COO (SoA and AoS) formats, and the ELL format (`rocsparse_Xellmv`); `rocsparse_Xbsrmv`; `rocsparse_Xbsrxmv`; `rocsparse_Xgemvi`; `rocsparse_Xgemmi`; and `rocsparse_spmm` with the blocked-ELL format. ### Resolved issues -* Fixed `rocsparse_Xsctr` and the generic `rocsparse_scatter` routine (which delegates to it), whose device kernel computed the element index with 32-bit arithmetic. For sparse vectors with `nnz` of `2^32` (~4.3 billion) or more entries, reachable with 64-bit indices, the block-index multiply overflowed, leaving elements past the overflow point unscattered. The index is now computed in the 64-bit index type, and the kernel uses a grid-stride loop with the launch grid clamped to the device limit. +* Fixed an overflow issue in `rocsparse_Xsctr` and the generic `rocsparse_scatter` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow and prevent some elements from being scattered. The kernel now computes element indices in 64-bit arithmetic and uses a grid-stride loop with the launch grid clamped to the device limit. * Fixed an integer overflow in `rocsparse_prune_dense2csr_by_percentage` and `rocsparse_prune_csr2csr_by_percentage`, which computed the matrix element count in 32-bit arithmetic. For matrices with more than `INT32_MAX` (~2.1 billion) elements the count overflowed to a negative value, resulting in out-of-bounds pointer construction and an invalid kernel launch grid. The element count is now computed in 64-bit arithmetic. * Fixed `rocsparse_spmm` with the segmented COO, atomic COO, segmented-atomic COO, and row-split CSR algorithms, which failed with `hipErrorInvalidConfiguration` for batch counts exceeding 65535 because the batch dimension of the kernel launch grid exceeded the maximum supported grid dimension. * Fixed an issue with `rocsparse_spmm` when using the nnz-split algorithm with the CSR or CSC format. The operation produced incorrect results because the segmented-block-reduction helper had shared-memory pointer parameters marked `__restrict__`, while threads in the block must read values written by other threads. The `__restrict__` attribute has now been removed. From c80d3c85309742744a697ad599e003cda15645ab Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 15 Sep 2026 11:35:47 -0600 Subject: [PATCH 5/6] test(rocsparse): make the relocated scatter_extra regression test runnable (AISPARSE-652) The large-nnz regression test built a sparse vector descriptor with size = 2 and nnz = 2^32 + 512. rocsparse_create_spvec_descr rejects nnz > size with rocsparse_status_invalid_size, so the test could never pass; it went unnoticed because its device_memory_gb guard kept it from being instantiated on every machine it had been run on so far. Size the dense vector to nnz so the descriptor is valid, and raise device_memory_gb to 70 to cover the now three nnz-sized device allocations (32 GiB indices + 16 GiB values + 16 GiB dense vector). Verified by running the test body at a scaled-down nnz: it passes with the fixed kernel and fails when the kernel is made to drop the tail of the sparse vector, so the probe is load-bearing. Co-authored-by: Cursor --- .../rocsparse/clients/testings/testing_scatter.cpp | 13 ++++++++----- projects/rocsparse/clients/tests/test_scatter.yaml | 2 +- 2 files changed, 9 insertions(+), 6 deletions(-) diff --git a/projects/rocsparse/clients/testings/testing_scatter.cpp b/projects/rocsparse/clients/testings/testing_scatter.cpp index 37c441761a62..fa3393283e1f 100644 --- a/projects/rocsparse/clients/testings/testing_scatter.cpp +++ b/projects/rocsparse/clients/testings/testing_scatter.cpp @@ -144,9 +144,9 @@ void testing_scatter_extra(const Arguments& arg) // // This drives the 64-bit-index path of rocsparse_scatter (which dispatches // to sctr_template) with nnz just past the 2^32 boundary and checks that an - // element beyond that boundary is actually scattered. To stay within a - // single device allocation everything is initialized on the device and a - // single element is probed. + // element beyond that boundary is actually scattered. Everything is + // initialized on the device and a single element is probed, so no host + // buffer of this size is ever needed. using I = int64_t; using T = float; @@ -154,8 +154,11 @@ void testing_scatter_extra(const Arguments& arg) // nnz just beyond 2^32 so at least one block has a block index whose // (blockIdx * BLOCKSIZE) product overflows 32-bit arithmetic. - const I nnz = two_pow_32 + 512; - const I size = 2; + const I nnz = two_pow_32 + 512; + + // rocsparse_create_spvec_descr rejects nnz > size, so the dense vector has + // to be at least as long as the sparse one. + const I size = nnz; const rocsparse_index_base base = rocsparse_index_base_zero; diff --git a/projects/rocsparse/clients/tests/test_scatter.yaml b/projects/rocsparse/clients/tests/test_scatter.yaml index d5dd8701de81..34f276daec45 100644 --- a/projects/rocsparse/clients/tests/test_scatter.yaml +++ b/projects/rocsparse/clients/tests/test_scatter.yaml @@ -154,4 +154,4 @@ Tests: category: nightly function: scatter_extra host_memory_gb: 1 - device_memory_gb: 55 + device_memory_gb: 70 From 132569904a63c4725d46c328e3dd8a55c597b15f Mon Sep 17 00:00:00 2001 From: kliegeois Date: Fri, 18 Sep 2026 17:00:16 +0000 Subject: [PATCH 6/6] docs(rocsparse): move the sctr overflow changelog entry to 5.2.0 (AISPARSE-652) The snap date for rocSPARSE 5.1.0 has passed, so this fix belongs in the next version. Open "## (Unreleased) rocSPARSE 5.2.0" and move only this PR's Resolved issues entry there. The 5.1.0 section is left otherwise unchanged so the net diff against develop stays additive. --- projects/rocsparse/CHANGELOG.md | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index b8a307b04bd4..3367d81af7ef 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -3,6 +3,11 @@ Documentation for rocSPARSE is available at [https://rocm.docs.amd.com/projects/rocSPARSE/en/latest/](https://rocm.docs.amd.com/projects/rocSPARSE/en/latest/). +## (Unreleased) rocSPARSE 5.2.0 + +### Resolved issues +* Fixed an overflow issue in `rocsparse_Xsctr` and the generic `rocsparse_scatter` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow and prevent some elements from being scattered. The kernel now computes element indices in 64-bit arithmetic and uses a grid-stride loop with the launch grid clamped to the device limit. + ## (Unreleased) rocSPARSE 5.1.0 ### Added @@ -15,7 +20,6 @@ Documentation for rocSPARSE is available at * Optimized architecture-aware launch configurations for RDNA (wave32) and CDNA (wave64) GPUs, improving performance and performance portability for several sparse level 2 and level 3 routines without algorithmic or numerical changes. Affected routines include `rocsparse_spmv` for the CSR adaptive, nnz-split, and LRB algorithms, the COO (SoA and AoS) formats, and the ELL format (`rocsparse_Xellmv`); `rocsparse_Xbsrmv`; `rocsparse_Xbsrxmv`; `rocsparse_Xgemvi`; `rocsparse_Xgemmi`; and `rocsparse_spmm` with the blocked-ELL format. ### Resolved issues -* Fixed an overflow issue in `rocsparse_Xsctr` and the generic `rocsparse_scatter` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow and prevent some elements from being scattered. The kernel now computes element indices in 64-bit arithmetic and uses a grid-stride loop with the launch grid clamped to the device limit. * Fixed an integer overflow in `rocsparse_prune_dense2csr_by_percentage` and `rocsparse_prune_csr2csr_by_percentage`, which computed the matrix element count in 32-bit arithmetic. For matrices with more than `INT32_MAX` (~2.1 billion) elements the count overflowed to a negative value, resulting in out-of-bounds pointer construction and an invalid kernel launch grid. The element count is now computed in 64-bit arithmetic. * Fixed `rocsparse_spmm` with the segmented COO, atomic COO, segmented-atomic COO, and row-split CSR algorithms, which failed with `hipErrorInvalidConfiguration` for batch counts exceeding 65535 because the batch dimension of the kernel launch grid exceeded the maximum supported grid dimension. * Fixed an issue with `rocsparse_spmm` when using the nnz-split algorithm with the CSR or CSC format. The operation produced incorrect results because the segmented-block-reduction helper had shared-memory pointer parameters marked `__restrict__`, while threads in the block must read values written by other threads. The `__restrict__` attribute has now been removed.