diff --git a/common/cuda_hip/base/kernel_launch_reduction.hpp b/common/cuda_hip/base/kernel_launch_reduction.hpp
index 4c4fb366802..63148fa05f8 100644
--- a/common/cuda_hip/base/kernel_launch_reduction.hpp
+++ b/common/cuda_hip/base/kernel_launch_reduction.hpp
@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2026 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -294,26 +294,31 @@ __launch_bounds__(default_block_size) void generic_kernel_col_reduction_2d_block
     const auto block = group::this_thread_block();
     const auto warp = group::tiled_partition<warp_size>(block);
     const auto warp_rank = warp.thread_rank();
-    const auto col = warp_rank + static_cast<int64>(blockIdx.y) * warp_size;
-    auto partial = identity;
-    // accumulate within a thread
-    if (col < cols) {
-        for (auto row = warp_id; row < rows; row += warp_num) {
-            partial = op(partial, fn(row, col, args...));
+
+    for (auto block_col = static_cast<int64>(blockIdx.y) * warp_size;
+         block_col < cols;
+         block_col += static_cast<int64>(gridDim.y) * warp_size) {
+        const auto col = warp_rank + block_col;
+        auto partial = identity;
+        // accumulate within a thread
+        if (col < cols) {
+            for (auto row = warp_id; row < rows; row += warp_num) {
+                partial = op(partial, fn(row, col, args...));
+            }
         }
-    }
-    block_partial[threadIdx.x] = partial;
-    block.sync();
-    // in a single warp: accumulate the results
-    if (threadIdx.x < warp_size) {
-        partial = identity;
-        // accumulate the partial results within a thread
+        block_partial[threadIdx.x] = partial;
+        block.sync();
+        // in a single warp: accumulate the results
+        if (threadIdx.x < warp_size) {
+            partial = identity;
+            // accumulate the partial results within a thread
 #pragma unroll
-        for (int i = 0; i < default_block_size; i += warp_size) {
-            partial = op(partial, block_partial[i + warp_rank]);
-        }
-        if (col < cols) {
-            result[col + blockIdx.x * cols] = finalize(partial);
+            for (int i = 0; i < default_block_size; i += warp_size) {
+                partial = op(partial, block_partial[i + warp_rank]);
+            }
+            if (col < cols) {
+                result[col + blockIdx.x * cols] = finalize(partial);
+            }
         }
     }
 }
@@ -488,7 +493,10 @@ void run_kernel_col_reduction_cached(
             syn::value_list<int>(), syn::type_list<>(), max_blocks, exec, fn,
             op, finalize, identity, result, size, tmp, map_to_device(args)...);
     } else {
-        const auto col_blocks = ceildiv(cols, config::warp_size);
+        // cuda only accept up to 65545 for grid's y-axix
+        constexpr int64 max_grid_y = 65535;
+        const auto col_blocks =
+            std::min(ceildiv(cols, config::warp_size), max_grid_y);
         const auto row_blocks =
             ceildiv(std::min<int64>(
                         ceildiv(rows * config::warp_size, default_block_size),
diff --git a/test/base/kernel_launch_generic.cpp b/test/base/kernel_launch_generic.cpp
index 93eadb8d573..4c5f3df69f0 100644
--- a/test/base/kernel_launch_generic.cpp
+++ b/test/base/kernel_launch_generic.cpp
@@ -562,7 +562,9 @@ void run2d_col_reduction(std::shared_ptr<gko::EXEC_TYPE> exec)
     for (auto num_rows : {0, 10, 100, 1000, 10000}) {
         // check different edge cases: subwarp sizes, blocked mode
         for (auto num_cols :
-             {0, 1, 2, 3, 4, 5, 7, 8, 9, 16, 31, 32, 63, 127, 128, 129}) {
+             {0, 1, 2, 3, 4, 5, 7, 8, 9, 16, 31, 32, 63, 127, 128, 129,
+              65535 * 128 +
+                  1 /* check we do not exceed the gridDim.y limit */}) {
             SCOPED_TRACE(std::to_string(num_rows) + " rows, " +
                          std::to_string(num_cols) + " cols");
             gko::array<int64> host_ref{exec->get_master(),