diff --git a/projects/rocsparse/CHANGELOG.md b/projects/rocsparse/CHANGELOG.md index 37579c3dbae8..d283e4a03ca1 100644 --- a/projects/rocsparse/CHANGELOG.md +++ b/projects/rocsparse/CHANGELOG.md @@ -36,6 +36,7 @@ Documentation for rocSPARSE is available at * 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`. * Fixed an invalid kernel launch in legacy `rocsparse_Xbsr2csr` and the 64-bit BSR-to-CSR conversion, where the 1024-thread kernels overflowed the 32-bit work-item grid limit at about 4.2 million block rows because the grid was not clamped. +* Fixed an integer overflow in the SELL SpMV kernels used by `rocsparse_v2_spmv`. With 32-bit indices, the row index of each thread was computed in 32-bit arithmetic, so for matrices with close to `2^31` rows the threads past the last row wrapped to a negative row index that passed the `row < m` check. Those threads could read past the end of the slice offsets, and the non-transposed kernel could write `y` at a negative index. The row index is now computed in 64-bit arithmetic. The transposed kernel for slice sizes up to 128 also no longer reads `x` for threads past the last row. * Fixed an overflow issue in the COO atomic SpMV kernels used by `rocsparse_spmv` and `rocsparse_v2_spmv` for COO and COO AoS matrices. When using 64-bit indices and `nnz` >= `2^32`, a 32-bit element-ID calculation could overflow and prevent some nonzeros from being accumulated. The kernels now compute element IDs in 64-bit arithmetic. ## rocSPARSE 5.1.0 for ROCm 10.1 diff --git a/projects/rocsparse/library/src/level2/sellmv_device.h b/projects/rocsparse/library/src/level2/sellmv_device.h index 714ae8791031..b31e7a152123 100644 --- a/projects/rocsparse/library/src/level2/sellmv_device.h +++ b/projects/rocsparse/library/src/level2/sellmv_device.h @@ -1,6 +1,6 @@ /*! \file */ /* ************************************************************************ - * Copyright (C) 2025 Advanced Micro Devices, Inc. All rights Reserved. + * Copyright (C) 2025-2026 Advanced Micro Devices, Inc. All rights Reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -60,7 +60,7 @@ namespace rocsparse const uint32_t sliceid = hipGridDim_x * hipBlockIdx_y + hipBlockIdx_x; - const J row = sell_slice_size * sliceid + tidx; + const int64_t row = static_cast(sell_slice_size) * sliceid + tidx; const I start = (row < m) ? sell_slice_offsets[sliceid] - idx_base : 0; const I end = (row < m) ? sell_slice_offsets[sliceid + 1] - idx_base : 0; @@ -133,7 +133,7 @@ namespace rocsparse { const J local_row = (BLOCKSIZE * p + tid); - const J row = sell_slice_size * sliceid + (BLOCKSIZE * p + tid); + const int64_t row = static_cast(sell_slice_size) * sliceid + local_row; const bool row_in_range = (row < m && local_row < sell_slice_size); @@ -195,12 +195,12 @@ namespace rocsparse const uint32_t sliceid = hipGridDim_x * hipBlockIdx_y + hipBlockIdx_x; - const J row = sell_slice_size * sliceid + tidx; + const int64_t row = static_cast(sell_slice_size) * sliceid + tidx; const I start = (row < m) ? sell_slice_offsets[sliceid] - idx_base : 0; const I end = (row < m) ? sell_slice_offsets[sliceid + 1] - idx_base : 0; - T row_val = alpha * x[row]; + const T row_val = (row < m) ? alpha * x[row] : static_cast(0); for(I j = start + idx; j < end; j += (sell_slice_size * THREADS_PER_ROW)) { @@ -251,7 +251,7 @@ namespace rocsparse { const J local_row = (BLOCKSIZE * p + tid); - const J row = sell_slice_size * sliceid + (BLOCKSIZE * p + tid); + const int64_t row = static_cast(sell_slice_size) * sliceid + local_row; const bool row_in_range = (row < m && local_row < sell_slice_size);