kliegeois · Oct 8, 2026 · Sep 24, 2026 · Sep 29, 2026 · Sep 29, 2026 · Sep 29, 2026
@@ -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.
 
 ## (Unreleased) rocSPARSE 5.1.0
 

@@ -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<void>(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();
             }

@@ -175,21 +175,24 @@ 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;
 
     static constexpr int64_t two_pow_32 = static_cast<int64_t>(1) << 32;
 
     // nnz just beyond 2^32 so at least one block has a block index whose
     // (blockIdx * BLOCKSIZE) product overflows 32-bit arithmetic.
-    const I nnz  = two_pow_32 + 512;
-    const I size = 2;
+    const 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,20 +218,38 @@ void testing_axpby_extra(const Arguments& arg)
     CHECK_HIP_ERROR(
         hipMemcpy(static_cast<T*>(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<T>(1);
+    const T tail_in  = static_cast<T>(5);
+    CHECK_HIP_ERROR(
+        hipMemcpy(static_cast<T*>(dy) + probe_ind, &y_in, sizeof(T), hipMemcpyHostToDevice));
+    CHECK_HIP_ERROR(
+        hipMemcpy(static_cast<T*>(dy) + tail_idx, &tail_in, sizeof(T), hipMemcpyHostToDevice));
+
     const T halpha = static_cast<T>(2);
-    const T hbeta  = static_cast<T>(1);
+    const T hbeta  = static_cast<T>(3);
 
     rocsparse_local_spvec x(size, nnz, dx_ind, dx_val, get_indextype<I>(), base, get_datatype<T>());
     rocsparse_local_dnvec y(size, dy, get_datatype<T>());
 
     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<T>(0);
-    CHECK_HIP_ERROR(hipMemcpy(&y_out, static_cast<T*>(dy) + 1, sizeof(T), hipMemcpyDeviceToHost));
+    CHECK_HIP_ERROR(
+        hipMemcpy(&y_out, static_cast<T*>(dy) + probe_ind, sizeof(T), hipMemcpyDeviceToHost));
+
+    unit_check_scalar<T>(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<T>(0);
+    CHECK_HIP_ERROR(
+        hipMemcpy(&tail_out, static_cast<T*>(dy) + tail_idx, sizeof(T), hipMemcpyDeviceToHost));
 
-    unit_check_scalar<T>(static_cast<T>(2), y_out);
+    unit_check_scalar<T>(hbeta * tail_in, tail_out);
 }
@@ -185,18 +185,21 @@ 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;
 
     static constexpr int64_t two_pow_32 = static_cast<int64_t>(1) << 32;
 
     // nnz just beyond 2^32 so at least one block has a block index whose
     // (blockIdx * BLOCKSIZE) product overflows 32-bit arithmetic.
-    const I nnz  = two_pow_32 + 512;
-    const I size = 2;
+    const 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;
 

@@ -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
@@ -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

@@ -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
@@ -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
@@ -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
@@ -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
@@ -100,11 +100,11 @@ namespace
         rocsparse::gthr_device<BS, I, T>(nnz, y, x_val, x_ind, base);
     }
 
-    template <typename T, typename I>
+    template <typename T, typename I, bool GRID_STRIDE = false>
     __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<BS, I, T>(nnz, x_val, x_ind, y, c, s, base);
+        rocsparse::roti_device<BS, GRID_STRIDE, I, T>(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<T> x_val(nnz);
+    std::vector<I> x_ind(nnz);
+    std::vector<T> y(nnz, 0);
+    for(I i = 0; i < nnz; ++i)
+    {
+        x_val[i] = static_cast<T>(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<T, I, true>),
+                       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<T> want_xval(nnz);
+    std::vector<T> 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)
 {