kliegeois · Oct 13, 2026 · Sep 17, 2026 · Sep 18, 2026 · Sep 18, 2026 · Sep 22, 2026
@@ -16,6 +16,9 @@ Documentation for rocSPARSE is available at
 * 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 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 launch grid can hold. 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 binding grid.x limit is `(2^32 - 1) / blockDim.x` blocks, because `gridDim.x * blockDim.x` must stay below 2^32, rather than the 2147483647 that `maxGridSize[0]` reports. All BSR kernels now clamp the extent and recover the remaining block rows inside the kernel.
+* 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.
 * 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`.
 

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

@@ -150,6 +150,9 @@ set(ROCSPARSE_UNIT_TEST_DEVICE_SOURCES
   unit_test_internal_grid.cpp
   # dispatch_grid_stride_x variant selection in rocsparse_grid.hpp (append-only leaf)
   unit_test_internal_grid_dispatch.cpp
+  # grid_x_chunk and the incomplete factorizations that walk it [AISPARSE-692-694]
+  unit_test_internal_grid_x.cpp
+  unit_test_internal_precond_grid_x_chunk.cpp
   unit_test_internal_collectives_boost.cpp
   unit_test_internal_infostructs.cpp
   #10156: distinct anchor to avoid stacked-branch merge conflicts

@@ -0,0 +1,162 @@
+/*! \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.
+ *
+ * ************************************************************************ */
+
+//
+// Unit tests for the contiguous chunk partition (rocsparse::grid_x_chunk in
+// rocsparse_grid.hpp) that the incomplete factorizations use to recover the
+// work above the grid.x clamp.
+//
+// rocsparse_grid.hpp pulls in the handle and HIP, so these build into
+// rocsparse-unit-test-device. They call grid_x_chunk on the host only and
+// 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.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);
+}