Implement a kernel dispatcher raft::launch_kernel

cpp/include/raft/common/detail/scatter.cuh

@@ -1,12 +1,13 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
 #pragma once
 
 #include <raft/core/detail/macros.hpp>
 #include <raft/util/cuda_utils.cuh>
+#include <raft/util/kernel_launch.hpp>
 #include <raft/util/vectorized.cuh>
 
 namespace raft {
@@ -36,8 +37,8 @@ void scatterImpl(
   DataT* out, const DataT* in, const IdxT* idx, IdxT len, Lambda op, cudaStream_t stream)
 {
   const IdxT nblks = raft::ceildiv(VecLen ? len / VecLen : len, (IdxT)TPB);
-  scatterKernel<DataT, VecLen, Lambda, IdxT><<<nblks, TPB, 0, stream>>>(out, in, idx, len, op);
-  RAFT_CUDA_TRY(cudaGetLastError());
+  raft::launch_kernel(
+    stream, nblks, TPB, scatterKernel<DataT, VecLen, Lambda, IdxT>, out, in, idx, len, op);
 }
 
 }  // namespace detail

cpp/include/raft/core/bitset.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -15,6 +15,7 @@
 #include <raft/linalg/reduce.cuh>
 #include <raft/sparse/convert/csr.cuh>
 #include <raft/util/device_atomics.cuh>
+#include <raft/util/kernel_launch.hpp>
 #include <raft/util/popc.cuh>
 
 #include <rmm/device_scalar.hpp>
@@ -147,16 +148,21 @@ void bitset_repeat(raft::resources const& handle,
                    index_t repeat_times)
 {
   if (src_bit_len == 0 || repeat_times == 0) return;
-  auto stream = resource::get_cuda_stream(handle);
 
   constexpr index_t bits_per_element = sizeof(bitset_t) * 8;
   const index_t total_bits           = src_bit_len * repeat_times;
   const index_t output_size          = (total_bits + bits_per_element - 1) / bits_per_element;
 
   int threadsPerBlock = 128;
   int blocksPerGrid   = (output_size + threadsPerBlock - 1) / threadsPerBlock;
-  bitset_repeat_kernel<<<blocksPerGrid, threadsPerBlock, 0, stream>>>(
-    d_src, d_output, src_bit_len, repeat_times);
+  raft::launch_kernel(handle,
+                      blocksPerGrid,
+                      threadsPerBlock,
+                      bitset_repeat_kernel,
+                      d_src,
+                      d_output,
+                      src_bit_len,
+                      repeat_times);
 
   return;
 }

cpp/include/raft/core/detail/copy.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -24,6 +24,7 @@
 #ifdef __CUDACC__
 #include <raft/linalg/transpose.cuh>
 #include <raft/util/cuda_dev_essentials.cuh>
+#include <raft/util/kernel_launch.hpp>
 #endif
 #endif
 
@@ -503,7 +504,12 @@ mdspan_copyable_t<DstType, SrcType> copy(resources const& res, DstType&& dst, Sr
       raft::ceildiv(typename config::index_type(dst.size()),
                     typename config::index_type(mdspan_copy_tile_elems)));
     auto constexpr const threads = dim3{mdspan_copy_tile_dim, mdspan_copy_tile_dim, 1};
-    mdspan_copy_kernel<<<blocks, threads, 0, resource::get_cuda_stream(res)>>>(dst, src);
+    raft::launch_kernel(res,
+                        blocks,
+                        threads,
+                        mdspan_copy_kernel<typename config::dst_type, typename config::src_type>,
+                        dst,
+                        src);
 #else
     // Should never actually reach this because of enable_ifs. Included for
     // safety.

cpp/include/raft/label/detail/classlabels.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -10,6 +10,7 @@
 #include <raft/linalg/unary_op.cuh>
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/cudart_utils.hpp>
+#include <raft/util/kernel_launch.hpp>
 
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
@@ -163,8 +164,17 @@ void make_monotonic(
   rmm::device_uvector<Type> map_ids(0, stream);
   int num_clusters = getUniquelabels(map_ids, in, N, stream);
 
-  map_label_kernel<Type, TPB_X><<<blocks, threads, 0, stream>>>(
-    map_ids.data(), num_clusters, in, out, N, filter_op, zero_based);
+  raft::launch_kernel(stream,
+                      blocks,
+                      threads,
+                      map_label_kernel<Type, TPB_X, Lambda>,
+                      map_ids.data(),
+                      num_clusters,
+                      in,
+                      out,
+                      N,
+                      filter_op,
+                      zero_based);
 }
 
 /**

cpp/include/raft/label/detail/merge_labels.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -9,6 +9,7 @@
 #include <raft/linalg/init.cuh>
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/cudart_utils.hpp>
+#include <raft/util/kernel_launch.hpp>
 
 #include <math.h>
 
@@ -129,18 +130,31 @@ void merge_labels(value_idx* labels_a,
   do {
     RAFT_CUDA_TRY(cudaMemsetAsync(m, false, sizeof(bool), stream));
 
-    propagate_label_kernel<value_idx, TPB_X>
-      <<<blocks, threads, 0, stream>>>(labels_a, labels_b, R, mask, m, N);
-    RAFT_CUDA_TRY(cudaPeekAtLastError());
+    raft::launch_kernel(stream,
+                        blocks,
+                        threads,
+                        propagate_label_kernel<value_idx, TPB_X>,
+                        labels_a,
+                        labels_b,
+                        R,
+                        mask,
+                        m,
+                        N);
 
     raft::update_host(&host_m, m, 1, stream);
     RAFT_CUDA_TRY(cudaStreamSynchronize(stream));
   } while (host_m);
 
   // Step 2: re-assign minimum equivalent label
-  reassign_label_kernel<value_idx, TPB_X>
-    <<<blocks, threads, 0, stream>>>(labels_a, labels_b, R, N, MAX_LABEL);
-  RAFT_CUDA_TRY(cudaPeekAtLastError());
+  raft::launch_kernel(stream,
+                      blocks,
+                      threads,
+                      reassign_label_kernel<value_idx, TPB_X>,
+                      labels_a,
+                      labels_b,
+                      R,
+                      N,
+                      MAX_LABEL);
 }
 
 }  // namespace detail

cpp/include/raft/linalg/detail/add.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -10,6 +10,7 @@
 #include <raft/linalg/binary_op.cuh>
 #include <raft/linalg/unary_op.cuh>
 #include <raft/util/cuda_utils.cuh>
+#include <raft/util/kernel_launch.hpp>
 
 namespace raft {
 namespace linalg {
@@ -44,8 +45,8 @@ void addDevScalar(
   // TODO: block dimension has not been tuned
   dim3 block(256);
   dim3 grid(raft::ceildiv(len, (IdxType)block.x));
-  add_dev_scalar_kernel<<<grid, block, 0, stream>>>(outDev, inDev, singleScalarDev, len);
-  RAFT_CUDA_TRY(cudaPeekAtLastError());
+  raft::launch_kernel(
+    stream, grid, block, add_dev_scalar_kernel, outDev, inDev, singleScalarDev, len);
 }
 
 }  // namespace detail

cpp/include/raft/linalg/detail/coalesced_reduction-inl.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -9,6 +9,7 @@
 #include <raft/core/nvtx.hpp>
 #include <raft/core/operators.hpp>
 #include <raft/util/cuda_utils.cuh>
+#include <raft/util/kernel_launch.hpp>
 
 #include <rmm/device_uvector.hpp>
 
@@ -252,13 +253,33 @@ void coalescedReductionThin(OutType* dots,
   dim3 threads(Policy::LogicalWarpSize, Policy::NumLogicalWarps, 1);
   dim3 blocks(ceildiv<IdxType>(N, Policy::RowsPerBlock), 1, 1);
   if constexpr (std::is_same_v<ReduceLambda, raft::add_op>) {
-    coalescedSumThinKernel<Policy>
-      <<<blocks, threads, 0, stream>>>(dots, data, D, N, init, main_op, final_op, inplace);
+    raft::launch_kernel(stream,
+                        blocks,
+                        threads,
+                        coalescedSumThinKernel<Policy>,
+                        dots,
+                        data,
+                        D,
+                        N,
+                        init,
+                        main_op,
+                        final_op,
+                        inplace);
   } else {
-    coalescedReductionThinKernel<Policy><<<blocks, threads, 0, stream>>>(
-      dots, data, D, N, init, main_op, reduce_op, final_op, inplace);
+    raft::launch_kernel(stream,
+                        blocks,
+                        threads,
+                        coalescedReductionThinKernel<Policy>,
+                        dots,
+                        data,
+                        D,
+                        N,
+                        init,
+                        main_op,
+                        reduce_op,
+                        final_op,
+                        inplace);
   }
-  RAFT_CUDA_TRY(cudaPeekAtLastError());
 }
 
 template <typename InType,
@@ -398,13 +419,33 @@ void coalescedReductionMedium(OutType* dots,
 {
   common::nvtx::range<common::nvtx::domain::raft> fun_scope("coalescedReductionMedium<%d>", TPB);
   if constexpr (std::is_same_v<ReduceLambda, raft::add_op>) {
-    coalescedSumMediumKernel<TPB>
-      <<<N, TPB, 0, stream>>>(dots, data, D, N, init, main_op, final_op, inplace);
+    raft::launch_kernel(stream,
+                        N,
+                        TPB,
+                        coalescedSumMediumKernel<TPB>,
+                        dots,
+                        data,
+                        D,
+                        N,
+                        init,
+                        main_op,
+                        final_op,
+                        inplace);
   } else {
-    coalescedReductionMediumKernel<TPB>
-      <<<N, TPB, 0, stream>>>(dots, data, D, N, init, main_op, reduce_op, final_op, inplace);
+    raft::launch_kernel(stream,
+                        N,
+                        TPB,
+                        coalescedReductionMediumKernel<TPB>,
+                        dots,
+                        data,
+                        D,
+                        N,
+                        init,
+                        main_op,
+                        reduce_op,
+                        final_op,
+                        inplace);
   }
-  RAFT_CUDA_TRY(cudaPeekAtLastError());
 }
 
 template <typename InType,
@@ -525,13 +566,29 @@ void coalescedReductionThick(OutType* dots,
    *     main_op but applies final_op. If in-place, the existing and new values are reduced.
    */
   if constexpr (std::is_same_v<ReduceLambda, raft::add_op>) {
-    coalescedSumThickKernel<ThickPolicy>
-      <<<blocks, threads, 0, stream>>>(buffer.data(), data, D, N, init, main_op);
+    raft::launch_kernel(stream,
+                        blocks,
+                        threads,
+                        coalescedSumThickKernel<ThickPolicy>,
+                        buffer.data(),
+                        data,
+                        D,
+                        N,
+                        init,
+                        main_op);
   } else {
-    coalescedReductionThickKernel<ThickPolicy>
-      <<<blocks, threads, 0, stream>>>(buffer.data(), data, D, N, init, main_op, reduce_op);
+    raft::launch_kernel(stream,
+                        blocks,
+                        threads,
+                        coalescedReductionThickKernel<ThickPolicy>,
+                        buffer.data(),
+                        data,
+                        D,
+                        N,
+                        init,
+                        main_op,
+                        reduce_op);
   }
-  RAFT_CUDA_TRY(cudaPeekAtLastError());
 
   coalescedReductionThin<ThinPolicy>(dots,
                                      buffer.data(),

cpp/include/raft/linalg/detail/map.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -12,6 +12,7 @@
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/input_validation.hpp>
 #include <raft/util/integer_utils.hpp>
+#include <raft/util/kernel_launch.hpp>
 #include <raft/util/pow2_utils.cuh>
 #include <raft/util/vectorized.cuh>
 #include <raft/util/vectorized_kvp.cuh>
@@ -97,7 +98,8 @@ void map_call(rmm::cuda_stream_view stream, OutT* out_ptr, IdxT len, Func f, con
   const int threads =
     std::max<int>(WarpSize, std::min<IdxT>(raft::bound_by_power_of_two<IdxT>(len_vectorized), 256));
   const IdxT blocks = raft::div_rounding_up_unsafe<IdxT>(len_vectorized, threads);
-  map_kernel<R, PassOffset><<<blocks, threads, 0, stream>>>(out_ptr, len, f, in_ptrs...);
+  raft::launch_kernel(
+    stream, blocks, threads, map_kernel<R, PassOffset>, out_ptr, len, f, in_ptrs...);
 }
 
 constexpr int kCoalescedVectorSize = 16;

cpp/include/raft/linalg/detail/map_then_reduce.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -8,6 +8,7 @@
 #include <raft/core/detail/macros.hpp>
 #include <raft/core/resources.hpp>
 #include <raft/util/cuda_utils.cuh>
+#include <raft/util/kernel_launch.hpp>
 #include <raft/util/vectorized.cuh>
 
 #include <cub/block/block_reduce.cuh>
@@ -79,9 +80,18 @@ void mapThenReduceImpl(OutType* out,
 {
   raft::update_device(out, &neutral, 1, stream);
   const int nblks = raft::ceildiv(len, IdxType(TPB));
-  mapThenReduceKernel<InType, OutType, IdxType, MapOp, ReduceLambda, TPB, Args...>
-    <<<nblks, TPB, 0, stream>>>(out, len, neutral, map, op, in, args...);
-  RAFT_CUDA_TRY(cudaPeekAtLastError());
+  raft::launch_kernel(
+    stream,
+    nblks,
+    TPB,
+    mapThenReduceKernel<InType, OutType, IdxType, MapOp, ReduceLambda, TPB, Args...>,
+    out,
+    len,
+    neutral,
+    map,
+    op,
+    in,
+    args...);
 }
 
 };  // end namespace detail