fix the doc build error and cpp symbol-check tests issues

rapidsai · Feb 1, 2024 · 97e7beb · 97e7beb
1 parent 904fdd6
commit 97e7beb
Show file tree

Hide file tree

Showing 3 changed files with 16 additions and 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
@@ -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
@@ -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
@@ -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