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 #include +#include #include #include #include @@ -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(); result_distances_buffer[i] = utils::get_max_value(); } @@ -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 #include #include +#include #include #include #include @@ -448,14 +449,6 @@ __device__ inline void hashmap_restore(INDEX_T* const hashmap_ptr, } } -template -__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