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..3085767dad 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 @@ -209,7 +209,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,7 +351,7 @@ __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( "query, %d, thread, %d" 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..38484c20ad 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 @@ -448,14 +448,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