testintel: fix severe issue in sellmvt_device (created by kliegeoi (@kliegeois))

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

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<int64_t>(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<int64_t>(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<int64_t>(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<T>(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<int64_t>(sell_slice_size) * sliceid + local_row;
 
             const bool row_in_range = (row < m && local_row < sell_slice_size);