diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index 35a18b104d3b..827c7c7d7cfb 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -9,9 +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_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_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_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. * Fixed an integer overflow in the binary search that maps nonzeros to rows in the nnz-split algorithms of `rocsparse_spmv` (CSR) and `rocsparse_spmm` (CSR and CSC). With 32-bit indices and matrices with more than `2^30` rows, the search midpoint `(left + right) / 2` could overflow, causing out-of-bounds reads or incorrect row assignments. The midpoint is now computed as `left + (right - left) / 2`. ## rocSPARSE 5.1.0 for ROCm 10.1 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(); } diff --git a/projects/rocsparse/clients/testings/testing_axpby.cpp b/projects/rocsparse/clients/testings/testing_axpby.cpp index e5aa35b3bcaf..7f2d3220fde5 100644 --- a/projects/rocsparse/clients/testings/testing_axpby.cpp +++ b/projects/rocsparse/clients/testings/testing_axpby.cpp @@ -175,12 +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. 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. + // 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; @@ -188,8 +189,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; @@ -215,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()); @@ -225,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/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_axpby.yaml b/projects/rocsparse/clients/tests/test_axpby.yaml index ce754d22a3fe..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 (~52 GB of index/value 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: 55 + 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 9e6ec85b68a4..ff668151f06a 100644 --- a/projects/rocsparse/clients/tests/test_rot.yaml +++ b/projects/rocsparse/clients/tests/test_rot.yaml @@ -32,11 +32,16 @@ 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 + # 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 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/common/rocsparse_common.cpp b/projects/rocsparse/library/src/common/rocsparse_common.cpp index 368078acd3f4..5c1b282e4e53 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 @@ -48,23 +49,28 @@ namespace rocsparse array[lid + ld * wid] = value; } - template + 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; + if constexpr(!GRID_STRIDE) + { + break; + } } } @@ -156,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), @@ -166,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); } } @@ -262,16 +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((length - 1) / 256 + 1), - 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/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..8056aef7426b 100644 --- a/projects/rocsparse/library/src/level1/rocsparse_axpyi.cpp +++ b/projects/rocsparse/library/src/level1/rocsparse_axpyi.cpp @@ -25,6 +25,7 @@ #include "internal/level1/rocsparse_axpyi.h" #include "axpyi_device.h" #include "rocsparse_axpyi.hpp" +#include "rocsparse_grid.hpp" #include "rocsparse_utility.hpp" namespace rocsparse @@ -98,13 +99,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, (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_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/rocsparse_sctr_impl.cpp b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp index dce0d3d76fa4..e730c0ecc4df 100644 --- a/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp +++ b/projects/rocsparse/library/src/level1/rocsparse_sctr_impl.cpp @@ -26,6 +26,7 @@ #include "rocsparse_sctr.hpp" #include "rocsparse_control.hpp" +#include "rocsparse_grid.hpp" #include "rocsparse_utility.hpp" #include "sctr_device.h" @@ -68,12 +69,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, (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/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); + } } } } 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/coomv_device.h b/projects/rocsparse/library/src/level2/coomv_device.h index e58f9e2ec01a..6775ca37f3a2 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,46 @@ namespace rocsparse } } + // 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 + 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) + { + if constexpr(!GRID_STRIDE) + { + rocsparse::coomvn_aos_atomic_loops_tile_device( + hipBlockIdx_x, nnz, m, alpha, coo_ind, coo_val, x, y, idx_base); + } + else + { + 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 +757,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/ellmv_device.h b/projects/rocsparse/library/src/level2/ellmv_device.h index 76b0d2d94c82..b53f6015dd8e 100644 --- a/projects/rocsparse/library/src/level2/ellmv_device.h +++ b/projects/rocsparse/library/src/level2/ellmv_device.h @@ -41,45 +41,48 @@ 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, static_cast(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 - template + template ROCSPARSE_DEVICE_ILF void ellmvt_device(rocsparse_operation trans, I m, I n, @@ -91,34 +94,36 @@ 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, static_cast(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 + + 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 941046b6d1d2..e9788c9d82fe 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_coomv_aos_impl.cpp @@ -25,6 +25,7 @@ #include "rocsparse_common.h" #include "rocsparse_control.hpp" #include "rocsparse_coomv_aos.hpp" +#include "rocsparse_grid.hpp" #include "rocsparse_utility.hpp" #include "coomv_device.h" @@ -87,6 +88,7 @@ namespace rocsparse template ( + rocsparse::coomvn_aos_atomic_loops_device( nnz, m, alpha, coo_ind, coo_val, x, y, idx_base); } } @@ -159,21 +161,35 @@ 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); +#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: @@ -181,7 +197,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, @@ -340,7 +356,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, diff --git a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp index 20826b3718d6..6e544f978761 100644 --- a/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp +++ b/projects/rocsparse/library/src/level2/rocsparse_ellmv_impl.cpp @@ -28,6 +28,7 @@ #include "ellmv_device.h" #include "rocsparse_control.hpp" +#include "rocsparse_grid.hpp" #include "rocsparse_utility.hpp" namespace rocsparse @@ -55,7 +56,13 @@ namespace rocsparse } } - template + template ROCSPARSE_KERNEL(BLOCKSIZE) void ellmvt_kernel(rocsparse_operation trans, I m, @@ -72,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); } } @@ -99,23 +106,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, (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. @@ -161,23 +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((m - 1) / ELLMVT_DIM + 1), - 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 }