From ad4e5433f9309a7e20bd2154607cd8bda173204b Mon Sep 17 00:00:00 2001
From: Yinzuo Jiang <jiangyinzuo@foxmail.com>
Date: Thu, 27 Jun 2024 23:20:59 +0800
Subject: [PATCH] Fix compilation error when _CLK_BREAKDOWN is defined in
 cagra. (#2350)

PR #1740 forgot to rename `BLOCK_SIZE` in `#ifdef _CLK_BREAKDOWN` blocks.

The use of `RAFT_LOG_DEBUG` in kernel function results in compilation
errors, replace it with `printf`.

Also remove an unused function in search_single_cta_kernel-inl.cuh

After merging:

- [x] port to cuVS https://github.com/rapidsai/cuvs/pull/202

Authors:
  - Yinzuo Jiang (https://github.com/jiangyinzuo)
  - Tamas Bela Feher (https://github.com/tfeher)
  - Corey J. Nolet (https://github.com/cjnolet)

Approvers:
  - tsuki (https://github.com/enp1s0)
  - Tamas Bela Feher (https://github.com/tfeher)

URL: https://github.com/rapidsai/raft/pull/2350
---
 .../cagra/search_multi_cta_kernel-inl.cuh      | 12 ++++++++----
 .../cagra/search_single_cta_kernel-inl.cuh     | 18 +++++++-----------
 2 files changed, 15 insertions(+), 15 deletions(-)

diff --git a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh
index cfbb1e100c..16bb555aa4 100644
--- a/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh
+++ b/cpp/include/raft/neighbors/detail/cagra/search_multi_cta_kernel-inl.cuh
@@ -36,6 +36,7 @@
 
 #include <algorithm>
 #include <cassert>
+#include <cstdio>
 #include <iostream>
 #include <memory>
 #include <numeric>
@@ -209,7 +210,7 @@ __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel(
 
 #if 0
     /* debug */
-    for (unsigned i = threadIdx.x; i < result_buffer_size_32; i += BLOCK_SIZE) {
+    for (unsigned i = threadIdx.x; i < result_buffer_size_32; i += blockDim.x) {
         result_indices_buffer[i] = utils::get_max_value<INDEX_T>();
         result_distances_buffer[i] = utils::get_max_value<DISTANCE_T>();
     }
@@ -351,16 +352,19 @@ __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel(
   }
 
 #ifdef _CLK_BREAKDOWN
-  if ((threadIdx.x == 0 || threadIdx.x == BLOCK_SIZE - 1) && (blockIdx.x == 0) &&
+  if ((threadIdx.x == 0 || threadIdx.x == blockDim.x - 1) && (blockIdx.x == 0) &&
       ((query_id * 3) % gridDim.y < 3)) {
-    RAFT_LOG_DEBUG(
+    printf(
+      "%s:%d "
       "query, %d, thread, %d"
-      ", init, %d"
+      ", init, %lu"
       ", 1st_distance, %lu"
       ", topk, %lu"
       ", pickup_parents, %lu"
       ", distance, %lu"
       "\n",
+      __FILE__,
+      __LINE__,
       query_id,
       threadIdx.x,
       clk_init,
diff --git a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh
index e8104bd6f6..232dcb782a 100644
--- a/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh
+++ b/cpp/include/raft/neighbors/detail/cagra/search_single_cta_kernel-inl.cuh
@@ -40,6 +40,7 @@
 #include <algorithm>
 #include <cassert>
 #include <cstdint>
+#include <cstdio>
 #include <iostream>
 #include <memory>
 #include <numeric>
@@ -448,14 +449,6 @@ __device__ inline void hashmap_restore(INDEX_T* const hashmap_ptr,
   }
 }
 
-template <class T, unsigned BLOCK_SIZE>
-__device__ inline void set_value_device(T* const ptr, const T fill, const std::uint32_t count)
-{
-  for (std::uint32_t i = threadIdx.x; i < count; i += BLOCK_SIZE) {
-    ptr[i] = fill;
-  }
-}
-
 // One query one thread block
 template <uint32_t TEAM_SIZE,
           uint32_t DATASET_BLOCK_DIM,
@@ -791,10 +784,11 @@ __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel(
     num_executed_iterations[query_id] = iter + 1;
   }
 #ifdef _CLK_BREAKDOWN
-  if ((threadIdx.x == 0 || threadIdx.x == BLOCK_SIZE - 1) && ((query_id * 3) % gridDim.y < 3)) {
-    RAFT_LOG_DEBUG(
+  if ((threadIdx.x == 0 || threadIdx.x == blockDim.x - 1) && ((query_id * 3) % gridDim.y < 3)) {
+    printf(
+      "%s:%d "
       "query, %d, thread, %d"
-      ", init, %d"
+      ", init, %lu"
       ", 1st_distance, %lu"
       ", topk, %lu"
       ", reset_hash, %lu"
@@ -802,6 +796,8 @@ __launch_bounds__(1024, 1) RAFT_KERNEL search_kernel(
       ", restore_hash, %lu"
       ", distance, %lu"
       "\n",
+      __FILE__,
+      __LINE__,
       query_id,
       threadIdx.x,
       clk_init,