From a111ffa39b99fa1db0a778509fb45cc5a725fe53 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Thu, 24 Sep 2026 11:50:08 -0600 Subject: [PATCH 01/10] fix(rocsparse): clamp extra-nightly launches to the 32-bit dispatch limit 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 #12545 Co-authored-by: Cursor --- projects/rocsparse/CHANGELOG.md | 3 +- .../clients/testings/testing_axpby.cpp | 13 +++-- .../rocsparse/clients/tests/test_axpby.yaml | 4 +- .../library/src/include/rocsparse_common.hpp | 18 ++++++ .../library/src/level1/axpyi_device.h | 9 ++- .../library/src/level1/rocsparse_axpyi.cpp | 12 ++-- .../src/level1/rocsparse_sctr_impl.cpp | 11 ++-- .../library/src/level1/sctr_device.h | 9 ++- .../library/src/level2/ellmv_device.h | 47 ++++++++-------- .../src/level2/rocsparse_coomv_aos_impl.cpp | 56 ++++++++++++++----- .../src/level2/rocsparse_ellmv_impl.cpp | 35 ++++++------ 11 files changed, 128 insertions(+), 89 deletions(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index d9c6996392a3..28a737d6e57f 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -7,8 +7,9 @@ Documentation for rocSPARSE is available at ### Resolved issues * Fixed an overflow issue in `rocsparse_roti` and the generic `rocsparse_rot` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow, leaving some elements unrotated and causing low-index elements to be processed with incorrect data. 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 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 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 32-bit work-item dispatch limit. * Fixed an integer overflow in the ELL SpMV kernels used by `rocsparse_spmv` (ELL format) and the legacy `rocsparse_Xellmv` routines when operating on matrices with more than `2^32` rows and 64-bit index types. The overflow caused rows beyond the `2^32` boundary to be mapped to incorrect row indices, leaving some output values uncomputed. Row indices are now always computed using the matrix index type. +* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` / `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV, and COO AoS atomic SpMV now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop or an extra inner loop. ## (Unreleased) rocSPARSE 5.1.0 diff --git a/projects/rocsparse/clients/testings/testing_axpby.cpp b/projects/rocsparse/clients/testings/testing_axpby.cpp index e5aa35b3bcaf..5747dcc691c9 100644 --- a/projects/rocsparse/clients/testings/testing_axpby.cpp +++ b/projects/rocsparse/clients/testings/testing_axpby.cpp @@ -177,10 +177,9 @@ void testing_axpby_extra(const Arguments& arg) // // This drives the 64-bit-index path of rocsparse_axpby (which dispatches to // axpyi_template) with nnz just past the 2^32 boundary and checks that an - // element beyond that boundary is actually accumulated into y. To stay - // within a single device allocation (host mirrors of the full arrays would - // need tens of GB) everything is initialized on the device and a single - // element is probed. + // element beyond that boundary is actually accumulated into y. Everything + // is initialized on the device and a single element is probed, so no host + // buffer of this size is needed. using I = int64_t; using T = float; @@ -188,8 +187,10 @@ void testing_axpby_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; + + // A sparse-vector descriptor requires nnz <= size. + const I size = nnz; const rocsparse_index_base base = rocsparse_index_base_zero; diff --git a/projects/rocsparse/clients/tests/test_axpby.yaml b/projects/rocsparse/clients/tests/test_axpby.yaml index ce754d22a3fe..6d31fbd52ad9 100644 --- a/projects/rocsparse/clients/tests/test_axpby.yaml +++ b/projects/rocsparse/clients/tests/test_axpby.yaml @@ -131,10 +131,10 @@ Tests: # Regression test for AISPARSE-650: axpby with a 64-bit index type and nnz just # past 2^32 exercises the block-index overflow that was fixed in the axpyi -# device kernel. Needs a large-memory device (~52 GB of index/value data); +# device kernel. Needs a large-memory device (~68 GiB of device data); # skipped elsewhere. - name: axpby_extra category: nightly function: axpby_extra host_memory_gb: 1 - device_memory_gb: 55 + device_memory_gb: 70 diff --git a/projects/rocsparse/library/src/include/rocsparse_common.hpp b/projects/rocsparse/library/src/include/rocsparse_common.hpp index 285be259ba9f..3198c6a1477d 100644 --- a/projects/rocsparse/library/src/include/rocsparse_common.hpp +++ b/projects/rocsparse/library/src/include/rocsparse_common.hpp @@ -60,6 +60,24 @@ namespace rocsparse { + // The AQL dispatch packet stores grid extents as 32-bit work-item counts. + // Convert that limit to workgroups for the requested block size. + static constexpr int64_t max_grid_size_x(int64_t block_size) + { + return static_cast(UINT32_MAX) / block_size; + } + + template + static inline uint32_t + get_grid_size_x(const hipDeviceProp_t& properties, J count, int64_t block_size) + { + const int64_t requested = static_cast(count); + const int64_t device_cap = static_cast(properties.maxGridSize[0]); + const int64_t dispatch_cap + = (device_cap < max_grid_size_x(block_size)) ? device_cap : max_grid_size_x(block_size); + return static_cast((requested < dispatch_cap) ? requested : dispatch_cap); + } + template static uint16_t get_batch_grid_size(J batch_count) { diff --git a/projects/rocsparse/library/src/level1/axpyi_device.h b/projects/rocsparse/library/src/level1/axpyi_device.h index f2fa323c035e..16603a8af4a4 100644 --- a/projects/rocsparse/library/src/level1/axpyi_device.h +++ b/projects/rocsparse/library/src/level1/axpyi_device.h @@ -33,14 +33,13 @@ namespace rocsparse ROCSPARSE_DEVICE_ILF void axpyi_device( I nnz, T alpha, const X* x_val, const I* x_ind, Y* y, rocsparse_index_base idx_base) { - // Cast to the (possibly 64-bit) index type before the multiply so the - // computation does not overflow a 32-bit unsigned int when nnz is large. - I gid = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; - I stride = static_cast(hipGridDim_x) * BLOCKSIZE; + // Keep the grid-stride arithmetic wide even when I is 32-bit. + const int64_t gid = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; + const int64_t stride = static_cast(hipGridDim_x) * BLOCKSIZE; // Grid-stride loop so the full vector is processed even when the number // of required blocks exceeds the grid size. - for(I idx = gid; idx < nnz; idx += stride) + for(int64_t idx = gid; idx < nnz; idx += stride) { I i = x_ind[idx] - idx_base; y[i] = rocsparse::fma(alpha, x_val[idx], y[i]); diff --git a/projects/rocsparse/library/src/level1/rocsparse_axpyi.cpp b/projects/rocsparse/library/src/level1/rocsparse_axpyi.cpp index b8cdef70723b..d5419d7e1470 100644 --- a/projects/rocsparse/library/src/level1/rocsparse_axpyi.cpp +++ b/projects/rocsparse/library/src/level1/rocsparse_axpyi.cpp @@ -98,13 +98,11 @@ rocsparse_status rocsparse::axpyi_template(rocsparse_handle handle, hipStream_t stream = handle->stream; #define AXPYI_DIM 256 - // Number of blocks needed to cover nnz with one thread per element. Clamp - // against the device grid limit; the kernel uses a grid-stride loop so any - // remaining elements are still processed when the requested grid is capped. - const int64_t num_blocks - = rocsparse::min((static_cast(nnz) - 1) / AXPYI_DIM + 1, - static_cast(handle->properties.maxGridSize[0])); - dim3 axpyi_blocks(static_cast(num_blocks)); + // Clamp to both the device limit and the dispatch packet's work-item limit. + // The grid-stride kernel processes elements beyond the clamped grid. + const uint32_t num_blocks = rocsparse::get_grid_size_x( + handle->properties, (static_cast(nnz) - 1) / AXPYI_DIM + 1, AXPYI_DIM); + dim3 axpyi_blocks(num_blocks); dim3 axpyi_threads(AXPYI_DIM); RETURN_IF_HIPLAUNCHKERNELGGL_ERROR((rocsparse::axpyi_kernel), diff --git a/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp index dce0d3d76fa4..98dc3678707e 100644 --- a/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp +++ b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp @@ -68,12 +68,11 @@ rocsparse_status rocsparse::sctr_template(rocsparse_handle handle, hipStream_t stream = handle->stream; #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. - 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)); + // Clamp to both the device limit and the dispatch packet's work-item limit. + // The grid-stride kernel processes elements beyond the clamped grid. + const uint32_t num_blocks = rocsparse::get_grid_size_x( + handle->properties, (static_cast(nnz) - 1) / SCTR_DIM + 1, SCTR_DIM); + 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 f5cf3e95653c..48afed78c8d6 100644 --- a/projects/rocsparse/library/src/level1/sctr_device.h +++ b/projects/rocsparse/library/src/level1/sctr_device.h @@ -32,12 +32,11 @@ 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) { - // 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; + // Keep the grid-stride arithmetic wide even when I is 32-bit. + const int64_t stride = static_cast(hipGridDim_x) * BLOCKSIZE; + const int64_t gid = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; - for(I idx = gid; idx < nnz; idx += stride) + for(int64_t idx = gid; idx < nnz; idx += stride) { y[x_ind[idx] - idx_base] = x_val[idx]; } diff --git a/projects/rocsparse/library/src/level2/ellmv_device.h b/projects/rocsparse/library/src/level2/ellmv_device.h index 76b0d2d94c82..03cfb8db4c2e 100644 --- a/projects/rocsparse/library/src/level2/ellmv_device.h +++ b/projects/rocsparse/library/src/level2/ellmv_device.h @@ -41,41 +41,38 @@ namespace rocsparse Y* y, rocsparse_index_base idx_base) { - // Cast to the (possibly 64-bit) index type I before the multiply so the - // global thread id does not wrap at 2^32 when m exceeds the 32-bit range. - const I ai = static_cast(BLOCKSIZE) * hipBlockIdx_x + hipThreadIdx_x; + const int64_t first_row = static_cast(BLOCKSIZE) * hipBlockIdx_x + hipThreadIdx_x; + const int64_t row_stride = static_cast(BLOCKSIZE) * hipGridDim_x; - if(ai >= m) + for(int64_t ai = first_row; ai < m; ai += row_stride) { - return; - } + T sum = static_cast(0); + for(I p = 0; p < ell_width; ++p) + { - T sum = static_cast(0); - for(I p = 0; p < ell_width; ++p) - { + const int64_t idx = ELL_IND(ai, (int64_t)p, m, ell_width); + const I col = rocsparse::nontemporal_load(ell_col_ind + idx) - idx_base; + if(col >= 0 && col < n) + { + sum = rocsparse::fma( + rocsparse::nontemporal_load(ell_val + idx), rocsparse::ldg(x + col), sum); + } + else + { + break; + } + } - const int64_t idx = ELL_IND(ai, (int64_t)p, m, ell_width); - const I col = rocsparse::nontemporal_load(ell_col_ind + idx) - idx_base; - if(col >= 0 && col < n) + if(beta != static_cast(0)) { - sum = rocsparse::fma( - rocsparse::nontemporal_load(ell_val + idx), rocsparse::ldg(x + col), sum); + const Y yv = rocsparse::nontemporal_load(y + ai); + rocsparse::nontemporal_store(rocsparse::fma(beta, yv, alpha * sum), y + ai); } else { - break; + rocsparse::nontemporal_store(alpha * sum, y + ai); } } - - if(beta != static_cast(0)) - { - const Y yv = rocsparse::nontemporal_load(y + ai); - rocsparse::nontemporal_store(rocsparse::fma(beta, yv, alpha * sum), y + ai); - } - else - { - rocsparse::nontemporal_store(alpha * sum, y + ai); - } } // ELL SpMV for general, (conjugate) transposed matrices diff --git a/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp b/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp index 941046b6d1d2..b1caaa7b01f3 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp @@ -159,21 +159,47 @@ namespace rocsparse { case rocsparse_operation_none: { - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( - (rocsparse::coomvn_aos_atomic_loops<256, 1>), - dim3((nnz - 1) / 256 + 1), - dim3(256), - 0, - stream, - nnz, - m, - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), - coo_ind, - coo_val, - x, - y, - descr->base, - handle->pointer_mode == rocsparse_pointer_mode_host); + const int64_t one_loop_blocks = (nnz - 1) / 256 + 1; + const bool use_two_loops + = rocsparse::get_grid_size_x(handle->properties, one_loop_blocks, 256) + < one_loop_blocks; + + if(use_two_loops) + { + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( + (rocsparse::coomvn_aos_atomic_loops<256, 2>), + dim3((nnz - 1) / (2 * 256) + 1), + dim3(256), + 0, + stream, + nnz, + m, + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), + coo_ind, + coo_val, + x, + y, + descr->base, + handle->pointer_mode == rocsparse_pointer_mode_host); + } + else + { + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( + (rocsparse::coomvn_aos_atomic_loops<256, 1>), + dim3(one_loop_blocks), + dim3(256), + 0, + stream, + nnz, + m, + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), + coo_ind, + coo_val, + x, + y, + descr->base, + handle->pointer_mode == rocsparse_pointer_mode_host); + } break; } case rocsparse_operation_transpose: diff --git a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp index 20826b3718d6..d5e2d2947ec3 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp @@ -99,23 +99,24 @@ namespace rocsparse // Run different ellmv kernels if(trans == rocsparse_operation_none) { -#define LAUNCH_ELLMVN(DIM) \ - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( \ - (rocsparse::ellmvn_kernel), \ - dim3((m - 1) / (DIM) + 1), \ - dim3(DIM), \ - 0, \ - stream, \ - m, \ - n, \ - ell_width, \ - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), \ - ell_col_ind, \ - ell_val, \ - x, \ - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, beta_device_host), \ - y, \ - descr->base, \ +#define LAUNCH_ELLMVN(DIM) \ + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( \ + (rocsparse::ellmvn_kernel), \ + dim3(rocsparse::get_grid_size_x( \ + handle->properties, (static_cast(m) - 1) / (DIM) + 1, (DIM))), \ + dim3(DIM), \ + 0, \ + stream, \ + m, \ + n, \ + ell_width, \ + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), \ + ell_col_ind, \ + ell_val, \ + x, \ + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, beta_device_host), \ + y, \ + descr->base, \ handle->pointer_mode == rocsparse_pointer_mode_host) // Launch tuning for the one-thread-per-row non-transpose kernel. From 6ec92105ab2f3558279adc247a645e874a6df5e1 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 18:58:32 +0000 Subject: [PATCH 02/10] fix(rocsparse): clamp and grid-stride the scale_array kernel 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. --- projects/rocsparse/CHANGELOG.md | 2 +- .../clients/testings/testing_axpby.cpp | 42 ++++++++++++++----- .../library/src/common/rocsparse_common.cpp | 28 +++++++------ 3 files changed, 47 insertions(+), 25 deletions(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index c2de73ed2872..0ee441b25ca2 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -12,7 +12,7 @@ Documentation for rocSPARSE is available at * Fixed an overflow issue in `rocsparse_roti` and the generic `rocsparse_rot` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow, leaving some elements unrotated and causing low-index elements to be processed with incorrect data. 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 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 32-bit work-item dispatch limit. * Fixed an integer overflow in the ELL SpMV kernels used by `rocsparse_spmv` (ELL format) and the legacy `rocsparse_Xellmv` routines when operating on matrices with more than `2^32` rows and 64-bit index types. The overflow caused rows beyond the `2^32` boundary to be mapped to incorrect row indices, leaving some output values uncomputed. Row indices are now always computed using the matrix index type. -* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` / `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV, and COO AoS atomic SpMV now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop or an extra inner loop. +* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` (both the `beta` scaling of `y` and the `axpyi` update), `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV, and COO AoS atomic SpMV now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop or an extra inner loop. ## (Unreleased) rocSPARSE 5.1.0 diff --git a/projects/rocsparse/clients/testings/testing_axpby.cpp b/projects/rocsparse/clients/testings/testing_axpby.cpp index 5747dcc691c9..7f2d3220fde5 100644 --- a/projects/rocsparse/clients/testings/testing_axpby.cpp +++ b/projects/rocsparse/clients/testings/testing_axpby.cpp @@ -175,11 +175,13 @@ void testing_axpby_extra(const Arguments& arg) // touched. 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_axpby (which dispatches to - // axpyi_template) with nnz just past the 2^32 boundary and checks that an - // element beyond that boundary is actually accumulated into y. Everything - // is initialized on the device and a single element is probed, so no host - // buffer of this size is needed. + // This drives the 64-bit-index path of rocsparse_axpby with nnz and size + // just past the 2^32 boundary. beta is neither 0 nor 1, so axpby first + // launches the scale kernel over all of y and then axpyi_template over x. + // The test checks that an element beyond the boundary is accumulated into + // y and that the last element of y is scaled. Everything is initialized on + // the device and only a few elements are probed, so no host buffer of this + // size is needed. using I = int64_t; using T = float; @@ -216,9 +218,18 @@ void testing_axpby_extra(const Arguments& arg) CHECK_HIP_ERROR( hipMemcpy(static_cast(dx_val) + probe_idx, &x_in, sizeof(T), hipMemcpyHostToDevice)); - // beta == 1 leaves the existing y untouched; alpha scales the gathered x. + // y[1] is the probe's output; y[size - 1] is past the 2^32 boundary and is + // only touched by the beta scaling. + const I tail_idx = size - 1; + const T y_in = static_cast(1); + const T tail_in = static_cast(5); + CHECK_HIP_ERROR( + hipMemcpy(static_cast(dy) + probe_ind, &y_in, sizeof(T), hipMemcpyHostToDevice)); + CHECK_HIP_ERROR( + hipMemcpy(static_cast(dy) + tail_idx, &tail_in, sizeof(T), hipMemcpyHostToDevice)); + const T halpha = static_cast(2); - const T hbeta = static_cast(1); + const T hbeta = static_cast(3); rocsparse_local_spvec x(size, nnz, dx_ind, dx_val, get_indextype(), base, get_datatype()); rocsparse_local_dnvec y(size, dy, get_datatype()); @@ -226,10 +237,19 @@ void testing_axpby_extra(const Arguments& arg) CHECK_ROCSPARSE_ERROR(rocsparse_set_pointer_mode(handle, rocsparse_pointer_mode_host)); CHECK_ROCSPARSE_ERROR(testing::rocsparse_axpby(handle, &halpha, x, &hbeta, y)); - // y[1] = beta * 0 + alpha * x_in = 2. Before the fix the wrapped block index - // leaves the probe element unprocessed, so y[1] stays 0. + // y[1] = beta * y_in + alpha * x_in = 5. Before the fix the wrapped block + // index leaves the probe element unprocessed. T y_out = static_cast(0); - CHECK_HIP_ERROR(hipMemcpy(&y_out, static_cast(dy) + 1, sizeof(T), hipMemcpyDeviceToHost)); + CHECK_HIP_ERROR( + hipMemcpy(&y_out, static_cast(dy) + probe_ind, sizeof(T), hipMemcpyDeviceToHost)); + + unit_check_scalar(hbeta * y_in + halpha * x_in, y_out); + + // y[size - 1] = beta * tail_in = 15. An unclamped scale grid is rejected at + // launch, and a clamped grid without a grid-stride loop leaves it unscaled. + T tail_out = static_cast(0); + CHECK_HIP_ERROR( + hipMemcpy(&tail_out, static_cast(dy) + tail_idx, sizeof(T), hipMemcpyDeviceToHost)); - unit_check_scalar(static_cast(2), y_out); + unit_check_scalar(hbeta * tail_in, tail_out); } diff --git a/projects/rocsparse/library/src/common/rocsparse_common.cpp b/projects/rocsparse/library/src/common/rocsparse_common.cpp index 368078acd3f4..8d1aed8cf1a1 100644 --- a/projects/rocsparse/library/src/common/rocsparse_common.cpp +++ b/projects/rocsparse/library/src/common/rocsparse_common.cpp @@ -24,6 +24,7 @@ #include "rocsparse_common.h" #include "rocsparse_common.hpp" +#include "rocsparse_grid.hpp" #include "rocsparse_utility.hpp" #include @@ -51,20 +52,20 @@ namespace rocsparse template ROCSPARSE_DEVICE_ILF void scale_device(I length, T scalar, A* __restrict__ array) { - const I gid = hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x; + const int64_t stride = static_cast(hipGridDim_x) * BLOCKSIZE; - if(gid >= length) + for(int64_t gid = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; + gid < length; + gid += stride) { - return; - } - - if(scalar == static_cast(0)) - { - array[gid] = static_cast(0); - } - else - { - array[gid] *= scalar; + if(scalar == static_cast(0)) + { + array[gid] = static_cast(0); + } + else + { + array[gid] *= scalar; + } } } @@ -264,7 +265,8 @@ rocsparse_status rocsparse::scale_array(rocsparse_handle handle, { RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( (rocsparse::scale_kernel<256>), - dim3((length - 1) / 256 + 1), + dim3(rocsparse::get_grid_size_x( + handle, (static_cast(length) - 1) / 256 + 1, 256)), dim3(256), 0, handle->stream, From fd8882b3cbf4b5841c1f58e5fcbdd3378e23faaf Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 18:58:33 +0000 Subject: [PATCH 03/10] fix(rocsparse): keep COO AoS SpMV grids within the dispatch limit 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. --- projects/rocsparse/CHANGELOG.md | 2 +- .../library/src/level2/coomv_device.h | 71 +++++++++++++------ .../src/level2/rocsparse_coomv_aos_impl.cpp | 59 +++++---------- 3 files changed, 69 insertions(+), 63 deletions(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index 0ee441b25ca2..1d6d2b815df2 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -12,7 +12,7 @@ Documentation for rocSPARSE is available at * Fixed an overflow issue in `rocsparse_roti` and the generic `rocsparse_rot` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow, leaving some elements unrotated and causing low-index elements to be processed with incorrect data. 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 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 32-bit work-item dispatch limit. * Fixed an integer overflow in the ELL SpMV kernels used by `rocsparse_spmv` (ELL format) and the legacy `rocsparse_Xellmv` routines when operating on matrices with more than `2^32` rows and 64-bit index types. The overflow caused rows beyond the `2^32` boundary to be mapped to incorrect row indices, leaving some output values uncomputed. Row indices are now always computed using the matrix index type. -* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` (both the `beta` scaling of `y` and the `axpyi` update), `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV, and COO AoS atomic SpMV now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop or an extra inner loop. +* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` (both the `beta` scaling of `y` and the `axpyi` update), `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV, and COO AoS SpMV (atomic and transposed) now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop. ## (Unreleased) rocSPARSE 5.1.0 diff --git a/projects/rocsparse/library/src/level2/coomv_device.h b/projects/rocsparse/library/src/level2/coomv_device.h index e58f9e2ec01a..016be32d0cc3 100644 --- a/projects/rocsparse/library/src/level2/coomv_device.h +++ b/projects/rocsparse/library/src/level2/coomv_device.h @@ -542,14 +542,15 @@ namespace rocsparse typename X, typename Y, typename T> - ROCSPARSE_DEVICE_ILF void coomvn_aos_atomic_loops_device(int64_t nnz, - I m, - T alpha, - const I* __restrict__ coo_ind, - const A* __restrict__ coo_val, - const X* __restrict__ x, - Y* __restrict__ y, - rocsparse_index_base idx_base) + ROCSPARSE_DEVICE_ILF void coomvn_aos_atomic_loops_tile_device(int64_t tile, + int64_t nnz, + I m, + T alpha, + const I* __restrict__ coo_ind, + const A* __restrict__ coo_val, + const X* __restrict__ x, + Y* __restrict__ y, + rocsparse_index_base idx_base) { const int tid = hipThreadIdx_x; @@ -561,7 +562,7 @@ namespace rocsparse T val; // Current threads index into COO structure - int64_t idx = static_cast(hipBlockIdx_x) * LOOPS * BLOCKSIZE + tid; + int64_t idx = tile * LOOPS * BLOCKSIZE + tid; if(idx < nnz) { @@ -677,6 +678,36 @@ namespace rocsparse } } + // Each tile covers LOOPS * BLOCKSIZE entries. Blocks stride over the tiles, + // so a grid clamped below nnz / (LOOPS * BLOCKSIZE) still covers every entry. + template + ROCSPARSE_DEVICE_ILF void coomvn_aos_atomic_loops_device(int64_t nnz, + I m, + T alpha, + const I* __restrict__ coo_ind, + const A* __restrict__ coo_val, + const X* __restrict__ x, + Y* __restrict__ y, + rocsparse_index_base idx_base) + { + const int64_t tile_size = static_cast(LOOPS) * BLOCKSIZE; + + for(int64_t tile = hipBlockIdx_x; tile * tile_size < nnz; tile += hipGridDim_x) + { + rocsparse::coomvn_aos_atomic_loops_tile_device( + tile, nnz, m, alpha, coo_ind, coo_val, x, y, idx_base); + + // The next tile overwrites the shared buffers this tile still reads. + __syncthreads(); + } + } + template ROCSPARSE_DEVICE_ILF void coomvt_device(rocsparse_operation trans, int64_t nnz, @@ -716,19 +747,19 @@ namespace rocsparse Y* y, rocsparse_index_base idx_base) { - const int64_t gid = static_cast(hipBlockIdx_x) * hipBlockDim_x + hipThreadIdx_x; + const int64_t stride = static_cast(hipGridDim_x) * hipBlockDim_x; - if(gid >= nnz) + for(int64_t gid = static_cast(hipBlockIdx_x) * hipBlockDim_x + hipThreadIdx_x; + gid < nnz; + gid += stride) { - return; - } - - const I row = coo_ind[2 * gid] - idx_base; - const I col = coo_ind[2 * gid + 1] - idx_base; - const A val = (trans == rocsparse_operation_conjugate_transpose) - ? rocsparse::conj(coo_val[gid]) - : coo_val[gid]; + const I row = coo_ind[2 * gid] - idx_base; + const I col = coo_ind[2 * gid + 1] - idx_base; + const A val = (trans == rocsparse_operation_conjugate_transpose) + ? rocsparse::conj(coo_val[gid]) + : coo_val[gid]; - rocsparse::atomic_add(y, col, n, alpha * val * x[row]); + rocsparse::atomic_add(y, col, n, alpha * val * x[row]); + } } } diff --git a/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp b/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp index 56e2d2bd33f7..d7cda03df604 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp @@ -160,46 +160,21 @@ namespace rocsparse { case rocsparse_operation_none: { - const int64_t one_loop_blocks = (nnz - 1) / 256 + 1; - const bool use_two_loops - = rocsparse::get_grid_size_x(handle, one_loop_blocks, 256) < one_loop_blocks; - - if(use_two_loops) - { - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( - (rocsparse::coomvn_aos_atomic_loops<256, 2>), - dim3((nnz - 1) / (2 * 256) + 1), - dim3(256), - 0, - stream, - nnz, - m, - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), - coo_ind, - coo_val, - x, - y, - descr->base, - handle->pointer_mode == rocsparse_pointer_mode_host); - } - else - { - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( - (rocsparse::coomvn_aos_atomic_loops<256, 1>), - dim3(one_loop_blocks), - dim3(256), - 0, - stream, - nnz, - m, - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), - coo_ind, - coo_val, - x, - y, - descr->base, - handle->pointer_mode == rocsparse_pointer_mode_host); - } + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( + (rocsparse::coomvn_aos_atomic_loops<256, 1>), + dim3(rocsparse::get_grid_size_x(handle, (nnz - 1) / 256 + 1, 256)), + dim3(256), + 0, + stream, + nnz, + m, + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), + coo_ind, + coo_val, + x, + y, + descr->base, + handle->pointer_mode == rocsparse_pointer_mode_host); break; } case rocsparse_operation_transpose: @@ -207,7 +182,7 @@ namespace rocsparse { RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( (rocsparse::coomvt_aos_kernel<1024>), - dim3((nnz - 1) / 1024 + 1), + dim3(rocsparse::get_grid_size_x(handle, (nnz - 1) / 1024 + 1, 1024)), dim3(1024), 0, handle->stream, @@ -366,7 +341,7 @@ namespace rocsparse { RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( (rocsparse::coomvt_aos_kernel<1024>), - dim3((nnz - 1) / 1024 + 1), + dim3(rocsparse::get_grid_size_x(handle, (nnz - 1) / 1024 + 1, 1024)), dim3(1024), 0, handle->stream, From 8180515781b2580893f0cdfa78314826e8090c6b Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 18:58:34 +0000 Subject: [PATCH 04/10] fix(rocsparse): clamp and grid-stride transposed ELL SpMV 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. --- projects/rocsparse/CHANGELOG.md | 2 +- .../library/src/level2/ellmv_device.h | 41 +++++++++---------- .../src/level2/rocsparse_ellmv_impl.cpp | 3 +- 3 files changed, 22 insertions(+), 24 deletions(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index 1d6d2b815df2..066569d772e2 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -12,7 +12,7 @@ Documentation for rocSPARSE is available at * Fixed an overflow issue in `rocsparse_roti` and the generic `rocsparse_rot` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow, leaving some elements unrotated and causing low-index elements to be processed with incorrect data. 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 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 32-bit work-item dispatch limit. * Fixed an integer overflow in the ELL SpMV kernels used by `rocsparse_spmv` (ELL format) and the legacy `rocsparse_Xellmv` routines when operating on matrices with more than `2^32` rows and 64-bit index types. The overflow caused rows beyond the `2^32` boundary to be mapped to incorrect row indices, leaving some output values uncomputed. Row indices are now always computed using the matrix index type. -* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` (both the `beta` scaling of `y` and the `axpyi` update), `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV, and COO AoS SpMV (atomic and transposed) now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop. +* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` (both the `beta` scaling of `y` and the `axpyi` update), `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV (non-transposed and transposed), and COO AoS SpMV (atomic and transposed) now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop. ## (Unreleased) rocSPARSE 5.1.0 diff --git a/projects/rocsparse/library/src/level2/ellmv_device.h b/projects/rocsparse/library/src/level2/ellmv_device.h index 03cfb8db4c2e..0c0a8a1edfb7 100644 --- a/projects/rocsparse/library/src/level2/ellmv_device.h +++ b/projects/rocsparse/library/src/level2/ellmv_device.h @@ -88,36 +88,33 @@ namespace rocsparse Y* y, rocsparse_index_base idx_base) { - // Cast to the (possibly 64-bit) index type I before the multiply so the - // global thread id does not wrap at 2^32 when m exceeds the 32-bit range. - const I ai = static_cast(BLOCKSIZE) * hipBlockIdx_x + hipThreadIdx_x; + const int64_t first_row = static_cast(BLOCKSIZE) * hipBlockIdx_x + hipThreadIdx_x; + const int64_t row_stride = static_cast(BLOCKSIZE) * hipGridDim_x; - if(ai >= m) + for(int64_t ai = first_row; ai < m; ai += row_stride) { - return; - } + const T row_val = alpha * rocsparse::ldg(x + ai); - const T row_val = alpha * rocsparse::ldg(x + ai); + for(I p = 0; p < ell_width; ++p) + { + const int64_t idx = ELL_IND(ai, (int64_t)p, m, ell_width); + const I col = rocsparse::nontemporal_load(ell_col_ind + idx) - idx_base; - for(I p = 0; p < ell_width; ++p) - { - const int64_t idx = ELL_IND(ai, (int64_t)p, m, ell_width); - const I col = rocsparse::nontemporal_load(ell_col_ind + idx) - idx_base; + if(col >= 0 && col < n) + { + A val = rocsparse::nontemporal_load(ell_val + idx); - if(col >= 0 && col < n) - { - A val = rocsparse::nontemporal_load(ell_val + idx); + if(trans == rocsparse_operation_conjugate_transpose) + { + val = rocsparse::conj(val); + } - if(trans == rocsparse_operation_conjugate_transpose) + rocsparse::atomic_add(y, col, n, row_val * val); + } + else { - val = rocsparse::conj(val); + break; } - - rocsparse::atomic_add(y, col, n, row_val * val); - } - else - { - break; } } } diff --git a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp index 6bece1b48d36..05c3a9fe93e9 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp @@ -165,7 +165,8 @@ namespace rocsparse RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( (rocsparse::ellmvt_kernel), - dim3((m - 1) / ELLMVT_DIM + 1), + dim3(rocsparse::get_grid_size_x( + handle, (static_cast(m) - 1) / ELLMVT_DIM + 1, ELLMVT_DIM)), dim3(ELLMVT_DIM), 0, stream, From 3b7cb14da4b226561e617c46cf9c5cd8c65d6352 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 19:46:50 +0000 Subject: [PATCH 05/10] style(rocsparse): use static_cast instead of C-style casts --- projects/rocsparse/library/src/level2/ellmv_device.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/rocsparse/library/src/level2/ellmv_device.h b/projects/rocsparse/library/src/level2/ellmv_device.h index 0c0a8a1edfb7..64b566842959 100644 --- a/projects/rocsparse/library/src/level2/ellmv_device.h +++ b/projects/rocsparse/library/src/level2/ellmv_device.h @@ -50,7 +50,7 @@ namespace rocsparse for(I p = 0; p < ell_width; ++p) { - const int64_t idx = ELL_IND(ai, (int64_t)p, m, ell_width); + const int64_t idx = ELL_IND(ai, static_cast(p), m, ell_width); const I col = rocsparse::nontemporal_load(ell_col_ind + idx) - idx_base; if(col >= 0 && col < n) { @@ -97,7 +97,7 @@ namespace rocsparse for(I p = 0; p < ell_width; ++p) { - const int64_t idx = ELL_IND(ai, (int64_t)p, m, ell_width); + const int64_t idx = ELL_IND(ai, static_cast(p), m, ell_width); const I col = rocsparse::nontemporal_load(ell_col_ind + idx) - idx_base; if(col >= 0 && col < n) From 0013cae814c51518e577ae3a29044e537c703cad Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 22:23:55 +0000 Subject: [PATCH 06/10] fix(rocsparse): clamp the roti grid to the 32-bit dispatch limit 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. --- projects/rocsparse/CHANGELOG.md | 4 +- .../unittests/unit_test_device_level1.cpp | 58 ++++++++++++++++++- .../library/src/level1/rocsparse_roti.cpp | 55 +++++++++++------- .../library/src/level1/roti_device.h | 44 ++++++++++---- 4 files changed, 123 insertions(+), 38 deletions(-) diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index 066569d772e2..5c6ea628419d 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -9,10 +9,10 @@ Documentation for rocSPARSE is available at * Added support for the `gfx1250-strict` architecture. ### Resolved issues -* Fixed an overflow issue in `rocsparse_roti` and the generic `rocsparse_rot` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow, leaving some elements unrotated and causing low-index elements to be processed with incorrect data. 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 overflow issue in `rocsparse_roti` and the generic `rocsparse_rot` routine. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-index calculation could overflow, leaving some elements unrotated and causing low-index elements to be processed with incorrect data. The kernel now computes element indices in 64-bit arithmetic and uses a grid-stride loop with the launch grid clamped to the 32-bit work-item dispatch 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 32-bit work-item dispatch limit. * Fixed an integer overflow in the ELL SpMV kernels used by `rocsparse_spmv` (ELL format) and the legacy `rocsparse_Xellmv` routines when operating on matrices with more than `2^32` rows and 64-bit index types. The overflow caused rows beyond the `2^32` boundary to be mapped to incorrect row indices, leaving some output values uncomputed. Row indices are now always computed using the matrix index type. -* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` (both the `beta` scaling of `y` and the `axpyi` update), `rocsparse_axpyi`, `rocsparse_scatter`, ELL SpMV (non-transposed and transposed), and COO AoS SpMV (atomic and transposed) now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop. +* Fixed kernel launches whose `grid.x * blockDim.x` exceeded the 32-bit work-item dispatch limit. CDNA rejects those launches (`hipErrorInvalidConfiguration` / `rocsparse_status_internal_error`). `rocsparse_axpby` (both the `beta` scaling of `y` and the `axpyi` update), `rocsparse_axpyi`, `rocsparse_roti` and `rocsparse_rot`, `rocsparse_scatter`, ELL SpMV (non-transposed and transposed), and COO AoS SpMV (atomic and transposed) now clamp `grid.x` to `(2^32 - 1) / blockDim.x` and cover the remainder with a grid-stride loop. ## (Unreleased) rocSPARSE 5.1.0 diff --git a/projects/rocsparse/clients/unittests/unit_test_device_level1.cpp b/projects/rocsparse/clients/unittests/unit_test_device_level1.cpp index 4018002bec01..37e22fd91774 100644 --- a/projects/rocsparse/clients/unittests/unit_test_device_level1.cpp +++ b/projects/rocsparse/clients/unittests/unit_test_device_level1.cpp @@ -100,11 +100,11 @@ namespace rocsparse::gthr_device(nnz, y, x_val, x_ind, base); } - template + template __global__ void k_roti(I nnz, T* x_val, const I* x_ind, T* y, T c, T s, rocsparse_index_base base) { - rocsparse::roti_device(nnz, x_val, x_ind, y, c, s, base); + rocsparse::roti_device(nnz, x_val, x_ind, y, c, s, base); } } // namespace @@ -353,6 +353,60 @@ TEST(device_level1_roti, rotation_90deg) device_free(d_y); } +// The grid-stride variant must rotate every element even when launched with +// fewer blocks than ceil(nnz / BS). +TEST(device_level1_roti, grid_stride) +{ + using T = float; + using I = int64_t; + + const I nnz = 2 * BS + 3; + std::vector x_val(nnz); + std::vector x_ind(nnz); + std::vector y(nnz, 0); + for(I i = 0; i < nnz; ++i) + { + x_val[i] = static_cast(i + 1); + x_ind[i] = i; + } + + T* d_xval = device_from(x_val); + I* d_xind = device_from(x_ind); + T* d_y = device_from(y); + ASSERT_NE(d_xval, nullptr); + ASSERT_NE(d_xind, nullptr); + ASSERT_NE(d_y, nullptr); + + // c=2, s=1, y=0 -> x' = 2*x ; y' = -x + hipLaunchKernelGGL((k_roti), + dim3(1), + dim3(BS), + 0, + 0, + nnz, + d_xval, + d_xind, + d_y, + (T)2, + (T)1, + rocsparse_index_base_zero); + CHECK_HIP(hipDeviceSynchronize()); + + std::vector want_xval(nnz); + std::vector want_y(nnz); + for(I i = 0; i < nnz; ++i) + { + want_xval[i] = 2 * x_val[i]; + want_y[i] = -x_val[i]; + } + EXPECT_EQ(host_from(d_xval, x_val.size()), want_xval); + EXPECT_EQ(host_from(d_y, y.size()), want_y); + + device_free(d_xval); + device_free(d_xind); + device_free(d_y); +} + // Identity rotation (c=1, s=0) must leave both vectors unchanged. TEST(device_level1_roti, identity) { diff --git a/projects/rocsparse/library/src/level1/rocsparse_roti.cpp b/projects/rocsparse/library/src/level1/rocsparse_roti.cpp index 06a8fc6807a8..8e0f6257aba8 100644 --- a/projects/rocsparse/library/src/level1/rocsparse_roti.cpp +++ b/projects/rocsparse/library/src/level1/rocsparse_roti.cpp @@ -23,12 +23,13 @@ * ************************************************************************ */ #include "internal/level1/rocsparse_roti.h" +#include "rocsparse_grid.hpp" #include "rocsparse_roti.hpp" #include "roti_device.h" namespace rocsparse { - template + template ROCSPARSE_KERNEL(BLOCKSIZE) void roti_kernel(I nnz, T* __restrict__ x_val, @@ -45,7 +46,7 @@ namespace rocsparse { return; } - rocsparse::roti_device(nnz, x_val, x_ind, y, c, s, idx_base); + rocsparse::roti_device(nnz, x_val, x_ind, y, c, s, idx_base); } } @@ -94,13 +95,12 @@ rocsparse_status rocsparse::roti_template(rocsparse_handle handle, //0 hipStream_t stream = handle->stream; #define ROTI_DIM 512 - // Clamp the grid size against the device limit; the kernel uses a grid-stride - // loop so a clamped grid still covers all nnz elements (supports large nnz). - const int64_t num_blocks - = rocsparse::min((static_cast(nnz) - 1) / ROTI_DIM + 1, - static_cast(handle->properties.maxGridSize[0])); - dim3 roti_blocks(static_cast(num_blocks)); - dim3 roti_threads(ROTI_DIM); + // Clamp to both the device limit and the dispatch packet's work-item limit. + // Only a clamped grid needs the grid-stride kernel. + const int64_t num_blocks_needed = (static_cast(nnz) - 1) / ROTI_DIM + 1; + const uint32_t num_blocks = rocsparse::get_grid_size_x(handle, num_blocks_needed, ROTI_DIM); + dim3 roti_blocks(num_blocks); + dim3 roti_threads(ROTI_DIM); const bool on_host = (handle->pointer_mode == rocsparse_pointer_mode_host); if(on_host && (*c == static_cast(1) && *s == static_cast(0))) @@ -108,19 +108,30 @@ rocsparse_status rocsparse::roti_template(rocsparse_handle handle, //0 return rocsparse_status_success; } - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR((rocsparse::roti_kernel), - roti_blocks, - roti_threads, - 0, - stream, - nnz, - x_val, - x_ind, - y, - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, c), - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, s), - idx_base, - handle->pointer_mode == rocsparse_pointer_mode_host); +#define ROTI_LAUNCH(GRID_STRIDE) \ + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR((rocsparse::roti_kernel), \ + roti_blocks, \ + roti_threads, \ + 0, \ + stream, \ + nnz, \ + x_val, \ + x_ind, \ + y, \ + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, c), \ + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, s), \ + idx_base, \ + handle->pointer_mode == rocsparse_pointer_mode_host) + + if(num_blocks < num_blocks_needed) + { + ROTI_LAUNCH(true); + } + else + { + ROTI_LAUNCH(false); + } +#undef ROTI_LAUNCH #undef ROTI_DIM return rocsparse_status_success; } diff --git a/projects/rocsparse/library/src/level1/roti_device.h b/projects/rocsparse/library/src/level1/roti_device.h index e61588860221..9fd782c892aa 100644 --- a/projects/rocsparse/library/src/level1/roti_device.h +++ b/projects/rocsparse/library/src/level1/roti_device.h @@ -28,24 +28,44 @@ namespace rocsparse { - template + template ROCSPARSE_DEVICE_ILF void - roti_device(I nnz, T* x_val, const I* x_ind, T* y, T c, T s, rocsparse_index_base idx_base) + roti_element(K idx, T* x_val, const I* x_ind, T* y, T c, T s, rocsparse_index_base idx_base) { - // Cast to the (possibly 64-bit) index type I before the multiply so the - // element index does not overflow 32-bit arithmetic for large nnz. - const I stride = static_cast(hipGridDim_x) * BLOCKSIZE; - const I gid = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; + const I i = x_ind[idx] - idx_base; + + const T xr = x_val[idx]; + const T yr = y[i]; - for(I idx = gid; idx < nnz; idx += stride) + x_val[idx] = rocsparse::fma(c, xr, s * yr); + y[i] = rocsparse::fma(c, yr, -s * xr); + } + + // GRID_STRIDE must be true whenever the grid was clamped below + // ceil(nnz / BLOCKSIZE) blocks. Otherwise the grid holds at most 2^32 - 1 + // work-items, so nnz and the element index both fit in 32 bits. + template + ROCSPARSE_DEVICE_ILF void + roti_device(I nnz, T* x_val, const I* x_ind, T* y, T c, T s, rocsparse_index_base idx_base) + { + if constexpr(GRID_STRIDE) { - const I i = x_ind[idx] - idx_base; + const int64_t stride = static_cast(hipGridDim_x) * BLOCKSIZE; + const int64_t gid = static_cast(hipBlockIdx_x) * BLOCKSIZE + hipThreadIdx_x; - const T xr = x_val[idx]; - const T yr = y[i]; + for(int64_t idx = gid; idx < nnz; idx += stride) + { + rocsparse::roti_element(idx, x_val, x_ind, y, c, s, idx_base); + } + } + else + { + const uint32_t idx = hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x; - x_val[idx] = rocsparse::fma(c, xr, s * yr); - y[i] = rocsparse::fma(c, yr, -s * xr); + if(idx < static_cast(nnz)) + { + rocsparse::roti_element(idx, x_val, x_ind, y, c, s, idx_base); + } } } } From 6b526d15c3c7d0fb337a4a8d4dc612350d96e1e5 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 22:23:56 +0000 Subject: [PATCH 07/10] test(rocsparse): size the rot_extra dense vector to nnz 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. --- projects/rocsparse/clients/testings/testing_rot.cpp | 13 ++++++++----- projects/rocsparse/clients/tests/test_rot.yaml | 4 +++- 2 files changed, 11 insertions(+), 6 deletions(-) diff --git a/projects/rocsparse/clients/testings/testing_rot.cpp b/projects/rocsparse/clients/testings/testing_rot.cpp index 7d721855cfb2..58fd1f66b6aa 100644 --- a/projects/rocsparse/clients/testings/testing_rot.cpp +++ b/projects/rocsparse/clients/testings/testing_rot.cpp @@ -185,9 +185,10 @@ void testing_rot_extra(const Arguments& arg) // // This drives the 64-bit-index path of rocsparse_rot with nnz just past the // 2^32 boundary and checks that an element beyond that boundary is actually - // rotated. To stay within a single device allocation (host mirrors of the - // full arrays would need tens of GB) everything is initialized on the device - // and a single element is probed. + // rotated. At this nnz the unclamped grid exceeds the 32-bit work-item + // dispatch limit, so the launch must be clamped and grid-strided. Host + // mirrors of the full arrays would need tens of GB, so everything is + // initialized on the device and a single element is probed. using I = int64_t; using T = float; @@ -195,8 +196,10 @@ void testing_rot_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; + + // A sparse-vector descriptor requires nnz <= size. + const I size = nnz; const rocsparse_index_base base = rocsparse_index_base_zero; diff --git a/projects/rocsparse/clients/tests/test_rot.yaml b/projects/rocsparse/clients/tests/test_rot.yaml index 9e6ec85b68a4..969eba6269c5 100644 --- a/projects/rocsparse/clients/tests/test_rot.yaml +++ b/projects/rocsparse/clients/tests/test_rot.yaml @@ -32,11 +32,13 @@ Tests: indextype: *i32_i64 precision: *single_double_precisions_complex_real +# rot with a 64-bit index type and nnz just past 2^32. Needs a large-memory +# device (~64 GiB of device data); skipped elsewhere. - name: rot_extra category: nightly function: rot_extra host_memory_gb: 1 - device_memory_gb: 55 + device_memory_gb: 70 - name: rot category: quick From 35d874c47d387fe9f1f5412d70d5df089fa60df0 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 22:23:57 +0000 Subject: [PATCH 08/10] test(rocsparse): skip the 2^32-element extras on gfx1151 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. --- projects/rocsparse/clients/tests/test_axpby.yaml | 5 ++++- projects/rocsparse/clients/tests/test_rot.yaml | 3 +++ projects/rocsparse/clients/tests/test_scatter.yaml | 3 +++ projects/rocsparse/clients/tests/test_spmv_coo.yaml | 3 +++ projects/rocsparse/clients/tests/test_spmv_coo_aos.yaml | 3 +++ projects/rocsparse/clients/tests/test_spmv_ell.yaml | 3 +++ 6 files changed, 19 insertions(+), 1 deletion(-) diff --git a/projects/rocsparse/clients/tests/test_axpby.yaml b/projects/rocsparse/clients/tests/test_axpby.yaml index 6d31fbd52ad9..53e45bbd3959 100644 --- a/projects/rocsparse/clients/tests/test_axpby.yaml +++ b/projects/rocsparse/clients/tests/test_axpby.yaml @@ -131,10 +131,13 @@ Tests: # Regression test for AISPARSE-650: axpby with a 64-bit index type and nnz just # past 2^32 exercises the block-index overflow that was fixed in the axpyi -# device kernel. Needs a large-memory device (~68 GiB of device data); +# device kernel. Needs a large-memory device (~64 GiB of device data); # skipped elsewhere. - name: axpby_extra category: nightly function: axpby_extra host_memory_gb: 1 device_memory_gb: 70 + # gfx1151 reports shared system memory as device memory, which passes the + # memory guard, but hipMalloc cannot reliably provide 64 GiB there. + skip_hardware: gfx1151 diff --git a/projects/rocsparse/clients/tests/test_rot.yaml b/projects/rocsparse/clients/tests/test_rot.yaml index 969eba6269c5..ff668151f06a 100644 --- a/projects/rocsparse/clients/tests/test_rot.yaml +++ b/projects/rocsparse/clients/tests/test_rot.yaml @@ -39,6 +39,9 @@ Tests: function: rot_extra host_memory_gb: 1 device_memory_gb: 70 + # gfx1151 reports shared system memory as device memory, which passes the + # memory guard, but hipMalloc cannot reliably provide 64 GiB there. + skip_hardware: gfx1151 - name: rot category: quick diff --git a/projects/rocsparse/clients/tests/test_scatter.yaml b/projects/rocsparse/clients/tests/test_scatter.yaml index 34f276daec45..b50f14c2f656 100644 --- a/projects/rocsparse/clients/tests/test_scatter.yaml +++ b/projects/rocsparse/clients/tests/test_scatter.yaml @@ -155,3 +155,6 @@ Tests: function: scatter_extra host_memory_gb: 1 device_memory_gb: 70 + # gfx1151 reports shared system memory as device memory, which passes the + # memory guard, but hipMalloc cannot reliably provide 64 GiB there. + skip_hardware: gfx1151 diff --git a/projects/rocsparse/clients/tests/test_spmv_coo.yaml b/projects/rocsparse/clients/tests/test_spmv_coo.yaml index 7b46832bd359..1f9061184e88 100644 --- a/projects/rocsparse/clients/tests/test_spmv_coo.yaml +++ b/projects/rocsparse/clients/tests/test_spmv_coo.yaml @@ -452,3 +452,6 @@ Tests: function: spmv_coo_extra host_memory_gb: 1 device_memory_gb: 90 + # gfx1151 reports shared system memory as device memory, which passes the + # memory guard, but hipMalloc cannot reliably provide 64 GiB there. + skip_hardware: gfx1151 diff --git a/projects/rocsparse/clients/tests/test_spmv_coo_aos.yaml b/projects/rocsparse/clients/tests/test_spmv_coo_aos.yaml index 8d69881ad8c4..72307e1baa2b 100644 --- a/projects/rocsparse/clients/tests/test_spmv_coo_aos.yaml +++ b/projects/rocsparse/clients/tests/test_spmv_coo_aos.yaml @@ -413,3 +413,6 @@ Tests: function: spmv_coo_aos_extra host_memory_gb: 1 device_memory_gb: 90 + # gfx1151 reports shared system memory as device memory, which passes the + # memory guard, but hipMalloc cannot reliably provide 64 GiB there. + skip_hardware: gfx1151 diff --git a/projects/rocsparse/clients/tests/test_spmv_ell.yaml b/projects/rocsparse/clients/tests/test_spmv_ell.yaml index 9feacf62053a..4f8a28bf14f1 100644 --- a/projects/rocsparse/clients/tests/test_spmv_ell.yaml +++ b/projects/rocsparse/clients/tests/test_spmv_ell.yaml @@ -362,3 +362,6 @@ Tests: function: spmv_ell_extra host_memory_gb: 1 device_memory_gb: 72 + # gfx1151 reports shared system memory as device memory, which passes the + # memory guard, but hipMalloc cannot reliably provide 64 GiB there. + skip_hardware: gfx1151 From 2ca6a28563decd213a0db2865709b03e0d511469 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 22:23:57 +0000 Subject: [PATCH 09/10] fix(rocsparse): clear the HIP error after a failed client allocation 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. --- .../clients/include/rocsparse_allocator.hpp | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/projects/rocsparse/clients/include/rocsparse_allocator.hpp b/projects/rocsparse/clients/include/rocsparse_allocator.hpp index 24f123d22d1a..c242cf6c5006 100644 --- a/projects/rocsparse/clients/include/rocsparse_allocator.hpp +++ b/projects/rocsparse/clients/include/rocsparse_allocator.hpp @@ -269,6 +269,15 @@ struct rocsparse_allocator } #endif + static void report_allocation_failure(size_t nbytes) + { + fprintf(stderr, "Error allocating %zu bytes (%zu GB)\n", nbytes, nbytes >> 30); + + // Consume the error so that it is not reported by the next hipGetLastError + // call, e.g. the one in rocsparse_create_handle of the next test. + static_cast(hipGetLastError()); + } + public: static T* malloc(size_t size) { @@ -284,7 +293,7 @@ struct rocsparse_allocator { if(rocsparse_hipHostMalloc(&d, nbytes) != hipSuccess) { - fprintf(stderr, "Error allocating %'zu bytes (%zu GB)\n", nbytes, nbytes >> 30); + report_allocation_failure(nbytes); d = nullptr; throw std::bad_alloc(); } @@ -294,7 +303,7 @@ struct rocsparse_allocator { if(rocsparse_hipMalloc(&d, nbytes) != hipSuccess) { - fprintf(stderr, "Error allocating %'zu bytes (%zu GB)\n", nbytes, nbytes >> 30); + report_allocation_failure(nbytes); d = nullptr; throw std::bad_alloc(); } @@ -305,7 +314,7 @@ struct rocsparse_allocator if(rocsparse_hipMallocManaged(&d, nbytes) != hipSuccess) { - fprintf(stderr, "Error allocating %'zu bytes (%zu GB)\n", nbytes, nbytes >> 30); + report_allocation_failure(nbytes); d = nullptr; throw std::bad_alloc(); } From 8ab10f96a2ef8c2afdc9d9ac5fac63de9eb8c597 Mon Sep 17 00:00:00 2001 From: kliegeois Date: Tue, 29 Sep 2026 21:50:47 +0000 Subject: [PATCH 10/10] perf(rocsparse): straight-line fast paths for COO AoS SpMV and scale_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. --- .../library/src/common/rocsparse_common.cpp | 46 ++++++++++----- .../library/src/level2/coomv_device.h | 26 +++++--- .../library/src/level2/ellmv_device.h | 13 +++- .../src/level2/rocsparse_coomv_aos_impl.cpp | 47 ++++++++++----- .../src/level2/rocsparse_ellmv_impl.cpp | 59 ++++++++++++------- 5 files changed, 132 insertions(+), 59 deletions(-) diff --git a/projects/rocsparse/library/src/common/rocsparse_common.cpp b/projects/rocsparse/library/src/common/rocsparse_common.cpp index 8d1aed8cf1a1..5c1b282e4e53 100644 --- a/projects/rocsparse/library/src/common/rocsparse_common.cpp +++ b/projects/rocsparse/library/src/common/rocsparse_common.cpp @@ -49,7 +49,7 @@ namespace rocsparse array[lid + ld * wid] = value; } - template + template ROCSPARSE_DEVICE_ILF void scale_device(I length, T scalar, A* __restrict__ array) { const int64_t stride = static_cast(hipGridDim_x) * BLOCKSIZE; @@ -66,6 +66,11 @@ namespace rocsparse { array[gid] *= scalar; } + + if constexpr(!GRID_STRIDE) + { + break; + } } } @@ -157,7 +162,7 @@ namespace rocsparse rocsparse::valset_2d_device(m, n, ld, value, array, order); } - template + template ROCSPARSE_KERNEL(BLOCKSIZE) void scale_kernel(I length, ROCSPARSE_DEVICE_HOST_SCALAR_PARAMS(T, scalar), @@ -167,7 +172,7 @@ namespace rocsparse ROCSPARSE_DEVICE_HOST_SCALAR_GET(scalar); if(scalar != static_cast(1)) { - rocsparse::scale_device(length, scalar, array); + rocsparse::scale_device(length, scalar, array); } } @@ -263,17 +268,30 @@ rocsparse_status rocsparse::scale_array(rocsparse_handle handle, } else if((on_host && *scalar_device_host != static_cast(1)) || on_host == false) { - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( - (rocsparse::scale_kernel<256>), - dim3(rocsparse::get_grid_size_x( - handle, (static_cast(length) - 1) / 256 + 1, 256)), - dim3(256), - 0, - handle->stream, - length, - ROCSPARSE_SCALAR_HOST_DEVICE_ARGUMENT(pointer_mode, scalar_device_host), - array, - on_host); +#define LAUNCH_SCALE(GRID_STRIDE) \ + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( \ + (rocsparse::scale_kernel<256, GRID_STRIDE>), \ + dim3(grid_size), \ + dim3(256), \ + 0, \ + handle->stream, \ + length, \ + ROCSPARSE_SCALAR_HOST_DEVICE_ARGUMENT(pointer_mode, scalar_device_host), \ + array, \ + on_host) + + // Only a clamped grid needs the grid-stride loop. + const int64_t num_blocks = (static_cast(length) - 1) / 256 + 1; + const uint32_t grid_size = rocsparse::get_grid_size_x(handle, num_blocks, 256); + if(grid_size < num_blocks) + { + LAUNCH_SCALE(true); + } + else + { + LAUNCH_SCALE(false); + } +#undef LAUNCH_SCALE } } return rocsparse_status_success; diff --git a/projects/rocsparse/library/src/level2/coomv_device.h b/projects/rocsparse/library/src/level2/coomv_device.h index 016be32d0cc3..6775ca37f3a2 100644 --- a/projects/rocsparse/library/src/level2/coomv_device.h +++ b/projects/rocsparse/library/src/level2/coomv_device.h @@ -678,10 +678,12 @@ namespace rocsparse } } - // Each tile covers LOOPS * BLOCKSIZE entries. Blocks stride over the tiles, - // so a grid clamped below nnz / (LOOPS * BLOCKSIZE) still covers every entry. + // Each tile covers LOOPS * BLOCKSIZE entries. With GRID_STRIDE, blocks stride + // over the tiles, so a grid clamped below nnz / (LOOPS * BLOCKSIZE) still + // covers every entry. Without it, the grid must have one block per tile. template (LOOPS) * BLOCKSIZE; - - for(int64_t tile = hipBlockIdx_x; tile * tile_size < nnz; tile += hipGridDim_x) + if constexpr(!GRID_STRIDE) { rocsparse::coomvn_aos_atomic_loops_tile_device( - tile, nnz, m, alpha, coo_ind, coo_val, x, y, idx_base); + hipBlockIdx_x, nnz, m, alpha, coo_ind, coo_val, x, y, idx_base); + } + else + { + const int64_t tile_size = static_cast(LOOPS) * BLOCKSIZE; - // The next tile overwrites the shared buffers this tile still reads. - __syncthreads(); + for(int64_t tile = hipBlockIdx_x; tile * tile_size < nnz; tile += hipGridDim_x) + { + rocsparse::coomvn_aos_atomic_loops_tile_device( + tile, nnz, m, alpha, coo_ind, coo_val, x, y, idx_base); + + // The next tile overwrites the shared buffers this tile still reads. + __syncthreads(); + } } } diff --git a/projects/rocsparse/library/src/level2/ellmv_device.h b/projects/rocsparse/library/src/level2/ellmv_device.h index 64b566842959..b53f6015dd8e 100644 --- a/projects/rocsparse/library/src/level2/ellmv_device.h +++ b/projects/rocsparse/library/src/level2/ellmv_device.h @@ -76,7 +76,13 @@ namespace rocsparse } // ELL SpMV for general, (conjugate) transposed matrices - template + template ROCSPARSE_DEVICE_ILF void ellmvt_device(rocsparse_operation trans, I m, I n, @@ -116,6 +122,11 @@ namespace rocsparse break; } } + + if constexpr(!GRID_STRIDE) + { + break; + } } } } diff --git a/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp b/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp index d7cda03df604..e9788c9d82fe 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp @@ -88,6 +88,7 @@ namespace rocsparse template ( + rocsparse::coomvn_aos_atomic_loops_device( nnz, m, alpha, coo_ind, coo_val, x, y, idx_base); } } @@ -160,21 +161,35 @@ namespace rocsparse { case rocsparse_operation_none: { - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( - (rocsparse::coomvn_aos_atomic_loops<256, 1>), - dim3(rocsparse::get_grid_size_x(handle, (nnz - 1) / 256 + 1, 256)), - dim3(256), - 0, - stream, - nnz, - m, - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), - coo_ind, - coo_val, - x, - y, - descr->base, - handle->pointer_mode == rocsparse_pointer_mode_host); +#define LAUNCH_COOMVN_AOS_ATOMIC(GRID_STRIDE) \ + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( \ + (rocsparse::coomvn_aos_atomic_loops<256, 1, GRID_STRIDE>), \ + dim3(grid_size), \ + dim3(256), \ + 0, \ + stream, \ + nnz, \ + m, \ + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), \ + coo_ind, \ + coo_val, \ + x, \ + y, \ + descr->base, \ + handle->pointer_mode == rocsparse_pointer_mode_host) + + // Only a clamped grid needs the tile loop. + const int64_t num_blocks = (nnz - 1) / 256 + 1; + const uint32_t grid_size = rocsparse::get_grid_size_x(handle, num_blocks, 256); + if(grid_size < num_blocks) + { + LAUNCH_COOMVN_AOS_ATOMIC(true); + } + else + { + LAUNCH_COOMVN_AOS_ATOMIC(false); + } +#undef LAUNCH_COOMVN_AOS_ATOMIC break; } case rocsparse_operation_transpose: diff --git a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp index 05c3a9fe93e9..6e544f978761 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp @@ -56,7 +56,13 @@ namespace rocsparse } } - template + template ROCSPARSE_KERNEL(BLOCKSIZE) void ellmvt_kernel(rocsparse_operation trans, I m, @@ -73,7 +79,7 @@ namespace rocsparse ROCSPARSE_DEVICE_HOST_SCALAR_GET(alpha); if(alpha != 0) { - rocsparse::ellmvt_device( + rocsparse::ellmvt_device( trans, m, n, ell_width, alpha, ell_col_ind, ell_val, x, y, idx_base); } } @@ -163,24 +169,37 @@ namespace rocsparse // Scale y with beta RETURN_IF_ROCSPARSE_ERROR(rocsparse::scale_array(handle, n, beta_device_host, y)); - RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( - (rocsparse::ellmvt_kernel), - dim3(rocsparse::get_grid_size_x( - handle, (static_cast(m) - 1) / ELLMVT_DIM + 1, ELLMVT_DIM)), - dim3(ELLMVT_DIM), - 0, - stream, - trans, - m, - n, - ell_width, - ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), - ell_col_ind, - ell_val, - x, - y, - descr->base, - handle->pointer_mode == rocsparse_pointer_mode_host); +#define LAUNCH_ELLMVT(GRID_STRIDE) \ + RETURN_IF_HIPLAUNCHKERNELGGL_ERROR( \ + (rocsparse::ellmvt_kernel), \ + dim3(grid_size), \ + dim3(ELLMVT_DIM), \ + 0, \ + stream, \ + trans, \ + m, \ + n, \ + ell_width, \ + ROCSPARSE_DEVICE_HOST_SCALAR_ARGS(handle, alpha_device_host), \ + ell_col_ind, \ + ell_val, \ + x, \ + y, \ + descr->base, \ + handle->pointer_mode == rocsparse_pointer_mode_host) + + // Only a clamped grid needs the grid-stride loop. + const int64_t num_blocks = (static_cast(m) - 1) / ELLMVT_DIM + 1; + const uint32_t grid_size = rocsparse::get_grid_size_x(handle, num_blocks, ELLMVT_DIM); + if(grid_size < num_blocks) + { + LAUNCH_ELLMVT(true); + } + else + { + LAUNCH_ELLMVT(false); + } +#undef LAUNCH_ELLMVT #undef ELLMVT_DIM }