AISPARSE-692/693/694: gtsv offsets, zero-batch guard and BSR row-map bounds checks

projects/rocsparse/CHANGELOG.md

@@ -33,6 +33,10 @@ Documentation for rocSPARSE is available at
 * Fixed `rocsparse_spmm` with the segmented COO, atomic COO, segmented-atomic COO, and row-split CSR algorithms, which failed with `hipErrorInvalidConfiguration` for batch counts exceeding 65535 because the batch dimension of the kernel launch grid exceeded the maximum supported grid dimension.
 * Fixed an out-of-bounds write caused by an incorrectly sized temporary buffer in the CSR-to-CSC conversion performed by `rocsparse_sparse_to_sparse` when the row pointer and column index types differ (for example,  64-bit row pointers with 32-bit column indices).
 * Fixed an issue with `rocsparse_spmm` when using the nnz-split algorithm with the CSR or CSC format. The operation produced incorrect results because the segmented-block-reduction helper had shared-memory pointer parameters marked `__restrict__`, while threads in the block must read values written by other threads. The `__restrict__` attribute has now been removed.
+* Fixed an out-of-bounds read of the row map in the block-row incomplete factorization kernels used by `rocsparse_spic0` and `rocsparse_spilu0`. Eight of the nine BSR kernels indexed the map with the raw block index and never compared it against the block row count, so a launch grid rounded up past that count read outside the map allocation. All nine now bounds check, which is what the CSR kernels in the same family already did.
+* Fixed `rocsparse_spic0` and `rocsparse_spilu0` silently truncating the launch grid for matrices with more block rows than the hardware grid.x maximum of 2147483647. The block row count is taken from the sparse matrix descriptor as a 64-bit value and was narrowed into the 32-bit grid extent without a clamp. The general BSR kernels now clamp the extent and recover the remaining block rows inside the kernel.
+* Fixed a signed integer overflow in the `rocsparse_csritilu0` async-inplace preprocessing, which computed four launch grid extents as `(target_size + BLOCKSIZE - 1) / BLOCKSIZE`. The addition overflowed 32-bit `rocsparse_int` before the division for target sizes above 2147482625, which is undefined behaviour. The extents are now computed as `(target_size - 1) / BLOCKSIZE + 1`, the form used everywhere else in the directory, which cannot overflow for any representable positive size.
+* Fixed an integer overflow in `rocsparse_Xgtsv_interleaved_batch`, which produced incorrect results when `batch_count * m` or `batch_stride * m` exceeded `INT32_MAX`. The interleaved layout makes every array `batch_stride * m` elements long, but the solver kernels computed the row offsets into those arrays in 32-bit arithmetic. The offsets are now computed in 64-bit arithmetic.
 
 ## rocSPARSE 5.0.0 for ROCm 10.0
 

projects/rocsparse/clients/tests/test_gtsv_interleaved_batch.yaml

@@ -57,6 +57,25 @@ Tests:
   gtsv_interleaved_alg: [rocsparse_gtsv_interleaved_alg_thomas]
   matrix: [rocsparse_matrix_random]
 
+# The interleaved kernels address the batch_count-strided workspaces and the
+# batch_stride-strided user arrays with two different multipliers of the same row
+# index. batch_stride is deliberately well above batch_count here so that a
+# widened offset that picks up the wrong multiplier, or that truncates on the way
+# back into a 32-bit local, lands outside the batch it belongs to. M is past the
+# m > 2 and m > 3 special cases all three algorithms carry, so every one of the
+# forward, backward and pivoting offsets is exercised.
+- name: gtsv_interleaved_batch
+  category: quick
+  function: gtsv_interleaved_batch
+  precision: *single_double_precisions_complex_real
+  M: [513]
+  batch_count: [1000]
+  batch_stride: [2048]
+  gtsv_interleaved_alg: [rocsparse_gtsv_interleaved_alg_thomas, rocsparse_gtsv_interleaved_alg_lu, rocsparse_gtsv_interleaved_alg_qr]
+  matrix: [rocsparse_matrix_random]
+  host_memory_gb: 1
+  device_memory_gb: 1
+
 # Pre-checkin
 - name: gtsv_interleaved_batch_bad_arg
   category: pre_checkin

projects/rocsparse/clients/unittests/CMakeLists.txt

@@ -48,6 +48,7 @@ set(ROCSPARSE_UNIT_TEST_SOURCES
   unit_test_indextype_utils.cpp
   unit_test_datatype_utils.cpp
   unit_test_csrmv_adaptive_analysis.cpp
+  unit_test_internal_grid_x.cpp
 )
 
 # Host-pure library translation units under test. These are compiled directly

projects/rocsparse/clients/unittests/unit_test_internal_grid_x.cpp

@@ -0,0 +1,160 @@
+/*! \file */
+/* ************************************************************************
+ * Copyright (C) 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ *
+ * ************************************************************************ */
+
+//
+// Host unit tests for the contiguous chunk partition that the incomplete
+// factorizations use to recover the work above the grid.x clamp.
+//
+// These live in the host-only rocsparse-unit-test binary rather than in
+// rocsparse-unit-test-device, and they declare no memory requirement, so
+// nothing can drop them at instantiation time. The whole row count range the
+// clamp exists for is far past what any available device can hold, so the
+// properties the kernels depend on are checked here arithmetically instead:
+//
+//   - the chunks are a partition: every item in [0, count) is owned by exactly
+//     one block;
+//   - the chunks are monotone in the block id, which is the property that keeps
+//     the done-flag spin in the factorization kernels from deadlocking;
+//   - at any unclamped extent the partition degenerates to one item per block,
+//     so the recovery loop is a no-op for every reachable problem size.
+//
+#include "rocsparse_grid_x.hpp"
+
+#include <cstdint>
+#include <gtest/gtest.h>
+#include <vector>
+
+TEST(grid_x, grid_x_chunk_is_a_partition)
+{
+    for(int64_t count : {int64_t(1),
+                         int64_t(2),
+                         int64_t(5),
+                         int64_t(16),
+                         int64_t(97),
+                         int64_t(1024),
+                         int64_t(4097)})
+    {
+        for(uint32_t grid_size = 1; grid_size <= static_cast<uint32_t>(count); ++grid_size)
+        {
+            std::vector<int> owners(static_cast<size_t>(count), 0);
+
+            for(uint32_t block_id = 0; block_id < grid_size; ++block_id)
+            {
+                int64_t first, last;
+                rocsparse::grid_x_chunk(count, grid_size, block_id, first, last);
+
+                ASSERT_LE(first, last) << "count " << count << " grid_size " << grid_size;
+
+                for(int64_t i = first; i < last; ++i)
+                {
+                    ASSERT_GE(i, 0);
+                    ASSERT_LT(i, count);
+                    ++owners[static_cast<size_t>(i)];
+                }
+            }
+
+            for(int64_t i = 0; i < count; ++i)
+            {
+                ASSERT_EQ(owners[static_cast<size_t>(i)], 1)
+                    << "item " << i << " of " << count << " with grid_size " << grid_size;
+            }
+        }
+    }
+}
+
+TEST(grid_x, grid_x_chunk_is_monotone_in_the_block_id)
+{
+    // A block only ever waits on items belonging to itself or to a lower
+    // numbered block. A grid-stride loop would not have this property and would
+    // deadlock the done-flag spin once the grid exceeds what stays resident.
+    for(int64_t count : {int64_t(7), int64_t(100), int64_t(4097)})
+    {
+        for(uint32_t grid_size = 1; grid_size <= static_cast<uint32_t>(count); ++grid_size)
+        {
+            int64_t previous_last = 0;
+
+            for(uint32_t block_id = 0; block_id < grid_size; ++block_id)
+            {
+                int64_t first, last;
+                rocsparse::grid_x_chunk(count, grid_size, block_id, first, last);
+
+                if(first < last)
+                {
+                    ASSERT_GE(first, previous_last)
+                        << "count " << count << " grid_size " << grid_size << " block " << block_id;
+                    previous_last = last;
+                }
+            }
+
+            ASSERT_EQ(previous_last, count);
+        }
+    }
+}
+
+TEST(grid_x, grid_x_chunk_degenerates_to_one_item_per_block_when_unclamped)
+{
+    // Every problem size that fits on hardware today launches grid_size ==
+    // count, and this is what makes the recovery loop a no-op there.
+    for(int64_t count : {int64_t(1), int64_t(3), int64_t(64), int64_t(65537), int64_t(100000)})
+    {
+        const uint32_t grid_size = static_cast<uint32_t>(count);
+
+        for(uint32_t block_id = 0; block_id < grid_size; ++block_id)
+        {
+            int64_t first, last;
+            rocsparse::grid_x_chunk(count, grid_size, block_id, first, last);
+
+            ASSERT_EQ(first, static_cast<int64_t>(block_id));
+            ASSERT_EQ(last, static_cast<int64_t>(block_id) + 1);
+        }
+    }
+}
+
+TEST(grid_x, grid_x_chunk_covers_a_count_above_the_clamp)
+{
+    // The case the clamp exists for: more items than the grid can express.
+    // grid_size is what rocsparse::get_grid_size_x returns at 256 threads per
+    // block, (2^32 - 1) / 256. Check the leading blocks and the final block
+    // without walking every item.
+    const int64_t  count     = 5000000000LL;
+    const uint32_t grid_size = 16777215u;
+
+    int64_t first, last;
+
+    rocsparse::grid_x_chunk(count, grid_size, 0, first, last);
+    EXPECT_EQ(first, 0);
+    EXPECT_EQ(last, 299);
+
+    int64_t total = 0;
+    for(uint32_t block_id = 0; block_id < 1000; ++block_id)
+    {
+        rocsparse::grid_x_chunk(count, grid_size, block_id, first, last);
+        total += last - first;
+    }
+    EXPECT_EQ(total, 299000);
+
+    rocsparse::grid_x_chunk(count, grid_size, grid_size - 1, first, last);
+    EXPECT_EQ(last, count);
+    EXPECT_LE(first, last);
+}

projects/rocsparse/library/src/include/rocsparse_grid_x.hpp

@@ -0,0 +1,78 @@
+/*! \file */
+/* ************************************************************************
+ * Copyright (C) 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ *
+ * ************************************************************************ */
+
+#pragma once
+
+#include <stdint.h>
+
+//
+// Kept in its own header, and deliberately free of any rocsparse or HIP
+// include, so that the host-only rocsparse-unit-test binary can compile it as
+// plain C++ without pulling in the HIP toolchain. The launch clamps grid.x with
+// rocsparse::get_grid_size_x (rocsparse_grid.hpp); grid_x_chunk recovers the
+// work above the clamp inside the kernel.
+//
+#if defined(__HIPCC__)
+#define ROCSPARSE_GRID_X_ILF static __host__ __device__ __forceinline__
+#else
+#define ROCSPARSE_GRID_X_ILF static inline
+#endif
+
+namespace rocsparse
+{
+    //
+    // Half-open range [first, last) of the count work items owned by block
+    // block_id when the items are split into grid_size contiguous chunks. The
+    // range is empty when the block has nothing to do, which happens for the
+    // trailing blocks when grid_size does not divide count.
+    //
+    // Contiguous chunks, not a grid-stride loop, and that is load bearing for
+    // the incomplete factorizations. Those kernels spin on the done flags of the
+    // rows they depend on, and the row map is a topological order, so an item
+    // only ever waits on items earlier in that order. With contiguous chunks
+    // every item a block waits on belongs either to itself, earlier in its own
+    // chunk, or to a lower numbered block, which the dispatcher started first.
+    // That is the same property the unclamped launch relies on today, and it is
+    // what keeps the spin from deadlocking when the grid is larger than the
+    // device can hold resident.
+    //
+    // A grid-stride loop does not have that property. Block 0 on its second item
+    // waits on the first item of block grid_size - 1, which cannot start until
+    // some block retires, and no block retires while it is spinning.
+    //
+    ROCSPARSE_GRID_X_ILF void grid_x_chunk(
+        int64_t count, uint32_t grid_size, uint32_t block_id, int64_t& first, int64_t& last)
+    {
+        const int64_t chunk = (count - 1) / static_cast<int64_t>(grid_size) + 1;
+        const int64_t start = static_cast<int64_t>(block_id) * chunk;
+
+        //
+        // Both ends are clamped, not just the upper one. grid_size does not have
+        // to divide count, so a trailing block can start past the end and must
+        // come back with an empty range rather than an inverted one.
+        //
+        first = (start < count) ? start : count;
+        last  = ((first + chunk) < count) ? (first + chunk) : count;
+    }
+}

projects/rocsparse/library/src/precond/bsric0/rocsparse_bsric0_kernel_17_32.cpp

@@ -59,7 +59,20 @@ namespace rocsparse
         __shared__ T values[BSRDIM][BSRDIM + 1];
 
         // Current block row this wavefront is working on
-        auto block_row = block_map[hipBlockIdx_x];
+        const J idx = hipBlockIdx_x;
+
+        //
+        // Do not run out of bounds. The row map is indexed by the block id, and
+        // bsrilu0_kernel_general is the only kernel in the family that checked
+        // it; the CSR kernels all do. Cheap, and the only thing standing between
+        // a grid rounded up past mb and a read outside the map allocation.
+        //
+        if(idx >= mb)
+        {
+            return;
+        }
+
+        auto block_row = block_map[idx];
 
         // Block diagonal entry point of the current block row
         I block_row_diag = bsr_diag_ind[block_row];

projects/rocsparse/library/src/precond/bsric0/rocsparse_bsric0_kernel_2_8.cpp

@@ -56,8 +56,21 @@ namespace rocsparse
         __shared__ T values[BSRDIM][BSRDIM + 1];
         __shared__ T local_values[BSRDIM][BSRDIM + 1];
 
+        const J idx = hipBlockIdx_x;
+
+        //
+        // Do not run out of bounds. The row map is indexed by the block id, and
+        // bsrilu0_kernel_general is the only kernel in the family that checked
+        // it; the CSR kernels all do. Cheap, and the only thing standing between
+        // a grid rounded up past mb and a read outside the map allocation.
+        //
+        if(idx >= mb)
+        {
+            return;
+        }
+
         // Current block row this wavefront is working on
-        J block_row = block_map[hipBlockIdx_x];
+        J block_row = block_map[idx];
 
         // Block diagonal entry point of the current block row
         I block_row_diag = bsr_diag_ind[block_row];

projects/rocsparse/library/src/precond/bsric0/rocsparse_bsric0_kernel_2_8_unrolled.cpp

@@ -57,7 +57,20 @@ namespace rocsparse
         __shared__ T local_values[BSRDIM][BSRDIM + 1];
 
         // Current block row this wavefront is working on
-        J block_row = block_map[hipBlockIdx_x];
+        const J idx = hipBlockIdx_x;
+
+        //
+        // Do not run out of bounds. The row map is indexed by the block id, and
+        // bsrilu0_kernel_general is the only kernel in the family that checked
+        // it; the CSR kernels all do. Cheap, and the only thing standing between
+        // a grid rounded up past mb and a read outside the map allocation.
+        //
+        if(idx >= mb)
+        {
+            return;
+        }
+
+        J block_row = block_map[idx];
 
         // Block diagonal entry point of the current block row
         I block_row_diag = bsr_diag_ind[block_row];

projects/rocsparse/library/src/precond/bsric0/rocsparse_bsric0_kernel_9_16.cpp

@@ -61,7 +61,20 @@ namespace rocsparse
         __shared__ T local_values[BSRDIM][BSRDIM + 1];
 
         // Current block row this wavefront is working on
-        J block_row = block_map[hipBlockIdx_x];
+        const J idx = hipBlockIdx_x;
+
+        //
+        // Do not run out of bounds. The row map is indexed by the block id, and
+        // bsrilu0_kernel_general is the only kernel in the family that checked
+        // it; the CSR kernels all do. Cheap, and the only thing standing between
+        // a grid rounded up past mb and a read outside the map allocation.
+        //
+        if(idx >= mb)
+        {
+            return;
+        }
+
+        J block_row = block_map[idx];
 
         // Block diagonal entry point of the current block row
         I block_row_diag = bsr_diag_ind[block_row];