diff --git a/cpp/internal/raft_internal/neighbors/ivf_pq_compute_similarity_filters_test-ext.cuh b/cpp/internal/raft_internal/neighbors/ivf_pq_compute_similarity_filters_test-ext.cuh
deleted file mode 100644
index aa14ab19b8..0000000000
--- a/cpp/internal/raft_internal/neighbors/ivf_pq_compute_similarity_filters_test-ext.cuh
+++ /dev/null
@@ -1,181 +0,0 @@
-/*
- * Copyright (c) 2024, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#include <raft/core/detail/macros.hpp>       // RAFT_WEAK_FUNCTION
-#include <raft/distance/distance_types.hpp>  // raft::distance::DistanceType
-#include <raft/neighbors/detail/ivf_pq_compute_similarity.cuh>
-#include <raft/neighbors/detail/ivf_pq_fp_8bit.cuh>  // raft::neighbors::ivf_pq::detail::fp_8bit
-#include <raft/neighbors/sample_filter.cuh>          // none_ivf_sample_filter
-#include <raft/neighbors/sample_filter_types.hpp>    // none_ivf_sample_filter
-
-#include <rmm/cuda_stream_view.hpp>  // rmm::cuda_stream_view
-
-#include <cuda_fp16.h>  // __half
-
-#define instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(                 \
-  OutT, LutT, IvfSampleFilterT)                                                             \
-  extern template auto                                                                      \
-  raft::neighbors::ivf_pq::detail::compute_similarity_select<OutT, LutT, IvfSampleFilterT>( \
-    const cudaDeviceProp& dev_props,                                                        \
-    bool manage_local_topk,                                                                 \
-    int locality_hint,                                                                      \
-    double preferred_shmem_carveout,                                                        \
-    uint32_t pq_bits,                                                                       \
-    uint32_t pq_dim,                                                                        \
-    uint32_t precomp_data_count,                                                            \
-    uint32_t n_queries,                                                                     \
-    uint32_t n_probes,                                                                      \
-    uint32_t topk)                                                                          \
-    ->raft::neighbors::ivf_pq::detail::selected<OutT, LutT, IvfSampleFilterT>;              \
-                                                                                            \
-  extern template void                                                                      \
-  raft::neighbors::ivf_pq::detail::compute_similarity_run<OutT, LutT, IvfSampleFilterT>(    \
-    raft::neighbors::ivf_pq::detail::selected<OutT, LutT, IvfSampleFilterT> s,              \
-    rmm::cuda_stream_view stream,                                                           \
-    uint32_t dim,                                                                           \
-    uint32_t n_probes,                                                                      \
-    uint32_t pq_dim,                                                                        \
-    uint32_t n_queries,                                                                     \
-    uint32_t queries_offset,                                                                \
-    raft::distance::DistanceType metric,                                                    \
-    raft::neighbors::ivf_pq::codebook_gen codebook_kind,                                    \
-    uint32_t topk,                                                                          \
-    uint32_t max_samples,                                                                   \
-    const float* cluster_centers,                                                           \
-    const float* pq_centers,                                                                \
-    const uint8_t* const* pq_dataset,                                                       \
-    const uint32_t* cluster_labels,                                                         \
-    const uint32_t* _chunk_indices,                                                         \
-    const float* queries,                                                                   \
-    const uint32_t* index_list,                                                             \
-    float* query_kths,                                                                      \
-    IvfSampleFilterT sample_filter,                                                         \
-    LutT* lut_scores,                                                                       \
-    OutT* _out_scores,                                                                      \
-    uint32_t* _out_indices);
-
-#define COMMA ,
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA false>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::none_ivf_sample_filter>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA true>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::none_ivf_sample_filter>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  half,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::none_ivf_sample_filter>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  half,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::none_ivf_sample_filter>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  float,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::none_ivf_sample_filter>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA false>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::none_ivf_sample_filter>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA true>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::none_ivf_sample_filter>);
-
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA false>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA true>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  half,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  half,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  float,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA false>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA true>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    uint32_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA false>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    int64_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA true>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    int64_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  half,
-  half,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    int64_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  half,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    int64_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  float,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    int64_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA false>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    int64_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>>);
-instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select(
-  float,
-  raft::neighbors::ivf_pq::detail::fp_8bit<5u COMMA true>,
-  raft::neighbors::filtering::ivf_to_sample_filter<
-    int64_t COMMA raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>>);
-#undef COMMA
-
-#undef instantiate_raft_neighbors_ivf_pq_detail_compute_similarity_select
diff --git a/cpp/internal/raft_internal/neighbors/ivf_pq_search_test-ext.cuh b/cpp/internal/raft_internal/neighbors/ivf_pq_search_test-ext.cuh
deleted file mode 100644
index 1e6f4f9976..0000000000
--- a/cpp/internal/raft_internal/neighbors/ivf_pq_search_test-ext.cuh
+++ /dev/null
@@ -1,89 +0,0 @@
-/*
- * Copyright (c) 2024, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#include <raft/core/device_mdspan.hpp>  // raft::device_matrix_view
-#include <raft/core/resources.hpp>      // raft::resources
-#include <raft/neighbors/ivf_pq.cuh>
-#include <raft/neighbors/ivf_pq_types.hpp>  // raft::neighbors::ivf_pq::index
-#include <raft/neighbors/sample_filter.cuh>
-#include <raft/neighbors/sample_filter_types.hpp>
-
-#include <raft_internal/neighbors/ivf_pq_compute_similarity_filters_test-ext.cuh>
-
-#include <rmm/resource_ref.hpp>
-
-#include <cstdint>  // int64_t
-
-#define instantiate_raft_neighbors_ivf_pq_search(T, IdxT)            \
-  extern template void raft::neighbors::ivf_pq::search<T, IdxT>(     \
-    raft::resources const& handle,                                   \
-    const raft::neighbors::ivf_pq::search_params& params,            \
-    const raft::neighbors::ivf_pq::index<IdxT>& idx,                 \
-    raft::device_matrix_view<const T, uint32_t, row_major> queries,  \
-    raft::device_matrix_view<IdxT, uint32_t, row_major> neighbors,   \
-    raft::device_matrix_view<float, uint32_t, row_major> distances); \
-                                                                     \
-  extern template void raft::neighbors::ivf_pq::search<T, IdxT>(     \
-    raft::resources const& handle,                                   \
-    const raft::neighbors::ivf_pq::search_params& params,            \
-    const raft::neighbors::ivf_pq::index<IdxT>& idx,                 \
-    const T* queries,                                                \
-    uint32_t n_queries,                                              \
-    uint32_t k,                                                      \
-    IdxT* neighbors,                                                 \
-    float* distances);                                               \
-                                                                     \
-  extern template void raft::neighbors::ivf_pq::search<T, IdxT>(     \
-    raft::resources const& handle,                                   \
-    const raft::neighbors::ivf_pq::search_params& params,            \
-    const raft::neighbors::ivf_pq::index<IdxT>& idx,                 \
-    const T* queries,                                                \
-    uint32_t n_queries,                                              \
-    uint32_t k,                                                      \
-    IdxT* neighbors,                                                 \
-    float* distances)
-
-instantiate_raft_neighbors_ivf_pq_search(float, uint32_t);
-
-#undef instantiate_raft_neighbors_ivf_pq_search
-
-#define instantiate_raft_neighbors_ivf_pq_search_with_filtering(T, IdxT, FilterT)        \
-  extern template void raft::neighbors::ivf_pq::search_with_filtering<T, IdxT, FilterT>( \
-    raft::resources const& handle,                                                       \
-    const search_params& params,                                                         \
-    const index<IdxT>& idx,                                                              \
-    raft::device_matrix_view<const T, uint32_t, row_major> queries,                      \
-    raft::device_matrix_view<IdxT, uint32_t, row_major> neighbors,                       \
-    raft::device_matrix_view<float, uint32_t, row_major> distances,                      \
-    FilterT sample_filter)
-
-#define COMMA ,
-instantiate_raft_neighbors_ivf_pq_search_with_filtering(
-  float, uint32_t, raft::neighbors::filtering::bitset_filter<uint32_t COMMA uint32_t>);
-
-instantiate_raft_neighbors_ivf_pq_search_with_filtering(
-  float, uint32_t, raft::neighbors::filtering::none_ivf_sample_filter);
-
-instantiate_raft_neighbors_ivf_pq_search_with_filtering(
-  float, int64_t, raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>);
-
-instantiate_raft_neighbors_ivf_pq_search_with_filtering(
-  int8_t, int64_t, raft::neighbors::filtering::bitset_filter<uint32_t COMMA int64_t>);
-
-#undef COMMA
-#undef instantiate_raft_neighbors_ivf_pq_search_with_filtering
diff --git a/cpp/internal/raft_internal/neighbors/naive_knn.cuh b/cpp/internal/raft_internal/neighbors/naive_knn.cuh
deleted file mode 100644
index c14a8e3e9f..0000000000
--- a/cpp/internal/raft_internal/neighbors/naive_knn.cuh
+++ /dev/null
@@ -1,124 +0,0 @@
-/*
- * Copyright (c) 2023-2024, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#include <raft/core/resource/cuda_stream.hpp>
-#include <raft/core/resource/device_memory_resource.hpp>
-#include <raft/distance/distance_types.hpp>
-#include <raft/matrix/detail/select_k.cuh>
-#include <raft/spatial/knn/detail/ann_utils.cuh>
-#include <raft/util/cuda_utils.cuh>
-
-#include <rmm/device_uvector.hpp>
-
-namespace raft::neighbors {
-
-template <typename EvalT, typename DataT, typename IdxT>
-RAFT_KERNEL naive_distance_kernel(EvalT* dist,
-                                  const DataT* x,
-                                  const DataT* y,
-                                  IdxT m,
-                                  IdxT n,
-                                  IdxT k,
-                                  raft::distance::DistanceType metric)
-{
-  IdxT midx = IdxT(threadIdx.x) + IdxT(blockIdx.x) * IdxT(blockDim.x);
-  if (midx >= m) return;
-  IdxT grid_size = IdxT(blockDim.y) * IdxT(gridDim.y);
-  for (IdxT nidx = threadIdx.y + blockIdx.y * blockDim.y; nidx < n; nidx += grid_size) {
-    EvalT acc = EvalT(0);
-    for (IdxT i = 0; i < k; ++i) {
-      IdxT xidx = i + midx * k;
-      IdxT yidx = i + nidx * k;
-      auto xv   = EvalT(x[xidx]);
-      auto yv   = EvalT(y[yidx]);
-      switch (metric) {
-        case raft::distance::DistanceType::InnerProduct: {
-          acc += xv * yv;
-        } break;
-        case raft::distance::DistanceType::L2SqrtExpanded:
-        case raft::distance::DistanceType::L2SqrtUnexpanded:
-        case raft::distance::DistanceType::L2Expanded:
-        case raft::distance::DistanceType::L2Unexpanded: {
-          auto diff = xv - yv;
-          acc += diff * diff;
-        } break;
-        default: break;
-      }
-    }
-    switch (metric) {
-      case raft::distance::DistanceType::L2SqrtExpanded:
-      case raft::distance::DistanceType::L2SqrtUnexpanded: {
-        acc = raft::sqrt(acc);
-      } break;
-      default: break;
-    }
-    dist[midx * n + nidx] = acc;
-  }
-}
-
-/**
- * Naive, but flexible bruteforce KNN search.
- *
- * TODO: either replace this with brute_force_knn or with distance+select_k
- *       when either distance or brute_force_knn support 8-bit int inputs.
- */
-template <typename EvalT, typename DataT, typename IdxT>
-void naive_knn(raft::resources const& handle,
-               EvalT* dist_topk,
-               IdxT* indices_topk,
-               const DataT* x,
-               const DataT* y,
-               size_t n_inputs,
-               size_t input_len,
-               size_t dim,
-               uint32_t k,
-               raft::distance::DistanceType type)
-{
-  auto mr     = resource::get_workspace_resource(handle);
-  auto stream = raft::resource::get_cuda_stream(handle);
-  dim3 block_dim(16, 32, 1);
-  // maximum reasonable grid size in `y` direction
-  auto grid_y =
-    static_cast<uint16_t>(std::min<size_t>(raft::ceildiv<size_t>(input_len, block_dim.y), 32768));
-
-  // bound the memory used by this function
-  size_t max_batch_size =
-    std::min<size_t>(n_inputs, raft::ceildiv<size_t>(size_t(1) << size_t(27), input_len));
-  rmm::device_uvector<EvalT> dist(max_batch_size * input_len, stream, mr);
-
-  for (size_t offset = 0; offset < n_inputs; offset += max_batch_size) {
-    size_t batch_size = std::min(max_batch_size, n_inputs - offset);
-    dim3 grid_dim(raft::ceildiv<size_t>(batch_size, block_dim.x), grid_y, 1);
-
-    naive_distance_kernel<EvalT, DataT, IdxT><<<grid_dim, block_dim, 0, stream>>>(
-      dist.data(), x + offset * dim, y, batch_size, input_len, dim, type);
-
-    matrix::detail::select_k<EvalT, IdxT>(handle,
-                                          dist.data(),
-                                          nullptr,
-                                          batch_size,
-                                          input_len,
-                                          static_cast<int>(k),
-                                          dist_topk + offset * k,
-                                          indices_topk + offset * k,
-                                          type != raft::distance::DistanceType::InnerProduct);
-  }
-  RAFT_CUDA_TRY(cudaStreamSynchronize(stream));
-}
-
-}  // namespace raft::neighbors
diff --git a/cpp/internal/raft_internal/neighbors/refine_helper.cuh b/cpp/internal/raft_internal/neighbors/refine_helper.cuh
deleted file mode 100644
index 665ec23d8e..0000000000
--- a/cpp/internal/raft_internal/neighbors/refine_helper.cuh
+++ /dev/null
@@ -1,158 +0,0 @@
-/*
- * Copyright (c) 2022-2024, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-#pragma once
-
-#include <raft/core/device_mdarray.hpp>
-#include <raft/core/device_mdspan.hpp>
-#include <raft/core/host_mdarray.hpp>
-#include <raft/core/host_mdspan.hpp>
-#include <raft/core/resource/cuda_stream.hpp>
-#include <raft/core/resources.hpp>
-#include <raft/distance/distance_types.hpp>
-#include <raft/random/rng.cuh>
-
-#include <raft_internal/neighbors/naive_knn.cuh>
-
-#include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_uvector.hpp>
-
-namespace raft::neighbors {
-
-template <typename IdxT>
-struct RefineInputs {
-  IdxT n_queries;
-  IdxT n_rows;
-  IdxT dim;
-  IdxT k;   // after refinement
-  IdxT k0;  // initial k before refinement (k0 >= k).
-  raft::distance::DistanceType metric;
-  bool host_data;
-};
-
-/** Helper class to allocate arrays and generate input data for refinement test and benchmark. */
-template <typename DataT, typename DistanceT, typename IdxT>
-class RefineHelper {
- public:
-  RefineHelper(const raft::resources& handle, RefineInputs<IdxT> params)
-    : handle_(handle),
-      stream_(resource::get_cuda_stream(handle)),
-      p(params),
-      dataset(handle),
-      queries(handle),
-      refined_distances(handle),
-      refined_indices(handle),
-      candidates(handle),
-      dataset_host(handle),
-      queries_host(handle),
-      candidates_host(handle),
-      refined_distances_host(handle),
-      refined_indices_host(handle)
-  {
-    raft::random::RngState rng(1234ULL);
-
-    dataset = raft::make_device_matrix<DataT, IdxT>(handle_, p.n_rows, p.dim);
-    queries = raft::make_device_matrix<DataT, IdxT>(handle_, p.n_queries, p.dim);
-    if constexpr (std::is_same<DataT, float>{}) {
-      raft::random::uniform(
-        handle, rng, dataset.data_handle(), dataset.size(), DataT(-10.0), DataT(10.0));
-      raft::random::uniform(
-        handle, rng, queries.data_handle(), queries.size(), DataT(-10.0), DataT(10.0));
-    } else {
-      raft::random::uniformInt(
-        handle, rng, dataset.data_handle(), dataset.size(), DataT(1), DataT(20));
-      raft::random::uniformInt(
-        handle, rng, queries.data_handle(), queries.size(), DataT(1), DataT(20));
-    }
-
-    refined_distances = raft::make_device_matrix<DistanceT, IdxT>(handle_, p.n_queries, p.k);
-    refined_indices   = raft::make_device_matrix<IdxT, IdxT>(handle_, p.n_queries, p.k);
-
-    // Generate candidate vectors
-    {
-      candidates = raft::make_device_matrix<IdxT, IdxT>(handle_, p.n_queries, p.k0);
-      rmm::device_uvector<DistanceT> distances_tmp(p.n_queries * p.k0, stream_);
-      naive_knn<DistanceT, DataT, IdxT>(handle_,
-                                        distances_tmp.data(),
-                                        candidates.data_handle(),
-                                        queries.data_handle(),
-                                        dataset.data_handle(),
-                                        p.n_queries,
-                                        p.n_rows,
-                                        p.dim,
-                                        p.k0,
-                                        p.metric);
-      resource::sync_stream(handle_, stream_);
-    }
-
-    if (p.host_data) {
-      dataset_host    = raft::make_host_matrix<DataT, IdxT>(p.n_rows, p.dim);
-      queries_host    = raft::make_host_matrix<DataT, IdxT>(p.n_queries, p.dim);
-      candidates_host = raft::make_host_matrix<IdxT, IdxT>(p.n_queries, p.k0);
-
-      raft::copy(dataset_host.data_handle(), dataset.data_handle(), dataset.size(), stream_);
-      raft::copy(queries_host.data_handle(), queries.data_handle(), queries.size(), stream_);
-      raft::copy(
-        candidates_host.data_handle(), candidates.data_handle(), candidates.size(), stream_);
-
-      refined_distances_host = raft::make_host_matrix<DistanceT, IdxT>(p.n_queries, p.k);
-      refined_indices_host   = raft::make_host_matrix<IdxT, IdxT>(p.n_queries, p.k);
-      resource::sync_stream(handle_, stream_);
-    }
-
-    // Generate ground thruth for testing.
-    {
-      rmm::device_uvector<DistanceT> distances_dev(p.n_queries * p.k, stream_);
-      rmm::device_uvector<IdxT> indices_dev(p.n_queries * p.k, stream_);
-      naive_knn<DistanceT, DataT, IdxT>(handle_,
-                                        distances_dev.data(),
-                                        indices_dev.data(),
-                                        queries.data_handle(),
-                                        dataset.data_handle(),
-                                        p.n_queries,
-                                        p.n_rows,
-                                        p.dim,
-                                        p.k,
-                                        p.metric);
-      true_refined_distances_host.resize(p.n_queries * p.k);
-      true_refined_indices_host.resize(p.n_queries * p.k);
-      raft::copy(true_refined_indices_host.data(), indices_dev.data(), indices_dev.size(), stream_);
-      raft::copy(
-        true_refined_distances_host.data(), distances_dev.data(), distances_dev.size(), stream_);
-      resource::sync_stream(handle_, stream_);
-    }
-  }
-
- public:
-  RefineInputs<IdxT> p;
-  const raft::resources& handle_;
-  rmm::cuda_stream_view stream_;
-
-  raft::device_matrix<DataT, IdxT, row_major> dataset;
-  raft::device_matrix<DataT, IdxT, row_major> queries;
-  raft::device_matrix<IdxT, IdxT, row_major> candidates;  // Neighbor candidate indices
-  raft::device_matrix<IdxT, IdxT, row_major> refined_indices;
-  raft::device_matrix<DistanceT, IdxT, row_major> refined_distances;
-
-  raft::host_matrix<DataT, IdxT, row_major> dataset_host;
-  raft::host_matrix<DataT, IdxT, row_major> queries_host;
-  raft::host_matrix<IdxT, IdxT, row_major> candidates_host;
-  raft::host_matrix<IdxT, IdxT, row_major> refined_indices_host;
-  raft::host_matrix<DistanceT, IdxT, row_major> refined_distances_host;
-
-  std::vector<IdxT> true_refined_indices_host;
-  std::vector<DistanceT> true_refined_distances_host;
-};
-}  // namespace raft::neighbors