From 97e7bebb62c41f1b5ecb80d51098c582e89b6a39 Mon Sep 17 00:00:00 2001
From: Mahesh Doijade <mdoijade@nvidia.com>
Date: Thu, 1 Feb 2024 09:14:03 +0000
Subject: [PATCH] fix the doc build error and cpp symbol-check tests issues

---
 .../detail/fused_distance_nn/cutlass_base.cuh |  2 +-
 .../fused_distance_nn/helper_structs.cuh      |  2 +-
 .../detail/fused_distance_nn/simt_kernel.cuh  | 28 +++++++++----------
 3 files changed, 16 insertions(+), 16 deletions(-)

diff --git a/cpp/include/raft/distance/detail/fused_distance_nn/cutlass_base.cuh b/cpp/include/raft/distance/detail/fused_distance_nn/cutlass_base.cuh
index cf13bafe85..be1c8891f5 100644
--- a/cpp/include/raft/distance/detail/fused_distance_nn/cutlass_base.cuh
+++ b/cpp/include/raft/distance/detail/fused_distance_nn/cutlass_base.cuh
@@ -48,7 +48,7 @@ namespace distance {
 namespace detail {
 
 template <typename IdxT>
-__global__ void initBinMutexKernel(cuda::binary_semaphore<cuda::thread_scope_device>* mut, IdxT m)
+RAFT_KERNEL initBinMutexKernel(cuda::binary_semaphore<cuda::thread_scope_device>* mut, IdxT m)
 {
   auto tid = IdxT(blockIdx.x) * blockDim.x + threadIdx.x;
 
diff --git a/cpp/include/raft/distance/detail/fused_distance_nn/helper_structs.cuh b/cpp/include/raft/distance/detail/fused_distance_nn/helper_structs.cuh
index 209c9ad8c8..e88ea9cfc8 100644
--- a/cpp/include/raft/distance/detail/fused_distance_nn/helper_structs.cuh
+++ b/cpp/include/raft/distance/detail/fused_distance_nn/helper_structs.cuh
@@ -107,7 +107,7 @@ struct MinReduceOpImpl {
 };
 
 template <typename DataT, typename OutT, typename IdxT, typename ReduceOpT>
-__global__ void initKernel(OutT* min, IdxT m, DataT maxVal, ReduceOpT redOp)
+RAFT_KERNEL initKernel(OutT* min, IdxT m, DataT maxVal, ReduceOpT redOp)
 {
   auto tid = IdxT(blockIdx.x) * blockDim.x + threadIdx.x;
   if (tid < m) { redOp.init(min + tid, maxVal); }
diff --git a/cpp/include/raft/distance/detail/fused_distance_nn/simt_kernel.cuh b/cpp/include/raft/distance/detail/fused_distance_nn/simt_kernel.cuh
index 242548fbfd..f5e4c725d6 100644
--- a/cpp/include/raft/distance/detail/fused_distance_nn/simt_kernel.cuh
+++ b/cpp/include/raft/distance/detail/fused_distance_nn/simt_kernel.cuh
@@ -66,20 +66,20 @@ template <typename DataT,
           typename KVPReduceOpT,
           typename OpT,
           typename FinalLambda>
-__global__ __launch_bounds__(P::Nthreads, 2) void fusedDistanceNNkernel(OutT* min,
-                                                                        const DataT* x,
-                                                                        const DataT* y,
-                                                                        const DataT* xn,
-                                                                        const DataT* yn,
-                                                                        IdxT m,
-                                                                        IdxT n,
-                                                                        IdxT k,
-                                                                        DataT maxVal,
-                                                                        int* mutex,
-                                                                        ReduceOpT redOp,
-                                                                        KVPReduceOpT pairRedOp,
-                                                                        OpT distance_op,
-                                                                        FinalLambda fin_op)
+__launch_bounds__(P::Nthreads, 2) RAFT_KERNEL fusedDistanceNNkernel(OutT* min,
+                                                                    const DataT* x,
+                                                                    const DataT* y,
+                                                                    const DataT* xn,
+                                                                    const DataT* yn,
+                                                                    IdxT m,
+                                                                    IdxT n,
+                                                                    IdxT k,
+                                                                    DataT maxVal,
+                                                                    int* mutex,
+                                                                    ReduceOpT redOp,
+                                                                    KVPReduceOpT pairRedOp,
+                                                                    OpT distance_op,
+                                                                    FinalLambda fin_op)
 {
 // compile only if below non-ampere arch.
 #if __CUDA_ARCH__ < 800