AISPARSE-687: clamp grid.x in the dense conversion paths and add grid-stride loops

projects/rocsparse/clients/testings/testing_dense_to_sparse_bell.cpp

@@ -299,4 +299,139 @@ INSTANTIATE(int64_t, float);
 INSTANTIATE(int64_t, double);
 INSTANTIATE(int64_t, rocsparse_float_complex);
 INSTANTIATE(int64_t, rocsparse_double_complex);
-void testing_dense_to_sparse_bell_extra(const Arguments& arg) {}
+
+template <typename I, typename T>
+void testing_dense_to_sparse_bell_extra_687(const Arguments& arg)
+{
+    if(!arg.unit_check)
+    {
+        return;
+    }
+
+    int             device;
+    hipDeviceProp_t prop;
+    CHECK_HIP_ERROR(hipGetDevice(&device));
+    CHECK_HIP_ERROR(hipGetDeviceProperties(&prop, device));
+
+    // The dense2bell kernels launch one 256-thread block per block-row, and their grid is
+    // clamped to min(maxGridSize[0], (2^32 - 1) / 256) blocks, i.e. 16,777,215 blocks. With a
+    // block size of one, m exceeds that clamp by exactly one block-row. The single trailing
+    // block-row is only reached if the grid is clamped and the kernels stride.
+    const int64_t block_threads  = 256;
+    const int64_t grid_x_max     = std::min(static_cast<int64_t>(prop.maxGridSize[0]),
+                                        static_cast<int64_t>(UINT32_MAX) / block_threads);
+    const int64_t ell_block_size = 1;
+    const int64_t m              = grid_x_max + 1;
+    const int64_t n              = 1;
+    const int64_t ld             = m;
+
+    const int64_t mb = (m - 1) / ell_block_size + 1;
+    const int64_t nb = (n - 1) / ell_block_size + 1;
+
+    const rocsparse_index_base          base  = rocsparse_index_base_one;
+    const rocsparse_order               order = rocsparse_order_column;
+    const rocsparse_dense_to_sparse_alg alg   = rocsparse_dense_to_sparse_alg_default;
+
+    const rocsparse_indextype itype = get_indextype<I>();
+    const rocsparse_datatype  ttype = get_datatype<T>();
+
+    rocsparse_local_handle handle;
+
+    // The dense matrix is zero everywhere but in its very last row, so that the converted
+    // blocked ELL matrix has a single occupied slot and that slot sits in the trailing
+    // block-row that an unclamped or non-striding grid does not reach.
+    const T h_last_val = static_cast<T>(3);
+
+    device_vector<T> d_dense_val(ld * n);
+    CHECK_HIP_ERROR(hipMemset(d_dense_val, 0, sizeof(T) * ld * n));
+    CHECK_HIP_ERROR(
+        hipMemcpy(d_dense_val.data() + (m - 1), &h_last_val, sizeof(T), hipMemcpyHostToDevice));
+
+    rocsparse_local_dnmat mat_dense(m, n, ld, d_dense_val, ttype, order);
+
+    rocsparse_local_spmat mat_sparse(mb * ell_block_size,
+                                     nb * ell_block_size,
+                                     rocsparse_direction_row,
+                                     ell_block_size,
+                                     0,
+                                     nullptr,
+                                     nullptr,
+                                     itype,
+                                     base,
+                                     ttype);
+
+    size_t buffer_size;
+    CHECK_ROCSPARSE_ERROR(
+        rocsparse_dense_to_sparse(handle, mat_dense, mat_sparse, alg, &buffer_size, nullptr));
+
+    device_vector<int64_t> d_temp_buffer(buffer_size / sizeof(int64_t));
+
+    CHECK_ROCSPARSE_ERROR(
+        rocsparse_dense_to_sparse(handle, mat_dense, mat_sparse, alg, nullptr, d_temp_buffer));
+
+    int64_t              rows_tmp;
+    int64_t              cols_tmp;
+    int64_t              ell_block_dim_tmp;
+    int64_t              ell_cols;
+    rocsparse_direction  ell_block_dir_tmp;
+    rocsparse_indextype  itype_tmp;
+    rocsparse_index_base base_tmp;
+    rocsparse_datatype   ttype_tmp;
+    void*                ell_col_ind_tmp;
+    void*                ell_val_tmp;
+
+    CHECK_ROCSPARSE_ERROR(rocsparse_bell_get(mat_sparse,
+                                             &rows_tmp,
+                                             &cols_tmp,
+                                             &ell_block_dir_tmp,
+                                             &ell_block_dim_tmp,
+                                             &ell_cols,
+                                             &ell_col_ind_tmp,
+                                             &ell_val_tmp,
+                                             &itype_tmp,
+                                             &base_tmp,
+                                             &ttype_tmp));
+
+    unit_check_scalar<int64_t>(ell_cols, ell_block_size);
+
+    const int64_t ell_block_width = ell_cols / ell_block_size;
+
+    device_vector<I> d_bell_col_ind(mb * ell_block_width);
+    device_vector<T> d_bell_val(m * ell_cols);
+
+    // Poison the column indices so that a block-row the fill kernel never visits cannot be
+    // mistaken for either an occupied or a padded slot.
+    CHECK_HIP_ERROR(hipMemset(d_bell_col_ind, 0xff, sizeof(I) * mb * ell_block_width));
+
+    CHECK_ROCSPARSE_ERROR(rocsparse_bell_set_pointers(mat_sparse, d_bell_col_ind, d_bell_val));
+
+    CHECK_ROCSPARSE_ERROR(
+        rocsparse_dense_to_sparse(handle, mat_dense, mat_sparse, alg, &buffer_size, d_temp_buffer));
+
+    I h_col_ind[2];
+    T h_val[2];
+
+    // Leading block-row: empty, so its single slot is padded and its value is zero.
+    CHECK_HIP_ERROR(
+        hipMemcpy(&h_col_ind[0], d_bell_col_ind.data(), sizeof(I), hipMemcpyDeviceToHost));
+    CHECK_HIP_ERROR(hipMemcpy(&h_val[0], d_bell_val.data(), sizeof(T), hipMemcpyDeviceToHost));
+
+    // Trailing block-row: holds the only non-zero of the matrix in block-column zero.
+    CHECK_HIP_ERROR(hipMemcpy(&h_col_ind[1],
+                              d_bell_col_ind.data() + (mb - 1) * ell_block_width,
+                              sizeof(I),
+                              hipMemcpyDeviceToHost));
+    CHECK_HIP_ERROR(hipMemcpy(
+        &h_val[1], d_bell_val.data() + (m - 1) * ell_cols, sizeof(T), hipMemcpyDeviceToHost));
+
+    unit_check_scalar<I>(h_col_ind[0], static_cast<I>(base) - 1);
+    unit_check_scalar<T>(h_val[0], static_cast<T>(0));
+
+    unit_check_scalar<I>(h_col_ind[1], static_cast<I>(base));
+    unit_check_scalar<T>(h_val[1], h_last_val);
+}
+
+void testing_dense_to_sparse_bell_extra(const Arguments& arg)
+{
+    testing_dense_to_sparse_bell_extra_687<int32_t, float>(arg);
+}

projects/rocsparse/clients/tests/test_dense_to_sparse_bell.yaml

@@ -131,3 +131,11 @@ Tests:
   baseA: [rocsparse_index_base_zero]
   order: [rocsparse_order_row]
 
+- name: dense_to_sparse_bell_extra
+  category: nightly
+  function: dense_to_sparse_bell_extra
+  indextype: *i32
+  precision: *single_only_precisions
+  host_memory_gb: 1
+  device_memory_gb: 1
+

projects/rocsparse/library/src/conversion/coo2dense_device.h

@@ -41,24 +41,24 @@ namespace rocsparse
                           T*                   A,
                           rocsparse_order      order)
     {
-        int64_t gid = hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x;
+        const int64_t NUM_THREADS = static_cast<int64_t>(BLOCKSIZE) * hipGridDim_x;
 
-        if(gid >= nnz)
-        {
-            return;
-        }
+        const int64_t gid = static_cast<int64_t>(BLOCKSIZE) * hipBlockIdx_x + hipThreadIdx_x;
 
-        I row = coo_row_ind[gid] - base;
-        I col = coo_col_ind[gid] - base;
-        T val = coo_val[gid];
-
-        if(order == rocsparse_order_column)
-        {
-            A[lda * col + row] = val;
-        }
-        else
+        for(int64_t idx = gid; idx < nnz; idx += NUM_THREADS)
         {
-            A[lda * row + col] = val;
+            I row = coo_row_ind[idx] - base;
+            I col = coo_col_ind[idx] - base;
+            T val = coo_val[idx];
+
+            if(order == rocsparse_order_column)
+            {
+                A[lda * col + row] = val;
+            }
+            else
+            {
+                A[lda * row + col] = val;
+            }
         }
     }
 
@@ -74,11 +74,11 @@ namespace rocsparse
                               T*                   A,
                               rocsparse_order      order)
     {
-        const auto NUM_THREADS = hipGridDim_x * BLOCKSIZE;
+        const int64_t NUM_THREADS = static_cast<int64_t>(BLOCKSIZE) * hipGridDim_x;
 
-        const auto gid = hipBlockIdx_x * BLOCKSIZE + hipThreadIdx_x;
+        const int64_t gid = static_cast<int64_t>(BLOCKSIZE) * hipBlockIdx_x + hipThreadIdx_x;
 
-        for(auto idx = gid; idx < nnz; idx += NUM_THREADS)
+        for(int64_t idx = gid; idx < nnz; idx += NUM_THREADS)
         {
             I row = coo_ind[2 * idx] - base;
             I col = coo_ind[2 * idx + 1] - base;

projects/rocsparse/library/src/conversion/csx2dense_device.h

@@ -46,9 +46,12 @@ namespace rocsparse
     {
         const rocsparse_int wavefront_index = hipThreadIdx_x / WF_SIZE,
                             lane_index      = hipThreadIdx_x % WF_SIZE;
-        const J row_index                   = NUMROWS_PER_BLOCK * hipBlockIdx_x + wavefront_index;
+        const int64_t stride = static_cast<int64_t>(NUMROWS_PER_BLOCK) * hipGridDim_x;
 
-        if(row_index < m)
+        for(int64_t row_index
+            = static_cast<int64_t>(NUMROWS_PER_BLOCK) * hipBlockIdx_x + wavefront_index;
+            row_index < m;
+            row_index += stride)
         {
             const I shift    = csr_row_ptr[row_index] - base;
             const I size_row = csr_row_ptr[row_index + 1] - base - shift;
@@ -90,9 +93,12 @@ namespace rocsparse
     {
         const rocsparse_int wavefront_index = hipThreadIdx_x / WF_SIZE,
                             lane_index      = hipThreadIdx_x % WF_SIZE;
-        const J column_index = NUMCOLUMNS_PER_BLOCK * hipBlockIdx_x + wavefront_index;
+        const int64_t stride = static_cast<int64_t>(NUMCOLUMNS_PER_BLOCK) * hipGridDim_x;
 
-        if(column_index < n)
+        for(int64_t column_index
+            = static_cast<int64_t>(NUMCOLUMNS_PER_BLOCK) * hipBlockIdx_x + wavefront_index;
+            column_index < n;
+            column_index += stride)
         {
             const I shift       = csc_col_ptr[column_index] - base;
             const I size_column = csc_col_ptr[column_index + 1] - base - shift;