diff --git a/CMakeLists.txt b/CMakeLists.txt
index 0b7184100..157d6903d 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -12,8 +12,7 @@
 # License for the specific language governing permissions and limitations under
 # the License
 
-cmake_minimum_required(VERSION 3.23.0 FATAL_ERROR)
-project(knowhere CXX C)
+cmake_minimum_required(VERSION 3.26.4 FATAL_ERROR)
 
 set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
 list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/")
@@ -21,10 +20,20 @@ include(GNUInstallDirs)
 include(ExternalProject)
 include(cmake/utils/utils.cmake)
 
+knowhere_option(WITH_RAFT "Build with RAFT indexes" OFF)
+if (WITH_RAFT)
+  if("${CMAKE_CUDA_ARCHITECTURES}" STREQUAL "")
+    set(CMAKE_CUDA_ARCHITECTURES RAPIDS)
+  endif()
+  include(cmake/libs/librapids.cmake)
+  project(knowhere CXX C CUDA)
+else()
+  project(knowhere CXX C)
+endif()
+
 knowhere_option(WITH_UT "Build with UT test" OFF)
 knowhere_option(WITH_ASAN "Build with ASAN" OFF)
 knowhere_option(WITH_DISKANN "Build with diskann index" OFF)
-knowhere_option(WITH_RAFT "Build with RAFT indexes" OFF)
 knowhere_option(WITH_BENCHMARK "Build with benchmark" OFF)
 knowhere_option(WITH_COVERAGE "Build with coverage" OFF)
 knowhere_option(WITH_CCACHE "Build with ccache" ON)
@@ -64,18 +73,6 @@ endif()
 
 list( APPEND CMAKE_MODULE_PATH ${CMAKE_BINARY_DIR}/)
 
-if(WITH_RAFT)
-  if("${CMAKE_CUDA_ARCHITECTURES}" STREQUAL "")
-    set(CMAKE_CUDA_ARCHITECTURES 86;80;75;70;61)
-  endif()
-  enable_language(CUDA)
-  find_package(CUDAToolkit REQUIRED)
-  if(${CUDAToolkit_VERSION_MAJOR} GREATER 10)
-    # cuda11 support --threads for compile some large .cu more efficient
-    add_compile_options($<$<COMPILE_LANGUAGE:CUDA>:--threads=4>)
-  endif()
-endif()
-
 add_definitions(-DNOT_COMPILE_FOR_SWIG)
 
 include(cmake/utils/compile_flags.cmake)
@@ -113,8 +110,7 @@ if(WITH_COVERAGE)
 endif()
 
 knowhere_file_glob(GLOB_RECURSE KNOWHERE_SRCS src/common/*.cc src/index/*.cc
-                   src/io/*.cc src/index/*.cu src/common/raft/*.cu
-                   src/common/raft/*.cc)
+                   src/io/*.cc src/common/*.cu src/index/*.cu src/io/*.cu)
 
 set(KNOWHERE_LINKER_LIBS "")
 
@@ -127,13 +123,13 @@ else()
 endif()
 
 knowhere_file_glob(GLOB_RECURSE KNOWHERE_GPU_SRCS src/index/gpu/flat_gpu/*.cc
-                     src/index/gpu/ivf_gpu/*.cc src/index/cagra/*.cu)
+                     src/index/gpu/ivf_gpu/*.cc)
 list(REMOVE_ITEM KNOWHERE_SRCS ${KNOWHERE_GPU_SRCS})
 
 if(NOT WITH_RAFT)
-  knowhere_file_glob(GLOB_RECURSE KNOWHERE_RAFT_SRCS src/index/ivf_raft/*.cc
-                     src/index/ivf_raft/*.cu src/index/cagra/*.cu
-                     src/common/raft/*.cu src/common/raft/*.cc)
+  knowhere_file_glob(GLOB_RECURSE KNOWHERE_RAFT_SRCS
+                     src/common/raft/*.cu src/common/raft/*.cc
+                     src/index/gpu_raft/*.cc)
   list(REMOVE_ITEM KNOWHERE_SRCS ${KNOWHERE_RAFT_SRCS})
 endif()
 
@@ -150,7 +146,7 @@ list(APPEND KNOWHERE_LINKER_LIBS ${FOLLY_LIBRARIES})
 add_library(knowhere SHARED ${KNOWHERE_SRCS})
 add_dependencies(knowhere ${KNOWHERE_LINKER_LIBS})
 if(WITH_RAFT)
-  list(APPEND KNOWHERE_LINKER_LIBS raft::raft)
+  list(APPEND KNOWHERE_LINKER_LIBS raft::raft CUDA::cublas CUDA::cusparse CUDA::cusolver)
 endif()
 target_link_libraries(knowhere PUBLIC ${KNOWHERE_LINKER_LIBS})
 target_include_directories(
diff --git a/benchmark/hdf5/benchmark_float_qps.cpp b/benchmark/hdf5/benchmark_float_qps.cpp
index 2022f4468..bc87e2a44 100644
--- a/benchmark/hdf5/benchmark_float_qps.cpp
+++ b/benchmark/hdf5/benchmark_float_qps.cpp
@@ -73,6 +73,60 @@ class Benchmark_float_qps : public Benchmark_knowhere, public ::testing::Test {
         }
     }
 
+    void
+    test_cagra(const knowhere::Json& cfg) {
+        auto conf = cfg;
+
+        auto find_smallest_max_iters = [&](float expected_recall) -> int32_t {
+            auto ds_ptr = knowhere::GenDataSet(nq_, dim_, xq_);
+            auto left = 32;
+            auto right = 256;
+            auto max_iterations = left;
+
+            float recall;
+            while (left <= right) {
+                max_iterations = left + (right - left) / 2;
+                conf[knowhere::indexparam::MAX_ITERATIONS] = max_iterations;
+
+                auto result = index_.Search(*ds_ptr, conf, nullptr);
+                recall = CalcRecall(result.value()->GetIds(), nq_, topk_);
+                printf(
+                    "[%0.3f s] iterate CAGRA param for recall %.4f: max_iterations=%d, k=%d, "
+                    "R@=%.4f\n",
+                    get_time_diff(), expected_recall, max_iterations, topk_, recall);
+                std::fflush(stdout);
+                if (std::abs(recall - expected_recall) <= 0.0001) {
+                    return max_iterations;
+                }
+                if (recall < expected_recall) {
+                    left = max_iterations + 1;
+                } else {
+                    right = max_iterations - 1;
+                }
+            }
+            return left;
+        };
+
+        for (auto expected_recall : EXPECTED_RECALLs_) {
+            conf[knowhere::indexparam::ITOPK_SIZE] = ((int{topk_} + 32 - 1) / 32) * 32;
+            conf[knowhere::meta::TOPK] = topk_;
+            conf[knowhere::indexparam::MAX_ITERATIONS] = find_smallest_max_iters(expected_recall);
+
+            printf(
+                "\n[%0.3f s] %s | %s | k=%d, "
+                "R@=%.4f\n",
+                get_time_diff(), ann_test_name_.c_str(), index_type_.c_str(), topk_, expected_recall);
+            printf("================================================================================\n");
+            for (auto thread_num : THREAD_NUMs_) {
+                CALC_TIME_SPAN(task(conf, thread_num, nq_));
+                printf("  thread_num = %2d, elapse = %6.3fs, VPS = %.3f\n", thread_num, t_diff, nq_ / t_diff);
+                std::fflush(stdout);
+            }
+            printf("================================================================================\n");
+            printf("[%.3f s] Test '%s/%s' done\n\n", get_time_diff(), ann_test_name_.c_str(), index_type_.c_str());
+        }
+    }
+
     void
     test_hnsw(const knowhere::Json& cfg) {
         auto conf = cfg;
@@ -221,6 +275,9 @@ class Benchmark_float_qps : public Benchmark_knowhere, public ::testing::Test {
 #ifdef KNOWHERE_WITH_GPU
         knowhere::KnowhereConfig::InitGPUResource(GPU_DEVICE_ID, 2);
         cfg_[knowhere::meta::DEVICE_ID] = GPU_DEVICE_ID;
+#endif
+#ifdef KNOWHERE_WITH_RAFT
+        knowhere::KnowhereConfig::SetRaftMemPool();
 #endif
     }
 
@@ -251,6 +308,9 @@ class Benchmark_float_qps : public Benchmark_knowhere, public ::testing::Test {
     // SCANN index params
     const std::vector<int32_t> SCANN_REORDER_K = {256, 512, 768, 1024};
     const std::vector<bool> SCANN_WITH_RAW_DATA = {true};
+
+    // CAGRA index params
+    const std::vector<int32_t> GRAPH_DEGREE_ = {32, 64};
 };
 
 TEST_F(Benchmark_float_qps, TEST_IVF_FLAT) {
@@ -344,3 +404,15 @@ TEST_F(Benchmark_float_qps, TEST_SCANN) {
         }
     }
 }
+TEST_F(Benchmark_float_qps, TEST_CAGRA) {
+    index_type_ = knowhere::IndexEnum::INDEX_RAFT_CAGRA;
+    knowhere::Json conf = cfg_;
+    for (auto gd : GRAPH_DEGREE_) {
+        conf[knowhere::indexparam::GRAPH_DEGREE] = gd;
+        conf[knowhere::indexparam::INTERMEDIATE_GRAPH_DEGREE] = gd;
+        conf[knowhere::indexparam::MAX_ITERATIONS] = 64;
+        std::string index_file_name = get_index_name({gd});
+        create_index(index_file_name, conf);
+        test_cagra(conf);
+    }
+}
diff --git a/benchmark/hdf5/benchmark_knowhere.h b/benchmark/hdf5/benchmark_knowhere.h
index 0505d87e2..909130381 100644
--- a/benchmark/hdf5/benchmark_knowhere.h
+++ b/benchmark/hdf5/benchmark_knowhere.h
@@ -81,7 +81,6 @@ class Benchmark_knowhere : public Benchmark_hdf5 {
             bin->size = dim_ * nb_ * sizeof(float);
             binary_set.Append("RAW_DATA", bin);
         }
-
         index.Deserialize(binary_set, conf);
     }
 
diff --git a/cmake/libs/libraft.cmake b/cmake/libs/libraft.cmake
index 3de1cfaec..514793933 100644
--- a/cmake/libs/libraft.cmake
+++ b/cmake/libs/libraft.cmake
@@ -14,22 +14,15 @@
 # the License.
 
 add_definitions(-DKNOWHERE_WITH_RAFT)
-include(cmake/utils/fetch_rapids.cmake)
-include(rapids-cmake)
-include(rapids-cpm)
-include(rapids-cuda)
-include(rapids-export)
-include(rapids-find)
-
-rapids_cpm_init()
-
-set(CMAKE_CUDA_FLAGS
-    "${CMAKE_CUDA_FLAGS} --expt-extended-lambda --expt-relaxed-constexpr")
-
-set(RAPIDS_VERSION 23.04)
 set(RAFT_VERSION "${RAPIDS_VERSION}")
 set(RAFT_FORK "rapidsai")
-set(RAFT_PINNED_TAG "branch-${RAPIDS_VERSION}")
+set(RAFT_PINNED_TAG "branch-23.12")
+
+
+rapids_find_package(CUDAToolkit REQUIRED
+  BUILD_EXPORT_SET knowhere-exports
+  INSTALL_EXPORT_SET knowhere-exports
+)
 
 function(find_and_configure_raft)
   set(oneValueArgs VERSION FORK PINNED_TAG)
diff --git a/cmake/utils/fetch_rapids.cmake b/cmake/libs/librapids.cmake
similarity index 71%
rename from cmake/utils/fetch_rapids.cmake
rename to cmake/libs/librapids.cmake
index 56899f2c5..315c4c3d2 100644
--- a/cmake/utils/fetch_rapids.cmake
+++ b/cmake/libs/librapids.cmake
@@ -13,7 +13,7 @@
 # License for the specific language governing permissions and limitations under
 # the License.
 
-set(RAPIDS_VERSION "23.04")
+set(RAPIDS_VERSION 23.12)
 
 if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/RAPIDS.cmake)
   file(
@@ -22,3 +22,15 @@ if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/RAPIDS.cmake)
     ${CMAKE_CURRENT_BINARY_DIR}/RAPIDS.cmake)
 endif()
 include(${CMAKE_CURRENT_BINARY_DIR}/RAPIDS.cmake)
+
+include(rapids-cpm)  # Dependency tracking
+include(rapids-find) # Wrappers for finding packages
+include(rapids-cuda) # Common CMake CUDA logic
+
+rapids_cuda_init_architectures(knowhere)
+message(STATUS "INIT: ${CMAKE_CUDA_ARCHITECTURES}")
+
+rapids_cpm_init()
+
+set(CMAKE_CUDA_FLAGS
+    "${CMAKE_CUDA_FLAGS} --expt-extended-lambda --expt-relaxed-constexpr")
diff --git a/include/knowhere/comp/index_param.h b/include/knowhere/comp/index_param.h
index 675ccc35c..82f748145 100644
--- a/include/knowhere/comp/index_param.h
+++ b/include/knowhere/comp/index_param.h
@@ -81,6 +81,35 @@ constexpr const char* M = "m";          // PQ param for IVFPQ
 constexpr const char* SSIZE = "ssize";
 constexpr const char* REORDER_K = "reorder_k";
 constexpr const char* WITH_RAW_DATA = "with_raw_data";
+// RAFT Params
+constexpr const char* REFINE_RATIO = "refine_ratio";
+// RAFT-specific IVF Params
+constexpr const char* KMEANS_N_ITERS = "kmeans_n_iters";
+constexpr const char* KMEANS_TRAINSET_FRACTION = "kmeans_trainset_fraction";
+constexpr const char* ADAPTIVE_CENTERS = "adaptive_centers";                              // IVF FLAT
+constexpr const char* CODEBOOK_KIND = "codebook_kind";                                    // IVF PQ
+constexpr const char* FORCE_RANDOM_ROTATION = "force_random_rotation";                    // IVF PQ
+constexpr const char* CONSERVATIVE_MEMORY_ALLOCATION = "conservative_memory_allocation";  // IVF PQ
+constexpr const char* LUT_DTYPE = "lut_dtype";                                            // IVF PQ
+constexpr const char* INTERNAL_DISTANCE_DTYPE = "internal_distance_dtype";                // IVF PQ
+constexpr const char* PREFERRED_SHMEM_CARVEOUT = "preferred_shmem_carveout";              // IVF PQ
+
+// CAGRA Params
+constexpr const char* INTERMEDIATE_GRAPH_DEGREE = "intermediate_graph_degree";
+constexpr const char* GRAPH_DEGREE = "graph_degree";
+constexpr const char* ITOPK_SIZE = "itopk_size";
+constexpr const char* MAX_QUERIES = "max_queries";
+constexpr const char* BUILD_ALGO = "build_algo";
+constexpr const char* SEARCH_ALGO = "search_algo";
+constexpr const char* TEAM_SIZE = "team_size";
+constexpr const char* SEARCH_WIDTH = "search_width";
+constexpr const char* MIN_ITERATIONS = "min_iterations";
+constexpr const char* MAX_ITERATIONS = "max_iterations";
+constexpr const char* THREAD_BLOCK_SIZE = "thread_block_size";
+constexpr const char* HASHMAP_MODE = "hashmap_mode";
+constexpr const char* HASHMAP_MIN_BITLEN = "hashmap_min_bitlen";
+constexpr const char* HASHMAP_MAX_FILL_RATE = "hashmap_max_fill_rate";
+constexpr const char* NN_DESCENT_NITER = "nn_descent_niter";
 
 // HNSW Params
 constexpr const char* EFCONSTRUCTION = "efConstruction";
diff --git a/include/knowhere/comp/knowhere_config.h b/include/knowhere/comp/knowhere_config.h
index de845b6cf..ba224cca7 100644
--- a/include/knowhere/comp/knowhere_config.h
+++ b/include/knowhere/comp/knowhere_config.h
@@ -107,6 +107,12 @@ class KnowhereConfig {
     */
     static void
     SetRaftMemPool(size_t init_size, size_t max_size);
+
+    /**
+     * Initialize RAFT with defaults
+     */
+    static void
+    SetRaftMemPool();
 };
 
 }  // namespace knowhere
diff --git a/include/knowhere/device_bitset.h b/include/knowhere/device_bitset.h
deleted file mode 100644
index 12532dc4e..000000000
--- a/include/knowhere/device_bitset.h
+++ /dev/null
@@ -1,90 +0,0 @@
-// Copyright (C) 2023 Zilliz. All rights reserved.
-//
-// 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.
-
-#ifndef DEVICE_BITSET_H
-#define DEVICE_BITSET_H
-
-#include "knowhere/bitsetview.h"
-#include "raft/core/device_mdarray.hpp"
-#include "raft/core/device_resources.hpp"
-#include "raft/util/cudart_utils.hpp"
-
-namespace knowhere {
-
-struct DeviceBitsetView {
-    __device__ __host__
-    DeviceBitsetView(const DeviceBitsetView& other)
-        : bits_{other.data()}, num_bits_{other.size()} {
-    }
-    __device__ __host__
-    DeviceBitsetView(const uint8_t* data, size_t num_bits = size_t{})
-        : bits_{data}, num_bits_{num_bits} {
-    }
-
-    __device__ __host__ bool
-    empty() const {
-        return num_bits_ == 0;
-    }
-
-    __device__ __host__ size_t
-    size() const {
-        return num_bits_;
-    }
-
-    __device__ __host__ size_t
-    byte_size() const {
-        return (num_bits_ + 8 - 1) >> 3;
-    }
-
-    __device__ __host__ const uint8_t*
-    data() const {
-        return bits_;
-    }
-
-    __device__ bool
-    test(int64_t index) const {
-        auto result = false;
-        if (index < num_bits_) {
-            result = bits_[index >> 3] & (0x1 << (index & 0x7));
-        }
-        return result;
-    }
-
- private:
-    const uint8_t* bits_ = nullptr;
-    size_t num_bits_ = 0;
-};
-
-struct DeviceBitset {
-    DeviceBitset(raft::device_resources& res, BitsetView const& other)
-        : storage_{[&res, &other]() {
-              auto result = raft::make_device_vector<uint8_t, uint32_t>(res, other.byte_size());
-              if (!other.empty()) {
-                  raft::copy(result.data_handle(), other.data(), other.byte_size(), res.get_stream());
-              }
-              return result;
-          }()},
-          num_bits_{other.size()} {
-    }
-
-    auto
-    view() {
-        return DeviceBitsetView{storage_.data_handle(), num_bits_};
-    }
-
- private:
-    raft::device_vector<uint8_t> storage_;
-    size_t num_bits_;
-};
-
-}  // namespace knowhere
-
-#endif /* DEVICE_BITSET_H */
diff --git a/src/common/comp/knowhere_config.cc b/src/common/comp/knowhere_config.cc
index 126529d21..50bb23b85 100644
--- a/src/common/comp/knowhere_config.cc
+++ b/src/common/comp/knowhere_config.cc
@@ -23,10 +23,10 @@
 #ifdef KNOWHERE_WITH_GPU
 #include "index/gpu/gpu_res_mgr.h"
 #endif
-#include "simd/hook.h"
 #ifdef KNOWHERE_WITH_RAFT
-#include "common/raft/raft_utils.h"
+#include "common/raft/integration/raft_initialization.hpp"
 #endif
+#include "simd/hook.h"
 
 namespace knowhere {
 
@@ -163,7 +163,20 @@ KnowhereConfig::FreeGPUResource() {
 void
 KnowhereConfig::SetRaftMemPool(size_t init_size, size_t max_size) {
 #ifdef KNOWHERE_WITH_RAFT
-    raft_utils::set_mem_pool_size(init_size, max_size);
+    auto config = raft_knowhere::raft_configuration{};
+    config.init_mem_pool_size_mb = init_size;
+    config.max_mem_pool_size_mb = max_size;
+    // This should probably be a separate configuration option, but fine for now
+    config.max_workspace_size_mb = max_size;
+    raft_knowhere::initialize_raft(config);
+#endif
+}
+void
+KnowhereConfig::SetRaftMemPool() {
+    // Overload for default values
+#ifdef KNOWHERE_WITH_RAFT
+    auto config = raft_knowhere::raft_configuration{};
+    raft_knowhere::initialize_raft(config);
 #endif
 }
 
diff --git a/src/common/raft/integration/cagra_index.cu b/src/common/raft/integration/cagra_index.cu
new file mode 100644
index 000000000..4879f5c39
--- /dev/null
+++ b/src/common/raft/integration/cagra_index.cu
@@ -0,0 +1,21 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "common/raft/integration/raft_knowhere_index.cuh"
+namespace raft_knowhere {
+template struct raft_knowhere_index<raft_proto::raft_index_kind::cagra>;
+}  // namespace raft_knowhere
diff --git a/src/common/raft/integration/ivf_flat_index.cu b/src/common/raft/integration/ivf_flat_index.cu
new file mode 100644
index 000000000..fc759075d
--- /dev/null
+++ b/src/common/raft/integration/ivf_flat_index.cu
@@ -0,0 +1,21 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "common/raft/integration/raft_knowhere_index.cuh"
+namespace raft_knowhere {
+template struct raft_knowhere_index<raft_proto::raft_index_kind::ivf_flat>;
+}  // namespace raft_knowhere
diff --git a/src/common/raft/integration/ivf_pq_index.cu b/src/common/raft/integration/ivf_pq_index.cu
new file mode 100644
index 000000000..9284a0930
--- /dev/null
+++ b/src/common/raft/integration/ivf_pq_index.cu
@@ -0,0 +1,21 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "common/raft/integration/raft_knowhere_index.cuh"
+namespace raft_knowhere {
+template struct raft_knowhere_index<raft_proto::raft_index_kind::ivf_pq>;
+}  // namespace raft_knowhere
diff --git a/src/common/raft/integration/raft_initialization.cc b/src/common/raft/integration/raft_initialization.cc
new file mode 100644
index 000000000..59bb6f842
--- /dev/null
+++ b/src/common/raft/integration/raft_initialization.cc
@@ -0,0 +1,79 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#include "common/raft/integration/raft_initialization.hpp"
+
+#include <cuda_runtime_api.h>
+
+#include <cstddef>
+#include <raft/core/device_resources_manager.hpp>
+#include <raft/core/device_setter.hpp>
+#include <raft/core/resource/device_memory_resource.hpp>
+namespace raft_knowhere {
+
+void
+initialize_raft(raft_configuration const& config) {
+    auto static initialization_flag = std::once_flag{};
+    std::call_once(initialization_flag, [&config]() {
+        raft::device_resources_manager::set_streams_per_device(config.streams_per_device);
+        if (config.stream_pool_size) {
+            raft::device_resources_manager::set_stream_pools_per_device(config.stream_pools_per_device,
+                                                                        *(config.stream_pool_size));
+        } else {
+            raft::device_resources_manager::set_stream_pools_per_device(config.stream_pools_per_device);
+        }
+        if (config.init_mem_pool_size_mb) {
+            raft::device_resources_manager::set_init_mem_pool_size(*(config.init_mem_pool_size_mb) << 20);
+        }
+        if (config.max_mem_pool_size_mb) {
+            if (*config.max_mem_pool_size_mb > 0) {
+                raft::device_resources_manager::set_max_mem_pool_size(*(config.max_mem_pool_size_mb) << 20);
+            }
+        } else {
+            raft::device_resources_manager::set_max_mem_pool_size(std::nullopt);
+        }
+        if (config.max_workspace_size_mb) {
+            raft::device_resources_manager::set_workspace_allocation_limit(*(config.max_workspace_size_mb) << 20);
+        }
+        auto device_count = []() {
+            auto result = 0;
+            RAFT_CUDA_TRY(cudaGetDeviceCount(&result));
+            RAFT_EXPECTS(result != 0, "No CUDA devices found");
+            return result;
+        }();
+
+        for (auto device_id = 0; device_id < device_count; ++device_id) {
+            auto scoped_device = raft::device_setter{device_id};
+            auto workspace_size = std::size_t{};
+            if (config.max_workspace_size_mb) {
+                workspace_size = *(config.max_workspace_size_mb) << 20;
+            } else {
+                auto free_mem = std::size_t{};
+                auto total_mem = std::size_t{};
+                RAFT_CUDA_TRY_NO_THROW(cudaMemGetInfo(&free_mem, &total_mem));
+                // Heuristic: If workspace size is not explicitly specified, use half of free memory or a quarter of
+                // total memory, whichever is larger
+                workspace_size = std::max(free_mem / std::size_t{2}, total_mem / std::size_t{4});
+            }
+            if (workspace_size > std::size_t{}) {
+                raft::device_resources_manager::set_workspace_memory_resource(
+                    raft::resource::workspace_resource_factory::default_pool_resource(workspace_size), device_id);
+            }
+        }
+    });
+}
+
+}  // namespace raft_knowhere
diff --git a/src/common/raft/integration/raft_initialization.hpp b/src/common/raft/integration/raft_initialization.hpp
new file mode 100644
index 000000000..a26f48e4c
--- /dev/null
+++ b/src/common/raft/integration/raft_initialization.hpp
@@ -0,0 +1,31 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#include <cstddef>
+#include <optional>
+namespace raft_knowhere {
+struct raft_configuration {
+    std::size_t streams_per_device = std::size_t{16};
+    std::size_t stream_pools_per_device = std::size_t{};
+    std::optional<std::size_t> stream_pool_size = std::nullopt;
+    std::optional<std::size_t> init_mem_pool_size_mb = std::nullopt;
+    std::optional<std::size_t> max_mem_pool_size_mb = std::nullopt;
+    std::optional<std::size_t> max_workspace_size_mb = std::nullopt;
+};
+
+void
+initialize_raft(raft_configuration const& config);
+}  // namespace raft_knowhere
diff --git a/src/common/raft/integration/raft_knowhere_config.hpp b/src/common/raft/integration/raft_knowhere_config.hpp
new file mode 100644
index 000000000..e29e452ef
--- /dev/null
+++ b/src/common/raft/integration/raft_knowhere_config.hpp
@@ -0,0 +1,120 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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 <optional>
+#include <string>
+
+#include "common/raft/proto/raft_index_kind.hpp"
+
+namespace raft_knowhere {
+// This struct includes all parameters that may be passed to underlying RAFT
+// indexes. It is designed to not expose ANY RAFT types in order to cleanly
+// separate RAFT from knowhere headers.
+struct raft_knowhere_config {
+    raft_proto::raft_index_kind index_type;
+    int k = 10;
+
+    // Common Parameters
+    std::string metric_type = std::string{"L2"};
+    float metric_arg = 2.0f;
+    bool add_data_on_build = true;
+    float refine_ratio = 1.0f;
+
+    // Shared IVF Parameters
+    std::optional<int> nlist = std::nullopt;
+    std::optional<int> nprobe = std::nullopt;
+    std::optional<int> kmeans_n_iters = std::nullopt;
+    std::optional<float> kmeans_trainset_fraction = std::nullopt;
+
+    // IVF Flat only Parameters
+    std::optional<bool> adaptive_centers = std::nullopt;
+
+    // IVFPQ only Parameters
+    std::optional<int> m = std::nullopt;
+    std::optional<int> nbits = std::nullopt;
+    std::optional<std::string> codebook_kind = std::nullopt;
+    std::optional<bool> force_random_rotation = std::nullopt;
+    std::optional<bool> conservative_memory_allocation = std::nullopt;
+    std::optional<std::string> lookup_table_dtype = std::nullopt;
+    std::optional<std::string> internal_distance_dtype = std::nullopt;
+    std::optional<float> preferred_shmem_carveout = std::nullopt;
+
+    // CAGRA Parameters
+    std::optional<int> intermediate_graph_degree = std::nullopt;
+    std::optional<int> graph_degree = std::nullopt;
+    std::optional<int> itopk_size = std::nullopt;
+    std::optional<int> max_queries = std::nullopt;
+    std::optional<std::string> build_algo = std::nullopt;
+    std::optional<std::string> search_algo = std::nullopt;
+    std::optional<int> team_size = std::nullopt;
+    std::optional<int> search_width = std::nullopt;
+    std::optional<int> min_iterations = std::nullopt;
+    std::optional<int> max_iterations = std::nullopt;
+    std::optional<int> thread_block_size = std::nullopt;
+    std::optional<std::string> hashmap_mode = std::nullopt;
+    std::optional<int> hashmap_min_bitlen = std::nullopt;
+    std::optional<float> hashmap_max_fill_rate = std::nullopt;
+    std::optional<int> nn_descent_niter = std::nullopt;
+};
+
+// The following function provides a single source of truth for default values
+// of RAFT index configurations.
+[[nodiscard]] inline auto
+validate_raft_knowhere_config(raft_knowhere_config config) {
+    if (config.index_type == raft_proto::raft_index_kind::ivf_flat ||
+        config.index_type == raft_proto::raft_index_kind::ivf_pq) {
+        config.add_data_on_build = true;
+        config.nlist = config.nlist.value_or(128);
+        config.nprobe = config.nprobe.value_or(8);
+        config.kmeans_n_iters = config.kmeans_n_iters.value_or(20);
+        config.kmeans_trainset_fraction = config.kmeans_trainset_fraction.value_or(0.5f);
+    }
+    if (config.index_type == raft_proto::raft_index_kind::ivf_flat) {
+        config.adaptive_centers = config.adaptive_centers.value_or(false);
+    }
+    if (config.index_type == raft_proto::raft_index_kind::ivf_pq) {
+        config.m = config.m.value_or(0);
+        config.nbits = config.nbits.value_or(8);
+        config.codebook_kind = config.codebook_kind.value_or("PER_SUBSPACE");
+        config.force_random_rotation = config.force_random_rotation.value_or(false);
+        config.conservative_memory_allocation = config.conservative_memory_allocation.value_or(false);
+        config.lookup_table_dtype = config.lookup_table_dtype.value_or("CUDA_R_32F");
+        config.internal_distance_dtype = config.internal_distance_dtype.value_or("CUDA_R_32F");
+        config.preferred_shmem_carveout = config.preferred_shmem_carveout.value_or(1.0f);
+    }
+    if (config.index_type == raft_proto::raft_index_kind::cagra) {
+        config.add_data_on_build = true;
+        config.intermediate_graph_degree = config.intermediate_graph_degree.value_or(128);
+        config.graph_degree = config.graph_degree.value_or(64);
+        config.itopk_size = config.itopk_size.value_or(64);
+        config.max_queries = config.max_queries.value_or(0);
+        config.build_algo = config.build_algo.value_or("IVF_PQ");
+        config.search_algo = config.search_algo.value_or("AUTO");
+        config.team_size = config.team_size.value_or(0);
+        config.search_width = config.search_width.value_or(1);
+        config.min_iterations = config.min_iterations.value_or(0);
+        config.max_iterations = config.max_iterations.value_or(0);
+        config.thread_block_size = config.thread_block_size.value_or(0);
+        config.hashmap_mode = config.hashmap_mode.value_or("AUTO");
+        config.hashmap_min_bitlen = config.hashmap_min_bitlen.value_or(0);
+        config.hashmap_max_fill_rate = config.hashmap_max_fill_rate.value_or(0.5f);
+        config.nn_descent_niter = config.nn_descent_niter.value_or(20);
+    }
+    return config;
+}
+
+}  // namespace raft_knowhere
diff --git a/src/common/raft/integration/raft_knowhere_index.cuh b/src/common/raft/integration/raft_knowhere_index.cuh
new file mode 100644
index 000000000..805c72730
--- /dev/null
+++ b/src/common/raft/integration/raft_knowhere_index.cuh
@@ -0,0 +1,620 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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 <cstdint>
+#include <ostream>
+#include <istream>
+#include <tuple>
+#include <type_traits>
+#include <raft/core/bitset.cuh>
+#include <raft/core/copy.cuh>
+#include <raft/core/device_resources_manager.hpp>
+#include <raft/core/device_setter.hpp>
+#include <raft/distance/distance_types.hpp>
+#include <raft/neighbors/sample_filter.cuh>
+#include <raft/neighbors/ivf_pq_types.hpp>
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "common/raft/proto/raft_index.cuh"
+#include "common/raft/integration/raft_knowhere_index.hpp"
+
+namespace raft_knowhere {
+namespace detail {
+
+// This helper struct maps the generic type of RAFT index to the specific
+// instantiation of that index used within knowhere.
+template <bool B, raft_proto::raft_index_kind IndexKind>
+struct raft_index_type_mapper : std::false_type {
+};
+
+template <>
+struct raft_index_type_mapper<true, raft_proto::raft_index_kind::ivf_flat> : std::true_type {
+  using data_type = raft_data_t<raft_proto::raft_index_kind::ivf_flat>;
+  using indexing_type = raft_indexing_t<raft_proto::raft_index_kind::ivf_flat>;
+  using type = raft_proto::raft_index<raft::neighbors::ivf_flat::index, data_type, indexing_type>;
+  using underlying_index_type = typename type::vector_index_type;
+  using index_params_type = typename type::index_params_type;
+  using search_params_type = typename type::search_params_type;
+};
+template <>
+struct raft_index_type_mapper<true, raft_proto::raft_index_kind::ivf_pq> : std::true_type {
+  using data_type = raft_data_t<raft_proto::raft_index_kind::ivf_pq>;
+  using indexing_type = raft_indexing_t<raft_proto::raft_index_kind::ivf_pq>;
+  using type = raft_proto::raft_index<raft::neighbors::ivf_pq::index, indexing_type>;
+  using underlying_index_type = typename type::vector_index_type;
+  using index_params_type = typename type::index_params_type;
+  using search_params_type = typename type::search_params_type;
+};
+template <>
+struct raft_index_type_mapper<true, raft_proto::raft_index_kind::cagra> : std::true_type {
+  using data_type = raft_data_t<raft_proto::raft_index_kind::cagra>;
+  using indexing_type = raft_indexing_t<raft_proto::raft_index_kind::cagra>;
+  using type = raft_proto::raft_index<raft::neighbors::cagra::index, data_type, indexing_type>;
+  using underlying_index_type = typename type::vector_index_type;
+  using index_params_type = typename type::index_params_type;
+  using search_params_type = typename type::search_params_type;
+};
+
+} // namespace detail
+
+template <raft_proto::raft_index_kind IndexKind>
+using raft_index_t = typename detail::raft_index_type_mapper<true, IndexKind>::type;
+
+template <raft_proto::raft_index_kind IndexKind>
+using raft_index_params_t = typename detail::raft_index_type_mapper<true, IndexKind>::index_params_type;
+template <raft_proto::raft_index_kind IndexKind>
+using raft_search_params_t = typename detail::raft_index_type_mapper<true, IndexKind>::search_params_type;
+
+// Metrics are passed between knowhere and RAFT as strings to avoid tight
+// coupling between the implementation details of either one.
+[[nodiscard]] inline auto metric_string_to_raft_distance_type(std::string const& metric_string) {
+  auto result = raft::distance::DistanceType::L2Expanded;
+  if (metric_string == "L2") {
+    result = raft::distance::DistanceType::L2Expanded;
+  } else if (metric_string == "L2SqrtExpanded") {
+    result = raft::distance::DistanceType::L2SqrtExpanded;
+  } else if (metric_string == "CosineExpanded") {
+   result = raft::distance::DistanceType::CosineExpanded;
+  } else if (metric_string == "L1") {
+    result = raft::distance::DistanceType::L1;
+  } else if (metric_string == "L2Unexpanded") {
+    result = raft::distance::DistanceType::L2Unexpanded;
+  } else if (metric_string == "L2SqrtUnexpanded") {
+    result = raft::distance::DistanceType::L2SqrtUnexpanded;
+  } else if (metric_string == "IP") {
+    result = raft::distance::DistanceType::InnerProduct;
+  } else if (metric_string == "Linf") {
+    result = raft::distance::DistanceType::Linf;
+  } else if (metric_string == "Canberra") {
+    result = raft::distance::DistanceType::Canberra;
+  } else if (metric_string == "LpUnexpanded") {
+    result = raft::distance::DistanceType::LpUnexpanded;
+  } else if (metric_string == "CorrelationExpanded") {
+    result = raft::distance::DistanceType::CorrelationExpanded;
+  } else if (metric_string == "JACCARD") {
+    result = raft::distance::DistanceType::JaccardExpanded;
+  } else if (metric_string == "HellingerExpanded") {
+    result = raft::distance::DistanceType::HellingerExpanded;
+  } else if (metric_string == "Haversine") {
+    result = raft::distance::DistanceType::Haversine;
+  } else if (metric_string == "BrayCurtis") {
+    result = raft::distance::DistanceType::BrayCurtis;
+  } else if (metric_string == "JensenShannon") {
+    result = raft::distance::DistanceType::JensenShannon;
+  } else if (metric_string == "HAMMING") {
+    result = raft::distance::DistanceType::HammingUnexpanded;
+  } else if (metric_string == "KLDivergence") {
+    result = raft::distance::DistanceType::KLDivergence;
+  } else if (metric_string == "RusselRaoExpanded") {
+    result = raft::distance::DistanceType::RusselRaoExpanded;
+  } else if (metric_string == "DiceExpanded") {
+    result = raft::distance::DistanceType::DiceExpanded;
+  } else if (metric_string == "Precomputed") {
+    result = raft::distance::DistanceType::Precomputed;
+  } else {
+    RAFT_FAIL("Unrecognized metric type %s", metric_string.c_str());
+  }
+  return result;
+}
+
+[[nodiscard]] inline auto codebook_string_to_raft_codebook_gen(std::string const& codebook_string) {
+  auto result = raft::neighbors::ivf_pq::codebook_gen::PER_SUBSPACE;
+  if (codebook_string == "PER_SUBSPACE") {
+    result = raft::neighbors::ivf_pq::codebook_gen::PER_SUBSPACE;
+  } else if (codebook_string == "PER_CLUSTER") {
+    result = raft::neighbors::ivf_pq::codebook_gen::PER_CLUSTER;
+  } else {
+    RAFT_FAIL("Unrecognized codebook type %s", codebook_string.c_str());
+  }
+  return result;
+}
+[[nodiscard]] inline auto build_algo_string_to_cagra_build_algo(std::string
+    const& algo_string) {
+  auto result = raft::neighbors::cagra::graph_build_algo::IVF_PQ;
+  if (algo_string == "IVF_PQ") {
+    result = raft::neighbors::cagra::graph_build_algo::IVF_PQ;
+  } else if (algo_string == "NN_DESCENT") {
+    result = raft::neighbors::cagra::graph_build_algo::NN_DESCENT;
+  } else {
+    RAFT_FAIL("Unrecognized CAGRA build algo %s", algo_string.c_str());
+  }
+  return result;
+}
+
+[[nodiscard]] inline auto search_algo_string_to_cagra_search_algo(std::string
+    const& algo_string) {
+  auto result = raft::neighbors::cagra::search_algo::AUTO;
+  if (algo_string == "SINGLE_CTA") {
+    result = raft::neighbors::cagra::search_algo::SINGLE_CTA;
+  } else if (algo_string == "MULTI_CTA") {
+    result = raft::neighbors::cagra::search_algo::MULTI_CTA;
+  } else if (algo_string == "MULTI_KERNEL") {
+    result = raft::neighbors::cagra::search_algo::MULTI_KERNEL;
+  } else if (algo_string == "AUTO") {
+    result = raft::neighbors::cagra::search_algo::AUTO;
+  } else {
+    RAFT_FAIL("Unrecognized CAGRA search algo %s", algo_string.c_str());
+  }
+  return result;
+}
+
+[[nodiscard]] inline auto hashmap_mode_string_to_cagra_hashmap_mode(std::string
+    const& mode_string) {
+  auto result = raft::neighbors::cagra::hash_mode::AUTO;
+  if (mode_string == "HASH") {
+    result = raft::neighbors::cagra::hash_mode::HASH;
+  } else if (mode_string == "SMALL") {
+    result = raft::neighbors::cagra::hash_mode::SMALL;
+  } else if (mode_string == "AUTO") {
+    result = raft::neighbors::cagra::hash_mode::AUTO;
+  } else {
+    RAFT_FAIL("Unrecognized CAGRA hash mode %s", mode_string.c_str());
+  }
+  return result;
+}
+
+[[nodiscard]] inline auto dtype_string_to_cuda_dtype(std::string
+    const& dtype_string) {
+  auto result = CUDA_R_32F;
+  if (dtype_string == "CUDA_R_16F") {
+    result = CUDA_R_16F;
+  } else if (dtype_string == "CUDA_C_16F") {
+    result = CUDA_C_16F;
+  } else if (dtype_string == "CUDA_R_16BF") {
+    result = CUDA_R_16BF;
+  } else if (dtype_string == "CUDA_R_32F") {
+    result = CUDA_R_32F;
+  } else if (dtype_string == "CUDA_C_32F") {
+    result = CUDA_C_32F;
+  } else if (dtype_string == "CUDA_R_64F") {
+    result = CUDA_R_64F;
+  } else if (dtype_string == "CUDA_C_64F") {
+    result = CUDA_C_64F;
+  } else if (dtype_string == "CUDA_R_8I") {
+    result = CUDA_R_8I;
+  } else if (dtype_string == "CUDA_C_8I") {
+    result = CUDA_C_8I;
+  } else if (dtype_string == "CUDA_R_8U") {
+    result = CUDA_R_8U;
+  } else if (dtype_string == "CUDA_C_8U") {
+    result = CUDA_C_8U;
+  } else if (dtype_string == "CUDA_R_32I") {
+    result = CUDA_R_32I;
+  } else if (dtype_string == "CUDA_C_32I") {
+    result = CUDA_C_32I;
+  } else if (dtype_string == "CUDA_R_8F_E4M3") {
+    result = CUDA_R_8F_E4M3;
+  } else if (dtype_string == "CUDA_R_8F_E5M2") {
+    result = CUDA_R_8F_E5M2;
+  } else {
+    RAFT_FAIL("Unrecognized dtype %s", dtype_string.c_str());
+  }
+  return result;
+}
+
+// Given a generic config without RAFT symbols, convert to RAFT index build
+// parameters
+template <raft_proto::raft_index_kind IndexKind>
+[[nodiscard]] auto config_to_index_params(raft_knowhere_config const& raw_config) {
+  RAFT_EXPECTS(raw_config.index_type == IndexKind, "Incorrect index type for this index");
+  auto config = validate_raft_knowhere_config(raw_config);
+  auto result = raft_index_params_t<IndexKind>{};
+
+  result.metric = metric_string_to_raft_distance_type(config.metric_type);
+  result.metric_arg = config.metric_arg;
+  result.add_data_on_build = config.add_data_on_build;
+
+  if constexpr (IndexKind == raft_proto::raft_index_kind::ivf_flat || IndexKind ==
+      raft_proto::raft_index_kind::ivf_pq) {
+    result.n_lists = *(config.nlist);
+    result.kmeans_n_iters = *(config.kmeans_n_iters);
+    result.kmeans_trainset_fraction = *(config.kmeans_trainset_fraction);
+    result.conservative_memory_allocation = *(config.conservative_memory_allocation);
+  }
+  if constexpr (IndexKind == raft_proto::raft_index_kind::ivf_flat) {
+    result.adaptive_centers = *(config.adaptive_centers);
+  }
+  if constexpr (IndexKind == raft_proto::raft_index_kind::ivf_pq) {
+    result.pq_dim = *(config.m);
+    result.pq_bits = *(config.nbits);
+    result.codebook_kind = codebook_string_to_raft_codebook_gen(*(config.codebook_kind));
+    result.force_random_rotation = *(config.force_random_rotation);
+  }
+  if constexpr (IndexKind == raft_proto::raft_index_kind::cagra) {
+    result.intermediate_graph_degree = *(config.intermediate_graph_degree);
+    result.graph_degree = *(config.graph_degree);
+    result.build_algo = build_algo_string_to_cagra_build_algo(*(config.build_algo));
+    result.nn_descent_niter = *(config.nn_descent_niter);
+  }
+  return result;
+}
+
+// Given a generic config without RAFT symbols, convert to RAFT index search
+// parameters
+template <raft_proto::raft_index_kind IndexKind>
+[[nodiscard]] auto config_to_search_params(raft_knowhere_config const& raw_config) {
+  RAFT_EXPECTS(raw_config.index_type == IndexKind, "Incorrect index type for this index");
+  auto config = validate_raft_knowhere_config(raw_config);
+  auto result = raft_search_params_t<IndexKind>{};
+  if constexpr (IndexKind == raft_proto::raft_index_kind::ivf_flat || IndexKind
+      == raft_proto::raft_index_kind::ivf_pq) {
+    result.n_probes = *(config.nprobe);
+  }
+  if constexpr (IndexKind == raft_proto::raft_index_kind::ivf_pq) {
+    result.lut_dtype = dtype_string_to_cuda_dtype(*(config.lookup_table_dtype));
+    result.internal_distance_dtype = dtype_string_to_cuda_dtype(*(config.internal_distance_dtype));
+    result.preferred_shmem_carveout = *(config.preferred_shmem_carveout);
+  }
+  if constexpr (IndexKind == raft_proto::raft_index_kind::cagra) {
+    result.max_queries = *(config.max_queries);
+    result.itopk_size = *(config.itopk_size);
+    result.max_iterations = *(config.max_iterations);
+    result.algo = search_algo_string_to_cagra_search_algo(*(config.search_algo));
+    result.team_size = *(config.team_size);
+    result.search_width = *(config.search_width);
+    result.min_iterations = *(config.min_iterations);
+    result.thread_block_size = *(config.thread_block_size);
+    result.hashmap_mode = hashmap_mode_string_to_cagra_hashmap_mode(*(config.hashmap_mode));
+    result.hashmap_min_bitlen = *(config.hashmap_min_bitlen);
+    result.hashmap_max_fill_rate = *(config.hashmap_max_fill_rate);
+  }
+  return result;
+}
+
+inline auto select_device_id() {
+  auto static device_count = []() {
+    auto result = 0;
+    RAFT_CUDA_TRY(cudaGetDeviceCount(&result));
+    RAFT_EXPECTS(result != 0, "No CUDA devices found");
+    return result;
+  }();
+  auto static index_counter = std::atomic<int>{0};
+  // Use round-robin assignment to distribute indexes across devices
+  auto result = index_counter.fetch_add(1) % device_count;
+  return result;
+}
+
+// This struct is used to connect knowhere to a RAFT index. The implementation
+// is provided here, but this header should never be directly included in
+// another knowhere header. This ensures that RAFT symbols are not exposed in
+// any knowhere header.
+template <raft_proto::raft_index_kind IndexKind>
+struct raft_knowhere_index<IndexKind>::impl {
+  auto static constexpr index_kind = IndexKind;
+  using data_type = raft_data_t<index_kind>;
+  using indexing_type = raft_indexing_t<index_kind>;
+  using input_indexing_type = raft_input_indexing_t<index_kind>;
+
+  impl() {}
+
+ private:
+  using raft_index_type = raft_index_t<index_kind>;
+
+ public:
+  auto is_trained() const {
+    return index_.has_value();
+  }
+  [[nodiscard]] auto size() const {
+    auto result = std::int64_t{};
+    if (is_trained()) {
+      result = index_->size();
+    }
+    return result;
+  }
+  [[nodiscard]] auto dim() const {
+    auto result = std::int64_t{};
+    if (is_trained()) {
+      result = index_->dim();
+    }
+    return result;
+  }
+
+  void train(raft_knowhere_config const& config, data_type const* data,
+      knowhere_indexing_type row_count, knowhere_indexing_type feature_count) {
+    auto scoped_device = raft::device_setter{device_id};
+    auto index_params = config_to_index_params<index_kind>(config);
+    auto const& res = raft::device_resources_manager::get_device_resources();
+    auto host_data = raft::make_host_matrix_view(data, row_count, feature_count);
+    device_dataset_storage = raft::make_device_matrix<data_type, input_indexing_type>(res, row_count, feature_count);
+    auto device_data = device_dataset_storage->view();
+    raft::copy(res, device_data, host_data);
+    index_ = raft_index_type::template build<data_type, indexing_type, input_indexing_type>(res, index_params, raft::make_const_mdspan(device_data));
+  }
+
+  void add(data_type const* data, knowhere_indexing_type row_count,
+      knowhere_indexing_type feature_count, knowhere_indexing_type const* new_ids) {
+    if constexpr (index_kind == raft_proto::raft_index_kind::cagra) {
+      if (index_) {
+        RAFT_FAIL("CAGRA does not support adding vectors after training");
+      }
+    } else if constexpr (index_kind == raft_proto::raft_index_kind::ivf_pq){
+      if (index_) {
+        RAFT_FAIL("IVFPQ does not support adding vectors after training");
+      }
+    } else {
+      if (index_) {
+        auto const& res = raft::device_resources_manager::get_device_resources();
+        raft::resource::set_workspace_to_pool_resource(res);
+        auto host_data = raft::make_host_matrix_view(data, row_count, feature_count);
+        device_dataset_storage = raft::make_device_matrix<data_type, input_indexing_type>(res, row_count, feature_count);
+        auto device_data = device_dataset_storage->view();
+        raft::copy(res, device_data, host_data);
+        auto device_ids_storage = std::optional<raft::device_vector<indexing_type,
+             input_indexing_type>>{};
+        if (new_ids != nullptr) {
+          auto host_ids = raft::make_host_vector_view(new_ids, row_count);
+          device_ids_storage = raft::make_device_vector<indexing_type,
+                             input_indexing_type>(res, row_count);
+          raft::copy(res, device_ids_storage->view(), host_ids);
+        }
+
+        if (device_ids_storage) {
+          index_ = raft_index_type::extend(
+            res,
+            raft::make_const_mdspan(device_data),
+            std::make_optional(raft::make_const_mdspan(device_ids_storage->view())),
+            *index_
+          );
+        } else {
+          index_ = raft_index_type::extend(
+            res,
+            raft::make_const_mdspan(device_data),
+            std::optional<raft::device_vector_view<indexing_type const, input_indexing_type>>{},
+            *index_
+          );
+        }
+      } else {
+        RAFT_FAIL("Index has not yet been trained");
+      }
+    }
+  }
+
+  auto search(
+    raft_knowhere_config const& config,
+    data_type const* data,
+    knowhere_indexing_type row_count,
+    knowhere_indexing_type feature_count,
+    knowhere_bitset_data_type const * bitset_data,
+    knowhere_bitset_indexing_type bitset_byte_size,
+    knowhere_bitset_indexing_type bitset_size
+  ) const {
+    auto scoped_device = raft::device_setter{device_id};
+    auto const& res = raft::device_resources_manager::get_device_resources();
+    auto k = knowhere_indexing_type(config.k);
+    auto search_params = config_to_search_params<index_kind>(config);
+
+    auto host_data = raft::make_host_matrix_view(data, row_count, feature_count);
+    auto device_data_storage = raft::make_device_matrix<data_type,
+         input_indexing_type>(res, row_count, feature_count);
+    raft::copy(res, device_data_storage.view(), host_data);
+
+    auto device_bitset = std::optional<raft::core::bitset<knowhere_bitset_data_type, knowhere_bitset_indexing_type>>{};
+
+    if(bitset_data != nullptr && bitset_byte_size != 0) {
+      device_bitset = raft::core::bitset<knowhere_bitset_data_type,
+                    knowhere_bitset_indexing_type>(res, bitset_size);
+      raft::copy(res, device_bitset->to_mdspan(), raft::make_host_vector_view(bitset_data, bitset_byte_size));
+      device_bitset->flip(res);
+    }
+
+    auto output_size = row_count * k;
+    auto ids = std::unique_ptr<knowhere_indexing_type[]>(new knowhere_indexing_type[output_size]);
+    auto distances = std::unique_ptr<knowhere_data_type[]>(new knowhere_data_type[output_size]);
+
+    auto host_ids = raft::make_host_matrix_view(ids.get(), row_count, k);
+    auto host_distances = raft::make_host_matrix_view(distances.get(), row_count, k);
+
+    auto device_ids_storage = raft::make_device_matrix<indexing_type,
+         input_indexing_type>(res, row_count, k);
+    auto device_distances_storage = raft::make_device_matrix<data_type,
+         input_indexing_type>(res, row_count, k);
+    auto device_ids = device_ids_storage.view();
+    auto device_distances = device_distances_storage.view();
+
+    RAFT_EXPECTS(index_, "Index has not yet been trained");
+    auto dataset_view = device_dataset_storage ?
+      std::make_optional(device_dataset_storage->view()) :
+      std::optional<raft::device_matrix_view<const data_type, input_indexing_type>>{};
+
+
+    if (device_bitset) {
+      raft_index_type::search(
+        res,
+        *index_,
+        search_params,
+        raft::make_const_mdspan(device_data_storage.view()),
+        device_ids,
+        device_distances,
+        config.refine_ratio,
+        input_indexing_type{},
+        dataset_view,
+        raft::neighbors::filtering::bitset_filter<knowhere_bitset_data_type, knowhere_bitset_indexing_type>{
+          device_bitset->view()
+        }
+      );
+    } else {
+      raft_index_type::search(
+        res,
+        *index_,
+        search_params,
+        raft::make_const_mdspan(device_data_storage.view()),
+        device_ids,
+        device_distances,
+        config.refine_ratio,
+        input_indexing_type{},
+        dataset_view
+      );
+    }
+    if constexpr (index_kind == raft_proto::raft_index_kind::ivf_pq) {
+      thrust::replace(
+        res.get_thrust_policy(),
+        thrust::device_ptr<indexing_type>(device_ids.data_handle()),
+        thrust::device_ptr<indexing_type>(device_ids.data_handle() + output_size),
+        raft::neighbors::ivf_pq::kOutOfBoundsRecord<indexing_type>,
+        indexing_type{-1}
+      );
+    }
+    raft::copy(res, host_ids, device_ids);
+    raft::copy(res, host_distances, device_distances);
+    return std::make_tuple(ids.release(), distances.release());
+  }
+  void range_search() const {
+    RAFT_FAIL("Range search not yet implemented for RAFT indexes");
+  }
+  void get_vector_by_id() const {
+    RAFT_FAIL("Vector reconstruction not yet implemented for RAFT indexes");
+  }
+  void serialize(
+      std::ostream& os
+  ) const {
+    auto scoped_device = raft::device_setter{device_id};
+    auto const& res = raft::device_resources_manager::get_device_resources();
+    RAFT_EXPECTS(index_, "Index has not yet been trained");
+    raft_index_type::template serialize<data_type, indexing_type>(res, os, *index_);
+  }
+  auto static deserialize(
+    std::istream& is
+  ) {
+    auto new_device_id = select_device_id();
+    auto scoped_device = raft::device_setter{new_device_id};
+    auto const& res = raft::device_resources_manager::get_device_resources();
+    return std::make_unique<typename raft_knowhere_index<index_kind>::impl>(
+      raft_index_type::template deserialize<data_type, indexing_type>(res, is),
+      new_device_id
+    );
+  }
+  void synchronize() const {
+    auto scoped_device = raft::device_setter{device_id};
+    raft::device_resources_manager::get_device_resources().sync_stream();
+  }
+  impl(raft_index_type&& index, int new_device_id) : index_{std::move(index)},
+    device_id{new_device_id} {}
+
+ private:
+  std::optional<raft_index_type> index_ = std::nullopt;
+  int device_id = select_device_id();
+  std::optional<raft::device_matrix<data_type, input_indexing_type>> device_dataset_storage = std::nullopt;
+};
+
+template <raft_proto::raft_index_kind IndexKind>
+raft_knowhere_index<IndexKind>::raft_knowhere_index() : pimpl{new
+  raft_knowhere_index<IndexKind>::impl()} {}
+
+template <raft_proto::raft_index_kind IndexKind>
+raft_knowhere_index<IndexKind>::~raft_knowhere_index<IndexKind>() = default;
+
+template <raft_proto::raft_index_kind IndexKind>
+raft_knowhere_index<IndexKind>::raft_knowhere_index(raft_knowhere_index<IndexKind>&& other) : pimpl{std::move(other.pimpl)} {}
+
+template <raft_proto::raft_index_kind IndexKind>
+raft_knowhere_index<IndexKind>& raft_knowhere_index<IndexKind>::operator=(raft_knowhere_index<IndexKind>&& other) {
+  pimpl = std::move(other.pimpl);
+  return *this;
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+bool raft_knowhere_index<IndexKind>::is_trained() const {
+  return pimpl->is_trained();
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+std::int64_t raft_knowhere_index<IndexKind>::size() const {
+  return pimpl->size();
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+std::int64_t raft_knowhere_index<IndexKind>::dim() const {
+  return pimpl->dim();
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+void raft_knowhere_index<IndexKind>::train(
+  raft_knowhere_config const& config,
+  data_type const* data,
+  knowhere_indexing_type row_count,
+  knowhere_indexing_type feature_count
+) {
+  return pimpl->train(config, data, row_count, feature_count);
+}
+template <raft_proto::raft_index_kind IndexKind>
+void raft_knowhere_index<IndexKind>::add(
+  data_type const* data,
+  knowhere_indexing_type row_count,
+  knowhere_indexing_type feature_count,
+  knowhere_indexing_type const* new_ids
+) {
+  return pimpl->add(data, row_count, feature_count, new_ids);
+}
+template <raft_proto::raft_index_kind IndexKind>
+std::tuple<knowhere_indexing_type*, knowhere_data_type*> raft_knowhere_index<IndexKind>::search(
+  raft_knowhere_config const& config,
+  data_type const* data,
+  knowhere_indexing_type row_count,
+  knowhere_indexing_type feature_count,
+  knowhere_bitset_data_type const* bitset_data,
+  knowhere_bitset_indexing_type bitset_byte_size,
+  knowhere_bitset_indexing_type bitset_size
+) const {
+  return pimpl->search(config, data, row_count, feature_count, bitset_data, bitset_byte_size, bitset_size);
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+void raft_knowhere_index<IndexKind>::range_search() const {
+  return pimpl->range_search();
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+void raft_knowhere_index<IndexKind>::get_vector_by_id() const {
+  return pimpl->get_vector_by_id();
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+void raft_knowhere_index<IndexKind>::serialize(std::ostream& os) const {
+  return pimpl->serialize(os);
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+raft_knowhere_index<IndexKind>
+raft_knowhere_index<IndexKind>::deserialize(std::istream& is) {
+  return raft_knowhere_index<IndexKind>(raft_knowhere_index<IndexKind>::impl::deserialize(is));
+}
+
+template <raft_proto::raft_index_kind IndexKind>
+void raft_knowhere_index<IndexKind>::synchronize() const {
+  return pimpl->synchronize();
+}
+
+} // namespace raft_knowhere
diff --git a/src/common/raft/integration/raft_knowhere_index.hpp b/src/common/raft/integration/raft_knowhere_index.hpp
new file mode 100644
index 000000000..f26a7f0c7
--- /dev/null
+++ b/src/common/raft/integration/raft_knowhere_index.hpp
@@ -0,0 +1,124 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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 <cstdint>
+
+#include "common/raft/integration/raft_knowhere_config.hpp"
+#include "common/raft/proto/raft_index_kind.hpp"
+
+namespace raft_knowhere {
+
+using knowhere_data_type = float;
+using knowhere_indexing_type = std::int64_t;
+using knowhere_bitset_data_type = std::uint8_t;
+using knowhere_bitset_indexing_type = std::uint32_t;
+
+namespace detail {
+
+template <bool B, raft_proto::raft_index_kind IndexKind>
+struct raft_io_type_mapper : std::false_type {};
+
+template <>
+struct raft_io_type_mapper<true, raft_proto::raft_index_kind::ivf_flat> : std::true_type {
+    using data_type = float;
+    using indexing_type = std::int64_t;
+    using input_indexing_type = std::int64_t;
+};
+
+template <>
+struct raft_io_type_mapper<true, raft_proto::raft_index_kind::ivf_pq> : std::true_type {
+    using data_type = float;
+    using indexing_type = std::int64_t;
+    using input_indexing_type = std::int64_t;
+};
+
+template <>
+struct raft_io_type_mapper<true, raft_proto::raft_index_kind::cagra> : std::true_type {
+    using data_type = float;
+    using indexing_type = std::uint32_t;
+    using input_indexing_type = std::int64_t;
+};
+
+}  // namespace detail
+
+template <raft_proto::raft_index_kind IndexKind>
+using raft_data_t = typename detail::raft_io_type_mapper<true, IndexKind>::data_type;
+
+template <raft_proto::raft_index_kind IndexKind>
+using raft_indexing_t = typename detail::raft_io_type_mapper<true, IndexKind>::indexing_type;
+
+template <raft_proto::raft_index_kind IndexKind>
+using raft_input_indexing_t = typename detail::raft_io_type_mapper<true, IndexKind>::input_indexing_type;
+
+template <raft_proto::raft_index_kind IndexKind>
+struct raft_knowhere_index {
+    auto static constexpr index_kind = IndexKind;
+
+    using data_type = raft_data_t<index_kind>;
+    using indexing_type = raft_indexing_t<index_kind>;
+    using input_indexing_type = raft_input_indexing_t<index_kind>;
+
+    raft_knowhere_index();
+    ~raft_knowhere_index();
+
+    raft_knowhere_index(raft_knowhere_index&& other);
+    raft_knowhere_index&
+    operator=(raft_knowhere_index&& other);
+
+    bool
+    is_trained() const;
+    std::int64_t
+    size() const;
+    std::int64_t
+    dim() const;
+
+    void
+    train(raft_knowhere_config const&, data_type const*, knowhere_indexing_type, knowhere_indexing_type);
+    void
+    add(data_type const* data, knowhere_indexing_type row_count, knowhere_indexing_type feature_count,
+        knowhere_indexing_type const* new_ids = nullptr);
+    std::tuple<knowhere_indexing_type*, knowhere_data_type*>
+    search(raft_knowhere_config const& config, data_type const* data, knowhere_indexing_type row_count,
+           knowhere_indexing_type feature_count, knowhere_bitset_data_type const* bitset_data = nullptr,
+           knowhere_bitset_indexing_type bitset_byte_size = knowhere_bitset_indexing_type{},
+           knowhere_bitset_indexing_type bitset_size = knowhere_bitset_indexing_type{}) const;
+    void
+    range_search() const;
+    void
+    get_vector_by_id() const;
+    void
+    serialize(std::ostream& os) const;
+    static raft_knowhere_index<IndexKind>
+    deserialize(std::istream& is);
+    void
+    synchronize() const;
+
+ private:
+    // Use a private implementation to completely separate knowhere headers from
+    // RAFT headers
+    struct impl;
+    std::unique_ptr<impl> pimpl;
+
+    raft_knowhere_index(std::unique_ptr<impl>&& new_pimpl) : pimpl{std::move(new_pimpl)} {
+    }
+};
+
+extern template struct raft_knowhere_index<raft_proto::raft_index_kind::ivf_flat>;
+extern template struct raft_knowhere_index<raft_proto::raft_index_kind::ivf_pq>;
+extern template struct raft_knowhere_index<raft_proto::raft_index_kind::cagra>;
+
+}  // namespace raft_knowhere
diff --git a/src/common/raft/proto/raft_index.cuh b/src/common/raft/proto/raft_index.cuh
new file mode 100644
index 000000000..ff25b7e04
--- /dev/null
+++ b/src/common/raft/proto/raft_index.cuh
@@ -0,0 +1,583 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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 <cstddef>
+#include <cstdint>
+#include <istream>
+#include <optional>
+#include <ostream>
+#include <type_traits>
+#include <raft/core/logger.hpp>
+#include <raft/neighbors/ivf_flat.cuh>
+#include <raft/neighbors/ivf_flat_serialize.cuh>
+#include <raft/neighbors/ivf_flat_types.hpp>
+#include <raft/neighbors/ivf_pq.cuh>
+#include <raft/neighbors/ivf_pq_serialize.cuh>
+#include <raft/neighbors/ivf_pq_types.hpp>
+#include <raft/neighbors/cagra.cuh>
+#include <raft/neighbors/cagra_serialize.cuh>
+#include <raft/neighbors/cagra_types.hpp>
+#include "common/raft/proto/raft_index_kind.hpp"
+
+namespace raft_proto {
+
+auto static const RAFT_NAME = raft::RAFT_NAME;
+
+namespace detail {
+template <raft_index_kind index_kind, template<typename...> typename index_template>
+struct template_matches_index_kind : std::false_type{};
+
+template<>
+struct template_matches_index_kind<raft_index_kind::ivf_flat, raft::neighbors::ivf_flat::index> : std::true_type{};
+
+template<>
+struct template_matches_index_kind<raft_index_kind::ivf_pq, raft::neighbors::ivf_pq::index> : std::true_type{};
+
+template<>
+struct template_matches_index_kind<raft_index_kind::cagra, raft::neighbors::cagra::index> : std::true_type{};
+
+template <raft_index_kind index_kind, template<typename...> typename index_template>
+auto static constexpr template_matches_index_kind_v = template_matches_index_kind<index_kind, index_template>::value;
+
+}
+
+template <template <typename...> typename underlying_index_type, typename... raft_index_args>
+struct raft_index {
+  using vector_index_type = underlying_index_type<raft_index_args...>;
+  auto static constexpr vector_index_kind = []() {
+    if constexpr (detail::template_matches_index_kind_v<raft_index_kind::ivf_flat, underlying_index_type>) {
+      return raft_index_kind::ivf_flat;
+    } else if constexpr (detail::template_matches_index_kind_v<raft_index_kind::ivf_pq, underlying_index_type>) {
+      return raft_index_kind::ivf_pq;
+    } else if constexpr (detail::template_matches_index_kind_v<raft_index_kind::cagra, underlying_index_type>) {
+      return raft_index_kind::cagra;
+    } else {
+      static_assert(detail::template_matches_index_kind_v<
+        raft_index_kind::ivf_flat, underlying_index_type
+      >, "Unsupported index template passed to raft_index");
+    }
+  }();
+
+  using index_params_type = std::conditional_t<
+    vector_index_kind == raft_index_kind::ivf_flat,
+    raft::neighbors::ivf_flat::index_params,
+    std::conditional_t<
+      vector_index_kind == raft_index_kind::ivf_pq,
+      raft::neighbors::ivf_pq::index_params,
+      std::conditional_t<
+        vector_index_kind == raft_index_kind::cagra,
+        raft::neighbors::cagra::index_params,
+        // Should never get here; precluded by static assertion above
+        raft::neighbors::ivf_flat::index_params
+      >
+    >
+  >;
+  using search_params_type = std::conditional_t<
+    vector_index_kind == raft_index_kind::ivf_flat,
+    raft::neighbors::ivf_flat::search_params,
+    std::conditional_t<
+      vector_index_kind == raft_index_kind::ivf_pq,
+      raft::neighbors::ivf_pq::search_params,
+      std::conditional_t<
+        vector_index_kind == raft_index_kind::cagra,
+        raft::neighbors::cagra::search_params,
+        // Should never get here; precluded by static assertion above
+        raft::neighbors::ivf_flat::search_params
+      >
+    >
+  >;
+
+  [[nodiscard]] auto& get_vector_index() { return vector_index_; }
+  [[nodiscard]] auto const& get_vector_index() const { return vector_index_; }
+  [[nodiscard]] auto size() const {
+    return vector_index_.size();
+  }
+  [[nodiscard]] auto dim() const {
+    return vector_index_.dim();
+  }
+
+  template <typename T, typename IdxT, typename InputIdxT>
+  auto static build(
+    raft::resources const& res,
+    index_params_type const& index_params,
+    raft::device_matrix_view<
+      T const, InputIdxT
+    > data
+  ) {
+    if constexpr (vector_index_kind == raft_index_kind::ivf_flat) {
+      return raft_index<underlying_index_type, raft_index_args...>{
+        raft::neighbors::ivf_flat::build<T, IdxT>(
+          res,
+          index_params,
+          data
+        )
+      };
+    } else if constexpr (vector_index_kind == raft_index_kind::ivf_pq) {
+      return raft_index<underlying_index_type, raft_index_args...>{
+        raft::neighbors::ivf_pq::build<T, IdxT>(
+          res,
+          index_params,
+          data
+        )
+      };
+    } else if constexpr (vector_index_kind == raft_index_kind::cagra) {
+      return raft_index<underlying_index_type, raft_index_args...>{
+        raft::neighbors::cagra::build<T>(
+          res,
+          index_params,
+          data
+        )
+      };
+    }
+  }
+
+  template <typename T, typename IdxT, typename InputIdxT, typename FilterT=std::nullptr_t>
+  auto static search(
+    raft::resources const& res,
+    raft_index<underlying_index_type, raft_index_args...> const& index,
+    search_params_type const& search_params,
+    raft::device_matrix_view<T const, InputIdxT> queries,
+    raft::device_matrix_view<IdxT, InputIdxT> neighbors,
+    raft::device_matrix_view<float, InputIdxT> distances,
+    float refine_ratio = 1.0f,
+    InputIdxT k_offset=InputIdxT{},
+    std::optional<raft::device_matrix_view<const T, InputIdxT>> dataset = std::nullopt,
+    FilterT filter = nullptr
+  ) {
+    auto const& underlying_index = index.get_vector_index();
+
+    auto k = neighbors.extent(1);
+    auto k_tmp = k + k_offset;
+    if (refine_ratio > 1.0f) {
+      k_tmp *= refine_ratio;
+    }
+
+    auto neighbors_tmp = neighbors;
+    auto distances_tmp = distances;
+    auto neighbors_storage = std::optional<raft::device_matrix<IdxT,
+         InputIdxT>>{};
+    auto distances_storage = std::optional<raft::device_matrix<float,
+         InputIdxT>>{};
+
+    if (k_tmp > k) {
+      neighbors_storage = raft::make_device_matrix<IdxT, InputIdxT>(
+        res,
+        queries.extent(0),
+        k_tmp
+      );
+      neighbors_tmp = neighbors_storage->view();
+      distances_storage = raft::make_device_matrix<float, InputIdxT>(
+        res,
+        queries.extent(0),
+        k_tmp
+      );
+      distances_tmp = distances_storage->view();
+    }
+
+    if constexpr (vector_index_kind == raft_index_kind::ivf_flat) {
+      if constexpr (std::is_same_v<FilterT, std::nullptr_t>){
+        raft::neighbors::ivf_flat::search<T, IdxT>(
+          res,
+          search_params,
+          underlying_index,
+          queries,
+          neighbors_tmp,
+          distances_tmp
+        );
+      } else {
+        raft::neighbors::ivf_flat::search_with_filtering<T, IdxT>(
+          res,
+          search_params,
+          underlying_index,
+          queries,
+          neighbors_tmp,
+          distances_tmp,
+          filter
+        );
+      }
+    } else if constexpr (vector_index_kind == raft_index_kind::ivf_pq) {
+      if constexpr (std::is_same_v<FilterT, std::nullptr_t>){
+        raft::neighbors::ivf_pq::search<T, IdxT>(
+          res,
+          search_params,
+          underlying_index,
+          queries,
+          neighbors_tmp,
+          distances_tmp
+        );
+      } else {
+        raft::neighbors::ivf_pq::search_with_filtering<T, IdxT>(
+          res,
+          search_params,
+          underlying_index,
+          queries,
+          neighbors_tmp,
+          distances_tmp,
+          filter
+        );
+      }
+    } else if constexpr (vector_index_kind == raft_index_kind::cagra) {
+      if constexpr (std::is_same_v<FilterT, std::nullptr_t>){
+        raft::neighbors::cagra::search<T, IdxT>(
+          res,
+          search_params,
+          underlying_index,
+          queries,
+          neighbors_tmp,
+          distances_tmp
+        );
+      } else {
+        raft::neighbors::cagra::search_with_filtering<T, IdxT>(
+          res,
+          search_params,
+          underlying_index,
+          queries,
+          neighbors_tmp,
+          distances_tmp,
+          filter
+        );
+      }
+    }
+    if (refine_ratio > 1.0f) {
+      if (dataset.has_value()) {
+        if constexpr (std::is_same_v<IdxT, InputIdxT>) {
+          raft::neighbors::refine(
+            res,
+            *dataset,
+            queries,
+            raft::make_const_mdspan(neighbors_tmp),
+            neighbors,
+            distances,
+            underlying_index.metric()
+          );
+        } else {
+          // https://github.com/rapidsai/raft/issues/1950
+          raft::neighbors::refine(
+            res,
+            raft::make_device_matrix_view(
+              dataset->data_handle(),
+              IdxT(dataset->extent(0)),
+              IdxT(dataset->extent(1))
+            ),
+            raft::make_device_matrix_view(
+              queries.data_handle(),
+              IdxT(queries.extent(0)),
+              IdxT(queries.extent(1))
+            ),
+            raft::make_const_mdspan(
+              raft::make_device_matrix_view(
+                neighbors_tmp.data_handle(),
+                IdxT(neighbors_tmp.extent(0)),
+                IdxT(neighbors_tmp.extent(1))
+              )
+            ),
+            raft::make_device_matrix_view(
+              neighbors.data_handle(),
+              IdxT(neighbors.extent(0)),
+              IdxT(neighbors.extent(1))
+            ),
+            raft::make_device_matrix_view(
+              distances.data_handle(),
+              IdxT(distances.extent(0)),
+              IdxT(distances.extent(1))
+            ),
+            underlying_index.metric()
+          );
+        }
+      } else {
+        RAFT_LOG_WARN(
+          "Refinement requested, but no dataset provided. Ignoring refinement request."
+        );
+      }
+    }
+  }
+
+  template <typename T, typename IdxT, typename InputIdxT>
+  auto static extend(
+    raft::resources const& res,
+    raft::device_matrix_view<T const, InputIdxT> new_vectors,
+    std::optional<raft::device_vector_view<IdxT const, InputIdxT>> new_ids,
+    raft_index<underlying_index_type, raft_index_args...>& index
+  ) {
+    auto const& underlying_index = index.get_vector_index();
+
+    if constexpr (vector_index_kind == raft_index_kind::ivf_flat) {
+      return raft_index{raft::neighbors::ivf_flat::extend<T, IdxT>(
+        res,
+        new_vectors,
+        new_ids,
+        underlying_index
+      )};
+    } else if constexpr (vector_index_kind == raft_index_kind::ivf_pq) {
+      return raft_index{raft::neighbors::ivf_pq::extend<T, IdxT>(
+        res,
+        new_vectors,
+        new_ids,
+        underlying_index
+      )};
+    } else if constexpr (vector_index_kind == raft_index_kind::cagra) {
+      RAFT_FAIL("CAGRA implements no extend method");
+    }
+  }
+
+  template <typename T, typename IdxT>
+  void static serialize(
+    raft::resources const& res,
+    std::ostream& os,
+    raft_index<underlying_index_type, raft_index_args...> const& index
+  ) {
+    auto const& underlying_index = index.get_vector_index();
+
+    if constexpr (vector_index_kind == raft_index_kind::ivf_flat) {
+      return raft::neighbors::ivf_flat::serialize<T, IdxT>(
+        res,
+        os,
+        underlying_index
+      );
+    } else if constexpr (vector_index_kind == raft_index_kind::ivf_pq) {
+      return raft::neighbors::ivf_pq::serialize<IdxT>(
+        res,
+        os,
+        underlying_index
+      );
+    } else if constexpr (vector_index_kind == raft_index_kind::cagra) {
+      return raft::neighbors::cagra::serialize<T, IdxT>(
+        res,
+        os,
+        underlying_index
+      );
+    }
+  }
+
+  template <typename T, typename IdxT>
+  auto static deserialize(
+    raft::resources const& res,
+    std::istream& is
+  ) {
+    if constexpr (vector_index_kind == raft_index_kind::ivf_flat) {
+      return raft_index{raft::neighbors::ivf_flat::deserialize<T, IdxT>(
+        res,
+        is
+      )};
+    } else if constexpr (vector_index_kind == raft_index_kind::ivf_pq) {
+      return raft_index{raft::neighbors::ivf_pq::deserialize<IdxT>(
+        res,
+        is
+      )};
+    } else if constexpr (vector_index_kind == raft_index_kind::cagra) {
+      return raft_index{raft::neighbors::cagra::deserialize<T, IdxT>(
+        res,
+        is
+      )};
+    }
+  }
+
+ private:
+  vector_index_type vector_index_;
+
+  explicit raft_index(vector_index_type&& vector_index)
+    : vector_index_{std::move(vector_index)} {}
+};
+
+/*template <typename T, typename IdxT, template <typename...> typename underlying_index, typename... raft_index_args>
+struct tiered_index {
+  using raft_index_type = raft_index<underlying_index, raft_index_args...>;
+  using data_type = T;
+  using index_type = IdxT;
+
+  tiered_index (
+    raft::device_resources const& res,
+    raft_index_type::index_params_type const& index_params,
+    raft::device_matrix_view<T const, IdxT> data,
+    IdxT max_staged_vectors = IdxT{}
+  ) : index_params_{index_params},
+      vector_index_{},
+      dataset_{
+        [&res, data]() {
+          auto result = raft::make_device_matrix<T, IdxT>(
+            res,
+            data.extent(0),
+            data.extent(1)
+          );
+          raft::copy(res, result.view(), data);
+          return result;
+        }()
+      },
+      staged_count{data.extent(0)},
+      max_staged_count{max_staged_vectors} {
+      if (staged_count >= max_staged_count) {
+        commit(res);
+      }
+    }
+
+  template <typename FilterT=std::nullptr_t>
+  auto search(
+    raft::device_resources const& res,
+    search_params_type const& search_params,
+    raft::device_matrix_view<T const, IdxT> queries,
+    raft::device_matrix_view<IdxT, IdxT> neighbors,
+    raft::device_matrix_view<float, IdxT> distances,
+    float refine_ratio = 1.0f,
+    FilterT filter = nullptr
+  ) const {
+    auto k = neighbors.extent(1);
+    auto k_committed = k;
+    if (refine_ratio > 1.0f) {
+      k_committed *= refine_ratio;
+    }
+    auto k_total = k_committed + staged_count;
+
+    auto neighbors_committed = neighbors;
+    auto distances_committed = distances;
+    auto neighbors_staged = neighbors;
+    auto distances_staged = distances;
+    auto neighbors_storage = std::optional<raft::device_matrix<IdxT, IdxT>>{};
+    auto distances_storage = std::optional<raft::device_matrix<float, IdxT>>{};
+    if (k_total > k) {
+      neighbors_storage = raft::make_device_matrix<IdxT, IdxT>(
+        res,
+        queries.extent(0),
+        k_total
+      );
+      neighbors_committed = raft::make_device_matrix_view(
+          new_dataset.data_handle(),
+          queries.extent(0),
+          dataset_.extent(1)
+        ),
+      distances_storage = raft::make_device_matrix<float, IdxT>(
+        res,
+        queries.extent(0),
+        k_total
+      );
+      distances_tmp = distances_storage.view();
+    }
+    raft::neighbors::refine(
+      res,
+      *dataset,
+      queries,
+      neighbors_tmp,
+      neighbors,
+      distances,
+      underlying_index.metric()
+    );
+    return raft_index_type::search(
+      res,
+      vector_index_,
+      search_params,
+      queries,
+      neighbors,
+      distances,
+      1.0f,
+      0,  // TODO (wphicks): Pass offset down when batching available
+      get_entire_dataset(),
+      filter
+    );
+  }
+
+  template <typename T, typename IdxT>
+  void extend(
+    raft::device_resources const& res,
+    raft::device_matrix_view<T const, IdxT> new_vectors,
+    std::optional<raft::device_matrix_view<IdxT const, IdxT>> new_ids
+  ) {
+    if (size() + new_vectors.extent(0) > capacity()) {
+      reserve(res, size() + std::max(size(), new_vectors.extent(0)));
+    }
+    staged_count += new_vectors.extent(0);
+    if (staged_count >= max_staged_vectors) {
+      commit(res);
+    }
+  }
+
+  auto size() const {
+    return committed_count + staged_count;
+  }
+
+  void reserve(raft::resources const& res, std::size_type new_capacity) {
+    if (new_capacity > capacity()) {
+      auto new_dataset = raft::make_device_matrix<T, IdxT>(
+        res,
+        new_capacity,
+        dataset_.extent(1)
+      );
+      raft::copy(
+        res,
+        raft::make_device_matrix_view(
+          new_dataset.data_handle(),
+          size(),
+          dataset_.extent(1)
+        ),
+        get_entire_dataset()
+      );
+    }
+  }
+
+ private:
+  raft_index_type::index_params_type index_params_;
+  std::optional<raft_index_type> vector_index_;
+
+  raft::device_matrix<T, IdxT> dataset_;
+  IdxT committed_count;
+  IdxT max_staged_count;
+  IdxT build_size = IdxT{};
+  IdxT staged_count = IdxT{};
+
+  auto commit(raft::resources const& res) {
+    if (
+      (!vector_index_ && staged_count > max_staged_count) ||
+      build_size * 2 < size()
+    ) {
+      vector_index_ = raft_index_type::build(res, index_params_, dataset_.view());
+      committed_count = dataset_.extent(0);
+      build_size = committed_count;
+    } else {
+      raft_index_type::extend(res, vector_index_, get_staged_dataset());
+      committed_count += staged_count;
+    }
+    staged_count = IdxT{};
+  }
+
+  auto capacity() const {
+    return dataset_.extent(0);
+  }
+
+  auto get_committed_dataset() {
+    return raft::make_device_matrix_view(
+      dataset_.data_handle(),
+      committed_count,
+      dataset_.extent(1)
+    );
+  }
+
+  auto get_staged_dataset() {
+    return raft::make_device_matrix_view(
+      dataset_.data_handle() + committed_count * dataset_.extent(1),
+      staged_count,
+      dataset_.extent(1)
+    );
+  }
+
+  auto get_entire_dataset() {
+    return raft::make_device_matrix_view(
+      dataset_.data_handle(),
+      size(),
+      dataset_.extent(1)
+    );
+  }
+}; */
+
+}  // namespace raft_proto
diff --git a/src/common/raft/proto/raft_index_kind.hpp b/src/common/raft/proto/raft_index_kind.hpp
new file mode 100644
index 000000000..68102a260
--- /dev/null
+++ b/src/common/raft/proto/raft_index_kind.hpp
@@ -0,0 +1,21 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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
+
+namespace raft_proto {
+enum struct raft_index_kind { ivf_flat, ivf_pq, cagra };
+}  // namespace raft_proto
diff --git a/src/common/raft/raft.cu b/src/common/raft/raft.cu
deleted file mode 100644
index 59635c162..000000000
--- a/src/common/raft/raft.cu
+++ /dev/null
@@ -1,39 +0,0 @@
-#include "../src/distance/specializations/fused_l2_nn_double_int.cu"
-#include "../src/distance/specializations/fused_l2_nn_double_int64.cu"
-#include "../src/distance/specializations/fused_l2_nn_float_int.cu"
-#include "../src/distance/specializations/fused_l2_nn_float_int64.cu"
-#include "../src/matrix/specializations/detail/select_k_float_int64_t.cu"
-#include "../src/matrix/specializations/detail/select_k_float_uint32_t.cu"
-#include "../src/matrix/specializations/detail/select_k_half_int64_t.cu"
-#include "../src/matrix/specializations/detail/select_k_half_uint32_t.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_float_fast.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_float_no_basediff.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_float_no_smem_lut.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_fp8s_fast.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_fp8s_no_basediff.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_fp8s_no_smem_lut.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_fp8u_fast.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_fp8u_no_basediff.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_fp8u_no_smem_lut.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_half_fast.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_half_no_basediff.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_float_half_no_smem_lut.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_fp8s_fast.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_fp8s_no_basediff.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_fp8s_no_smem_lut.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_fp8u_fast.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_fp8u_no_basediff.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_fp8u_no_smem_lut.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_half_fast.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_half_no_basediff.cu"
-#include "../src/neighbors/specializations/detail/compute_similarity_half_half_no_smem_lut.cu"
-#include "../src/neighbors/specializations/fused_l2_knn_int_float_false.cu"
-#include "../src/neighbors/specializations/fused_l2_knn_int_float_true.cu"
-#include "../src/neighbors/specializations/fused_l2_knn_long_float_false.cu"
-#include "../src/neighbors/specializations/fused_l2_knn_long_float_true.cu"
-#include "../src/neighbors/specializations/ivfflat_build_float_int64_t.cu"
-#include "../src/neighbors/specializations/ivfflat_extend_float_int64_t.cu"
-#include "../src/neighbors/specializations/ivfflat_search_float_int64_t.cu"
-#include "../src/neighbors/specializations/ivfpq_build_float_int64_t.cu"
-#include "../src/neighbors/specializations/ivfpq_extend_float_int64_t.cu"
-#include "../src/neighbors/specializations/ivfpq_search_float_int64_t.cu"
diff --git a/src/common/raft/raft_utils.cc b/src/common/raft/raft_utils.cc
deleted file mode 100644
index 14e906fdc..000000000
--- a/src/common/raft/raft_utils.cc
+++ /dev/null
@@ -1,46 +0,0 @@
-#include "raft_utils.h"
-
-namespace raft_utils {
-int
-gpu_device_manager::random_choose() const {
-    srand(time(NULL));
-    return rand() % memory_load_.size();
-}
-
-int
-gpu_device_manager::choose_with_load(size_t load) {
-    std::lock_guard<std::mutex> lock(mtx_);
-
-    auto it = std::min_element(memory_load_.begin(), memory_load_.end());
-    *it += load;
-    return std::distance(memory_load_.begin(), it);
-}
-
-void
-gpu_device_manager::release_load(int device_id, size_t load) {
-    if (size_t(device_id) < memory_load_.size()) {
-        std::lock_guard<std::mutex> lock(mtx_);
-        memory_load_[device_id] -= load;
-    } else {
-        LOG_KNOWHERE_WARNING_ << "please check device id " << device_id;
-    }
-}
-
-gpu_device_manager::gpu_device_manager() {
-    int device_counts;
-    try {
-        RAFT_CUDA_TRY(cudaGetDeviceCount(&device_counts));
-    } catch (const raft::exception& e) {
-        LOG_KNOWHERE_FATAL_ << e.what();
-    }
-    memory_load_.resize(device_counts);
-    std::fill(memory_load_.begin(), memory_load_.end(), 0);
-}
-
-gpu_device_manager&
-gpu_device_manager::instance() {
-    static gpu_device_manager mgr;
-    return mgr;
-}
-
-}  // namespace raft_utils
diff --git a/src/common/raft/raft_utils.h b/src/common/raft/raft_utils.h
deleted file mode 100644
index b7d0ae33e..000000000
--- a/src/common/raft/raft_utils.h
+++ /dev/null
@@ -1,200 +0,0 @@
-#include <cuda_runtime.h>
-
-#include <atomic>
-#include <cstddef>
-#include <cstdlib>
-#include <map>
-#include <memory>
-#include <mutex>
-#include <optional>
-#include <sstream>
-#include <vector>
-
-#include "knowhere/log.h"
-#include "raft/core/device_resources.hpp"
-#include "rmm/cuda_stream_pool.hpp"
-#include "rmm/mr/device/cuda_memory_resource.hpp"
-#include "rmm/mr/device/per_device_resource.hpp"
-#include "rmm/mr/device/pool_memory_resource.hpp"
-#include "thrust/optional.h"
-
-namespace raft_utils {
-
-inline auto
-get_current_device() {
-    auto result = int{};
-    RAFT_CUDA_TRY(cudaGetDevice(&result));
-    return result;
-}
-
-inline auto thread_counter = std::atomic<std::size_t>{};
-inline auto
-get_thread_id() {
-    thread_local std::size_t id = ++thread_counter;
-    return id;
-}
-
-// TODO(wphicks): Replace this with version from RAFT once merged
-struct device_setter {
-    device_setter(int new_device) : prev_device_{[]() { return get_current_device(); }()} {
-        RAFT_CUDA_TRY(cudaSetDevice(new_device));
-    }
-
-    ~device_setter() {
-        RAFT_CUDA_TRY_NO_THROW(cudaSetDevice(prev_device_));
-    }
-
- private:
-    int prev_device_;
-};
-
-inline auto raft_mutex = std::mutex{};
-
-struct gpu_resources {
-    gpu_resources(std::optional<std::size_t> streams_per_device = std::nullopt)
-        : streams_per_device_{streams_per_device.value_or([]() {
-              auto result = std::size_t{16};
-              if (auto* env_str = std::getenv("KNOWHERE_STREAMS_PER_GPU")) {
-                  auto str_stream = std::stringstream{env_str};
-                  str_stream >> result;
-                  if (str_stream.fail() || result == std::size_t{0}) {
-                      LOG_KNOWHERE_WARNING_ << "KNOWHERE_STREAMS_PER_GPU env variable should be a positive integer";
-                      result = std::size_t{16};
-                  } else {
-                      LOG_KNOWHERE_INFO_ << "streams per gpu set to " << result;
-                  }
-              }
-              return result;
-          }())},
-          stream_pools_{},
-          memory_resources_{},
-          upstream_mr_{},
-          init_mem_pool_size_{},
-          max_mem_pool_size_{} {
-    }
-    ~gpu_resources() {
-        memory_resources_.clear();
-    }
-
-    /* Set pool size in megabytes */
-    void
-    set_pool_size(std::size_t init_size, std::size_t max_size) {
-        init_mem_pool_size_ = init_size << 20;
-        max_mem_pool_size_ = max_size << 20;
-    }
-
-    void
-    init(int device_id = get_current_device()) {
-        auto lock = std::lock_guard{raft_mutex};
-        auto stream_iter = stream_pools_.find(device_id);
-
-        if (stream_iter == stream_pools_.end()) {
-            auto scoped_device = device_setter{device_id};
-            stream_pools_[device_id] = std::make_shared<rmm::cuda_stream_pool>(streams_per_device_);
-
-            // Set up device memory pool for this device
-            if (!init_mem_pool_size_.has_value() && !max_mem_pool_size_.has_value()) {
-                auto* env_str = getenv("KNOWHERE_GPU_MEM_POOL_SIZE");
-                if (env_str != NULL) {
-                    auto init_pool_size_tmp = std::size_t{};
-                    auto max_pool_size_tmp = std::size_t{};
-                    auto stat = sscanf(env_str, "%zu;%zu", &init_pool_size_tmp, &max_pool_size_tmp);
-                    if (stat == 2) {
-                        LOG_KNOWHERE_INFO_ << "Get Gpu Pool Size From env, init size: " << init_pool_size_tmp
-                                           << " MB, max size: " << max_pool_size_tmp << " MB";
-                        init_mem_pool_size_ = init_pool_size_tmp << 20;
-                        max_mem_pool_size_ = max_pool_size_tmp << 20;
-                    } else {
-                        LOG_KNOWHERE_WARNING_ << "please check env format";
-                    }
-                }
-            }
-            if (max_mem_pool_size_.value_or(std::size_t{1}) != std::size_t{}) {
-                memory_resources_[device_id] =
-                    std::make_unique<rmm::mr::pool_memory_resource<rmm::mr::cuda_memory_resource>>(
-                        &upstream_mr_, init_mem_pool_size_, max_mem_pool_size_);
-                rmm::mr::set_current_device_resource(memory_resources_[device_id].get());
-            }
-        }
-    }
-
-    auto
-    get_streams_per_device() const {
-        return streams_per_device_;
-    }
-
-    auto
-    get_stream_view(int device_id = get_current_device(), std::size_t thread_id = get_thread_id()) {
-        return stream_pools_[device_id]->get_stream(thread_id % streams_per_device_);
-    }
-
- private:
-    std::size_t streams_per_device_;
-    std::map<int, std::shared_ptr<rmm::cuda_stream_pool>> stream_pools_;
-    std::map<int, std::unique_ptr<rmm::mr::pool_memory_resource<rmm::mr::cuda_memory_resource>>> memory_resources_;
-    rmm::mr::cuda_memory_resource upstream_mr_;
-    thrust::optional<std::size_t> init_mem_pool_size_;
-    thrust::optional<std::size_t> max_mem_pool_size_;
-};
-
-inline auto&
-get_gpu_resources(std::optional<std::size_t> streams_per_device = std::nullopt) {
-    static auto resources = gpu_resources{streams_per_device};
-    return resources;
-}
-
-inline void
-init_gpu_resources(std::optional<std::size_t> streams_per_device = std::nullopt, int device_id = get_current_device()) {
-    get_gpu_resources(streams_per_device).init(device_id);
-}
-
-class gpu_device_manager {
- public:
-    static gpu_device_manager&
-    instance();
-    int
-    random_choose() const;
-    int
-    choose_with_load(size_t load);
-    void
-    release_load(int device_id, size_t load);
-
- private:
-    gpu_device_manager();
-    std::vector<size_t> memory_load_;
-    mutable std::mutex mtx_;
-};
-
-inline auto&
-get_raft_resources(int device_id = get_current_device()) {
-    thread_local auto all_resources = std::map<int, std::unique_ptr<raft::device_resources>>{};
-
-    auto iter = all_resources.find(device_id);
-    if (iter == all_resources.end()) {
-        auto scoped_device = device_setter{device_id};
-        all_resources[device_id] = std::make_unique<raft::device_resources>(
-            get_gpu_resources().get_stream_view(), nullptr, rmm::mr::get_current_device_resource());
-    }
-    return *all_resources[device_id];
-}
-
-inline void
-set_mem_pool_size(size_t init_size, size_t max_size) {
-    LOG_KNOWHERE_INFO_ << "Set GPU pool size: init size " << init_size << ", max size " << max_size;
-    get_gpu_resources().set_pool_size(init_size, max_size);
-}
-
-};  // namespace raft_utils
-
-#define RANDOM_CHOOSE_DEVICE_WITH_ASSIGN(x)                             \
-    do {                                                                \
-        x = raft_utils::gpu_device_manager::instance().random_choose(); \
-    } while (0)
-#define MIN_LOAD_CHOOSE_DEVICE_WITH_ASSIGN(x, load)                            \
-    do {                                                                       \
-        x = raft_utils::gpu_device_manager::instance().choose_with_load(load); \
-    } while (0)
-#define RELEASE_DEVICE(x, load)                                           \
-    do {                                                                  \
-        raft_utils::gpu_device_manager::instance().release_load(x, load); \
-    } while (0)
diff --git a/src/index/cagra/cagra.cu b/src/index/cagra/cagra.cu
deleted file mode 100644
index b1a971efe..000000000
--- a/src/index/cagra/cagra.cu
+++ /dev/null
@@ -1,210 +0,0 @@
-#include "cagra_config.h"
-#include "common/raft/raft_utils.h"
-#include "common/raft_metric.h"
-#include "knowhere/index_node.h"
-#include "knowhere/log.h"
-#include "raft/neighbors/cagra.cuh"
-#include "raft/neighbors/cagra_serialize.cuh"
-namespace knowhere {
-
-using idx_type = uint32_t;
-
-using cagra_index = raft::neighbors::experimental::cagra::index<float, idx_type>;
-
-class CagraIndexNode : public IndexNode {
- public:
-    CagraIndexNode(const int32_t& /*version*/, const Object& object) : devs_{}, gpu_index_{} {
-    }
-
-    Status
-    Train(const DataSet& dataset, const Config& cfg) override {
-        auto cagra_cfg = static_cast<const knowhere::CagraConfig&>(cfg);
-        if (gpu_index_) {
-            LOG_KNOWHERE_WARNING_ << "index is already trained";
-            return Status::index_already_trained;
-        }
-        if (cagra_cfg.gpu_ids.value().size() != 1) {
-            LOG_KNOWHERE_WARNING_ << "Cagra implementation is single-GPU only" << std::endl;
-            return Status::raft_inner_error;
-        }
-        auto metric = Str2RaftMetricType(cagra_cfg.metric_type.value());
-        if (!metric.has_value()) {
-            LOG_KNOWHERE_WARNING_ << "please check metric value: " << cagra_cfg.metric_type.value();
-            return metric.error();
-        }
-        if (metric.value() != raft::distance::DistanceType::L2Expanded) {
-            LOG_KNOWHERE_WARNING_ << "only support L2Expanded metric type";
-            return Status::invalid_metric_type;
-        }
-        devs_.insert(devs_.begin(), cagra_cfg.gpu_ids.value().begin(), cagra_cfg.gpu_ids.value().end());
-        auto scoped_device = raft_utils::device_setter{*cagra_cfg.gpu_ids.begin()};
-        auto build_params = raft::neighbors::experimental::cagra::index_params{};
-        build_params.intermediate_graph_degree = cagra_cfg.intermediate_graph_degree.value();
-        build_params.graph_degree = cagra_cfg.graph_degree.value();
-        build_params.metric = metric.value();
-        auto& res = raft_utils::get_raft_resources();
-        auto rows = dataset.GetRows();
-        auto dim = dataset.GetDim();
-        auto* data = reinterpret_cast<float const*>(dataset.GetTensor());
-        auto data_gpu = raft::make_device_matrix<float, idx_type>(res, rows, dim);
-        RAFT_CUDA_TRY(cudaMemcpyAsync(data_gpu.data_handle(), data, data_gpu.size() * sizeof(float), cudaMemcpyDefault,
-                                      res.get_stream().value()));
-        gpu_index_ = raft::neighbors::experimental::cagra::build(
-            res, build_params,
-            raft::make_device_matrix_view<const float, idx_type>((const float*)data_gpu.data_handle(), rows, dim));
-        this->dim_ = dim;
-        this->counts_ = rows;
-        res.sync_stream();
-    }
-
-    Status
-    Add(const DataSet& dataset, const Config& cfg) override {
-        return Status::success;
-    }
-
-    expected<DataSetPtr>
-    Search(const DataSet& dataset, const Config& cfg, const BitsetView& bitset) const override {
-        auto cagra_cfg = static_cast<const CagraConfig&>(cfg);
-        auto rows = dataset.GetRows();
-        auto dim = dataset.GetDim();
-        auto* data = reinterpret_cast<float const*>(dataset.GetTensor());
-        auto output_size = rows * cagra_cfg.k.value();
-        auto ids = std::unique_ptr<idx_type[]>(new idx_type[output_size]);
-        auto dis = std::unique_ptr<float[]>(new float[output_size]);
-        try {
-            auto scoped_device = raft_utils::device_setter{devs_[0]};
-            auto& res_ = raft_utils::get_raft_resources();
-
-            auto data_gpu = raft::make_device_matrix<float, idx_type>(res_, rows, dim);
-            raft::copy(data_gpu.data_handle(), data, data_gpu.size(), res_.get_stream());
-
-            auto search_params = raft::neighbors::experimental::cagra::search_params{};
-            search_params.max_queries = cagra_cfg.max_queries.value();
-            auto ids_dev = raft::make_device_matrix<idx_type, idx_type>(*res_, rows, cagra_cfg.k.value());
-            auto dis_dev = raft::make_device_matrix<float, idx_type>(*res_, rows, cagra_cfg.k.value());
-            raft::neighbors::experimental::cagra::search(res_, search_params, *gpu_index_,
-                                                         raft::make_const_mdspan(data_gpu.view()), ids_dev.view(),
-                                                         dis_dev.view());
-
-            raft::copy(ids.get(), ids_dev.data_handle(), output_size, res_.get_stream());
-            raft::copy(dis.get(), dis_dev.data_handle(), output_size, res_.get_stream());
-            res_.sync_stream();
-
-        } catch (std::exception& e) {
-            LOG_KNOWHERE_WARNING_ << "RAFT inner error, " << e.what();
-            return Status::raft_inner_error;
-        }
-        return Status::not_implemented;
-        //        return GenResultDataSet(rows, cagra_cfg.k, ids.release(), dis.release());
-    }
-
-    expected<DataSetPtr>
-    RangeSearch(const DataSet& dataset, const Config& cfg, const BitsetView& bitset) const override {
-        return Status::not_implemented;
-    }
-
-    expected<DataSetPtr>
-    GetVectorByIds(const DataSet& dataset) const override {
-        return Status::not_implemented;
-    }
-
-    bool
-    HasRawData(const std::string& metric_type) const override {
-        return false;
-    }
-
-    expected<DataSetPtr>
-    GetIndexMeta(const Config& cfg) const override {
-        return Status::not_implemented;
-    }
-
-    Status
-    Serialize(BinarySet& binset) const override {
-        if (!gpu_index_.has_value()) {
-            LOG_KNOWHERE_ERROR_ << "Can not serialize empty RaftCagraIndex.";
-            return Status::empty_index;
-        }
-        std::stringbuf buf;
-        std::ostream os(&buf);
-        os.write((char*)(&this->dim_), sizeof(this->dim_));
-        os.write((char*)(&this->counts_), sizeof(this->counts_));
-        os.write((char*)(&this->devs_[0]), sizeof(this->devs_[0]));
-
-        auto scoped_device = raft_utils::device_setter{devs_[0]};
-        rmm::mr::cuda_memory_resource mr;
-        rmm::cuda_stream stm;
-        raft::device_resources res(stm.view(), nullptr, &mr);
-
-        raft::neighbors::experimental::cagra::serialize<float, idx_type>(res, os, *gpu_index_);
-
-        os.flush();
-        std::shared_ptr<uint8_t[]> index_binary(new (std::nothrow) uint8_t[buf.str().size()]);
-
-        memcpy(index_binary.get(), buf.str().c_str(), buf.str().size());
-        binset.Append(this->Type(), index_binary, buf.str().size());
-        return Status::success;
-    }
-
-    Status
-    Deserialize(const BinarySet& binset, const Config& config) override {
-        std::stringbuf buf;
-        auto binary = binset.GetByName(this->Type());
-        if (binary == nullptr) {
-            LOG_KNOWHERE_ERROR_ << "Invalid binary set.";
-            return Status::invalid_binary_set;
-        }
-        buf.sputn((char*)binary->data.get(), binary->size);
-        std::istream is(&buf);
-
-        is.read((char*)(&this->dim_), sizeof(this->dim_));
-        is.read((char*)(&this->counts_), sizeof(this->counts_));
-        this->devs_.resize(1);
-        is.read((char*)(&this->devs_[0]), sizeof(this->devs_[0]));
-        auto scoped_device = raft_utils::device_setter{devs_[0]};
-
-        auto& res = raft_utils::get_raft_resources();
-
-        cagra_index index_ = raft::neighbors::experimental::cagra::deserialize<float, idx_type>(res, is);
-        is.sync();
-        gpu_index_ = cagra_index(std::move(index_));
-
-        return Status::success;
-    }
-
-    Status
-    DeserializeFromFile(const std::string& filename, const Config& config) override {
-    }
-
-    std::unique_ptr<BaseConfig>
-    CreateConfig() const override {
-        return std::make_unique<CagraConfig>();
-    }
-
-    int64_t
-    Dim() const override {
-        return dim_;
-    }
-
-    int64_t
-    Size() const override {
-        return 0;
-    }
-
-    int64_t
-    Count() const override {
-        return counts_;
-    }
-
-    std::string
-    Type() const override {
-        return knowhere::IndexEnum::INDEX_RAFT_IVFFLAT;
-    }
-
- private:
-    std::vector<int32_t> devs_;
-    int64_t dim_ = 0;
-    int64_t counts_ = 0;
-    std::optional<cagra_index> gpu_index_;
-};
-
-}  // namespace knowhere
diff --git a/src/index/cagra/cagra_config.h b/src/index/cagra/cagra_config.h
deleted file mode 100644
index 5c3d790c8..000000000
--- a/src/index/cagra/cagra_config.h
+++ /dev/null
@@ -1,47 +0,0 @@
-/**
- * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
- * SPDX-License-Identifier: Apache-2.0
- *
- * 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.
- */
-#include "knowhere/config.h"
-namespace knowhere {
-
-class CagraConfig : public BaseConfig {
- public:
-    CFG_INT intermediate_graph_degree;
-    CFG_INT graph_degree;
-    CFG_LIST gpu_ids;
-    CFG_INT max_queries;
-    KNOHWERE_DECLARE_CONFIG(CagraConfig) {
-        KNOWHERE_CONFIG_DECLARE_FIELD(intermediate_graph_degree)
-            .set_default(128)
-            .description("degree of input graph for pruning.")
-            .for_train()
-            .set_range(1, 65536);
-        KNOWHERE_CONFIG_DECLARE_FIELD(graph_degree)
-            .set_default(64)
-            .description("degree of output graph.")
-            .for_search()
-            .set_range(1, 65536);
-        KNOWHERE_CONFIG_DECLARE_FIELD(gpu_ids)
-            .description("gpu device ids")
-            .set_default({
-                0,
-            })
-            .for_train();
-        KNOWHERE_CONFIG_DECLARE_FIELD(max_queries).description("query batch size.").set_default(1).for_search();
-    }
-};
-
-}  // namespace knowhere
diff --git a/src/index/gpu_raft/gpu_raft.h b/src/index/gpu_raft/gpu_raft.h
new file mode 100644
index 000000000..628726491
--- /dev/null
+++ b/src/index/gpu_raft/gpu_raft.h
@@ -0,0 +1,268 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#ifndef GPU_RAFT_H
+#define GPU_RAFT_H
+
+#include <cstdint>
+#include <exception>
+#include <fstream>
+#include <istream>
+#include <numeric>
+#include <thread>
+#include <tuple>
+#include <vector>
+
+#include "common/raft/integration/raft_knowhere_config.hpp"
+#include "common/raft/integration/raft_knowhere_index.hpp"
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "index/gpu_raft/gpu_raft_cagra_config.h"
+#include "index/gpu_raft/gpu_raft_ivf_flat_config.h"
+#include "index/gpu_raft/gpu_raft_ivf_pq_config.h"
+#include "knowhere/comp/index_param.h"
+#include "knowhere/expected.h"
+#include "knowhere/factory.h"
+#include "knowhere/log.h"
+#include "knowhere/utils.h"
+
+namespace knowhere {
+
+auto static constexpr cuda_concurrent_size = std::uint32_t{32};
+
+template <raft_proto::raft_index_kind K>
+struct KnowhereConfigType {};
+
+template <>
+struct KnowhereConfigType<raft_proto::raft_index_kind::ivf_flat> {
+    using Type = GpuRaftIvfFlatConfig;
+};
+
+template <>
+struct KnowhereConfigType<raft_proto::raft_index_kind::ivf_pq> {
+    using Type = GpuRaftIvfPqConfig;
+};
+
+template <>
+struct KnowhereConfigType<raft_proto::raft_index_kind::cagra> {
+    using Type = GpuRaftCagraConfig;
+};
+
+template <raft_proto::raft_index_kind K>
+struct GpuRaftIndexNode : public IndexNode {
+    auto static constexpr index_kind = K;
+    using knowhere_config_type = typename KnowhereConfigType<index_kind>::Type;
+
+    GpuRaftIndexNode(int32_t, const Object& object) : index_{} {
+    }
+
+    Status
+    Train(const DataSet& dataset, const Config& cfg) override {
+        auto result = Status::success;
+        auto raft_cfg = raft_knowhere::raft_knowhere_config{};
+        try {
+            raft_cfg = to_raft_knowhere_config(static_cast<const knowhere_config_type&>(cfg));
+        } catch (const std::exception& e) {
+            LOG_KNOWHERE_ERROR_ << e.what();
+            result = Status::invalid_args;
+        }
+        if (index_.is_trained()) {
+            result = Status::index_already_trained;
+        }
+        if (result == Status::success) {
+            auto rows = dataset.GetRows();
+            auto dim = dataset.GetDim();
+            auto const* data = reinterpret_cast<float const*>(dataset.GetTensor());
+            try {
+                index_.train(raft_cfg, data, rows, dim);
+                index_.synchronize();
+            } catch (const std::exception& e) {
+                LOG_KNOWHERE_ERROR_ << e.what();
+                result = Status::raft_inner_error;
+            }
+        }
+        return result;
+    }
+
+    Status
+    Add(const DataSet& dataset, const Config& cfg) override {
+        return Status::success;
+    }
+
+    expected<DataSetPtr>
+    Search(const DataSet& dataset, const Config& cfg, const BitsetView& bitset) const override {
+        auto result = Status::success;
+        auto raft_cfg = raft_knowhere::raft_knowhere_config{};
+        auto err_msg = std::string{};
+        try {
+            raft_cfg = to_raft_knowhere_config(static_cast<const knowhere_config_type&>(cfg));
+        } catch (const std::exception& e) {
+            err_msg = std::string{e.what()};
+            LOG_KNOWHERE_ERROR_ << e.what();
+            result = Status::invalid_args;
+        }
+        if (result == Status::success) {
+            try {
+                auto rows = dataset.GetRows();
+                auto dim = dataset.GetDim();
+                auto const* data = reinterpret_cast<float const*>(dataset.GetTensor());
+                auto search_result =
+                    index_.search(raft_cfg, data, rows, dim, bitset.data(), bitset.byte_size(), bitset.size());
+                std::this_thread::yield();
+                index_.synchronize();
+                return GenResultDataSet(rows, raft_cfg.k, std::get<0>(search_result), std::get<1>(search_result));
+            } catch (const std::exception& e) {
+                err_msg = std::string{e.what()};
+                LOG_KNOWHERE_ERROR_ << e.what();
+                result = Status::raft_inner_error;
+            }
+        }
+        return expected<DataSetPtr>::Err(result, err_msg.c_str());
+    }
+
+    expected<DataSetPtr>
+    RangeSearch(const DataSet& dataset, const Config& cfg, const BitsetView& bitset) const override {
+        return expected<DataSetPtr>::Err(Status::not_implemented, "RangeSearch not implemented");
+    }
+
+    expected<DataSetPtr>
+    GetVectorByIds(const DataSet& dataset) const override {
+        return expected<DataSetPtr>::Err(Status::not_implemented, "GetVectorByIds not implemented");
+    }
+
+    bool
+    HasRawData(const std::string& metric_type) const override {
+        if constexpr (index_kind == raft_proto::raft_index_kind::ivf_flat) {
+            return !IsMetricType(metric_type, metric::COSINE);
+        } else {
+            return false;
+        }
+    }
+
+    Status
+    Serialize(BinarySet& binset) const override {
+        auto result = Status::success;
+        std::stringbuf buf;
+        if (!index_.is_trained()) {
+            result = Status::empty_index;
+        } else {
+            std::ostream os(&buf);
+
+            try {
+                index_.serialize(os);
+                index_.synchronize();
+            } catch (const std::exception& e) {
+                LOG_KNOWHERE_ERROR_ << e.what();
+                result = Status::raft_inner_error;
+            }
+            os.flush();
+        }
+        if (result == Status::success) {
+            std::shared_ptr<uint8_t[]> index_binary(new (std::nothrow) uint8_t[buf.str().size()]);
+            memcpy(index_binary.get(), buf.str().c_str(), buf.str().size());
+            binset.Append(this->Type(), index_binary, buf.str().size());
+        }
+        return result;
+    }
+
+    Status
+    Deserialize(const BinarySet& binset, const Config& config) override {
+        auto result = Status::success;
+        std::stringbuf buf;
+        auto binary = binset.GetByName(this->Type());
+        if (binary == nullptr) {
+            LOG_KNOWHERE_ERROR_ << "Invalid binary set.";
+            result = Status::invalid_binary_set;
+        } else {
+            buf.sputn((char*)binary->data.get(), binary->size);
+            std::istream is(&buf);
+            result = DeserializeFromStream(is);
+        }
+        return result;
+    }
+
+    Status
+    DeserializeFromFile(const std::string& filename, const Config& config) {
+        auto stream = std::ifstream{filename};
+        return DeserializeFromStream(stream);
+    }
+
+    std::unique_ptr<BaseConfig>
+    CreateConfig() const override {
+        return std::make_unique<knowhere_config_type>();
+    }
+
+    expected<DataSetPtr>
+    GetIndexMeta(const Config& cfg) const override {
+        return expected<DataSetPtr>::Err(Status::not_implemented, "GetIndexMeta not implemented");
+    }
+
+    int64_t
+    Dim() const override {
+        return index_.dim();
+    }
+
+    int64_t
+    Size() const override {
+        return 0;
+    }
+
+    int64_t
+    Count() const override {
+        return index_.size();
+    }
+
+    std::string
+    Type() const override {
+        if constexpr (index_kind == raft_proto::raft_index_kind::ivf_flat) {
+            return knowhere::IndexEnum::INDEX_RAFT_IVFFLAT;
+        } else if constexpr (index_kind == raft_proto::raft_index_kind::ivf_pq) {
+            return knowhere::IndexEnum::INDEX_RAFT_IVFPQ;
+        } else if constexpr (index_kind == raft_proto::raft_index_kind::cagra) {
+            return knowhere::IndexEnum::INDEX_RAFT_CAGRA;
+        }
+    }
+
+ private:
+    using raft_knowhere_index_type = typename raft_knowhere::raft_knowhere_index<K>;
+
+    raft_knowhere_index_type index_;
+
+    Status
+    DeserializeFromStream(std::istream& stream) {
+        auto result = Status::success;
+        try {
+            index_ = raft_knowhere_index_type::deserialize(stream);
+            index_.synchronize();
+        } catch (const std::exception& e) {
+            LOG_KNOWHERE_ERROR_ << e.what();
+            result = Status::raft_inner_error;
+        }
+        stream.sync();
+        return result;
+    }
+};
+
+extern template struct GpuRaftIndexNode<raft_proto::raft_index_kind::ivf_flat>;
+extern template struct GpuRaftIndexNode<raft_proto::raft_index_kind::ivf_pq>;
+extern template struct GpuRaftIndexNode<raft_proto::raft_index_kind::cagra>;
+
+using GpuRaftIvfFlatIndexNode = GpuRaftIndexNode<raft_proto::raft_index_kind::ivf_flat>;
+using GpuRaftIvfPqIndexNode = GpuRaftIndexNode<raft_proto::raft_index_kind::ivf_pq>;
+using GpuRaftCagraIndexNode = GpuRaftIndexNode<raft_proto::raft_index_kind::cagra>;
+
+}  // namespace knowhere
+
+#endif /* GPU_RAFT_H */
diff --git a/src/index/gpu_raft/gpu_raft_cagra.cc b/src/index/gpu_raft/gpu_raft_cagra.cc
new file mode 100644
index 000000000..c6ea08661
--- /dev/null
+++ b/src/index/gpu_raft/gpu_raft_cagra.cc
@@ -0,0 +1,33 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+
+#include <vector>
+
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "gpu_raft.h"
+#include "knowhere/factory.h"
+#include "knowhere/index_node_thread_pool_wrapper.h"
+
+namespace knowhere {
+template struct GpuRaftIndexNode<raft_proto::raft_index_kind::cagra>;
+
+KNOWHERE_REGISTER_GLOBAL(GPU_RAFT_CAGRA, [](const int32_t& version, const Object& object) {
+    return Index<IndexNodeThreadPoolWrapper>::Create(std::make_unique<GpuRaftCagraIndexNode>(version, object),
+                                                     cuda_concurrent_size);
+});
+
+}  // namespace knowhere
diff --git a/src/index/gpu_raft/gpu_raft_cagra_config.h b/src/index/gpu_raft/gpu_raft_cagra_config.h
new file mode 100644
index 000000000..371c70f9c
--- /dev/null
+++ b/src/index/gpu_raft/gpu_raft_cagra_config.h
@@ -0,0 +1,132 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#ifndef GPU_RAFT_CAGRA_CONFIG_H
+#define GPU_RAFT_CAGRA_CONFIG_H
+
+#include "common/raft/integration/raft_knowhere_config.hpp"
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "index/ivf/ivf_config.h"
+#include "knowhere/config.h"
+
+namespace knowhere {
+
+struct GpuRaftCagraConfig : public BaseConfig {
+    CFG_FLOAT refine_ratio;
+    CFG_INT intermediate_graph_degree;
+    CFG_INT graph_degree;
+    CFG_INT itopk_size;
+    CFG_INT max_queries;
+    CFG_STRING build_algo;
+    CFG_STRING search_algo;
+    CFG_INT team_size;
+    CFG_INT search_width;
+    CFG_INT min_iterations;
+    CFG_INT max_iterations;
+    CFG_INT thread_block_size;
+    CFG_STRING hashmap_mode;
+    CFG_INT hashmap_min_bitlen;
+    CFG_FLOAT hashmap_max_fill_rate;
+    CFG_INT nn_descent_niter;
+
+    KNOHWERE_DECLARE_CONFIG(GpuRaftCagraConfig) {
+        KNOWHERE_CONFIG_DECLARE_FIELD(refine_ratio)
+            .set_default(1.0f)
+            .description("search refine_ratio * k results then refine")
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(intermediate_graph_degree)
+            .description("degree of intermediate knn graph")
+            .set_default(128)
+            .for_train();
+        KNOWHERE_CONFIG_DECLARE_FIELD(graph_degree).description("degree of knn graph").set_default(64).for_train();
+        KNOWHERE_CONFIG_DECLARE_FIELD(itopk_size)
+            .description("intermediate results retained during search")
+            .set_default(64)
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(max_queries).description("maximum batch size").set_default(0).for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(build_algo)
+            .description("algorithm used to build knn graph")
+            .set_default("IVF_PQ")
+            .for_train();
+        KNOWHERE_CONFIG_DECLARE_FIELD(search_algo)
+            .description("algorithm used for search")
+            .set_default("AUTO")
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(team_size)
+            .description("threads used to calculate single distance")
+            .set_default(0)
+            .set_range(0, 32)
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(search_width)
+            .description("nodes to select as starting point in each iteration")
+            .set_default(1)
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(min_iterations)
+            .description("minimum number of search iterations")
+            .set_default(0)
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(max_iterations)
+            .description("maximum number of search iterations")
+            .set_default(0)
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(thread_block_size).description("threads per block").set_default(0).for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(hashmap_mode).description("hashmap mode").set_default("AUTO").for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(hashmap_min_bitlen)
+            .description("minimum bit length of hashmap")
+            .set_default(0)
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(hashmap_max_fill_rate)
+            .description("minimum bit length of hashmap")
+            .set_default(0.5f)
+            .set_range(0.1f, 0.9f)
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(nn_descent_niter)
+            .description("number of iterations for NN descent")
+            .set_default(20)
+            .for_train();
+    }
+};
+
+[[nodiscard]] inline auto
+to_raft_knowhere_config(GpuRaftCagraConfig const& cfg) {
+    auto result = raft_knowhere::raft_knowhere_config{raft_proto::raft_index_kind::cagra};
+
+    result.metric_type = cfg.metric_type.value();
+    result.refine_ratio = cfg.refine_ratio.value();
+    result.k = cfg.k.value();
+
+    result.intermediate_graph_degree = cfg.intermediate_graph_degree;
+    result.graph_degree = cfg.graph_degree;
+    result.itopk_size = cfg.itopk_size;
+    result.max_queries = cfg.max_queries;
+    result.build_algo = cfg.build_algo;
+    result.search_algo = cfg.search_algo;
+    result.team_size = cfg.team_size;
+    result.search_width = cfg.search_width;
+    result.min_iterations = cfg.min_iterations;
+    result.max_iterations = cfg.max_iterations;
+    result.thread_block_size = cfg.thread_block_size;
+    result.hashmap_mode = cfg.hashmap_mode;
+    result.hashmap_min_bitlen = cfg.hashmap_min_bitlen;
+    result.hashmap_max_fill_rate = cfg.hashmap_max_fill_rate;
+    result.nn_descent_niter = cfg.nn_descent_niter;
+
+    return result;
+}
+
+}  // namespace knowhere
+
+#endif /*GPU_RAFT_CAGRA_CONFIG_H*/
diff --git a/src/index/gpu_raft/gpu_raft_ivf_flat.cc b/src/index/gpu_raft/gpu_raft_ivf_flat.cc
new file mode 100644
index 000000000..ec39e1309
--- /dev/null
+++ b/src/index/gpu_raft/gpu_raft_ivf_flat.cc
@@ -0,0 +1,33 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+
+#include <vector>
+
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "gpu_raft.h"
+#include "knowhere/factory.h"
+#include "knowhere/index_node_thread_pool_wrapper.h"
+
+namespace knowhere {
+template struct GpuRaftIndexNode<raft_proto::raft_index_kind::ivf_flat>;
+
+KNOWHERE_REGISTER_GLOBAL(GPU_RAFT_IVF_FLAT, [](const int32_t& version, const Object& object) {
+    return Index<IndexNodeThreadPoolWrapper>::Create(std::make_unique<GpuRaftIvfFlatIndexNode>(version, object),
+                                                     cuda_concurrent_size);
+});
+
+}  // namespace knowhere
diff --git a/src/index/gpu_raft/gpu_raft_ivf_flat_config.h b/src/index/gpu_raft/gpu_raft_ivf_flat_config.h
new file mode 100644
index 000000000..2f6cb0def
--- /dev/null
+++ b/src/index/gpu_raft/gpu_raft_ivf_flat_config.h
@@ -0,0 +1,75 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+#ifndef GPU_RAFT_IVF_FLAT_CONFIG_H
+#define GPU_RAFT_IVF_FLAT_CONFIG_H
+
+#include "common/raft/integration/raft_knowhere_config.hpp"
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "index/ivf/ivf_config.h"
+#include "knowhere/config.h"
+
+namespace knowhere {
+
+struct GpuRaftIvfFlatConfig : public IvfFlatConfig {
+    CFG_FLOAT refine_ratio;
+    CFG_INT kmeans_n_iters;
+    CFG_FLOAT kmeans_trainset_fraction;
+    CFG_BOOL adaptive_centers;
+    KNOHWERE_DECLARE_CONFIG(GpuRaftIvfFlatConfig) {
+        KNOWHERE_CONFIG_DECLARE_FIELD(refine_ratio)
+            .set_default(1.0f)
+            .description("search refine_ratio * k results then refine")
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(k)
+            .set_default(10)
+            .description("search for top k similar vector.")
+            .set_range(1, 256)  // Declared in base but limited to 256
+            .for_search();
+        KNOWHERE_CONFIG_DECLARE_FIELD(kmeans_n_iters)
+            .description("iterations to search for kmeans centers")
+            .set_default(20)
+            .for_train();
+        KNOWHERE_CONFIG_DECLARE_FIELD(kmeans_trainset_fraction)
+            .description("fraction of data to use in kmeans building")
+            .set_default(0.5)
+            .for_train();
+        KNOWHERE_CONFIG_DECLARE_FIELD(adaptive_centers)
+            .description("update centroids with new data")
+            .set_default(false)
+            .for_train();
+    }
+};
+
+[[nodiscard]] inline auto
+to_raft_knowhere_config(GpuRaftIvfFlatConfig const& cfg) {
+    auto result = raft_knowhere::raft_knowhere_config{raft_proto::raft_index_kind::ivf_flat};
+
+    result.metric_type = cfg.metric_type.value();
+    result.refine_ratio = cfg.refine_ratio.value();
+    result.k = cfg.k.value();
+    result.nlist = cfg.nlist;
+    result.nprobe = cfg.nprobe;
+    result.kmeans_n_iters = cfg.kmeans_n_iters;
+    result.kmeans_trainset_fraction = cfg.kmeans_trainset_fraction;
+    result.adaptive_centers = cfg.adaptive_centers;
+
+    return result;
+}
+
+}  // namespace knowhere
+
+#endif /*GPU_RAFT_IVF_FLAT_CONFIG_H*/
diff --git a/src/index/gpu_raft/gpu_raft_ivf_pq.cc b/src/index/gpu_raft/gpu_raft_ivf_pq.cc
new file mode 100644
index 000000000..d59e647ea
--- /dev/null
+++ b/src/index/gpu_raft/gpu_raft_ivf_pq.cc
@@ -0,0 +1,33 @@
+/**
+ * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * 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.
+ */
+
+#include <vector>
+
+#include "common/raft/proto/raft_index_kind.hpp"
+#include "gpu_raft.h"
+#include "knowhere/factory.h"
+#include "knowhere/index_node_thread_pool_wrapper.h"
+
+namespace knowhere {
+template struct GpuRaftIndexNode<raft_proto::raft_index_kind::ivf_pq>;
+
+KNOWHERE_REGISTER_GLOBAL(GPU_RAFT_IVF_PQ, [](const int32_t& version, const Object& object) {
+    return Index<IndexNodeThreadPoolWrapper>::Create(std::make_unique<GpuRaftIvfPqIndexNode>(version, object),
+                                                     cuda_concurrent_size);
+});
+
+}  // namespace knowhere
diff --git a/src/index/ivf_raft/ivf_raft_config.h b/src/index/gpu_raft/gpu_raft_ivf_pq_config.h
similarity index 54%
rename from src/index/ivf_raft/ivf_raft_config.h
rename to src/index/gpu_raft/gpu_raft_ivf_pq_config.h
index 620f36a0e..c689cf69f 100644
--- a/src/index/ivf_raft/ivf_raft_config.h
+++ b/src/index/gpu_raft/gpu_raft_ivf_pq_config.h
@@ -14,62 +14,43 @@
  * See the License for the specific language governing permissions and
  * limitations under the License.
  */
+#ifndef GPU_RAFT_IVF_PQ_CONFIG_H
+#define GPU_RAFT_IVF_PQ_CONFIG_H
 
-#ifndef IVF_RAFT_CONFIG_H
-#define IVF_RAFT_CONFIG_H
-
+#include "common/raft/integration/raft_knowhere_config.hpp"
+#include "common/raft/proto/raft_index_kind.hpp"
 #include "index/ivf/ivf_config.h"
 #include "knowhere/config.h"
 
 namespace knowhere {
 
-class RaftIvfFlatConfig : public IvfFlatConfig {
- public:
+struct GpuRaftIvfPqConfig : public IvfPqConfig {
+    CFG_FLOAT refine_ratio;
     CFG_INT kmeans_n_iters;
     CFG_FLOAT kmeans_trainset_fraction;
-    CFG_BOOL adaptive_centers;
-    KNOHWERE_DECLARE_CONFIG(RaftIvfFlatConfig) {
-        KNOWHERE_CONFIG_DECLARE_FIELD(k)
-            .set_default(10)
-            .description("search for top k similar vector.")
-            .set_range(1, 256)
-            .for_search();
-        KNOWHERE_CONFIG_DECLARE_FIELD(kmeans_n_iters)
-            .description("iterations to search for kmeans centers")
-            .set_default(20)
-            .for_train();
-        KNOWHERE_CONFIG_DECLARE_FIELD(kmeans_trainset_fraction)
-            .description("fraction of data to use in kmeans building")
-            .set_default(0.5)
-            .for_train();
-        KNOWHERE_CONFIG_DECLARE_FIELD(adaptive_centers)
-            .description("update centroids with new data")
-            .set_default(false)
-            .for_train();
-    }
-};
 
-class RaftIvfPqConfig : public IvfPqConfig {
- public:
-    CFG_INT streams_per_device;
-    CFG_INT kmeans_n_iters;
-    CFG_FLOAT kmeans_trainset_fraction;
-    CFG_INT m;
     CFG_STRING codebook_kind;
     CFG_BOOL force_random_rotation;
     CFG_BOOL conservative_memory_allocation;
     CFG_STRING lut_dtype;
     CFG_STRING internal_distance_dtype;
     CFG_FLOAT preferred_shmem_carveout;
-    KNOHWERE_DECLARE_CONFIG(RaftIvfPqConfig) {
+
+    KNOHWERE_DECLARE_CONFIG(GpuRaftIvfPqConfig) {
+        KNOWHERE_CONFIG_DECLARE_FIELD(refine_ratio)
+            .set_default(1.0f)
+            .description("search refine_ratio * k results then refine")
+            .for_search();
         KNOWHERE_CONFIG_DECLARE_FIELD(k)
             .set_default(10)
             .description("search for top k similar vector.")
-            .set_range(1, 1024)
+            .set_range(1, 1024)  // Declared in base but limited to 1024
             .for_search();
-        KNOWHERE_CONFIG_DECLARE_FIELD(streams_per_device)
-            .description("CUDA streams per GPU")
-            .set_default(1)
+        KNOWHERE_CONFIG_DECLARE_FIELD(m).set_default(8).description("m").set_range(0, 65536).for_train();
+        KNOWHERE_CONFIG_DECLARE_FIELD(nbits)
+            .set_default(0)
+            .description("nbits")
+            .set_range(0, 8)  // Declared in base but limited here
             .for_train();
         KNOWHERE_CONFIG_DECLARE_FIELD(kmeans_n_iters)
             .description("iterations to search for kmeans centers")
@@ -79,21 +60,21 @@ class RaftIvfPqConfig : public IvfPqConfig {
             .description("fraction of data to use in kmeans building")
             .set_default(0.5)
             .for_train();
-        KNOWHERE_CONFIG_DECLARE_FIELD(m).description("dimension after compression by PQ").set_default(0).for_train();
+
         KNOWHERE_CONFIG_DECLARE_FIELD(codebook_kind)
-            .description("how PQ codebooks are created")
+            .description("how PQ codebooks are generated")
             .set_default("PER_SUBSPACE")
             .for_train();
         KNOWHERE_CONFIG_DECLARE_FIELD(force_random_rotation)
-            .description("always perform random rotation")
+            .description("always perform random_rotation")
             .set_default(false)
             .for_train();
         KNOWHERE_CONFIG_DECLARE_FIELD(conservative_memory_allocation)
-            .description("use minimal possible GPU memory")
+            .description("use minimum GPU memory at the cost of speed")
             .set_default(false)
             .for_train();
         KNOWHERE_CONFIG_DECLARE_FIELD(lut_dtype)
-            .description("Data type for lookup table")
+            .description("data type for lookup table")
             .set_default("CUDA_R_32F")
             .for_search();
         KNOWHERE_CONFIG_DECLARE_FIELD(internal_distance_dtype)
@@ -102,10 +83,37 @@ class RaftIvfPqConfig : public IvfPqConfig {
             .for_search();
         KNOWHERE_CONFIG_DECLARE_FIELD(preferred_shmem_carveout)
             .description("preferred fraction of memory for shmem vs L1")
-            .set_default(1.0)
+            .set_range(0.0f, 1.0f)
+            .set_default(1.0f)
             .for_search();
     }
 };
 
+[[nodiscard]] inline auto
+to_raft_knowhere_config(GpuRaftIvfPqConfig const& cfg) {
+    auto result = raft_knowhere::raft_knowhere_config{raft_proto::raft_index_kind::ivf_pq};
+
+    result.metric_type = cfg.metric_type.value();
+    result.refine_ratio = cfg.refine_ratio.value();
+    result.k = cfg.k.value();
+
+    result.nlist = cfg.nlist;
+    result.nprobe = cfg.nprobe;
+    result.kmeans_n_iters = cfg.kmeans_n_iters;
+    result.kmeans_trainset_fraction = cfg.kmeans_trainset_fraction;
+
+    result.m = cfg.m;
+    result.nbits = cfg.nbits;
+    result.codebook_kind = cfg.codebook_kind;
+    result.force_random_rotation = cfg.force_random_rotation;
+    result.conservative_memory_allocation = cfg.conservative_memory_allocation;
+    result.lookup_table_dtype = cfg.lut_dtype;
+    result.internal_distance_dtype = cfg.internal_distance_dtype;
+    result.preferred_shmem_carveout = cfg.preferred_shmem_carveout;
+
+    return result;
+}
+
 }  // namespace knowhere
-#endif /* IVF_CONFIG_H */
+
+#endif /*GPU_RAFT_IVF_PQ_CONFIG_H*/
diff --git a/src/index/ivf_raft/ivf_raft.cu b/src/index/ivf_raft/ivf_raft.cu
deleted file mode 100644
index 622666c91..000000000
--- a/src/index/ivf_raft/ivf_raft.cu
+++ /dev/null
@@ -1,45 +0,0 @@
-/**
- * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
- * SPDX-License-Identifier: Apache-2.0
- *
- * 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.
- */
-
-#include "ivf_raft.cuh"
-#include "knowhere/factory.h"
-#include "knowhere/index_node_thread_pool_wrapper.h"
-
-constexpr uint32_t cuda_concurrent_size = 32;
-
-namespace knowhere {
-
-KNOWHERE_REGISTER_GLOBAL(GPU_RAFT_IVF_FLAT, [](const int32_t& version, const Object& object) {
-    return Index<IndexNodeThreadPoolWrapper>::Create(
-        std::make_unique<RaftIvfIndexNode<detail::raft_ivf_flat_index>>(version, object), cuda_concurrent_size);
-});
-
-KNOWHERE_REGISTER_GLOBAL(GPU_RAFT_IVF_PQ, [](const int32_t& version, const Object& object) {
-    return Index<IndexNodeThreadPoolWrapper>::Create(
-        std::make_unique<RaftIvfIndexNode<detail::raft_ivf_pq_index>>(version, object), cuda_concurrent_size);
-});
-
-KNOWHERE_REGISTER_GLOBAL(GPU_IVF_FLAT, [](const int32_t& version, const Object& object) {
-    return Index<IndexNodeThreadPoolWrapper>::Create(
-        std::make_unique<RaftIvfIndexNode<detail::raft_ivf_flat_index>>(version, object), cuda_concurrent_size);
-});
-
-KNOWHERE_REGISTER_GLOBAL(GPU_IVF_PQ, [](const int32_t& version, const Object& object) {
-    return Index<IndexNodeThreadPoolWrapper>::Create(
-        std::make_unique<RaftIvfIndexNode<detail::raft_ivf_pq_index>>(version, object), cuda_concurrent_size);
-});
-}  // namespace knowhere
diff --git a/src/index/ivf_raft/ivf_raft.cuh b/src/index/ivf_raft/ivf_raft.cuh
deleted file mode 100644
index 02eee67b7..000000000
--- a/src/index/ivf_raft/ivf_raft.cuh
+++ /dev/null
@@ -1,631 +0,0 @@
-/**
- * SPDX-FileCopyrightText: Copyright (c) 2023,NVIDIA CORPORATION & AFFILIATES. All rights reserved.
- * SPDX-License-Identifier: Apache-2.0
- *
- * 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.
- */
-
-#ifndef IVF_RAFT_CUH
-#define IVF_RAFT_CUH
-
-#include <cstddef>
-#include <cstdint>
-#include <limits>
-#include <memory>
-#include <optional>
-#include <raft/neighbors/specializations.cuh>
-
-#include "common/raft/raft_utils.h"
-#include "common/raft_metric.h"
-#include "fmt/core.h"
-#include "index/ivf_raft/ivf_raft_config.h"
-#include "knowhere/comp/index_param.h"
-#include "knowhere/device_bitset.h"
-#include "knowhere/factory.h"
-#include "knowhere/log.h"
-#include "knowhere/utils.h"
-#include "raft/core/device_resources.hpp"
-#include "thrust/execution_policy.h"
-#include "thrust/logical.h"
-#include "thrust/sequence.h"
-
-namespace knowhere {
-
-namespace raft_detail {
-struct raft_results {
-    raft_results(raft::device_resources& res)
-        : ids_{raft::make_device_matrix<std::int64_t, std::int64_t>(res, 0, 0)},
-          dists_{raft::make_device_matrix<float, std::int64_t>(res, 0, 0)} {
-    }
-    raft_results(raft::device_resources& res, std::int64_t rows, std::int64_t k)
-        : ids_{raft::make_device_matrix<std::int64_t, std::int64_t>(res, rows, k)},
-          dists_{raft::make_device_matrix<float, std::int64_t>(res, rows, k)} {
-    }
-    auto
-    ids() {
-        return ids_.view();
-    }
-    auto
-    dists() {
-        return dists_.view();
-    }
-    auto
-    ids_data() {
-        return ids_.data_handle();
-    }
-    auto
-    dists_data() {
-        return dists_.data_handle();
-    }
-
-    auto
-    rows() const {
-        return ids_.extent(0);
-    }
-    auto
-    k() const {
-        return ids_.extent(1);
-    }
-
- private:
-    raft::device_matrix<std::int64_t, std::int64_t> ids_;
-    raft::device_matrix<float, std::int64_t> dists_;
-};
-
-auto constexpr static const MAX_IVF_FLAT_K = 256;
-
-__global__ void
-postprocess_device_results(bool* enough_valid, int64_t* ids, float* dists, int64_t rows, int k, int target_k,
-                           DeviceBitsetView bitset) {
-    __shared__ int invalid_count[1];
-    for (auto row_index = blockIdx.x; row_index < rows; row_index += gridDim.x) {
-        if (threadIdx.x == 0) {
-            invalid_count[0] = 0;
-        }
-        __syncthreads();
-        // First, replace all invalid IDs with -1
-        for (auto col_index = threadIdx.x; col_index < k; col_index += blockDim.x) {
-            auto elem_index = row_index * k + col_index;
-            auto cur_id = ids[elem_index];
-            auto invalid_id = bitset.test(cur_id);
-            if (invalid_id) {  // TODO(wphicks): assuming the branch is worth it here;
-                               // should analyze perf.
-                atomicAdd_block(invalid_count, 1);
-            }
-            invalid_id |= cur_id == std::numeric_limits<int64_t>::max();
-            ids[elem_index] = int(invalid_id) * -1 + int(!invalid_id) * ids[elem_index];
-        }
-        __syncthreads();
-
-        // Now move all valid results to the front of the row
-        auto cur_valid_index = row_index * k;
-        for (auto col_index = 0; col_index < k; ++col_index) {
-            auto elem_index = row_index * k + col_index;
-            // Just do one row per block for now; can improve this later
-            if (ids[elem_index] != -1 && threadIdx.x == 0) {
-                // Swap valid id to an earlier place in the row
-                ids[cur_valid_index] = ids[elem_index];
-                if (elem_index != cur_valid_index) {
-                    ids[elem_index] = -1;
-                    // Only count elements invalidated by the bitset. These are the
-                    // elements that will require a swap as opposed to just appearing at
-                    // the end of the row
-                }
-                dists[cur_valid_index++] = dists[elem_index];
-            }
-        }
-        // Check if we have enough valid results
-        if (threadIdx.x == 0) {
-            enough_valid[row_index] = (k - invalid_count[0] >= target_k);
-        }
-    }
-}
-
-/* The following kernel is used to copy data from one row-major 2D array to
- * another. If the input array has more columns than the output array, the
- * rows will be truncated to the length of the output array. Similarly, the
- * number of rows will be truncated if the output has fewer rows. If the output
- * array has more rows/columns than the input, the output will be padded with
- * entries of -1. */
-template <typename T, typename IndexT>
-__global__ void
-slice(T* out, T* in, IndexT out_rows, IndexT out_cols, IndexT in_rows, IndexT in_cols) {
-    for (auto i = blockIdx.x * blockDim.x + threadIdx.x; i < out_rows * out_cols; i += blockDim.x * gridDim.x) {
-        auto row_index = i / out_cols;
-        auto col_index = i % out_cols;
-        if (row_index < in_rows && col_index < in_cols) {
-            out[i] = in[row_index * in_cols + col_index];
-        } else {
-            out[i] = -1;
-        }
-    }
-}
-
-}  // namespace raft_detail
-
-namespace detail {
-using raft_ivf_flat_index = raft::neighbors::ivf_flat::index<float, std::int64_t>;
-using raft_ivf_pq_index = raft::neighbors::ivf_pq::index<std::int64_t>;
-
-namespace codebook {
-auto static constexpr const PER_SUBSPACE = "PER_SUBSPACE";
-auto static constexpr const PER_CLUSTER = "PER_CLUSTER";
-}  // namespace codebook
-
-inline expected<raft::neighbors::ivf_pq::codebook_gen>
-str_to_codebook_gen(std::string const& str) {
-    static const std::unordered_map<std::string, raft::neighbors::ivf_pq::codebook_gen> name_map = {
-        {codebook::PER_SUBSPACE, raft::neighbors::ivf_pq::codebook_gen::PER_SUBSPACE},
-        {codebook::PER_CLUSTER, raft::neighbors::ivf_pq::codebook_gen::PER_CLUSTER},
-    };
-
-    auto it = name_map.find(str);
-    if (it == name_map.end())
-        return expected<raft::neighbors::ivf_pq::codebook_gen>::Err(Status::invalid_args,
-                                                                    fmt::format("invalid arguments: {}", str));
-    return it->second;
-}
-
-namespace cuda_type {
-auto static constexpr const CUDA_R_16F = "CUDA_R_16F";
-auto static constexpr const CUDA_C_16F = "CUDA_C_16F";
-auto static constexpr const CUDA_R_16BF = "CUDA_R_16BF";
-auto static constexpr const CUDA_C_16BF = "CUDA_C_16BF";
-auto static constexpr const CUDA_R_32F = "CUDA_R_32F";
-auto static constexpr const CUDA_C_32F = "CUDA_C_32F";
-auto static constexpr const CUDA_R_64F = "CUDA_R_64F";
-auto static constexpr const CUDA_C_64F = "CUDA_C_64F";
-auto static constexpr const CUDA_R_8I = "CUDA_R_8I";
-auto static constexpr const CUDA_C_8I = "CUDA_C_8I";
-auto static constexpr const CUDA_R_8U = "CUDA_R_8U";
-auto static constexpr const CUDA_C_8U = "CUDA_C_8U";
-auto static constexpr const CUDA_R_32I = "CUDA_R_32I";
-auto static constexpr const CUDA_C_32I = "CUDA_C_32I";
-auto static constexpr const CUDA_R_8F_E4M3 = "CUDA_R_8F_E4M3";
-auto static constexpr const CUDA_R_8F_E5M2 = "CUDA_R_8F_E5M2";
-}  // namespace cuda_type
-
-inline expected<cudaDataType_t>
-str_to_cuda_dtype(std::string const& str) {
-    static const std::unordered_map<std::string, cudaDataType_t> name_map = {
-        {cuda_type::CUDA_R_16F, CUDA_R_16F},   {cuda_type::CUDA_C_16F, CUDA_C_16F},
-        {cuda_type::CUDA_R_16BF, CUDA_R_16BF}, {cuda_type::CUDA_C_16BF, CUDA_C_16BF},
-        {cuda_type::CUDA_R_32F, CUDA_R_32F},   {cuda_type::CUDA_C_32F, CUDA_C_32F},
-        {cuda_type::CUDA_R_64F, CUDA_R_64F},   {cuda_type::CUDA_C_64F, CUDA_C_64F},
-        {cuda_type::CUDA_R_8I, CUDA_R_8I},     {cuda_type::CUDA_C_8I, CUDA_C_8I},
-        {cuda_type::CUDA_R_8U, CUDA_R_8U},     {cuda_type::CUDA_C_8U, CUDA_C_8U},
-        {cuda_type::CUDA_R_32I, CUDA_R_32I},   {cuda_type::CUDA_C_32I, CUDA_C_32I},
-        // not support, when we use cuda 11.6
-        //{cuda_type::CUDA_R_8F_E4M3, CUDA_R_8F_E4M3}, {cuda_type::CUDA_R_8F_E5M2, CUDA_R_8F_E5M2},
-
-    };
-
-    auto it = name_map.find(str);
-    if (it == name_map.end())
-        return expected<cudaDataType_t>::Err(Status::invalid_args, fmt::format("invalid arguments: {}", str));
-    return it->second;
-}
-
-}  // namespace detail
-
-template <typename T>
-struct KnowhereConfigType {};
-
-template <>
-struct KnowhereConfigType<detail::raft_ivf_flat_index> {
-    typedef RaftIvfFlatConfig Type;
-};
-template <>
-struct KnowhereConfigType<detail::raft_ivf_pq_index> {
-    typedef RaftIvfPqConfig Type;
-};
-
-template <typename T>
-class RaftIvfIndexNode : public IndexNode {
- public:
-    RaftIvfIndexNode(const int32_t& /*version*/, const Object& object) : device_id_{-1}, gpu_index_{} {
-    }
-
-    Status
-    Train(const DataSet& dataset, const Config& cfg) override {
-        auto ivf_raft_cfg = static_cast<const typename KnowhereConfigType<T>::Type&>(cfg);
-        if (gpu_index_) {
-            LOG_KNOWHERE_WARNING_ << "index is already trained";
-            return Status::index_already_trained;
-        }
-        try {
-            RANDOM_CHOOSE_DEVICE_WITH_ASSIGN(this->device_id_);
-            raft_utils::device_setter with_this_device(this->device_id_);
-            raft_utils::init_gpu_resources();
-
-            auto metric = Str2RaftMetricType(ivf_raft_cfg.metric_type.value());
-            if (!metric.has_value()) {
-                LOG_KNOWHERE_WARNING_ << "please check metric value: " << ivf_raft_cfg.metric_type.value();
-                return metric.error();
-            }
-            if (metric.value() != raft::distance::DistanceType::L2Expanded &&
-                metric.value() != raft::distance::DistanceType::InnerProduct) {
-                LOG_KNOWHERE_WARNING_ << "selected metric not supported in RAFT IVF indexes: "
-                                      << ivf_raft_cfg.metric_type.value();
-                return Status::invalid_metric_type;
-            }
-            auto& res = raft_utils::get_raft_resources();
-
-            auto rows = dataset.GetRows();
-            auto dim = dataset.GetDim();
-            auto* data = reinterpret_cast<float const*>(dataset.GetTensor());
-
-            auto data_gpu = raft::make_device_matrix<float, std::int64_t>(res, rows, dim);
-            RAFT_CUDA_TRY(cudaMemcpyAsync(data_gpu.data_handle(), data, data_gpu.size() * sizeof(float),
-                                          cudaMemcpyDefault, res.get_stream().value()));
-            if constexpr (std::is_same_v<detail::raft_ivf_flat_index, T>) {
-                auto build_params = raft::neighbors::ivf_flat::index_params{};
-                build_params.metric = metric.value();
-                build_params.n_lists = ivf_raft_cfg.nlist.value();
-                build_params.kmeans_n_iters = ivf_raft_cfg.kmeans_n_iters.value();
-                build_params.kmeans_trainset_fraction = ivf_raft_cfg.kmeans_trainset_fraction.value();
-                build_params.adaptive_centers = ivf_raft_cfg.adaptive_centers.value();
-                gpu_index_ = raft::neighbors::ivf_flat::build<float, std::int64_t>(res, build_params, data_gpu.view());
-            } else if constexpr (std::is_same_v<detail::raft_ivf_pq_index, T>) {
-                auto build_params = raft::neighbors::ivf_pq::index_params{};
-                build_params.metric = metric.value();
-                build_params.n_lists = ivf_raft_cfg.nlist.value();
-                build_params.pq_bits = ivf_raft_cfg.nbits.value();
-                build_params.kmeans_n_iters = ivf_raft_cfg.kmeans_n_iters.value();
-                build_params.kmeans_trainset_fraction = ivf_raft_cfg.kmeans_trainset_fraction.value();
-                build_params.pq_dim = ivf_raft_cfg.m.value();
-                auto codebook_kind = detail::str_to_codebook_gen(ivf_raft_cfg.codebook_kind.value());
-                if (!codebook_kind.has_value()) {
-                    LOG_KNOWHERE_WARNING_ << "please check codebook kind: " << ivf_raft_cfg.codebook_kind.value();
-                    return codebook_kind.error();
-                }
-                build_params.codebook_kind = codebook_kind.value();
-                build_params.force_random_rotation = ivf_raft_cfg.force_random_rotation.value();
-                gpu_index_ = raft::neighbors::ivf_pq::build<float, std::int64_t>(res, build_params, data_gpu.view());
-            }
-            dim_ = dim;
-            counts_ = rows;
-            res.sync_stream();
-
-        } catch (std::exception& e) {
-            LOG_KNOWHERE_WARNING_ << "RAFT inner error, " << e.what();
-            return Status::raft_inner_error;
-        }
-        return Status::success;
-    }
-
-    Status
-    Add(const DataSet& dataset, const Config& cfg) override {
-        if (!gpu_index_) {
-            return Status::index_not_trained;
-        }
-        try {
-            RAFT_EXPECTS(this->device_id_ != -1, "call data add before index train.");
-            raft_utils::device_setter with_this_device{this->device_id_};
-            auto rows = dataset.GetRows();
-            auto dim = dataset.GetDim();
-            auto* data = reinterpret_cast<float const*>(dataset.GetTensor());
-
-            raft_utils::init_gpu_resources();
-            auto& res = raft_utils::get_raft_resources();
-
-            // TODO(wphicks): Clean up transfer with raft
-            // buffer objects when available
-            auto data_gpu = raft::make_device_matrix<float, std::int64_t>(res, rows, dim);
-            RAFT_CUDA_TRY(cudaMemcpyAsync(data_gpu.data_handle(), data, data_gpu.size() * sizeof(float),
-                                          cudaMemcpyDefault, res.get_stream().value()));
-
-            auto indices = rmm::device_uvector<std::int64_t>(rows, res.get_stream());
-            thrust::sequence(res.get_thrust_policy(), indices.begin(), indices.end(), gpu_index_->size());
-
-            if constexpr (std::is_same_v<detail::raft_ivf_flat_index, T>) {
-                raft::neighbors::ivf_flat::extend<float, std::int64_t>(
-                    res, raft::make_const_mdspan(data_gpu.view()),
-                    std::make_optional(
-                        raft::make_device_vector_view<const std::int64_t, std::int64_t>(indices.data(), rows)),
-                    gpu_index_.value());
-            } else if constexpr (std::is_same_v<detail::raft_ivf_pq_index, T>) {
-                raft::neighbors::ivf_pq::extend<float, std::int64_t>(
-                    res, raft::make_const_mdspan(data_gpu.view()),
-                    std::make_optional(
-                        raft::make_device_matrix_view<const std::int64_t, std::int64_t>(indices.data(), rows, 1)),
-                    gpu_index_.value());
-            } else {
-                static_assert(std::is_same_v<detail::raft_ivf_flat_index, T>);
-            }
-            dim_ = dim;
-            counts_ = rows;
-        } catch (std::exception& e) {
-            LOG_KNOWHERE_WARNING_ << "RAFT inner error, " << e.what();
-            return Status::raft_inner_error;
-        }
-
-        return Status::success;
-    }
-
-    expected<DataSetPtr>
-    Search(const DataSet& dataset, const Config& cfg, const BitsetView& bitset) const override {
-        auto ivf_raft_cfg = static_cast<const typename KnowhereConfigType<T>::Type&>(cfg);
-        auto rows = dataset.GetRows();
-        auto dim = dataset.GetDim();
-        auto* data = reinterpret_cast<float const*>(dataset.GetTensor());
-        auto output_size = rows * ivf_raft_cfg.k.value();
-        auto ids = std::unique_ptr<std::int64_t[]>(new std::int64_t[output_size]);
-        auto dis = std::unique_ptr<float[]>(new float[output_size]);
-        try {
-            RAFT_EXPECTS(this->device_id_ != -1, "device id is -1, when call search");
-            raft_utils::device_setter with_this_device{this->device_id_};
-            auto& res_ = raft_utils::get_raft_resources();
-
-            // TODO(wphicks): Clean up transfer with raft
-            // buffer objects when available
-            auto data_gpu = raft::make_device_matrix<float, std::int64_t>(res_, rows, dim);
-            raft::copy(data_gpu.data_handle(), data, data_gpu.size(), res_.get_stream());
-
-            auto gpu_results = raft_detail::raft_results{res_};
-            auto gpu_bitset = DeviceBitset{res_, bitset};
-
-            if constexpr (std::is_same_v<detail::raft_ivf_flat_index, T>) {
-                auto search_params = raft::neighbors::ivf_flat::search_params{};
-                search_params.n_probes = std::min<uint32_t>(ivf_raft_cfg.nprobe.value(), gpu_index_->n_lists());
-                auto search_k = std::min(
-                    ivf_raft_cfg.k.value() + (bitset.count() * ivf_raft_cfg.k.value() / counts_),
-                    std::min(static_cast<uint64_t>(counts_), static_cast<uint64_t>(raft_detail::MAX_IVF_FLAT_K)));
-                gpu_results = RawSearch(res_, raft::make_const_mdspan(data_gpu.view()), search_params, search_k,
-                                        ivf_raft_cfg.k.value(), gpu_bitset.view());
-            } else if constexpr (std::is_same_v<detail::raft_ivf_pq_index, T>) {
-                auto search_params = raft::neighbors::ivf_pq::search_params{};
-                search_params.n_probes = std::min<uint32_t>(ivf_raft_cfg.nprobe.value(), gpu_index_->n_lists());
-                auto lut_dtype = detail::str_to_cuda_dtype(ivf_raft_cfg.lut_dtype.value());
-                if (!lut_dtype.has_value()) {
-                    LOG_KNOWHERE_WARNING_ << "please check lookup dtype: " << ivf_raft_cfg.lut_dtype.value();
-                    return expected<DataSetPtr>::Err(
-                        lut_dtype.error(), fmt::format("invalid lookup dtype: {}", ivf_raft_cfg.lut_dtype.value()));
-                }
-                if (lut_dtype.value() != CUDA_R_32F && lut_dtype.value() != CUDA_R_16F &&
-                    lut_dtype.value() != CUDA_R_8U) {
-                    LOG_KNOWHERE_WARNING_ << "selected lookup dtype not supported: " << ivf_raft_cfg.lut_dtype.value();
-                    return expected<DataSetPtr>::Err(Status::invalid_args,
-                                                     fmt::format("invalid lookup dtype: {}", lut_dtype.value()));
-                }
-                search_params.lut_dtype = lut_dtype.value();
-                auto internal_distance_dtype = detail::str_to_cuda_dtype(ivf_raft_cfg.internal_distance_dtype.value());
-                if (!internal_distance_dtype.has_value()) {
-                    LOG_KNOWHERE_WARNING_ << "please check internal distance dtype: "
-                                          << ivf_raft_cfg.internal_distance_dtype.value();
-                    return expected<DataSetPtr>::Err(internal_distance_dtype.error(),
-                                                     fmt::format("invalid internal distance dtype: {}",
-                                                                 ivf_raft_cfg.internal_distance_dtype.value()));
-                }
-                if (internal_distance_dtype.value() != CUDA_R_32F && internal_distance_dtype.value() != CUDA_R_16F) {
-                    LOG_KNOWHERE_WARNING_ << "selected internal distance dtype not supported: "
-                                          << ivf_raft_cfg.internal_distance_dtype.value();
-                    return expected<DataSetPtr>::Err(
-                        Status::invalid_args, fmt::format("invalid lookup dtype: {}", internal_distance_dtype.value()));
-                }
-                search_params.internal_distance_dtype = internal_distance_dtype.value();
-                search_params.preferred_shmem_carveout = search_params.preferred_shmem_carveout;
-                gpu_results = RawSearch(res_, raft::make_const_mdspan(data_gpu.view()), search_params,
-                                        ivf_raft_cfg.k.value(), ivf_raft_cfg.k.value(), gpu_bitset.view());
-            } else {
-                static_assert(std::is_same_v<detail::raft_ivf_flat_index, T>);
-            }
-
-            if (gpu_results.k() != ivf_raft_cfg.k.value()) {
-                auto new_gpu_results = raft_detail::raft_results{res_, gpu_results.rows(), ivf_raft_cfg.k.value()};
-                raft_detail::slice<<<1024, 256, 0, res_.get_stream().value()>>>(
-                    new_gpu_results.ids_data(), gpu_results.ids_data(), new_gpu_results.rows(), new_gpu_results.k(),
-                    gpu_results.rows(), gpu_results.k());
-                res_.sync_stream();
-                gpu_results = new_gpu_results;
-            }
-            raft::copy(ids.get(), gpu_results.ids_data(), output_size, res_.get_stream());
-            raft::copy(dis.get(), gpu_results.dists_data(), output_size, res_.get_stream());
-            res_.sync_stream();
-
-        } catch (std::exception& e) {
-            LOG_KNOWHERE_WARNING_ << "RAFT inner error, " << e.what();
-            return expected<DataSetPtr>::Err(Status::raft_inner_error, e.what());
-        }
-
-        return GenResultDataSet(rows, ivf_raft_cfg.k.value(), ids.release(), dis.release());
-    }
-
-    expected<DataSetPtr>
-    RangeSearch(const DataSet& dataset, const Config& cfg, const BitsetView& bitset) const override {
-        return expected<DataSetPtr>::Err(Status::not_implemented, "RangeSearch not implemented");
-    }
-
-    expected<DataSetPtr>
-    GetVectorByIds(const DataSet& dataset) const override {
-        return expected<DataSetPtr>::Err(Status::not_implemented, "GetVectorByIds not implemented");
-    }
-
-    bool
-    HasRawData(const std::string& metric_type) const override {
-        if constexpr (std::is_same_v<detail::raft_ivf_flat_index, T>) {
-            return !IsMetricType(metric_type, metric::COSINE);
-        }
-        if constexpr (std::is_same_v<detail::raft_ivf_pq_index, T>) {
-            return false;
-        }
-    }
-
-    expected<DataSetPtr>
-    GetIndexMeta(const Config& cfg) const override {
-        return expected<DataSetPtr>::Err(Status::not_implemented, "GetIndexMeta not implemented");
-    }
-
-    Status
-    Serialize(BinarySet& binset) const override {
-        if (!gpu_index_.has_value()) {
-            LOG_KNOWHERE_ERROR_ << "Can not serialize empty RaftIvfIndex.";
-            return Status::empty_index;
-        }
-
-        RAFT_EXPECTS(this->device_id_ != -1, "index serialize before trained.");
-        std::stringbuf buf;
-
-        std::ostream os(&buf);
-
-        os.write((char*)(&this->dim_), sizeof(this->dim_));
-        os.write((char*)(&this->counts_), sizeof(this->counts_));
-        os.write((char*)(&this->device_id_), sizeof(this->device_id_));
-
-        raft_utils::device_setter with_this_device{device_id_};
-        auto& res = raft_utils::get_raft_resources();
-
-        if constexpr (std::is_same_v<T, detail::raft_ivf_flat_index>) {
-            raft::neighbors::ivf_flat::serialize<float, std::int64_t>(res, os, *gpu_index_);
-        }
-        if constexpr (std::is_same_v<T, detail::raft_ivf_pq_index>) {
-            raft::neighbors::ivf_pq::serialize<std::int64_t>(res, os, *gpu_index_);
-        }
-        res.sync_stream();
-
-        os.flush();
-        std::shared_ptr<uint8_t[]> index_binary(new (std::nothrow) uint8_t[buf.str().size()]);
-
-        memcpy(index_binary.get(), buf.str().c_str(), buf.str().size());
-        binset.Append(this->Type(), index_binary, buf.str().size());
-        return Status::success;
-    }
-
-    Status
-    Deserialize(const BinarySet& binset, const Config& config) override {
-        std::stringbuf buf;
-        auto binary = binset.GetByName(this->Type());
-        if (binary == nullptr) {
-            LOG_KNOWHERE_ERROR_ << "Invalid binary set.";
-            return Status::invalid_binary_set;
-        }
-        buf.sputn((char*)binary->data.get(), binary->size);
-        std::istream is(&buf);
-
-        is.read((char*)(&this->dim_), sizeof(this->dim_));
-        is.read((char*)(&this->counts_), sizeof(this->counts_));
-        // device_id from binset is useless, will gen device id from global
-        // status
-        is.read((char*)(&this->device_id_), sizeof(this->device_id_));
-        MIN_LOAD_CHOOSE_DEVICE_WITH_ASSIGN(this->device_id_, binary->size);
-        load_ = binary->size;
-        raft_utils::device_setter with_this_device{this->device_id_};
-
-        raft_utils::init_gpu_resources();
-        auto& res = raft_utils::get_raft_resources();
-
-        if constexpr (std::is_same_v<T, detail::raft_ivf_flat_index>) {
-            T index_ = raft::neighbors::ivf_flat::deserialize<float, std::int64_t>(res, is);
-            res.sync_stream();
-            is.sync();
-            gpu_index_ = T(std::move(index_));
-        }
-        if constexpr (std::is_same_v<T, detail::raft_ivf_pq_index>) {
-            T index_ = raft::neighbors::ivf_pq::deserialize<std::int64_t>(res, is);
-            res.sync_stream();
-            is.sync();
-            gpu_index_ = T(std::move(index_));
-        }
-
-        return Status::success;
-    }
-
-    Status
-    DeserializeFromFile(const std::string& filename, const Config& config) {
-        LOG_KNOWHERE_ERROR_ << "RaftIvfIndex doesn't support Deserialization from file.";
-        return Status::not_implemented;
-    }
-
-    std::unique_ptr<BaseConfig>
-    CreateConfig() const override {
-        return std::make_unique<typename KnowhereConfigType<T>::Type>();
-    }
-
-    int64_t
-    Dim() const override {
-        return dim_;
-    }
-
-    int64_t
-    Size() const override {
-        return 0;
-    }
-
-    int64_t
-    Count() const override {
-        return counts_;
-    }
-
-    std::string
-    Type() const override {
-        if constexpr (std::is_same_v<detail::raft_ivf_flat_index, T>) {
-            return knowhere::IndexEnum::INDEX_RAFT_IVFFLAT;
-        }
-        if constexpr (std::is_same_v<detail::raft_ivf_pq_index, T>) {
-            return knowhere::IndexEnum::INDEX_RAFT_IVFPQ;
-        }
-    }
-    virtual ~RaftIvfIndexNode() {
-        if (device_id_ >= 0) {
-            RELEASE_DEVICE(this->device_id_, this->load_);
-        }
-    }
-
- private:
-    int device_id_ = -1;
-    int64_t dim_ = 0;
-    int64_t counts_ = 0;
-    size_t load_ = 0;
-    std::optional<T> gpu_index_;
-
-    template <typename raft_search_params_t>
-    raft_detail::raft_results
-    RawSearch(raft::device_resources& res, raft::device_matrix_view<const float, std::int64_t> queries,
-              raft_search_params_t const& search_params, int k, int target_k, DeviceBitsetView const& bitset) const {
-        auto max_k = counts_;
-        if constexpr (std::is_same_v<detail::raft_ivf_flat_index, T>) {
-            k = std::min(k, raft_detail::MAX_IVF_FLAT_K);
-            max_k = std::min(max_k, int64_t{raft_detail::MAX_IVF_FLAT_K});
-        }
-        auto result = raft_detail::raft_results{res, queries.extent(0), k};
-        if constexpr (std::is_same_v<detail::raft_ivf_flat_index, T>) {
-            raft::neighbors::ivf_flat::search<float, std::int64_t>(res, search_params, *gpu_index_, queries,
-                                                                   result.ids(), result.dists());
-        } else if constexpr (std::is_same_v<detail::raft_ivf_pq_index, T>) {
-            raft::neighbors::ivf_pq::search<float, std::int64_t>(res, search_params, *gpu_index_, queries, result.ids(),
-                                                                 result.dists());
-        } else {
-            static_assert(std::is_same_v<detail::raft_ivf_flat_index, T>);
-        }
-
-        auto blocks = std::min(queries.extent(0), int64_t{1024});
-        auto threads = std::min(k, 256);
-        auto warp_remainder = threads % 32;
-        if (warp_remainder != 0) {
-            threads += (32 - warp_remainder);
-        }
-        auto enough_valid = raft::make_device_vector<bool>(res, queries.extent(0));
-
-        raft_detail::postprocess_device_results<<<blocks, threads, 0, res.get_stream().value()>>>(
-            enough_valid.data_handle(), result.ids_data(), result.dists_data(), queries.extent(0), k, target_k, bitset);
-
-        if (k < max_k && !thrust::all_of(res.get_thrust_policy(), enough_valid.data_handle(),
-                                         enough_valid.data_handle() + queries.extent(0), thrust::identity<bool>())) {
-            result = RawSearch(res, queries, search_params, std::min(int64_t{k * 2}, max_k), target_k, bitset);
-        }
-        return result;
-    }
-};
-
-}  // namespace knowhere
-#endif /* IVF_RAFT_CUH */
diff --git a/tests/ut/test_gpu_search.cc b/tests/ut/test_gpu_search.cc
index 22c430fb4..6c4e7cb56 100644
--- a/tests/ut/test_gpu_search.cc
+++ b/tests/ut/test_gpu_search.cc
@@ -9,6 +9,8 @@
 // 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.
 
+#include <string>
+
 #include "catch2/catch_approx.hpp"
 #include "catch2/catch_test_macros.hpp"
 #include "catch2/generators/catch_generators.hpp"
@@ -45,8 +47,6 @@ TEST_CASE("Test All GPU Index", "[search]") {
         return json;
     };
 
-    auto ivfsq_gen = ivfflat_gen;
-
     auto ivfpq_gen = [&ivfflat_gen]() {
         knowhere::Json json = ivfflat_gen();
         json[knowhere::indexparam::M] = 0;
@@ -54,11 +54,22 @@ TEST_CASE("Test All GPU Index", "[search]") {
         return json;
     };
 
-    auto gpu_flat_gen = [&base_gen]() {
-        auto json = base_gen();
+    auto cagra_gen = [&base_gen]() {
+        knowhere::Json json = base_gen();
+        json[knowhere::indexparam::INTERMEDIATE_GRAPH_DEGREE] = 128;
+        json[knowhere::indexparam::GRAPH_DEGREE] = 64;
+        json[knowhere::indexparam::ITOPK_SIZE] = 128;
         return json;
     };
 
+    auto refined_gen = [](auto&& upstream_gen) {
+        return [upstream_gen]() {
+            knowhere::Json json = upstream_gen();
+            json[knowhere::indexparam::REFINE_RATIO] = 1.5;
+            return json;
+        };
+    };
+
     SECTION("Test Gpu Index Search") {
         using std::make_tuple;
         auto [name, gen] = GENERATE_REF(table<std::string, std::function<knowhere::Json()>>({
@@ -69,7 +80,10 @@ TEST_CASE("Test All GPU Index", "[search]") {
             // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFPQ, ivfpq_gen),
             // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFSQ8, ivfsq_gen),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, ivfflat_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, refined_gen(ivfflat_gen)),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, ivfpq_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, refined_gen(ivfpq_gen)),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_CAGRA, cagra_gen),
         }));
         auto idx = knowhere::IndexFactory::Instance().Create(name, version);
         auto cfg_json = gen().dump();
@@ -91,14 +105,11 @@ TEST_CASE("Test All GPU Index", "[search]") {
     SECTION("Test Gpu Index Search With Bitset") {
         using std::make_tuple;
         auto [name, gen] = GENERATE_REF(table<std::string, std::function<knowhere::Json()>>({
-            // GPU_FLAT cannot run this test is because its Train() and Add() actually run in CPU,
-            // "res_" in gpu_index_ is not set correctly
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IDMAP, gpu_flat_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFFLAT, ivfflat_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFPQ, ivfpq_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFSQ8, ivfsq_gen),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, ivfflat_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, refined_gen(ivfflat_gen)),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, ivfpq_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, refined_gen(ivfpq_gen)),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_CAGRA, cagra_gen),
         }));
         auto idx = knowhere::IndexFactory::Instance().Create(name, version);
         auto cfg_json = gen().dump();
@@ -133,14 +144,11 @@ TEST_CASE("Test All GPU Index", "[search]") {
     SECTION("Test Gpu Index Search TopK") {
         using std::make_tuple;
         auto [name, gen] = GENERATE_REF(table<std::string, std::function<knowhere::Json()>>({
-            // GPU_FLAT cannot run this test is because its Train() and Add() actually run in CPU,
-            // "res_" in gpu_index_ is not set correctly
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IDMAP, gpu_flat_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFFLAT, ivfflat_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFPQ, ivfpq_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFSQ8, ivfsq_gen),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, ivfflat_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, refined_gen(ivfflat_gen)),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, ivfpq_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, refined_gen(ivfpq_gen)),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_CAGRA, cagra_gen),
         }));
         auto idx = knowhere::IndexFactory::Instance().Create(name, version);
         auto cfg_json = gen().dump();
@@ -167,12 +175,11 @@ TEST_CASE("Test All GPU Index", "[search]") {
     SECTION("Test Gpu Index Serialize/Deserialize") {
         using std::make_tuple;
         auto [name, gen] = GENERATE_REF(table<std::string, std::function<knowhere::Json()>>({
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IDMAP, gpu_flat_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFFLAT, ivfflat_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFPQ, ivfpq_gen),
-            // make_tuple(knowhere::IndexEnum::INDEX_FAISS_GPU_IVFSQ8, ivfsq_gen),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, ivfflat_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, refined_gen(ivfflat_gen)),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, ivfpq_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, refined_gen(ivfpq_gen)),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_CAGRA, cagra_gen),
         }));
 
         auto idx = knowhere::IndexFactory::Instance().Create(name, version);
@@ -200,7 +207,10 @@ TEST_CASE("Test All GPU Index", "[search]") {
         using std::make_tuple;
         auto [name, gen] = GENERATE_REF(table<std::string, std::function<knowhere::Json()>>({
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, ivfflat_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFFLAT, refined_gen(ivfflat_gen)),
             make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, ivfpq_gen),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_IVFPQ, refined_gen(ivfpq_gen)),
+            make_tuple(knowhere::IndexEnum::INDEX_RAFT_CAGRA, cagra_gen),
         }));
         auto rows = 16;
         auto idx = knowhere::IndexFactory::Instance().Create(name, version);