diff --git a/c/include/cuvs/neighbors/ivf_pq.h b/c/include/cuvs/neighbors/ivf_pq.h
index 8d66662a09..5b95dfe7c5 100644
--- a/c/include/cuvs/neighbors/ivf_pq.h
+++ b/c/include/cuvs/neighbors/ivf_pq.h
@@ -23,9 +23,18 @@ extern "C" {
  * @brief A type for specifying how PQ codebooks are created
  *
  */
-enum codebook_gen {  // NOLINT
-  PER_SUBSPACE = 0,  // NOLINT
-  PER_CLUSTER  = 1,  // NOLINT
+enum cuvsIvfPqCodebookGen {
+  CUVS_IVF_PQ_CODEBOOK_GEN_PER_SUBSPACE = 0,
+  CUVS_IVF_PQ_CODEBOOK_GEN_PER_CLUSTER  = 1,
+};
+
+/**
+ * @brief A type for specifying the memory layout of IVF-PQ list data
+ *
+ */
+enum cuvsIvfPqListLayout {
+  CUVS_IVF_PQ_LIST_LAYOUT_FLAT        = 0,
+  CUVS_IVF_PQ_LIST_LAYOUT_INTERLEAVED = 1,
 };
 
 /**
@@ -80,7 +89,7 @@ struct cuvsIvfPqIndexParams {
    */
   uint32_t pq_dim;
   /** How PQ codebooks are created. */
-  enum codebook_gen codebook_kind;
+  enum cuvsIvfPqCodebookGen codebook_kind;
   /**
    * Apply a random rotation matrix on the input data and queries even if `dim % pq_dim == 0`.
    *
@@ -114,6 +123,14 @@ struct cuvsIvfPqIndexParams {
    * points to train each codebook.
    */
   uint32_t max_train_points_per_pq_code;
+  /**
+   * Memory layout of the IVF-PQ list data.
+   *
+   * - CUVS_IVF_PQ_LIST_LAYOUT_FLAT: Codes are stored contiguously, one vector's codes after another.
+   * - CUVS_IVF_PQ_LIST_LAYOUT_INTERLEAVED: Codes are interleaved for optimized search performance.
+   *   This is the default and recommended for search workloads.
+   */
+  enum cuvsIvfPqListLayout codes_layout;
 };
 
 typedef struct cuvsIvfPqIndexParams* cuvsIvfPqIndexParams_t;
@@ -294,8 +311,8 @@ cuvsError_t cuvsIvfPqIndexGetCentersPadded(cuvsIvfPqIndex_t index, DLManagedTens
 /**
  * @brief Get the PQ cluster centers
  *
- *   - codebook_gen::PER_SUBSPACE: [pq_dim , pq_len, pq_book_size]
- *   - codebook_gen::PER_CLUSTER:  [n_lists, pq_len, pq_book_size]
+ *   - CUVS_IVF_PQ_CODEBOOK_GEN_PER_SUBSPACE: [pq_dim , pq_len, pq_book_size]
+ *   - CUVS_IVF_PQ_CODEBOOK_GEN_PER_CLUSTER:  [n_lists, pq_len, pq_book_size]
  *
  * @param[in] index cuvsIvfPqIndex_t Built Ivf-Pq index
  * @param[out] pq_centers Output tensor that will be populated with a non-owning view of the data
@@ -443,8 +460,8 @@ cuvsError_t cuvsIvfPqBuild(cuvsResources_t res,
  * matrices)
  * @param[in] dim dimensionality of the input data
  * @param[in] pq_centers PQ codebook on device memory with required shape:
- *   - codebook_kind PER_SUBSPACE: [pq_dim, pq_len, pq_book_size]
- *   - codebook_kind PER_CLUSTER:  [n_lists, pq_len, pq_book_size]
+ *   - codebook_kind CUVS_IVF_PQ_CODEBOOK_GEN_PER_SUBSPACE: [pq_dim, pq_len, pq_book_size]
+ *   - codebook_kind CUVS_IVF_PQ_CODEBOOK_GEN_PER_CLUSTER:  [n_lists, pq_len, pq_book_size]
  * @param[in] centers Cluster centers in the original space [n_lists, dim_ext]
  *   where dim_ext = round_up(dim + 1, 8)
  * @param[in] centers_rot Rotated cluster centers [n_lists, rot_dim]
diff --git a/c/src/neighbors/cagra.cpp b/c/src/neighbors/cagra.cpp
index 611ef3e086..e344aed415 100644
--- a/c/src/neighbors/cagra.cpp
+++ b/c/src/neighbors/cagra.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -365,7 +365,7 @@ static void _populate_c_ivf_pq_params(cuvsIvfPqParams* c_ivf_pq,
   c_ivf_pq->ivf_pq_build_params->kmeans_trainset_fraction = bp.kmeans_trainset_fraction;
   c_ivf_pq->ivf_pq_build_params->pq_bits = bp.pq_bits;
   c_ivf_pq->ivf_pq_build_params->pq_dim = bp.pq_dim;
-  c_ivf_pq->ivf_pq_build_params->codebook_kind = static_cast<codebook_gen>(bp.codebook_kind);
+  c_ivf_pq->ivf_pq_build_params->codebook_kind = static_cast<cuvsIvfPqCodebookGen>(bp.codebook_kind);
   c_ivf_pq->ivf_pq_build_params->force_random_rotation = bp.force_random_rotation;
   c_ivf_pq->ivf_pq_build_params->conservative_memory_allocation = bp.conservative_memory_allocation;
   c_ivf_pq->ivf_pq_build_params->max_train_points_per_pq_code = bp.max_train_points_per_pq_code;
diff --git a/c/src/neighbors/ivf_pq.cpp b/c/src/neighbors/ivf_pq.cpp
index f99180697a..4ffff4bfaa 100644
--- a/c/src/neighbors/ivf_pq.cpp
+++ b/c/src/neighbors/ivf_pq.cpp
@@ -34,6 +34,7 @@ void convert_c_index_params(cuvsIvfPqIndexParams params, cuvs::neighbors::ivf_pq
   out->force_random_rotation          = params.force_random_rotation;
   out->conservative_memory_allocation = params.conservative_memory_allocation;
   out->max_train_points_per_pq_code   = params.max_train_points_per_pq_code;
+  out->codes_layout = static_cast<cuvs::neighbors::ivf_pq::list_layout>((int)params.codes_layout);
 }
 void convert_c_search_params(cuvsIvfPqSearchParams params,
                              cuvs::neighbors::ivf_pq::search_params* out)
@@ -218,8 +219,16 @@ void _get_list_indices(cuvsIvfPqIndex index,
                        uint32_t label,
                        DLManagedTensor* out_labels)
 {
-  auto index_ptr    = reinterpret_cast<cuvs::neighbors::ivf_pq::index<IdxT>*>(index.addr);
-  cuvs::core::to_dlpack(index_ptr->lists()[label]->indices.view(), out_labels);
+  auto index_ptr = reinterpret_cast<cuvs::neighbors::ivf_pq::index<IdxT>*>(index.addr);
+  if (index_ptr->codes_layout() == cuvs::neighbors::ivf_pq::list_layout::FLAT) {
+    auto& list =
+      static_cast<cuvs::neighbors::ivf_pq::list_data_flat<IdxT>&>(*index_ptr->lists()[label]);
+    cuvs::core::to_dlpack(list.indices.view(), out_labels);
+  } else {
+    auto& list = static_cast<cuvs::neighbors::ivf_pq::list_data_interleaved<IdxT>&>(
+      *index_ptr->lists()[label]);
+    cuvs::core::to_dlpack(list.indices.view(), out_labels);
+  }
 }
 }  // namespace
 
@@ -325,10 +334,11 @@ extern "C" cuvsError_t cuvsIvfPqIndexParamsCreate(cuvsIvfPqIndexParams_t* params
                                        .kmeans_trainset_fraction       = 0.5,
                                        .pq_bits                        = 8,
                                        .pq_dim                         = 0,
-                                       .codebook_kind                  = codebook_gen::PER_SUBSPACE,
+                                       .codebook_kind                  = CUVS_IVF_PQ_CODEBOOK_GEN_PER_SUBSPACE,
                                        .force_random_rotation          = false,
                                        .conservative_memory_allocation = false,
-                                       .max_train_points_per_pq_code   = 256};
+                                       .max_train_points_per_pq_code   = 256,
+                                       .codes_layout                   = CUVS_IVF_PQ_LIST_LAYOUT_INTERLEAVED};
   });
 }
 
diff --git a/cpp/cmake/patches/faiss-1.13-cuvs-26.02.diff b/cpp/cmake/patches/faiss-1.13-cuvs-26.02.diff
index 1e4443416a..50199e0975 100644
--- a/cpp/cmake/patches/faiss-1.13-cuvs-26.02.diff
+++ b/cpp/cmake/patches/faiss-1.13-cuvs-26.02.diff
@@ -1,5 +1,5 @@
 diff --git a/faiss/gpu/impl/CuvsIVFPQ.cu b/faiss/gpu/impl/CuvsIVFPQ.cu
-index 1e2fef225..35b388147 100644
+index 1e2fef225..2ee40da46 100644
 --- a/faiss/gpu/impl/CuvsIVFPQ.cu
 +++ b/faiss/gpu/impl/CuvsIVFPQ.cu
 @@ -129,8 +129,14 @@ void CuvsIVFPQ::updateQuantizer(Index* quantizer) {
@@ -122,7 +122,35 @@ index 1e2fef225..35b388147 100644
      }
 
      setPQCentroids_();
-@@ -520,7 +583,7 @@ void CuvsIVFPQ::setPQCentroids_() {
+@@ -404,10 +467,11 @@ void CuvsIVFPQ::copyInvertedListsFrom(const InvertedLists* ivf) {
+     auto& cuvs_index_lists = cuvs_index->lists();
+
+     // conservative memory alloc for cloning cpu inverted lists
+-    cuvs::neighbors::ivf_pq::list_spec<uint32_t, idx_t> ivf_list_spec{
+-            static_cast<uint32_t>(bitsPerSubQuantizer_),
+-            static_cast<uint32_t>(numSubQuantizers_),
+-            true};
++    cuvs::neighbors::ivf_pq::list_spec_interleaved<uint32_t, idx_t>
++            ivf_list_spec{
++                    static_cast<uint32_t>(bitsPerSubQuantizer_),
++                    static_cast<uint32_t>(numSubQuantizers_),
++                    true};
+
+     for (size_t i = 0; i < nlist; ++i) {
+         size_t listSize = ivf->list_size(i);
+@@ -426,9 +490,9 @@ void CuvsIVFPQ::copyInvertedListsFrom(const InvertedLists* ivf) {
+         // This cuVS list must currently be empty
+         FAISS_ASSERT(getListLength(i) == 0);
+
+-        cuvs::neighbors::ivf::resize_list(
++        cuvs::neighbors::ivf_pq::helpers::resize_list(
+                 raft_handle,
+-                cuvs_index_lists[i],
++                cuvs_index_lists[i],
+                 ivf_list_spec,
+                 static_cast<uint32_t>(listSize),
+                 static_cast<uint32_t>(0));
+@@ -520,7 +587,7 @@ void CuvsIVFPQ::setPQCentroids_() {
      auto stream = resources_->getDefaultStreamCurrentDevice();
 
      raft::copy(
diff --git a/cpp/include/cuvs/neighbors/common.hpp b/cpp/include/cuvs/neighbors/common.hpp
index 4f697b3604..0d23989a33 100644
--- a/cpp/include/cuvs/neighbors/common.hpp
+++ b/cpp/include/cuvs/neighbors/common.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -683,11 +683,52 @@ template <typename IdxT>
 constexpr static IdxT kInvalidRecord =
   (std::is_signed_v<IdxT> ? IdxT{0} : std::numeric_limits<IdxT>::max()) - 1;
 
+/**
+ * Abstract base class for IVF list data.
+ * This allows polymorphic access to list data regardless of the underlying layout.
+ *
+ * @tparam ValueT The data element type (e.g., uint8_t for PQ codes, float for raw vectors)
+ * @tparam IdxT The index type for source indices
+ * @tparam SizeT The size type
+ *
+ * TODO: Make this struct internal (tracking issue: https://github.com/rapidsai/cuvs/issues/1726)
+ */
+template <typename ValueT, typename IdxT, typename SizeT = uint32_t>
+struct list_base {
+  using value_type = ValueT;
+  using index_type = IdxT;
+  using size_type  = SizeT;
+
+  virtual ~list_base() = default;
+
+  /** Get the raw data pointer. */
+  virtual value_type* data_ptr() noexcept             = 0;
+  virtual const value_type* data_ptr() const noexcept = 0;
+
+  /** Get the indices pointer. */
+  virtual index_type* indices_ptr() noexcept             = 0;
+  virtual const index_type* indices_ptr() const noexcept = 0;
+
+  /** Get the current size (number of records). */
+  virtual size_type get_size() const noexcept = 0;
+
+  /** Set the current size (number of records). */
+  virtual void set_size(size_type new_size) noexcept = 0;
+
+  /** Get the total size of the data array in bytes. */
+  virtual size_t data_byte_size() const noexcept = 0;
+
+  /** Get the capacity (number of indices that can be stored). */
+  virtual size_type indices_capacity() const noexcept = 0;
+};
+
 /** The data for a single IVF list. */
 template <template <typename, typename...> typename SpecT,
           typename SizeT,
           typename... SpecExtraArgs>
-struct list {
+struct list : public list_base<typename SpecT<SizeT, SpecExtraArgs...>::value_type,
+                               typename SpecT<SizeT, SpecExtraArgs...>::index_type,
+                               SizeT> {
   using size_type    = SizeT;
   using spec_type    = SpecT<size_type, SpecExtraArgs...>;
   using value_type   = typename spec_type::value_type;
@@ -703,6 +744,18 @@ struct list {
 
   /** Allocate a new list capable of holding at least `n_rows` data records and indices. */
   list(raft::resources const& res, const spec_type& spec, size_type n_rows);
+
+  value_type* data_ptr() noexcept override { return data.data_handle(); }
+  const value_type* data_ptr() const noexcept override { return data.data_handle(); }
+
+  index_type* indices_ptr() noexcept override { return indices.data_handle(); }
+  const index_type* indices_ptr() const noexcept override { return indices.data_handle(); }
+
+  size_type get_size() const noexcept override { return size.load(); }
+  void set_size(size_type new_size) noexcept override { size.store(new_size); }
+
+  size_t data_byte_size() const noexcept override { return data.size() * sizeof(value_type); }
+  size_type indices_capacity() const noexcept override { return indices.extent(0); }
 };
 
 template <typename ListT, class T = void>
@@ -727,6 +780,10 @@ using enable_if_valid_list_t = typename enable_if_valid_list<ListT, T>::type;
 /**
  * Resize a list by the given id, so that it can contain the given number of records;
  * copy the data if necessary.
+ *
+ * @note This is an internal function that requires the concrete list type.
+ *       For IVF-PQ indexes, prefer using the helper functions in
+ *       `cuvs::neighbors::ivf_pq::helpers::resize_list` which handle type casting internally.
  */
 template <typename ListT>
 void resize_list(raft::resources const& res,
@@ -735,6 +792,15 @@ void resize_list(raft::resources const& res,
                  typename ListT::size_type new_used_size,
                  typename ListT::size_type old_used_size);
 
+/**
+ * Serialize a list to an output stream.
+ *
+ * @note This function requires the concrete list type (not the base class) because:
+ *       1. It needs access to the spec_type to determine the data layout for serialization
+ *       2. The serialized format depends on the spec's make_list_extents() method
+ *       When calling from code that only has a base class pointer, use std::static_pointer_cast
+ *       to obtain the typed pointer first.
+ */
 template <typename ListT>
 enable_if_valid_list_t<ListT> serialize_list(
   const raft::resources& handle,
@@ -742,6 +808,7 @@ enable_if_valid_list_t<ListT> serialize_list(
   const ListT& ld,
   const typename ListT::spec_type& store_spec,
   std::optional<typename ListT::size_type> size_override = std::nullopt);
+
 template <typename ListT>
 enable_if_valid_list_t<ListT> serialize_list(
   const raft::resources& handle,
diff --git a/cpp/include/cuvs/neighbors/ivf_flat.hpp b/cpp/include/cuvs/neighbors/ivf_flat.hpp
index 9928a0e3ca..b28c01de04 100644
--- a/cpp/include/cuvs/neighbors/ivf_flat.hpp
+++ b/cpp/include/cuvs/neighbors/ivf_flat.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -2933,17 +2933,17 @@ void reset_index(const raft::resources& res, index<uint8_t, int64_t>* index);
  *   ivf_flat::index<uint8_t, int64_t> index(res, index_params, D);
  *   ivf_flat::helpers::reset_index(res, &index);
  *   // resize the first IVF list to hold 5 records
- *   auto spec = list_spec<uint32_t, uint8_t, int64_t>{
- *     index->dim(), index->conservative_memory_allocation()};
+ *   auto spec = list_spec<uint32_t, float, int64_t>{
+ *     index.dim(), index.conservative_memory_allocation()};
  *   uint32_t new_size = 5;
- *   ivf::resize_list(res, list, spec, new_size, 0);
+ *   ivf::resize_list(res, index.lists()[0], spec, new_size, 0);
  *   raft::update_device(index.list_sizes(), &new_size, 1, stream);
  *   // recompute the internal state of the index
  *   ivf_flat::helpers::recompute_internal_state(res, index);
  * @endcode
  *
  * @param[in] res raft resource
- * @param[inout] index pointer to IVF-PQ index
+ * @param[inout] index pointer to IVF-Flat index
  */
 void recompute_internal_state(const raft::resources& res, index<float, int64_t>* index);
 
@@ -2961,17 +2961,17 @@ void recompute_internal_state(const raft::resources& res, index<float, int64_t>*
  *   ivf_flat::index<uint8_t, int64_t> index(res, index_params, D);
  *   ivf_flat::helpers::reset_index(res, &index);
  *   // resize the first IVF list to hold 5 records
- *   auto spec = list_spec<uint32_t, uint8_t, int64_t>{
- *     index->dim(), index->conservative_memory_allocation()};
+ *   auto spec = list_spec<uint32_t, half, int64_t>{
+ *     index.dim(), index.conservative_memory_allocation()};
  *   uint32_t new_size = 5;
- *   ivf::resize_list(res, list, spec, new_size, 0);
+ *   ivf::resize_list(res, index.lists()[0], spec, new_size, 0);
  *   raft::update_device(index.list_sizes(), &new_size, 1, stream);
  *   // recompute the internal state of the index
  *   ivf_flat::helpers::recompute_internal_state(res, index);
  * @endcode
  *
  * @param[in] res raft resource
- * @param[inout] index pointer to IVF-PQ index
+ * @param[inout] index pointer to IVF-Flat index
  */
 void recompute_internal_state(const raft::resources& res, index<half, int64_t>* index);
 
@@ -2989,17 +2989,17 @@ void recompute_internal_state(const raft::resources& res, index<half, int64_t>*
  *   ivf_flat::index<uint8_t, int64_t> index(res, index_params, D);
  *   ivf_flat::helpers::reset_index(res, &index);
  *   // resize the first IVF list to hold 5 records
- *   auto spec = list_spec<uint32_t, uint8_t, int64_t>{
- *     index->dim(), index->conservative_memory_allocation()};
+ *   auto spec = list_spec<uint32_t, int8_t, int64_t>{
+ *     index.dim(), index.conservative_memory_allocation()};
  *   uint32_t new_size = 5;
- *   ivf::resize_list(res, list, spec, new_size, 0);
+ *   ivf::resize_list(res, index.lists()[0], spec, new_size, 0);
  *   raft::update_device(index.list_sizes(), &new_size, 1, stream);
  *   // recompute the internal state of the index
  *   ivf_flat::helpers::recompute_internal_state(res, index);
  * @endcode
  *
  * @param[in] res raft resource
- * @param[inout] index pointer to IVF-PQ index
+ * @param[inout] index pointer to IVF-Flat index
  */
 void recompute_internal_state(const raft::resources& res, index<int8_t, int64_t>* index);
 
@@ -3018,9 +3018,9 @@ void recompute_internal_state(const raft::resources& res, index<int8_t, int64_t>
  *   ivf_flat::helpers::reset_index(res, &index);
  *   // resize the first IVF list to hold 5 records
  *   auto spec = list_spec<uint32_t, uint8_t, int64_t>{
- *     index->dim(), index->conservative_memory_allocation()};
+ *     index.dim(), index.conservative_memory_allocation()};
  *   uint32_t new_size = 5;
- *   ivf::resize_list(res, list, spec, new_size, 0);
+ *   ivf::resize_list(res, index.lists()[0], spec, new_size, 0);
  *   raft::update_device(index.list_sizes(), &new_size, 1, stream);
  *   // recompute the internal state of the index
  *   ivf_flat::helpers::recompute_internal_state(res, index);
diff --git a/cpp/include/cuvs/neighbors/ivf_pq.hpp b/cpp/include/cuvs/neighbors/ivf_pq.hpp
index 53b84f24ed..4ae46fd1f5 100644
--- a/cpp/include/cuvs/neighbors/ivf_pq.hpp
+++ b/cpp/include/cuvs/neighbors/ivf_pq.hpp
@@ -33,6 +33,14 @@ enum class codebook_gen {  // NOLINT
   PER_CLUSTER  = 1,        // NOLINT
 };
 
+/** A type for specifying the memory layout of PQ codes in IVF lists. */
+enum class list_layout {  // NOLINT
+  /** Flat layout: each vector's PQ codes stored contiguously [n_rows, bytes_per_vector]. */
+  FLAT = 0,  // NOLINT
+  /** Interleaved layout: codes from multiple vectors interleaved for coalesced memory access. */
+  INTERLEAVED = 1,  // NOLINT
+};
+
 struct index_params : cuvs::neighbors::index_params {
   /**
    * The number of inverted lists (clusters)
@@ -69,6 +77,14 @@ struct index_params : cuvs::neighbors::index_params {
   uint32_t pq_dim = 0;
   /** How PQ codebooks are created. */
   codebook_gen codebook_kind = codebook_gen::PER_SUBSPACE;
+  /**
+   * Memory layout of PQ codes in IVF lists.
+   *
+   * - INTERLEAVED (default): Codes from multiple vectors are interleaved for coalesced GPU memory
+   *   access during search. This is optimized for search performance.
+   * - FLAT: Each vector's PQ codes are stored contiguously.
+   */
+  list_layout codes_layout = list_layout::INTERLEAVED;
   /**
    * Apply a random rotation matrix on the input data and queries even if `dim % pq_dim == 0`.
    *
@@ -213,7 +229,7 @@ template <typename IdxT>
 constexpr static IdxT kOutOfBoundsRecord = std::numeric_limits<IdxT>::max();
 
 template <typename SizeT, typename IdxT>
-struct list_spec {
+struct list_spec_interleaved {
   using value_type = uint8_t;
   using index_type = IdxT;
   /** PQ-encoded data stored in the interleaved format:
@@ -232,20 +248,22 @@ struct list_spec {
   uint32_t pq_bits;
   uint32_t pq_dim;
 
-  constexpr list_spec(uint32_t pq_bits, uint32_t pq_dim, bool conservative_memory_allocation);
+  constexpr list_spec_interleaved(uint32_t pq_bits,
+                                  uint32_t pq_dim,
+                                  bool conservative_memory_allocation);
 
   // Allow casting between different size-types (for safer size and offset calculations)
   template <typename OtherSizeT>
-  constexpr explicit list_spec(const list_spec<OtherSizeT, IdxT>& other_spec);
+  constexpr explicit list_spec_interleaved(
+    const list_spec_interleaved<OtherSizeT, IdxT>& other_spec);
 
   /** Determine the extents of an array enough to hold a given amount of data. */
   constexpr list_extents make_list_extents(SizeT n_rows) const;
 };
 
 template <typename SizeT, typename IdxT>
-constexpr list_spec<SizeT, IdxT>::list_spec(uint32_t pq_bits,
-                                            uint32_t pq_dim,
-                                            bool conservative_memory_allocation)
+constexpr list_spec_interleaved<SizeT, IdxT>::list_spec_interleaved(
+  uint32_t pq_bits, uint32_t pq_dim, bool conservative_memory_allocation)
   : pq_bits(pq_bits),
     pq_dim(pq_dim),
     align_min(kIndexGroupSize),
@@ -255,7 +273,8 @@ constexpr list_spec<SizeT, IdxT>::list_spec(uint32_t pq_bits,
 
 template <typename SizeT, typename IdxT>
 template <typename OtherSizeT>
-constexpr list_spec<SizeT, IdxT>::list_spec(const list_spec<OtherSizeT, IdxT>& other_spec)
+constexpr list_spec_interleaved<SizeT, IdxT>::list_spec_interleaved(
+  const list_spec_interleaved<OtherSizeT, IdxT>& other_spec)
   : pq_bits{other_spec.pq_bits},
     pq_dim{other_spec.pq_dim},
     align_min{other_spec.align_min},
@@ -264,8 +283,8 @@ constexpr list_spec<SizeT, IdxT>::list_spec(const list_spec<OtherSizeT, IdxT>& o
 }
 
 template <typename SizeT, typename IdxT>
-constexpr typename list_spec<SizeT, IdxT>::list_extents list_spec<SizeT, IdxT>::make_list_extents(
-  SizeT n_rows) const
+constexpr typename list_spec_interleaved<SizeT, IdxT>::list_extents
+list_spec_interleaved<SizeT, IdxT>::make_list_extents(SizeT n_rows) const
 {
   // how many elems of pq_dim fit into one kIndexGroupVecLen-byte chunk
   auto pq_chunk = (kIndexGroupVecLen * 8u) / pq_bits;
@@ -274,7 +293,68 @@ constexpr typename list_spec<SizeT, IdxT>::list_extents list_spec<SizeT, IdxT>::
 }
 
 template <typename IdxT, typename SizeT = uint32_t>
-using list_data = ivf::list<list_spec, SizeT, IdxT>;
+using list_data_interleaved = ivf::list<list_spec_interleaved, SizeT, IdxT>;
+
+/**
+ * Flat (non-interleaved) storage specification for PQ-encoded data.
+ *
+ * This stores each vector's PQ codes contiguously:
+ *   [n_rows, bytes_per_vector] where bytes_per_vector = ceildiv(pq_dim * pq_bits, 8)
+ */
+template <typename SizeT, typename IdxT>
+struct list_spec_flat {
+  using value_type   = uint8_t;
+  using index_type   = IdxT;
+  using list_extents = raft::matrix_extent<SizeT>;
+
+  SizeT align_max;
+  SizeT align_min;
+  uint32_t pq_bits;
+  uint32_t pq_dim;
+
+  constexpr list_spec_flat(uint32_t pq_bits, uint32_t pq_dim, bool conservative_memory_allocation)
+    : pq_bits(pq_bits), pq_dim(pq_dim), align_min(1), align_max(1)
+  {
+  }
+
+  // Allow casting between different size-types (for safer size and offset calculations)
+  template <typename OtherSizeT>
+  constexpr explicit list_spec_flat(const list_spec_flat<OtherSizeT, IdxT>& other_spec);
+
+  /** Number of bytes per encoded vector. */
+  constexpr SizeT bytes_per_vector() const
+  {
+    return raft::div_rounding_up_safe<SizeT>(pq_dim * pq_bits, 8u);
+  }
+
+  /** Determine the extents of an array enough to hold a given amount of data. */
+  constexpr list_extents make_list_extents(SizeT n_rows) const
+  {
+    return list_extents{n_rows, bytes_per_vector()};
+  }
+};
+
+template <typename SizeT, typename IdxT>
+template <typename OtherSizeT>
+constexpr list_spec_flat<SizeT, IdxT>::list_spec_flat(
+  const list_spec_flat<OtherSizeT, IdxT>& other_spec)
+  : pq_bits{other_spec.pq_bits},
+    pq_dim{other_spec.pq_dim},
+    align_min{other_spec.align_min},
+    align_max{other_spec.align_max}
+{
+}
+
+template <typename IdxT, typename SizeT = uint32_t>
+using list_data_flat = ivf::list<list_spec_flat, SizeT, IdxT>;
+
+/**
+ * Type alias for the polymorphic base class for IVF-PQ list data.
+ * IVF-PQ uses uint8_t for PQ-encoded data.
+ * Both list_data_interleaved and list_data_flat now inherit from this via ivf::list.
+ */
+template <typename IdxT, typename SizeT = uint32_t>
+using list_data_base = ivf::list_base<uint8_t, IdxT, SizeT>;
 
 using pq_centers_extents =
   raft::extents<uint32_t, raft::dynamic_extent, raft::dynamic_extent, raft::dynamic_extent>;
@@ -286,6 +366,7 @@ class index_iface {
 
   virtual cuvs::distance::DistanceType metric() const noexcept  = 0;
   virtual codebook_gen codebook_kind() const noexcept           = 0;
+  virtual list_layout codes_layout() const noexcept             = 0;
   virtual IdxT size() const noexcept                            = 0;
   virtual uint32_t dim() const noexcept                         = 0;
   virtual uint32_t dim_ext() const noexcept                     = 0;
@@ -298,8 +379,8 @@ class index_iface {
   virtual bool conservative_memory_allocation() const noexcept  = 0;
   virtual uint32_t get_list_size_in_bytes(uint32_t label) const = 0;
 
-  virtual std::vector<std::shared_ptr<list_data<IdxT>>>& lists() noexcept             = 0;
-  virtual const std::vector<std::shared_ptr<list_data<IdxT>>>& lists() const noexcept = 0;
+  virtual std::vector<std::shared_ptr<list_data_base<IdxT>>>& lists() noexcept             = 0;
+  virtual const std::vector<std::shared_ptr<list_data_base<IdxT>>>& lists() const noexcept = 0;
 
   virtual raft::device_vector_view<uint32_t, uint32_t, raft::row_major> list_sizes() noexcept = 0;
   virtual raft::device_vector_view<const uint32_t, uint32_t, raft::row_major> list_sizes()
@@ -479,6 +560,9 @@ class index : public index_iface<IdxT>, cuvs::neighbors::index {
   /** How PQ codebooks are created. */
   codebook_gen codebook_kind() const noexcept override;
 
+  /** Memory layout of PQ codes in IVF lists. */
+  list_layout codes_layout() const noexcept override;
+
   /** Number of clusters/inverted lists (first level quantization). */
   uint32_t n_lists() const noexcept;
 
@@ -497,9 +581,9 @@ class index : public index_iface<IdxT>, cuvs::neighbors::index {
   raft::device_mdspan<const float, pq_centers_extents, raft::row_major> pq_centers()
     const noexcept override;
 
-  /** Lists' data and indices. */
-  std::vector<std::shared_ptr<list_data<IdxT>>>& lists() noexcept override;
-  const std::vector<std::shared_ptr<list_data<IdxT>>>& lists() const noexcept override;
+  /** Lists' data and indices (polymorphic, works for both FLAT and INTERLEAVED layouts). */
+  std::vector<std::shared_ptr<list_data_base<IdxT>>>& lists() noexcept override;
+  const std::vector<std::shared_ptr<list_data_base<IdxT>>>& lists() const noexcept override;
 
   /** Pointers to the inverted lists (clusters) data  [n_lists]. */
   raft::device_vector_view<uint8_t*, uint32_t, raft::row_major> data_ptrs() noexcept override;
@@ -2465,13 +2549,13 @@ namespace codepacker {
  *   The length `n_take` defines how many records to unpack,
  *   it must be smaller than the list size.
  */
-void unpack(
-  raft::resources const& res,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  uint32_t pq_bits,
-  uint32_t offset,
-  raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes);
+void unpack(raft::resources const& res,
+            raft::device_mdspan<const uint8_t,
+                                list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                raft::row_major> list_data,
+            uint32_t pq_bits,
+            uint32_t offset,
+            raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes);
 
 /**
  * @brief Unpack `n_rows` consecutive records of a single list (cluster) in the compressed index
@@ -2505,15 +2589,15 @@ void unpack(
  *   The length `n_rows` defines how many records to unpack,
  *   it must be smaller than the list size.
  */
-void unpack_contiguous(
-  raft::resources const& res,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  uint32_t pq_bits,
-  uint32_t offset,
-  uint32_t n_rows,
-  uint32_t pq_dim,
-  uint8_t* codes);
+void unpack_contiguous(raft::resources const& res,
+                       raft::device_mdspan<const uint8_t,
+                                           list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                           raft::row_major> list_data,
+                       uint32_t pq_bits,
+                       uint32_t offset,
+                       uint32_t n_rows,
+                       uint32_t pq_dim,
+                       uint8_t* codes);
 
 /**
  * Write flat PQ codes into an existing list by the given offset.
@@ -2541,8 +2625,9 @@ void pack(raft::resources const& res,
           raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
           uint32_t pq_bits,
           uint32_t offset,
-          raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-            list_data);
+          raft::device_mdspan<uint8_t,
+                              list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                              raft::row_major> list_data);
 
 /**
  * Write flat PQ codes into an existing list by the given offset. The input codes of a single vector
@@ -2572,15 +2657,15 @@ void pack(raft::resources const& res,
  * @param[in] offset how many records to skip before writing the data into the list
  * @param[in] list_data block to write into
  */
-void pack_contiguous(
-  raft::resources const& res,
-  const uint8_t* codes,
-  uint32_t n_rows,
-  uint32_t pq_dim,
-  uint32_t pq_bits,
-  uint32_t offset,
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data);
+void pack_contiguous(raft::resources const& res,
+                     const uint8_t* codes,
+                     uint32_t n_rows,
+                     uint32_t pq_dim,
+                     uint32_t pq_bits,
+                     uint32_t offset,
+                     raft::device_mdspan<uint8_t,
+                                         list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                         raft::row_major> list_data);
 
 /**
  * Write flat PQ codes into an existing list by the given offset.
@@ -2909,6 +2994,45 @@ void extend_list_with_codes(
   raft::device_vector_view<const int64_t, uint32_t, raft::row_major> new_indices,
   uint32_t label);
 
+/**
+ * @brief Extend one list of the index in-place, by the list label, skipping the classification and
+ * encoding steps. Uses contiguous/packed codes format.
+ *
+ * This is similar to extend_list_with_codes but takes codes in contiguous packed format
+ * [n_rows, ceildiv(pq_dim * pq_bits, 8)] instead of unpacked format [n_rows, pq_dim].
+ * This works correctly with any pq_bits value.
+ *
+ * Usage example:
+ * @code{.cpp}
+ *   // We will extend the fourth cluster
+ *   uint32_t label = 3;
+ *   // We will fill 4 new vectors
+ *   uint32_t n_vec = 4;
+ *   // Indices of the new vectors
+ *   auto indices = raft::make_device_vector<int64_t>(res, n_vec);
+ *   ... fill the indices ...
+ *   // Allocate buffer for packed codes
+ *   uint32_t code_size = raft::ceildiv(index.pq_dim() * index.pq_bits(), 8u);
+ *   auto new_codes = raft::make_device_matrix<uint8_t, uint32_t, row_major>(res, n_vec, code_size);
+ *   ... fill codes ...
+ *   // extend list with new codes
+ *   ivf_pq::helpers::codepacker::extend_list_with_contiguous_codes(
+ *       res, &index, new_codes.view(), indices.view(), label);
+ * @endcode
+ *
+ * @param[in] res
+ * @param[inout] index
+ * @param[in] new_codes flat contiguous PQ codes [n_rows, ceildiv(pq_dim * pq_bits, 8)]
+ * @param[in] new_indices source indices [n_rows]
+ * @param[in] label the id of the target list (cluster).
+ */
+void extend_list_with_contiguous_codes(
+  raft::resources const& res,
+  index<int64_t>* index,
+  raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> new_codes,
+  raft::device_vector_view<const int64_t, uint32_t, raft::row_major> new_indices,
+  uint32_t label);
+
 /**
  * @brief Extend one list of the index in-place, by the list label, skipping the classification
  * step.
@@ -3081,13 +3205,13 @@ void extract_centers(raft::resources const& res,
  *   ivf_pq::index<int64_t> index(res, index_params, D);
  *   ivf_pq::helpers::reset_index(res, &index);
  *   // resize the first IVF list to hold 5 records
- *   auto spec = list_spec<uint32_t, int64_t>{
- *     index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
+ *   auto spec = ivf_pq::list_spec_interleaved<uint32_t, int64_t>{
+ *     index.pq_bits(), index.pq_dim(), index.conservative_memory_allocation()};
  *   uint32_t new_size = 5;
- *   ivf::resize_list(res, list, spec, new_size, 0);
- *   raft::update_device(index.list_sizes(), &new_size, 1, stream);
+ *   ivf_pq::helpers::resize_list(res, index.lists()[0], spec, new_size, 0);
+ *   raft::update_device(index.list_sizes().data_handle(), &new_size, 1, stream);
  *   // recompute the internal state of the index
- *   ivf_pq::helpers::recompute_internal_state(res, index);
+ *   ivf_pq::helpers::recompute_internal_state(res, &index);
  * @endcode
  *
  * @param[in] res raft resource
@@ -3124,6 +3248,68 @@ void make_rotation_matrix(
   raft::resources const& res,
   raft::device_matrix_view<float, uint32_t, raft::row_major> rotation_matrix,
   bool force_random_rotation);
+
+/**
+ * @brief Resize an IVF-PQ list with flat layout.
+ *
+ * This helper resizes an IVF list that uses the flat (non-interleaved) PQ code layout.
+ * If the new size exceeds the current capacity, a new list is allocated and existing
+ * data is copied. The function handles the type casting internally.
+ *
+ * Usage example:
+ * @code{.cpp}
+ *   using namespace cuvs::neighbors;
+ *   raft::resources res;
+ *   // Assuming index uses FLAT layout
+ *   auto spec = ivf_pq::list_spec_flat<uint32_t, int64_t>{
+ *     index.pq_bits(), index.pq_dim(), index.conservative_memory_allocation()};
+ *   uint32_t old_size = current_list_size;
+ *   uint32_t new_size = old_size + n_new_vectors;
+ *   ivf_pq::helpers::resize_list(res, index.lists()[label], spec, new_size, old_size);
+ * @endcode
+ *
+ * @param[in] res raft resource
+ * @param[inout] orig_list the list to resize (may be replaced with a new allocation)
+ * @param[in] spec the list specification containing pq_bits, pq_dim, and allocation settings
+ * @param[in] new_used_size the new size of the list (number of vectors)
+ * @param[in] old_used_size the current size of the list (data up to this size is preserved)
+ */
+void resize_list(raft::resources const& res,
+                 std::shared_ptr<list_data_base<int64_t, uint32_t>>& orig_list,
+                 const list_spec_flat<uint32_t, int64_t>& spec,
+                 uint32_t new_used_size,
+                 uint32_t old_used_size);
+
+/**
+ * @brief Resize an IVF-PQ list with interleaved layout.
+ *
+ * This helper resizes an IVF list that uses the interleaved PQ code layout (default).
+ * If the new size exceeds the current capacity, a new list is allocated and existing
+ * data is copied. The function handles the type casting internally.
+ *
+ * Usage example:
+ * @code{.cpp}
+ *   using namespace cuvs::neighbors;
+ *   raft::resources res;
+ *   // Assuming index uses INTERLEAVED layout (default)
+ *   auto spec = ivf_pq::list_spec_interleaved<uint32_t, int64_t>{
+ *     index.pq_bits(), index.pq_dim(), index.conservative_memory_allocation()};
+ *   uint32_t old_size = current_list_size;
+ *   uint32_t new_size = old_size + n_new_vectors;
+ *   ivf_pq::helpers::resize_list(res, index.lists()[label], spec, new_size, old_size);
+ * @endcode
+ *
+ * @param[in] res raft resource
+ * @param[inout] orig_list the list to resize (may be replaced with a new allocation)
+ * @param[in] spec the list specification containing pq_bits, pq_dim, and allocation settings
+ * @param[in] new_used_size the new size of the list (number of vectors)
+ * @param[in] old_used_size the current size of the list (data up to this size is preserved)
+ */
+void resize_list(raft::resources const& res,
+                 std::shared_ptr<list_data_base<int64_t, uint32_t>>& orig_list,
+                 const list_spec_interleaved<uint32_t, int64_t>& spec,
+                 uint32_t new_used_size,
+                 uint32_t old_used_size);
 /**
  * @}
  */
diff --git a/cpp/src/neighbors/ivf_common.cuh b/cpp/src/neighbors/ivf_common.cuh
index 9b10488c3f..d7a5da6e4e 100644
--- a/cpp/src/neighbors/ivf_common.cuh
+++ b/cpp/src/neighbors/ivf_common.cuh
@@ -240,8 +240,8 @@ void recompute_internal_state(const raft::resources& res, Index& index)
   auto inds_ptrs = index.inds_ptrs();
   for (uint32_t label = 0; label < index.n_lists(); label++) {
     auto& list          = index.lists()[label];
-    const auto data_ptr = list ? list->data.data_handle() : nullptr;
-    const auto inds_ptr = list ? list->indices.data_handle() : nullptr;
+    const auto data_ptr = list ? list->data_ptr() : nullptr;
+    const auto inds_ptr = list ? list->indices_ptr() : nullptr;
     raft::copy(&data_ptrs(label), &data_ptr, 1, stream);
     raft::copy(&inds_ptrs(label), &inds_ptr, 1, stream);
   }
diff --git a/cpp/src/neighbors/ivf_pq/detail/ivf_pq_contiguous_list_data.cu b/cpp/src/neighbors/ivf_pq/detail/ivf_pq_contiguous_list_data.cu
index 7ca05e7262..5fac491532 100644
--- a/cpp/src/neighbors/ivf_pq/detail/ivf_pq_contiguous_list_data.cu
+++ b/cpp/src/neighbors/ivf_pq/detail/ivf_pq_contiguous_list_data.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -9,8 +9,9 @@
 namespace cuvs::neighbors::ivf_pq::detail {
 void unpack_contiguous_list_data(
   uint8_t* codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   uint32_t n_rows,
   uint32_t pq_dim,
   std::variant<uint32_t, const uint32_t*> offset_or_indices,
@@ -22,8 +23,9 @@ void unpack_contiguous_list_data(
 };
 
 void pack_contiguous_list_data(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   const uint8_t* codes,
   uint32_t n_rows,
   uint32_t pq_dim,
diff --git a/cpp/src/neighbors/ivf_pq/detail/ivf_pq_list_data.cu b/cpp/src/neighbors/ivf_pq/detail/ivf_pq_list_data.cu
index acef4c516e..34a7244a67 100644
--- a/cpp/src/neighbors/ivf_pq/detail/ivf_pq_list_data.cu
+++ b/cpp/src/neighbors/ivf_pq/detail/ivf_pq_list_data.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -7,24 +7,24 @@
 #include <cuvs/neighbors/ivf_pq.hpp>
 
 namespace cuvs::neighbors::ivf_pq::detail {
-void unpack_list_data(
-  raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  std::variant<uint32_t, const uint32_t*> offset_or_indices,
-  uint32_t pq_bits,
-  rmm::cuda_stream_view stream)
+void unpack_list_data(raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes,
+                      raft::device_mdspan<const uint8_t,
+                                          list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                          raft::row_major> list_data,
+                      std::variant<uint32_t, const uint32_t*> offset_or_indices,
+                      uint32_t pq_bits,
+                      rmm::cuda_stream_view stream)
 {
   unpack_list_data_impl(codes, list_data, offset_or_indices, pq_bits, stream);
 };
 
-void pack_list_data(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
-  std::variant<uint32_t, const uint32_t*> offset_or_indices,
-  uint32_t pq_bits,
-  rmm::cuda_stream_view stream)
+void pack_list_data(raft::device_mdspan<uint8_t,
+                                        list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                        raft::row_major> list_data,
+                    raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
+                    std::variant<uint32_t, const uint32_t*> offset_or_indices,
+                    uint32_t pq_bits,
+                    rmm::cuda_stream_view stream)
 {
   pack_list_data_impl(list_data, codes, offset_or_indices, pq_bits, stream);
 };
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_build.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_build.cuh
index 8b63efd53f..a5caa5b965 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_build.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_build.cuh
@@ -552,8 +552,9 @@ struct reconstruct_vectors {
 template <uint32_t BlockSize, uint32_t PqBits>
 __launch_bounds__(BlockSize) static __global__ void reconstruct_list_data_kernel(
   raft::device_matrix_view<float, uint32_t, raft::row_major> out_vectors,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    in_list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> in_list_data,
   raft::device_mdspan<const float, raft::extent_3d<uint32_t>, raft::row_major> pq_centers,
   raft::device_matrix_view<const float, uint32_t, raft::row_major> centers_rot,
   codebook_gen codebook_kind,
@@ -577,12 +578,16 @@ void reconstruct_list_data(raft::resources const& res,
 {
   auto n_rows = out_vectors.extent(0);
   if (n_rows == 0) { return; }
-  auto& list = index.lists()[label];
+  // Currently only supports interleaved layout
+  RAFT_EXPECTS(index.codes_layout() == list_layout::INTERLEAVED,
+               "reconstruct_list_data currently only supports INTERLEAVED layout");
+  auto& list_data_base_ptr = index.lists()[label];
+  auto typed_list = std::static_pointer_cast<const list_data_interleaved<IdxT>>(list_data_base_ptr);
   if (std::holds_alternative<uint32_t>(offset_or_indices)) {
     auto n_skip = std::get<uint32_t>(offset_or_indices);
     // sic! I'm using the upper bound `list.size` instead of exact `list_sizes(label)`
     // to avoid an extra device-host data raft::copy and the stream sync.
-    RAFT_EXPECTS(n_skip + n_rows <= list->size.load(),
+    RAFT_EXPECTS(n_skip + n_rows <= typed_list->size.load(),
                  "offset + output size must be not bigger than the cluster size.");
   }
 
@@ -605,7 +610,7 @@ void reconstruct_list_data(raft::resources const& res,
     }
   }(index.pq_bits());
   kernel<<<blocks, threads, 0, raft::resource::get_cuda_stream(res)>>>(tmp.view(),
-                                                                       list->data.view(),
+                                                                       typed_list->data.view(),
                                                                        index.pq_centers(),
                                                                        index.centers_rot(),
                                                                        index.codebook_kind(),
@@ -652,9 +657,10 @@ void reconstruct_list_data(raft::resources const& res,
 }
 
 template <uint32_t BlockSize, uint32_t PqBits>
-__launch_bounds__(BlockSize) static __global__ void encode_list_data_kernel(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+__launch_bounds__(BlockSize) static __global__ void encode_list_data_interleaved_kernel(
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   raft::device_matrix_view<const float, uint32_t, raft::row_major> new_vectors,
   raft::device_mdspan<const float, raft::extent_3d<uint32_t>, raft::row_major> pq_centers,
   codebook_gen codebook_kind,
@@ -669,6 +675,24 @@ __launch_bounds__(BlockSize) static __global__ void encode_list_data_kernel(
     list_data, offset_or_indices, new_vectors.extent(0), pq_dim, encode_action);
 }
 
+template <uint32_t BlockSize, uint32_t PqBits>
+__launch_bounds__(BlockSize) static __global__ void encode_list_data_flat_kernel(
+  uint8_t* list_data,
+  raft::device_matrix_view<const float, uint32_t, raft::row_major> new_vectors,
+  raft::device_mdspan<const float, raft::extent_3d<uint32_t>, raft::row_major> pq_centers,
+  codebook_gen codebook_kind,
+  uint32_t cluster_ix,
+  std::variant<uint32_t, const uint32_t*> offset_or_indices,
+  uint32_t bytes_per_vector)
+{
+  constexpr uint32_t kSubWarpSize = std::min<uint32_t>(raft::WarpSize, 1u << PqBits);
+  const uint32_t pq_dim           = new_vectors.extent(1) / pq_centers.extent(1);
+  auto encode_action =
+    encode_vectors<kSubWarpSize, uint32_t>{pq_centers, new_vectors, codebook_kind, cluster_ix};
+  write_list_flat<PqBits, kSubWarpSize>(
+    list_data, offset_or_indices, new_vectors.extent(0), pq_dim, bytes_per_vector, encode_action);
+}
+
 template <typename T, typename IdxT>
 void encode_list_data(raft::resources const& res,
                       index<IdxT>* index,
@@ -698,23 +722,49 @@ void encode_list_data(raft::resources const& res,
   const uint32_t threads_per_vec = std::min<uint32_t>(raft::WarpSize, index->pq_book_size());
   dim3 blocks(raft::div_rounding_up_safe<uint32_t>(n_rows, kBlockSize / threads_per_vec), 1, 1);
   dim3 threads(kBlockSize, 1, 1);
-  auto kernel = [](uint32_t pq_bits) {
-    switch (pq_bits) {
-      case 4: return encode_list_data_kernel<kBlockSize, 4>;
-      case 5: return encode_list_data_kernel<kBlockSize, 5>;
-      case 6: return encode_list_data_kernel<kBlockSize, 6>;
-      case 7: return encode_list_data_kernel<kBlockSize, 7>;
-      case 8: return encode_list_data_kernel<kBlockSize, 8>;
-      default: RAFT_FAIL("Invalid pq_bits (%u), the value must be within [4, 8]", pq_bits);
-    }
-  }(index->pq_bits());
-  kernel<<<blocks, threads, 0, raft::resource::get_cuda_stream(res)>>>(
-    index->lists()[label]->data.view(),
-    new_vectors_residual.view(),
-    index->pq_centers(),
-    index->codebook_kind(),
-    label,
-    offset_or_indices);
+
+  if (index->codes_layout() == list_layout::FLAT) {
+    const uint32_t bytes_per_vector =
+      raft::div_rounding_up_safe(index->pq_dim() * index->pq_bits(), 8u);
+    auto kernel = [](uint32_t pq_bits) {
+      switch (pq_bits) {
+        case 4: return encode_list_data_flat_kernel<kBlockSize, 4>;
+        case 5: return encode_list_data_flat_kernel<kBlockSize, 5>;
+        case 6: return encode_list_data_flat_kernel<kBlockSize, 6>;
+        case 7: return encode_list_data_flat_kernel<kBlockSize, 7>;
+        case 8: return encode_list_data_flat_kernel<kBlockSize, 8>;
+        default: RAFT_FAIL("Invalid pq_bits (%u), the value must be within [4, 8]", pq_bits);
+      }
+    }(index->pq_bits());
+    kernel<<<blocks, threads, 0, raft::resource::get_cuda_stream(res)>>>(
+      index->lists()[label]->data_ptr(),
+      new_vectors_residual.view(),
+      index->pq_centers(),
+      index->codebook_kind(),
+      label,
+      offset_or_indices,
+      bytes_per_vector);
+  } else {
+    auto kernel = [](uint32_t pq_bits) {
+      switch (pq_bits) {
+        case 4: return encode_list_data_interleaved_kernel<kBlockSize, 4>;
+        case 5: return encode_list_data_interleaved_kernel<kBlockSize, 5>;
+        case 6: return encode_list_data_interleaved_kernel<kBlockSize, 6>;
+        case 7: return encode_list_data_interleaved_kernel<kBlockSize, 7>;
+        case 8: return encode_list_data_interleaved_kernel<kBlockSize, 8>;
+        default: RAFT_FAIL("Invalid pq_bits (%u), the value must be within [4, 8]", pq_bits);
+      }
+    }(index->pq_bits());
+    auto typed_list = std::static_pointer_cast<list_data_interleaved<IdxT>>(index->lists()[label]);
+    kernel<<<blocks, threads, 0, raft::resource::get_cuda_stream(res)>>>(
+      typed_list->data.view(),
+      new_vectors_residual.view(),
+      index->pq_centers(),
+      index->codebook_kind(),
+      label,
+      offset_or_indices);
+  }
+
   RAFT_CUDA_TRY(cudaPeekAtLastError());
 }
 
@@ -793,11 +843,17 @@ auto extend_list_prepare(
   uint32_t new_size = offset + n_rows;
   raft::copy(
     index->list_sizes().data_handle() + label, &new_size, 1, raft::resource::get_cuda_stream(res));
-  auto spec = list_spec<uint32_t, IdxT>{
-    index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
-  auto& list = index->lists()[label];
-  ivf::resize_list(res, list, spec, new_size, offset);
-  raft::copy(list->indices.data_handle() + offset,
+  auto& list_data_base_ptr = index->lists()[label];
+  if (index->codes_layout() == list_layout::FLAT) {
+    auto spec = list_spec_flat<uint32_t, IdxT>{
+      index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
+    cuvs::neighbors::ivf_pq::helpers::resize_list(res, list_data_base_ptr, spec, new_size, offset);
+  } else {
+    auto spec = list_spec_interleaved<uint32_t, IdxT>{
+      index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
+    cuvs::neighbors::ivf_pq::helpers::resize_list(res, list_data_base_ptr, spec, new_size, offset);
+  }
+  raft::copy(list_data_base_ptr->indices_ptr() + offset,
              new_indices.data_handle(),
              n_rows,
              raft::resource::get_cuda_stream(res));
@@ -825,6 +881,25 @@ void extend_list_with_codes(
   ivf::detail::recompute_internal_state(res, *index);
 }
 
+/**
+ * Extend one list of the index in-place, by the list label, skipping the classification and
+ * encoding steps. Uses contiguous/packed codes format.
+ * See the public interface for the api and usage.
+ */
+template <typename IdxT>
+void extend_list_with_contiguous_codes(
+  raft::resources const& res,
+  index<IdxT>* index,
+  raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> new_codes,
+  raft::device_vector_view<const IdxT, uint32_t, raft::row_major> new_indices,
+  uint32_t label)
+{
+  uint32_t n_rows = new_indices.extent(0);
+  auto offset     = extend_list_prepare(res, index, new_indices, label);
+  pack_contiguous_list_data(res, index, new_codes.data_handle(), n_rows, label, offset);
+  ivf::detail::recompute_internal_state(res, *index);
+}
+
 /**
  * Extend one list of the index in-place, by the list label, skipping the classification step.
  * See the public interface for the api and usage.
@@ -872,7 +947,8 @@ auto clone(const raft::resources& res, const index<IdxT>& source) -> index<IdxT>
                                                   source.dim(),
                                                   source.pq_bits(),
                                                   source.pq_dim(),
-                                                  source.conservative_memory_allocation());
+                                                  source.conservative_memory_allocation(),
+                                                  source.codes_layout());
 
   // raft::copy the independent parts using mutable accessors
   raft::copy(impl->list_sizes().data_handle(),
@@ -932,17 +1008,24 @@ void extend(raft::resources const& handle,
   rmm::device_async_resource_ref large_memory =
     raft::resource::get_large_workspace_resource(handle);
 
-  // The spec defines how the clusters look like
-  auto spec = list_spec<uint32_t, IdxT>{
-    index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
   // Try to allocate an index with the same parameters and the projected new size
-  // (which can be slightly larger than index->size() + n_rows, due to padding).
+  // (which can be slightly larger than index->size() + n_rows, due to padding for interleaved).
   // If this fails, the index would be too big to fit in the device anyway.
-  std::optional<list_data<IdxT, size_t>> placeholder_list(
-    std::in_place_t{},
-    handle,
-    list_spec<size_t, IdxT>{spec},
-    n_rows + (kIndexGroupSize - 1) * std::min<IdxT>(n_clusters, n_rows));
+  std::optional<list_data_interleaved<IdxT, size_t>> placeholder_list_interleaved;
+  std::optional<raft::device_matrix<uint8_t, IdxT, raft::row_major>> placeholder_list_flat;
+  if (index->codes_layout() == list_layout::FLAT) {
+    auto spec = list_spec_flat<uint32_t, IdxT>{
+      index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
+    placeholder_list_flat.emplace(raft::make_device_matrix<uint8_t, IdxT>(
+      handle, n_rows, static_cast<IdxT>(spec.bytes_per_vector())));
+  } else {
+    auto spec = list_spec_interleaved<uint32_t, IdxT>{
+      index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
+    placeholder_list_interleaved.emplace(
+      handle,
+      list_spec_interleaved<size_t, IdxT>{spec},
+      n_rows + (kIndexGroupSize - 1) * std::min<IdxT>(n_clusters, n_rows));
+  }
 
   // Available device memory
   size_t free_mem = raft::resource::get_workspace_free_bytes(handle);
@@ -1015,7 +1098,8 @@ void extend(raft::resources const& handle,
   // Release the placeholder memory, because we don't intend to allocate any more long-living
   // temporary buffers before we allocate the index data.
   // This memory could potentially speed up UVM accesses, if any.
-  placeholder_list.reset();
+  placeholder_list_interleaved.reset();
+  placeholder_list_flat.reset();
   {
     // The cluster centers in the index are stored padded, which is not acceptable by
     // the kmeans_balanced::predict. Thus, we need the restructuring raft::copy.
@@ -1070,9 +1154,20 @@ void extend(raft::resources const& handle,
     raft::copy(new_cluster_sizes.data(), list_sizes, n_clusters, stream);
     raft::copy(old_cluster_sizes.data(), orig_list_sizes.data(), n_clusters, stream);
     raft::resource::sync_stream(handle);
-    for (uint32_t label = 0; label < n_clusters; label++) {
-      ivf::resize_list(
-        handle, index->lists()[label], spec, new_cluster_sizes[label], old_cluster_sizes[label]);
+    if (index->codes_layout() == list_layout::FLAT) {
+      auto spec = list_spec_flat<uint32_t, IdxT>{
+        index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
+      for (uint32_t label = 0; label < n_clusters; label++) {
+        cuvs::neighbors::ivf_pq::helpers::resize_list(
+          handle, index->lists()[label], spec, new_cluster_sizes[label], old_cluster_sizes[label]);
+      }
+    } else {
+      auto spec = list_spec_interleaved<uint32_t, IdxT>{
+        index->pq_bits(), index->pq_dim(), index->conservative_memory_allocation()};
+      for (uint32_t label = 0; label < n_clusters; label++) {
+        cuvs::neighbors::ivf_pq::helpers::resize_list(
+          handle, index->lists()[label], spec, new_cluster_sizes[label], old_cluster_sizes[label]);
+      }
     }
   }
 
@@ -1147,7 +1242,8 @@ auto build(raft::resources const& handle,
     dim,
     params.pq_bits,
     params.pq_dim == 0 ? index<IdxT>::calculate_pq_dim(dim) : params.pq_dim,
-    params.conservative_memory_allocation);
+    params.conservative_memory_allocation,
+    params.codes_layout);
 
   auto stream = raft::resource::get_cuda_stream(handle);
   utils::memzero(
@@ -1352,7 +1448,8 @@ auto build(raft::resources const& handle,
                                                 pq_centers,
                                                 centers,
                                                 centers_rot,
-                                                rotation_matrix);
+                                                rotation_matrix,
+                                                index_params.codes_layout);
 
   index<IdxT> view_index(std::move(impl));
 
@@ -1402,7 +1499,7 @@ auto extend(
                 n_rows);
 }
 
-// In-place extend for base class pointer (clones, extends, moves back)
+// True in-place extend (does not clone the index)
 template <typename T, typename IdxT, typename accessor, typename accessor2>
 void extend(
   raft::resources const& handle,
@@ -1420,11 +1517,12 @@ void extend(
            "new_vectors and new_indices have different number of rows");
   }
 
-  *index = extend(handle,
-                  *index,
-                  new_vectors.data_handle(),
-                  new_indices.has_value() ? new_indices.value().data_handle() : nullptr,
-                  n_rows);
+  // Call the true in-place extend (line 972) instead of clone-and-replace
+  extend(handle,
+         index,
+         new_vectors.data_handle(),
+         new_indices.has_value() ? new_indices.value().data_handle() : nullptr,
+         n_rows);
 }
 
 template <typename IdxT>
@@ -1451,7 +1549,8 @@ auto build(
                                                   dim,
                                                   index_params.pq_bits,
                                                   pq_dim,
-                                                  index_params.conservative_memory_allocation);
+                                                  index_params.conservative_memory_allocation,
+                                                  index_params.codes_layout);
 
   utils::memzero(
     impl->accum_sorted_sizes().data_handle(), impl->accum_sorted_sizes().size(), stream);
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu b/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu
index b2177caee3..54cc5fee87 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_build_common.cu
@@ -15,26 +15,26 @@ namespace helpers {
 
 namespace codepacker {
 
-void unpack(
-  raft::resources const& res,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  uint32_t pq_bits,
-  uint32_t offset,
-  raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes)
+void unpack(raft::resources const& res,
+            raft::device_mdspan<const uint8_t,
+                                list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                raft::row_major> list_data,
+            uint32_t pq_bits,
+            uint32_t offset,
+            raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes)
 {
   detail::unpack_list_data(codes, list_data, offset, pq_bits, raft::resource::get_cuda_stream(res));
 }
 
-void unpack_contiguous(
-  raft::resources const& res,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  uint32_t pq_bits,
-  uint32_t offset,
-  uint32_t n_rows,
-  uint32_t pq_dim,
-  uint8_t* codes)
+void unpack_contiguous(raft::resources const& res,
+                       raft::device_mdspan<const uint8_t,
+                                           list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                           raft::row_major> list_data,
+                       uint32_t pq_bits,
+                       uint32_t offset,
+                       uint32_t n_rows,
+                       uint32_t pq_dim,
+                       uint8_t* codes)
 {
   detail::unpack_contiguous_list_data(
     codes, list_data, n_rows, pq_dim, offset, pq_bits, raft::resource::get_cuda_stream(res));
@@ -43,21 +43,22 @@ void pack(raft::resources const& res,
           raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
           uint32_t pq_bits,
           uint32_t offset,
-          raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-            list_data)
+          raft::device_mdspan<uint8_t,
+                              list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                              raft::row_major> list_data)
 {
   detail::pack_list_data(list_data, codes, offset, pq_bits, raft::resource::get_cuda_stream(res));
 }
 
-void pack_contiguous(
-  raft::resources const& res,
-  const uint8_t* codes,
-  uint32_t n_rows,
-  uint32_t pq_dim,
-  uint32_t pq_bits,
-  uint32_t offset,
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data)
+void pack_contiguous(raft::resources const& res,
+                     const uint8_t* codes,
+                     uint32_t n_rows,
+                     uint32_t pq_dim,
+                     uint32_t pq_bits,
+                     uint32_t offset,
+                     raft::device_mdspan<uint8_t,
+                                         list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                         raft::row_major> list_data)
 {
   detail::pack_contiguous_list_data(
     list_data, codes, n_rows, pq_dim, offset, pq_bits, raft::resource::get_cuda_stream(res));
@@ -174,6 +175,16 @@ void extend_list_with_codes(
   detail::extend_list_with_codes<int64_t>(res, index, new_codes, new_indices, label);
 }
 
+void extend_list_with_contiguous_codes(
+  raft::resources const& res,
+  index<int64_t>* index,
+  raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> new_codes,
+  raft::device_vector_view<const int64_t, uint32_t, raft::row_major> new_indices,
+  uint32_t label)
+{
+  detail::extend_list_with_contiguous_codes<int64_t>(res, index, new_codes, new_indices, label);
+}
+
 void extend_list(raft::resources const& res,
                  index<int64_t>* index,
                  raft::device_matrix_view<const float, uint32_t, raft::row_major> new_vectors,
@@ -375,6 +386,28 @@ void rotate_padded_centers(
                      stream);
 }
 
+void resize_list(raft::resources const& res,
+                 std::shared_ptr<list_data_base<int64_t, uint32_t>>& orig_list,
+                 const list_spec_flat<uint32_t, int64_t>& spec,
+                 uint32_t new_used_size,
+                 uint32_t old_used_size)
+{
+  auto typed_list = std::static_pointer_cast<list_data_flat<int64_t>>(orig_list);
+  cuvs::neighbors::ivf::resize_list(res, typed_list, spec, new_used_size, old_used_size);
+  orig_list = typed_list;
+}
+
+void resize_list(raft::resources const& res,
+                 std::shared_ptr<list_data_base<int64_t, uint32_t>>& orig_list,
+                 const list_spec_interleaved<uint32_t, int64_t>& spec,
+                 uint32_t new_used_size,
+                 uint32_t old_used_size)
+{
+  auto typed_list = std::static_pointer_cast<list_data_interleaved<int64_t>>(orig_list);
+  cuvs::neighbors::ivf::resize_list(res, typed_list, spec, new_used_size, old_used_size);
+  orig_list = typed_list;
+}
+
 }  // namespace helpers
 
 }  // namespace cuvs::neighbors::ivf_pq
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_codepacking.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_codepacking.cuh
index 669907577a..f917631db1 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_codepacking.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_codepacking.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -89,8 +89,9 @@ struct bitfield_view_t {
  */
 template <uint32_t PqBits, typename Action>
 __device__ void run_on_vector(
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    in_list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> in_list_data,
   uint32_t in_ix,
   uint32_t out_ix,
   uint32_t pq_dim,
@@ -131,9 +132,10 @@ __device__ void run_on_vector(
  *    type: (uint32_t in_ix, uint32_t j) -> uint8_t, where j = [0..pq_dim).
  */
 template <uint32_t PqBits, uint32_t SubWarpSize, typename IdxT, typename Action>
-__device__ void write_vector(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    out_list_data,
+__device__ void write_vector_interleaved(
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> out_list_data,
   uint32_t out_ix,
   IdxT in_ix,
   uint32_t pq_dim,
@@ -165,11 +167,62 @@ __device__ void write_vector(
   }
 }
 
+/**
+ * Process a single vector in a list.
+ *
+ * @tparam PqBits
+ * @tparam SubWarpSize how many threads work on the same ix (only the first thread writes data).
+ * @tparam IdxT type of the index passed to the action
+ * @tparam Action tells how to process a single vector (e.g. encode or just pack)
+ *
+ * @param[in] out_list_data the encoded cluster data.
+ * @param[in] out_ix in-cluster index of the vector to be processed (one-per-SubWarpSize threads).
+ * @param[in] in_ix the input index passed to the action (one-per-SubWarpSize threads).
+ * @param[in] pq_dim
+ * @param action a callable action to be invoked on each PQ code (component of the encoding)
+ *    type: (uint32_t in_ix, uint32_t j) -> uint8_t, where j = [0..pq_dim).
+ */
+/**
+ * Process a single vector for flat (non-interleaved) storage.
+ *
+ * @tparam PqBits
+ * @tparam SubWarpSize how many threads cooperatively compute each code
+ * @tparam IdxT type of the index passed to the action
+ * @tparam Action tells how to process a single vector (e.g. encode or just pack)
+ *
+ * @param[in] out_flat_codes flat PQ codes buffer [n_vectors, bytes_per_vector]
+ * @param[in] out_ix output vector index
+ * @param[in] in_ix the input index passed to the action
+ * @param[in] pq_dim
+ * @param[in] bytes_per_vector number of bytes per encoded vector = ceildiv(pq_dim * PqBits, 8)
+ * @param action a callable action to be invoked on each PQ code (component of the encoding)
+ *    type: (IdxT in_ix, uint32_t j) -> uint8_t, where j = [0..pq_dim).
+ */
+template <uint32_t PqBits, uint32_t SubWarpSize, typename IdxT, typename Action>
+__device__ void write_vector_flat(uint8_t* out_flat_codes,
+                                  uint32_t out_ix,
+                                  IdxT in_ix,
+                                  uint32_t pq_dim,
+                                  uint32_t bytes_per_vector,
+                                  Action action)
+{
+  const uint32_t lane_id = raft::Pow2<SubWarpSize>::mod(threadIdx.x);
+
+  uint8_t* out_vec_ptr = out_flat_codes + out_ix * bytes_per_vector;
+
+  bitfield_view_t<PqBits> code_view{out_vec_ptr};
+  for (uint32_t j = 0; j < pq_dim; j++) {
+    uint8_t code = action(in_ix, j);
+    if (lane_id == 0) { code_view[j] = code; }
+  }
+}
+
 /** Process the given indices or a block of a single list (cluster). */
 template <uint32_t PqBits, typename Action>
 __device__ void run_on_list(
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    in_list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> in_list_data,
   std::variant<uint32_t, const uint32_t*> offset_or_indices,
   uint32_t len,
   uint32_t pq_dim,
@@ -183,11 +236,12 @@ __device__ void run_on_list(
   }
 }
 
-/** Process the given indices or a block of a single list (cluster). */
+/** Process the given indices or a block of a single list (cluster) - interleaved layout. */
 template <uint32_t PqBits, uint32_t SubWarpSize, typename Action>
 __device__ void write_list(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    out_list_data,
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> out_list_data,
   std::variant<uint32_t, const uint32_t*> offset_or_indices,
   uint32_t len,
   uint32_t pq_dim,
@@ -200,7 +254,28 @@ __device__ void write_list(
     const uint32_t dst_ix = std::holds_alternative<uint32_t>(offset_or_indices)
                               ? std::get<uint32_t>(offset_or_indices) + ix
                               : std::get<const uint32_t*>(offset_or_indices)[ix];
-    write_vector<PqBits, SubWarpSize>(out_list_data, dst_ix, ix, pq_dim, action);
+    write_vector_interleaved<PqBits, SubWarpSize>(out_list_data, dst_ix, ix, pq_dim, action);
+  }
+}
+
+/** Process the given indices or a block of a single list (cluster) - flat layout. */
+template <uint32_t PqBits, uint32_t SubWarpSize, typename Action>
+__device__ void write_list_flat(uint8_t* out_flat_codes,
+                                std::variant<uint32_t, const uint32_t*> offset_or_indices,
+                                uint32_t len,
+                                uint32_t pq_dim,
+                                uint32_t bytes_per_vector,
+                                Action action)
+{
+  using subwarp_align = raft::Pow2<SubWarpSize>;
+  uint32_t stride     = subwarp_align::div(blockDim.x);
+  uint32_t ix         = subwarp_align::div(threadIdx.x + blockDim.x * blockIdx.x);
+  for (; ix < len; ix += stride) {
+    const uint32_t dst_ix = std::holds_alternative<uint32_t>(offset_or_indices)
+                              ? std::get<uint32_t>(offset_or_indices) + ix
+                              : std::get<const uint32_t*>(offset_or_indices)[ix];
+    write_vector_flat<PqBits, SubWarpSize>(
+      out_flat_codes, dst_ix, ix, pq_dim, bytes_per_vector, action);
   }
 }
 
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data.cuh
index c1b3ab083d..87c4c4ea79 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -13,8 +13,9 @@ namespace cuvs::neighbors::ivf_pq::detail {
 
 void unpack_contiguous_list_data(
   uint8_t* codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   uint32_t n_rows,
   uint32_t pq_dim,
   std::variant<uint32_t, const uint32_t*> offset_or_indices,
@@ -29,8 +30,13 @@ void unpack_contiguous_list_data(raft::resources const& res,
                                  uint32_t label,
                                  std::variant<uint32_t, const uint32_t*> offset_or_indices)
 {
+  // Currently only supports interleaved layout
+  RAFT_EXPECTS(index.codes_layout() == list_layout::INTERLEAVED,
+               "unpack_contiguous_list_data currently only supports INTERLEAVED layout");
+  auto typed_list =
+    std::static_pointer_cast<const list_data_interleaved<IdxT>>(index.lists()[label]);
   unpack_contiguous_list_data(out_codes,
-                              index.lists()[label]->data.view(),
+                              typed_list->data.view(),
                               n_rows,
                               index.pq_dim(),
                               offset_or_indices,
@@ -39,8 +45,9 @@ void unpack_contiguous_list_data(raft::resources const& res,
 };
 
 void pack_contiguous_list_data(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   const uint8_t* codes,
   uint32_t n_rows,
   uint32_t pq_dim,
@@ -56,7 +63,11 @@ void pack_contiguous_list_data(raft::resources const& res,
                                uint32_t label,
                                std::variant<uint32_t, const uint32_t*> offset_or_indices)
 {
-  pack_contiguous_list_data(index->lists()[label]->data.view(),
+  // Currently only supports interleaved layout
+  RAFT_EXPECTS(index->codes_layout() == list_layout::INTERLEAVED,
+               "pack_contiguous_list_data currently only supports INTERLEAVED layout");
+  auto typed_list = std::static_pointer_cast<list_data_interleaved<IdxT>>(index->lists()[label]);
+  pack_contiguous_list_data(typed_list->data.view(),
                             new_codes,
                             n_rows,
                             index->pq_dim(),
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data_impl.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data_impl.cuh
index 7387511ef9..bc477004b0 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data_impl.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_contiguous_list_data_impl.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -43,8 +43,9 @@ struct unpack_contiguous {
 template <uint32_t BlockSize, uint32_t PqBits>
 __launch_bounds__(BlockSize) static __global__ void unpack_contiguous_list_data_kernel(
   uint8_t* out_codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    in_list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> in_list_data,
   uint32_t n_rows,
   uint32_t pq_dim,
   std::variant<uint32_t, const uint32_t*> offset_or_indices)
@@ -64,8 +65,9 @@ __launch_bounds__(BlockSize) static __global__ void unpack_contiguous_list_data_
  */
 inline void unpack_contiguous_list_data_impl(
   uint8_t* codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   uint32_t n_rows,
   uint32_t pq_dim,
   std::variant<uint32_t, const uint32_t*> offset_or_indices,
@@ -120,8 +122,9 @@ struct pack_contiguous {
 
 template <uint32_t BlockSize, uint32_t PqBits>
 __launch_bounds__(BlockSize) static __global__ void pack_contiguous_list_data_kernel(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   const uint8_t* codes,
   uint32_t n_rows,
   uint32_t pq_dim,
@@ -143,8 +146,9 @@ __launch_bounds__(BlockSize) static __global__ void pack_contiguous_list_data_ke
  * @param[in] stream
  */
 inline void pack_contiguous_list_data_impl(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   const uint8_t* codes,
   uint32_t n_rows,
   uint32_t pq_dim,
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_list_data.hpp b/cpp/src/neighbors/ivf_pq/ivf_pq_list_data.hpp
index f9c9e28b3b..ac6bbec662 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_list_data.hpp
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_list_data.hpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -11,13 +11,22 @@
 
 namespace cuvs::neighbors::ivf_pq::detail {
 
-void unpack_list_data(
-  raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  std::variant<uint32_t, const uint32_t*> offset_or_indices,
-  uint32_t pq_bits,
-  rmm::cuda_stream_view stream);
+// Interleaved layout versions
+void unpack_list_data(raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes,
+                      raft::device_mdspan<const uint8_t,
+                                          list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                          raft::row_major> list_data,
+                      std::variant<uint32_t, const uint32_t*> offset_or_indices,
+                      uint32_t pq_bits,
+                      rmm::cuda_stream_view stream);
+
+void pack_list_data(raft::device_mdspan<uint8_t,
+                                        list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                                        raft::row_major> list_data,
+                    raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
+                    std::variant<uint32_t, const uint32_t*> offset_or_indices,
+                    uint32_t pq_bits,
+                    rmm::cuda_stream_view stream);
 
 /** Unpack the list data; see the public interface for the api and usage. */
 template <typename IdxT>
@@ -27,21 +36,15 @@ void unpack_list_data(raft::resources const& res,
                       uint32_t label,
                       std::variant<uint32_t, const uint32_t*> offset_or_indices)
 {
-  unpack_list_data(out_codes,
-                   index.lists()[label]->data.view(),
-                   offset_or_indices,
-                   index.pq_bits(),
-                   raft::resource::get_cuda_stream(res));
+  auto stream = raft::resource::get_cuda_stream(res);
+  // Currently only supports interleaved layout
+  RAFT_EXPECTS(index.codes_layout() == list_layout::INTERLEAVED,
+               "unpack_list_data currently only supports INTERLEAVED layout");
+  auto typed_list =
+    std::static_pointer_cast<const list_data_interleaved<IdxT>>(index.lists()[label]);
+  unpack_list_data(out_codes, typed_list->data.view(), offset_or_indices, index.pq_bits(), stream);
 }
 
-void pack_list_data(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
-  raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
-  std::variant<uint32_t, const uint32_t*> offset_or_indices,
-  uint32_t pq_bits,
-  rmm::cuda_stream_view stream);
-
 template <typename IdxT>
 void pack_list_data(raft::resources const& res,
                     index<IdxT>* index,
@@ -49,11 +52,12 @@ void pack_list_data(raft::resources const& res,
                     uint32_t label,
                     std::variant<uint32_t, const uint32_t*> offset_or_indices)
 {
-  pack_list_data(index->lists()[label]->data.view(),
-                 new_codes,
-                 offset_or_indices,
-                 index->pq_bits(),
-                 raft::resource::get_cuda_stream(res));
+  auto stream = raft::resource::get_cuda_stream(res);
+  // Currently only supports interleaved layout
+  RAFT_EXPECTS(index->codes_layout() == list_layout::INTERLEAVED,
+               "pack_list_data currently only supports INTERLEAVED layout");
+  auto typed_list = std::static_pointer_cast<list_data_interleaved<IdxT>>(index->lists()[label]);
+  pack_list_data(typed_list->data.view(), new_codes, offset_or_indices, index->pq_bits(), stream);
 }
 
 }  // namespace cuvs::neighbors::ivf_pq::detail
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_list_data_impl.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_list_data_impl.cuh
index c20857cb10..e03fb79398 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_list_data_impl.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_list_data_impl.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -42,8 +42,9 @@ struct unpack_codes {
 template <uint32_t BlockSize, uint32_t PqBits>
 __launch_bounds__(BlockSize) static __global__ void unpack_list_data_kernel(
   raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> out_codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    in_list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> in_list_data,
   std::variant<uint32_t, const uint32_t*> offset_or_indices)
 {
   const uint32_t pq_dim = out_codes.extent(1);
@@ -62,8 +63,9 @@ __launch_bounds__(BlockSize) static __global__ void unpack_list_data_kernel(
  */
 inline void unpack_list_data_impl(
   raft::device_matrix_view<uint8_t, uint32_t, raft::row_major> codes,
-  raft::device_mdspan<const uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<const uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   std::variant<uint32_t, const uint32_t*> offset_or_indices,
   uint32_t pq_bits,
   rmm::cuda_stream_view stream)
@@ -113,8 +115,9 @@ struct pass_codes {
 
 template <uint32_t BlockSize, uint32_t PqBits>
 __launch_bounds__(BlockSize) static __global__ void pack_list_data_kernel(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
   std::variant<uint32_t, const uint32_t*> offset_or_indices)
 {
@@ -134,8 +137,9 @@ __launch_bounds__(BlockSize) static __global__ void pack_list_data_kernel(
  * @param[in] stream
  */
 inline void pack_list_data_impl(
-  raft::device_mdspan<uint8_t, list_spec<uint32_t, uint32_t>::list_extents, raft::row_major>
-    list_data,
+  raft::device_mdspan<uint8_t,
+                      list_spec_interleaved<uint32_t, uint32_t>::list_extents,
+                      raft::row_major> list_data,
   raft::device_matrix_view<const uint8_t, uint32_t, raft::row_major> codes,
   std::variant<uint32_t, const uint32_t*> offset_or_indices,
   uint32_t pq_bits,
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_process_and_fill_codes_impl.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_process_and_fill_codes_impl.cuh
index ea56b79426..01f37dcd1a 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_process_and_fill_codes_impl.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_process_and_fill_codes_impl.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -24,7 +24,9 @@ __launch_bounds__(BlockSize) static __global__ void process_and_fill_codes_kerne
   raft::device_vector_view<IdxT*, uint32_t, raft::row_major> inds_ptrs,
   raft::device_vector_view<uint8_t*, uint32_t, raft::row_major> data_ptrs,
   raft::device_mdspan<const float, raft::extent_3d<uint32_t>, raft::row_major> pq_centers,
-  codebook_gen codebook_kind)
+  codebook_gen codebook_kind,
+  list_layout codes_layout,
+  uint32_t bytes_per_vector)
 {
   constexpr uint32_t kSubWarpSize = std::min<uint32_t>(raft::WarpSize, 1u << PqBits);
   using subwarp_align             = raft::Pow2<kSubWarpSize>;
@@ -37,7 +39,7 @@ __launch_bounds__(BlockSize) static __global__ void process_and_fill_codes_kerne
   if (lane_id == 0) { out_ix = atomicAdd(&list_sizes(cluster_ix), 1); }
   out_ix = raft::shfl(out_ix, 0, kSubWarpSize);
 
-  // write the label  (one record per subwarp)
+  // write the label (one record per subwarp)
   auto pq_indices = inds_ptrs(cluster_ix);
   if (lane_id == 0) {
     if (std::holds_alternative<IdxT>(src_offset_or_indices)) {
@@ -47,17 +49,22 @@ __launch_bounds__(BlockSize) static __global__ void process_and_fill_codes_kerne
     }
   }
 
-  // write the codes (one record per subwarp):
+  // write the codes (one record per subwarp)
   const uint32_t pq_dim = new_vectors.extent(1) / pq_centers.extent(1);
-  auto pq_extents = list_spec<uint32_t, IdxT>{PqBits, pq_dim, true}.make_list_extents(out_ix + 1);
-  auto pq_dataset = raft::make_mdspan<uint8_t, uint32_t, raft::row_major, false, true>(
-    data_ptrs[cluster_ix], pq_extents);
-  write_vector<PqBits, kSubWarpSize>(
-    pq_dataset,
-    out_ix,
-    row_ix,
-    pq_dim,
-    encode_vectors<kSubWarpSize, IdxT>{pq_centers, new_vectors, codebook_kind, cluster_ix});
+  auto encode_action =
+    encode_vectors<kSubWarpSize, IdxT>{pq_centers, new_vectors, codebook_kind, cluster_ix};
+
+  if (codes_layout == list_layout::FLAT) {
+    write_vector_flat<PqBits, kSubWarpSize>(
+      data_ptrs[cluster_ix], out_ix, row_ix, pq_dim, bytes_per_vector, encode_action);
+  } else {
+    auto pq_extents =
+      list_spec_interleaved<uint32_t, IdxT>{PqBits, pq_dim, true}.make_list_extents(out_ix + 1);
+    auto pq_dataset = raft::make_mdspan<uint8_t, uint32_t, raft::row_major, false, true>(
+      data_ptrs[cluster_ix], pq_extents);
+    write_vector_interleaved<PqBits, kSubWarpSize>(
+      pq_dataset, out_ix, row_ix, pq_dim, encode_action);
+  }
 }
 
 template <typename IdxT>
@@ -73,6 +80,10 @@ void launch_process_and_fill_codes_kernel(
   const uint32_t threads_per_vec = std::min<uint32_t>(raft::WarpSize, index.pq_book_size());
   dim3 blocks(raft::div_rounding_up_safe<IdxT>(n_rows, kBlockSize / threads_per_vec), 1, 1);
   dim3 threads(kBlockSize, 1, 1);
+
+  const uint32_t bytes_per_vector =
+    raft::div_rounding_up_safe(index.pq_dim() * index.pq_bits(), 8u);
+
   auto kernel = [](uint32_t pq_bits) {
     switch (pq_bits) {
       case 4: return process_and_fill_codes_kernel<kBlockSize, 4, IdxT>;
@@ -91,7 +102,9 @@ void launch_process_and_fill_codes_kernel(
                                                                           index.inds_ptrs(),
                                                                           index.data_ptrs(),
                                                                           index.pq_centers(),
-                                                                          index.codebook_kind());
+                                                                          index.codebook_kind(),
+                                                                          index.codes_layout(),
+                                                                          bytes_per_vector);
 
   RAFT_CUDA_TRY(cudaPeekAtLastError());
 }
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_search.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_search.cuh
index 8086cecfbf..79732642f8 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_search.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_search.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -859,6 +859,9 @@ inline void search(raft::resources const& handle,
     static_cast<uint64_t>(index.size()));
   RAFT_EXPECTS(params.n_probes > 0,
                "n_probes (number of clusters to probe in the search) must be positive.");
+  RAFT_EXPECTS(index.codes_layout() == list_layout::INTERLEAVED,
+               "IVF-PQ search requires INTERLEAVED codes layout. FLAT layout is not supported for "
+               "GPU search.");
 
   switch (utils::check_pointer_residency(queries, neighbors, distances)) {
     case utils::pointer_residency::device_only:
diff --git a/cpp/src/neighbors/ivf_pq/ivf_pq_serialize.cuh b/cpp/src/neighbors/ivf_pq/ivf_pq_serialize.cuh
index 0eecc862be..dbb4095dce 100644
--- a/cpp/src/neighbors/ivf_pq/ivf_pq_serialize.cuh
+++ b/cpp/src/neighbors/ivf_pq/ivf_pq_serialize.cuh
@@ -22,11 +22,8 @@
 namespace cuvs::neighbors::ivf_pq::detail {
 
 // Serialization version
-// No backward compatibility yet; that is, can't add additional fields without breaking
-// backward compatibility.
-// TODO(hcho3) Implement next-gen serializer for IVF that allows for expansion in a backward
-//             compatible fashion.
-constexpr int kSerializationVersion = 3;
+// Version 4 adds codes_layout field
+constexpr int kSerializationVersion = 4;
 
 /**
  * Write the index to an output stream
@@ -56,6 +53,7 @@ void serialize(raft::resources const& handle_, std::ostream& os, const index<Idx
 
   raft::serialize_scalar(handle_, os, index.metric());
   raft::serialize_scalar(handle_, os, index.codebook_kind());
+  raft::serialize_scalar(handle_, os, index.codes_layout());
   raft::serialize_scalar(handle_, os, index.n_lists());
 
   raft::serialize_mdspan(handle_, os, index.pq_centers());
@@ -71,9 +69,21 @@ void serialize(raft::resources const& handle_, std::ostream& os, const index<Idx
              raft::resource::get_cuda_stream(handle_));
   raft::resource::sync_stream(handle_);
   raft::serialize_mdspan(handle_, os, sizes_host.view());
-  auto list_store_spec = list_spec<uint32_t, IdxT>{index.pq_bits(), index.pq_dim(), true};
-  for (uint32_t label = 0; label < index.n_lists(); label++) {
-    ivf::serialize_list(handle_, os, index.lists()[label], list_store_spec, sizes_host(label));
+  // NOTE: We use static_cast here because serialize_list requires the concrete list type
+  // to access the spec_type for determining the serialized data layout.
+  if (index.codes_layout() == list_layout::FLAT) {
+    auto list_store_spec = list_spec_flat<uint32_t, IdxT>{index.pq_bits(), index.pq_dim(), true};
+    for (uint32_t label = 0; label < index.n_lists(); label++) {
+      auto& typed_list = static_cast<const list_data_flat<IdxT>&>(*index.lists()[label]);
+      ivf::serialize_list(handle_, os, typed_list, list_store_spec, sizes_host(label));
+    }
+  } else {
+    auto list_store_spec =
+      list_spec_interleaved<uint32_t, IdxT>{index.pq_bits(), index.pq_dim(), true};
+    for (uint32_t label = 0; label < index.n_lists(); label++) {
+      auto& typed_list = static_cast<const list_data_interleaved<IdxT>&>(*index.lists()[label]);
+      ivf::serialize_list(handle_, os, typed_list, list_store_spec, sizes_host(label));
+    }
   }
 }
 
@@ -126,6 +136,7 @@ auto deserialize(raft::resources const& handle_, std::istream& is) -> index<IdxT
 
   auto metric        = raft::deserialize_scalar<cuvs::distance::DistanceType>(handle_, is);
   auto codebook_kind = raft::deserialize_scalar<cuvs::neighbors::ivf_pq::codebook_gen>(handle_, is);
+  auto codes_layout  = raft::deserialize_scalar<cuvs::neighbors::ivf_pq::list_layout>(handle_, is);
   auto n_lists       = raft::deserialize_scalar<std::uint32_t>(handle_, is);
 
   RAFT_LOG_DEBUG("n_rows %zu, dim %d, pq_dim %d, pq_bits %d, n_lists %d",
@@ -137,7 +148,7 @@ auto deserialize(raft::resources const& handle_, std::istream& is) -> index<IdxT
 
   // Create owning_impl directly to get mutable access for deserialization
   auto impl = std::make_unique<owning_impl<IdxT>>(
-    handle_, metric, codebook_kind, n_lists, dim, pq_bits, pq_dim, cma);
+    handle_, metric, codebook_kind, n_lists, dim, pq_bits, pq_dim, cma, codes_layout);
 
   // Deserialize center/matrix data using mutable accessors
   raft::deserialize_mdspan(handle_, is, impl->pq_centers());
@@ -145,10 +156,22 @@ auto deserialize(raft::resources const& handle_, std::istream& is) -> index<IdxT
   raft::deserialize_mdspan(handle_, is, impl->centers_rot());
   raft::deserialize_mdspan(handle_, is, impl->rotation_matrix());
   raft::deserialize_mdspan(handle_, is, impl->list_sizes());
-  auto list_device_spec = list_spec<uint32_t, IdxT>{pq_bits, pq_dim, cma};
-  auto list_store_spec  = list_spec<uint32_t, IdxT>{pq_bits, pq_dim, true};
-  for (auto& list : impl->lists()) {
-    ivf::deserialize_list(handle_, is, list, list_store_spec, list_device_spec);
+  if (codes_layout == list_layout::FLAT) {
+    auto list_device_spec = list_spec_flat<uint32_t, IdxT>{pq_bits, pq_dim, cma};
+    auto list_store_spec  = list_spec_flat<uint32_t, IdxT>{pq_bits, pq_dim, true};
+    for (auto& list_data_base_ptr : impl->lists()) {
+      std::shared_ptr<list_data_flat<IdxT>> typed_list;
+      ivf::deserialize_list(handle_, is, typed_list, list_store_spec, list_device_spec);
+      list_data_base_ptr = typed_list;
+    }
+  } else {
+    auto list_device_spec = list_spec_interleaved<uint32_t, IdxT>{pq_bits, pq_dim, cma};
+    auto list_store_spec  = list_spec_interleaved<uint32_t, IdxT>{pq_bits, pq_dim, true};
+    for (auto& list_data_base_ptr : impl->lists()) {
+      std::shared_ptr<list_data_interleaved<IdxT>> typed_list;
+      ivf::deserialize_list(handle_, is, typed_list, list_store_spec, list_device_spec);
+      list_data_base_ptr = typed_list;
+    }
   }
 
   index<IdxT> idx(std::move(impl));
diff --git a/cpp/src/neighbors/ivf_pq_impl.hpp b/cpp/src/neighbors/ivf_pq_impl.hpp
index 67b4d27f4c..5c96755808 100644
--- a/cpp/src/neighbors/ivf_pq_impl.hpp
+++ b/cpp/src/neighbors/ivf_pq_impl.hpp
@@ -19,7 +19,8 @@ class index_impl : public index_iface<IdxT> {
              uint32_t dim,
              uint32_t pq_bits,
              uint32_t pq_dim,
-             bool conservative_memory_allocation);
+             bool conservative_memory_allocation,
+             list_layout codes_layout = list_layout::INTERLEAVED);
 
   ~index_impl()                            = default;
   index_impl(index_impl&&)                 = default;
@@ -29,6 +30,7 @@ class index_impl : public index_iface<IdxT> {
 
   cuvs::distance::DistanceType metric() const noexcept override;
   codebook_gen codebook_kind() const noexcept override;
+  list_layout codes_layout() const noexcept override;
   IdxT size() const noexcept override;
   uint32_t dim() const noexcept override;
   uint32_t dim_ext() const noexcept override;
@@ -40,8 +42,8 @@ class index_impl : public index_iface<IdxT> {
   uint32_t n_lists() const noexcept override;
   bool conservative_memory_allocation() const noexcept override;
 
-  std::vector<std::shared_ptr<list_data<IdxT>>>& lists() noexcept override;
-  const std::vector<std::shared_ptr<list_data<IdxT>>>& lists() const noexcept override;
+  std::vector<std::shared_ptr<list_data_base<IdxT>>>& lists() noexcept override;
+  const std::vector<std::shared_ptr<list_data_base<IdxT>>>& lists() const noexcept override;
 
   raft::device_vector_view<uint32_t, uint32_t, raft::row_major> list_sizes() noexcept override;
   raft::device_vector_view<const uint32_t, uint32_t, raft::row_major> list_sizes()
@@ -73,13 +75,14 @@ class index_impl : public index_iface<IdxT> {
  protected:
   cuvs::distance::DistanceType metric_;
   codebook_gen codebook_kind_;
+  list_layout codes_layout_;
   uint32_t dim_;
   uint32_t pq_bits_;
   uint32_t pq_dim_;
   bool conservative_memory_allocation_;
 
   // Primary data members
-  std::vector<std::shared_ptr<list_data<IdxT>>> lists_;
+  std::vector<std::shared_ptr<list_data_base<IdxT>>> lists_;
   raft::device_vector<uint32_t, uint32_t, raft::row_major> list_sizes_;
 
   // Lazy-initialized low-precision variants of index members - for low-precision coarse search.
@@ -109,7 +112,8 @@ class owning_impl : public index_impl<IdxT> {
               uint32_t dim,
               uint32_t pq_bits,
               uint32_t pq_dim,
-              bool conservative_memory_allocation);
+              bool conservative_memory_allocation,
+              list_layout codes_layout = list_layout::INTERLEAVED);
 
   ~owning_impl()                             = default;
   owning_impl(owning_impl&&)                 = default;
@@ -154,7 +158,8 @@ class view_impl : public index_impl<IdxT> {
             raft::device_mdspan<const float, pq_centers_extents, raft::row_major> pq_centers_view,
             raft::device_matrix_view<const float, uint32_t, raft::row_major> centers_view,
             raft::device_matrix_view<const float, uint32_t, raft::row_major> centers_rot_view,
-            raft::device_matrix_view<const float, uint32_t, raft::row_major> rotation_matrix_view);
+            raft::device_matrix_view<const float, uint32_t, raft::row_major> rotation_matrix_view,
+            list_layout codes_layout = list_layout::INTERLEAVED);
 
   ~view_impl()                           = default;
   view_impl(view_impl&&)                 = default;
diff --git a/cpp/src/neighbors/ivf_pq_index.cu b/cpp/src/neighbors/ivf_pq_index.cu
index 5ae14b7e28..fe230cd8aa 100644
--- a/cpp/src/neighbors/ivf_pq_index.cu
+++ b/cpp/src/neighbors/ivf_pq_index.cu
@@ -24,9 +24,11 @@ index_impl<IdxT>::index_impl(raft::resources const& handle,
                              uint32_t dim,
                              uint32_t pq_bits,
                              uint32_t pq_dim,
-                             bool conservative_memory_allocation)
+                             bool conservative_memory_allocation,
+                             list_layout codes_layout)
   : metric_(metric),
     codebook_kind_(codebook_kind),
+    codes_layout_(codes_layout),
     dim_(dim),
     pq_bits_(pq_bits),
     pq_dim_(pq_dim == 0 ? index<IdxT>::calculate_pq_dim(dim) : pq_dim),
@@ -53,6 +55,12 @@ codebook_gen index_impl<IdxT>::codebook_kind() const noexcept
   return codebook_kind_;
 }
 
+template <typename IdxT>
+list_layout index_impl<IdxT>::codes_layout() const noexcept
+{
+  return codes_layout_;
+}
+
 template <typename IdxT>
 IdxT index_impl<IdxT>::size() const noexcept
 {
@@ -114,13 +122,13 @@ bool index_impl<IdxT>::conservative_memory_allocation() const noexcept
 }
 
 template <typename IdxT>
-std::vector<std::shared_ptr<list_data<IdxT>>>& index_impl<IdxT>::lists() noexcept
+std::vector<std::shared_ptr<list_data_base<IdxT>>>& index_impl<IdxT>::lists() noexcept
 {
   return lists_;
 }
 
 template <typename IdxT>
-const std::vector<std::shared_ptr<list_data<IdxT>>>& index_impl<IdxT>::lists() const noexcept
+const std::vector<std::shared_ptr<list_data_base<IdxT>>>& index_impl<IdxT>::lists() const noexcept
 {
   return lists_;
 }
@@ -188,9 +196,17 @@ owning_impl<IdxT>::owning_impl(raft::resources const& handle,
                                uint32_t dim,
                                uint32_t pq_bits,
                                uint32_t pq_dim,
-                               bool conservative_memory_allocation)
-  : index_impl<IdxT>(
-      handle, metric, codebook_kind, n_lists, dim, pq_bits, pq_dim, conservative_memory_allocation),
+                               bool conservative_memory_allocation,
+                               list_layout codes_layout)
+  : index_impl<IdxT>(handle,
+                     metric,
+                     codebook_kind,
+                     n_lists,
+                     dim,
+                     pq_bits,
+                     pq_dim,
+                     conservative_memory_allocation,
+                     codes_layout),
     pq_centers_{raft::make_device_mdarray<float>(
       handle, index<IdxT>::make_pq_centers_extents(dim, pq_dim, pq_bits, codebook_kind, n_lists))},
     centers_{
@@ -230,9 +246,17 @@ view_impl<IdxT>::view_impl(
   raft::device_mdspan<const float, pq_centers_extents, raft::row_major> pq_centers_view,
   raft::device_matrix_view<const float, uint32_t, raft::row_major> centers_view,
   raft::device_matrix_view<const float, uint32_t, raft::row_major> centers_rot_view,
-  raft::device_matrix_view<const float, uint32_t, raft::row_major> rotation_matrix_view)
-  : index_impl<IdxT>(
-      handle, metric, codebook_kind, n_lists, dim, pq_bits, pq_dim, conservative_memory_allocation),
+  raft::device_matrix_view<const float, uint32_t, raft::row_major> rotation_matrix_view,
+  list_layout codes_layout)
+  : index_impl<IdxT>(handle,
+                     metric,
+                     codebook_kind,
+                     n_lists,
+                     dim,
+                     pq_bits,
+                     pq_dim,
+                     conservative_memory_allocation,
+                     codes_layout),
     pq_centers_view_(pq_centers_view),
     centers_view_(centers_view),
     centers_rot_view_(centers_rot_view),
@@ -452,6 +476,12 @@ codebook_gen index<IdxT>::codebook_kind() const noexcept
   return impl_->codebook_kind();
 }
 
+template <typename IdxT>
+list_layout index<IdxT>::codes_layout() const noexcept
+{
+  return impl_->codes_layout();
+}
+
 template <typename IdxT>
 uint32_t index<IdxT>::n_lists() const noexcept
 {
@@ -493,13 +523,13 @@ raft::device_matrix_view<const float, uint32_t, raft::row_major> index<IdxT>::ro
 }
 
 template <typename IdxT>
-std::vector<std::shared_ptr<list_data<IdxT>>>& index<IdxT>::lists() noexcept
+std::vector<std::shared_ptr<list_data_base<IdxT>>>& index<IdxT>::lists() noexcept
 {
   return impl_->lists();
 }
 
 template <typename IdxT>
-const std::vector<std::shared_ptr<list_data<IdxT>>>& index<IdxT>::lists() const noexcept
+const std::vector<std::shared_ptr<list_data_base<IdxT>>>& index<IdxT>::lists() const noexcept
 {
   return impl_->lists();
 }
@@ -595,8 +625,7 @@ uint32_t index_impl<IdxT>::get_list_size_in_bytes(uint32_t label) const
 {
   RAFT_EXPECTS(label < lists_.size(),
                "Expected label to be less than number of lists in the index");
-  auto& list_data = lists_[label]->data;
-  return list_data.size();
+  return lists_[label]->data_byte_size();
 }
 
 template <typename IdxT>
diff --git a/cpp/tests/neighbors/ann_ivf_pq.cuh b/cpp/tests/neighbors/ann_ivf_pq.cuh
index 140b569242..4ce9c96077 100644
--- a/cpp/tests/neighbors/ann_ivf_pq.cuh
+++ b/cpp/tests/neighbors/ann_ivf_pq.cuh
@@ -322,12 +322,13 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
   {
     // the original data cannot be reconstructed since the dataset was normalized
     if (index.metric() == cuvs::distance::DistanceType::CosineExpanded) { return; }
-    auto& rec_list = index.lists()[label];
+    auto& rec_list_base = index.lists()[label];
     // If the data is unbalanced the list might be empty, which is actually nullptr
-    if (!rec_list) { return; }
-    auto dim = index.dim();
-    n_take   = std::min<uint32_t>(n_take, rec_list->size.load());
-    n_skip   = std::min<uint32_t>(n_skip, rec_list->size.load() - n_take);
+    if (!rec_list_base) { return; }
+    auto rec_list = std::static_pointer_cast<list_data_interleaved<IdxT>>(rec_list_base);
+    auto dim      = index.dim();
+    n_take        = std::min<uint32_t>(n_take, rec_list->size.load());
+    n_skip        = std::min<uint32_t>(n_skip, rec_list->size.load() - n_take);
 
     if (n_take == 0) { return; }
 
@@ -352,10 +353,11 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
   {
     // NB: this is not reference, the list is retained; the index will have to create a new list on
     // `erase_list` op.
-    auto old_list = index->lists()[label];
+    auto old_list_base = index->lists()[label];
     // If the data is unbalanced the list might be empty, which is actually nullptr
-    if (!old_list) { return; }
-    auto n_rows = old_list->size.load();
+    if (!old_list_base) { return; }
+    auto old_list = std::static_pointer_cast<list_data_interleaved<IdxT>>(old_list_base);
+    auto n_rows   = old_list->size.load();
     if (n_rows == 0) { return; }
     if (index->metric() == cuvs::distance::DistanceType::CosineExpanded) { return; }
 
@@ -388,10 +390,11 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
 
   void check_packing(index<IdxT>* index, uint32_t label)
   {
-    auto old_list = index->lists()[label];
+    auto old_list_base = index->lists()[label];
     // If the data is unbalanced the list might be empty, which is actually nullptr
-    if (!old_list) { return; }
-    auto n_rows = old_list->size.load();
+    if (!old_list_base) { return; }
+    auto old_list = std::static_pointer_cast<list_data_interleaved<IdxT>>(old_list_base);
+    auto n_rows   = old_list->size.load();
 
     if (n_rows == 0) { return; }
 
@@ -405,19 +408,19 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
     ivf_pq::helpers::codepacker::extend_list_with_codes(
       handle_, index, codes.view(), indices.view(), label);
 
-    auto& new_list = index->lists()[label];
+    auto new_list = std::static_pointer_cast<list_data_interleaved<IdxT>>(index->lists()[label]);
     ASSERT_NE(old_list.get(), new_list.get())
       << "The old list should have been shared and retained after ivf_pq index has erased the "
          "corresponding cluster.";
-    auto list_data_size = (n_rows / cuvs::neighbors::ivf_pq::kIndexGroupSize) *
-                          new_list->data.extent(1) * new_list->data.extent(2) *
-                          new_list->data.extent(3);
+    auto list_data_byte_size = (n_rows / cuvs::neighbors::ivf_pq::kIndexGroupSize) *
+                               new_list->data.extent(1) * new_list->data.extent(2) *
+                               new_list->data.extent(3);
 
-    ASSERT_TRUE(old_list->data.size() >= list_data_size);
-    ASSERT_TRUE(new_list->data.size() >= list_data_size);
+    ASSERT_TRUE(old_list->data.size() >= list_data_byte_size);
+    ASSERT_TRUE(new_list->data.size() >= list_data_byte_size);
     ASSERT_TRUE(cuvs::devArrMatch(old_list->data.data_handle(),
                                   new_list->data.data_handle(),
-                                  list_data_size,
+                                  list_data_byte_size,
                                   cuvs::Compare<uint8_t>{}));
 
     // Pack a few vectors back to the list.
@@ -438,12 +441,13 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
       handle_, index, codes_to_pack, label, uint32_t(row_offset));
     ASSERT_TRUE(cuvs::devArrMatch(old_list->data.data_handle(),
                                   new_list->data.data_handle(),
-                                  list_data_size,
+                                  list_data_byte_size,
                                   cuvs::Compare<uint8_t>{}));
 
     // Another test with the API that take list_data directly
-    [[maybe_unused]] auto list_data = index->lists()[label]->data.view();
-    uint32_t n_take                 = 4;
+    [[maybe_unused]] auto list_data_view =
+      std::static_pointer_cast<list_data_interleaved<IdxT>>(index->lists()[label])->data.view();
+    uint32_t n_take = 4;
     if (static_cast<decltype(n_rows)>(row_offset + n_take) > n_rows) {
       RAFT_LOG_INFO(
         "Skipping IVF-PQ check_packing/take test for label %u due to insufficient data (%u "
@@ -454,17 +458,137 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
     }
     auto codes2 = raft::make_device_matrix<uint8_t>(handle_, n_take, index->pq_dim());
     ivf_pq::helpers::codepacker::unpack(
-      handle_, list_data, index->pq_bits(), row_offset, codes2.view());
+      handle_, list_data_view, index->pq_bits(), row_offset, codes2.view());
 
     // Write it back
     ivf_pq::helpers::codepacker::pack(
-      handle_, make_const_mdspan(codes2.view()), index->pq_bits(), row_offset, list_data);
+      handle_, make_const_mdspan(codes2.view()), index->pq_bits(), row_offset, list_data_view);
     ASSERT_TRUE(cuvs::devArrMatch(old_list->data.data_handle(),
                                   new_list->data.data_handle(),
-                                  list_data_size,
+                                  list_data_byte_size,
                                   cuvs::Compare<uint8_t>{}));
   }
 
+  void check_flat_layout_codes()
+  {
+    auto database_view =
+      raft::make_device_matrix_view<const DataT, int64_t>(database.data(), ps.num_db_vecs, ps.dim);
+
+    // Step 1: Build INTERLEAVED index WITHOUT data (just train the codebooks)
+    auto ipams_interleaved              = ps.index_params;
+    ipams_interleaved.add_data_on_build = false;
+    ipams_interleaved.codes_layout      = list_layout::INTERLEAVED;
+    auto index_interleaved =
+      cuvs::neighbors::ivf_pq::build(handle_, ipams_interleaved, database_view);
+
+    // Step 2: Build FLAT index using the SAME precomputed centers from the interleaved index
+    auto ipams_flat              = ps.index_params;
+    ipams_flat.add_data_on_build = false;
+    ipams_flat.codes_layout      = list_layout::FLAT;
+    auto index_flat              = cuvs::neighbors::ivf_pq::build(handle_,
+                                                     ipams_flat,
+                                                     index_interleaved.dim(),
+                                                     index_interleaved.pq_centers(),
+                                                     index_interleaved.centers(),
+                                                     index_interleaved.centers_rot(),
+                                                     index_interleaved.rotation_matrix());
+
+    ASSERT_EQ(index_interleaved.codes_layout(), list_layout::INTERLEAVED);
+    ASSERT_EQ(index_flat.codes_layout(), list_layout::FLAT);
+
+    // Step 3: Extend both indexes with the same data
+    auto indices = raft::make_device_vector<IdxT, int64_t>(handle_, ps.num_db_vecs);
+    raft::linalg::map_offset(handle_, indices.view(), raft::identity_op{});
+    cuvs::neighbors::ivf_pq::extend(
+      handle_, database_view, std::make_optional(indices.view()), &index_interleaved);
+    cuvs::neighbors::ivf_pq::extend(
+      handle_, database_view, std::make_optional(indices.view()), &index_flat);
+
+    ASSERT_TRUE(cuvs::devArrMatch(index_interleaved.list_sizes().data_handle(),
+                                  index_flat.list_sizes().data_handle(),
+                                  index_interleaved.n_lists(),
+                                  cuvs::Compare<uint32_t>{}));
+
+    auto ipams_from_flat              = ps.index_params;
+    ipams_from_flat.add_data_on_build = false;
+    ipams_from_flat.codes_layout      = list_layout::INTERLEAVED;
+    auto index_from_flat              = cuvs::neighbors::ivf_pq::build(handle_,
+                                                          ipams_from_flat,
+                                                          index_interleaved.dim(),
+                                                          index_interleaved.pq_centers(),
+                                                          index_interleaved.centers(),
+                                                          index_interleaved.centers_rot(),
+                                                          index_interleaved.rotation_matrix());
+
+    std::vector<uint32_t> list_sizes_host(index_flat.n_lists());
+    raft::update_host(
+      list_sizes_host.data(), index_flat.list_sizes().data_handle(), index_flat.n_lists(), stream_);
+    raft::resource::sync_stream(handle_);
+
+    // Step 4: Pack flat codes into the new interleaved index
+    uint32_t bytes_per_vec =
+      raft::div_rounding_up_safe(index_flat.pq_dim() * index_flat.pq_bits(), 8u);
+
+    for (uint32_t label = 0; label < index_flat.n_lists(); label++) {
+      uint32_t n_rows = list_sizes_host[label];
+      if (n_rows == 0) { continue; }
+
+      auto& list      = index_flat.lists()[label];
+      auto codes_view = raft::make_device_matrix_view<const uint8_t, uint32_t>(
+        list->data_ptr(), n_rows, bytes_per_vec);
+      auto indices_view =
+        raft::make_device_vector_view<const IdxT, uint32_t>(list->indices_ptr(), n_rows);
+
+      helpers::codepacker::extend_list_with_contiguous_codes(
+        handle_, &index_from_flat, codes_view, indices_view, label);
+    }
+
+    size_t queries_size = ps.num_queries * ps.k;
+    rmm::device_uvector<IdxT> indices_original(queries_size, stream_);
+    rmm::device_uvector<EvalT> distances_original(queries_size, stream_);
+    rmm::device_uvector<IdxT> indices_from_flat(queries_size, stream_);
+    rmm::device_uvector<EvalT> distances_from_flat(queries_size, stream_);
+
+    auto query_view =
+      raft::make_device_matrix_view<DataT, uint32_t>(search_queries.data(), ps.num_queries, ps.dim);
+    auto inds_orig_view =
+      raft::make_device_matrix_view<IdxT, uint32_t>(indices_original.data(), ps.num_queries, ps.k);
+    auto dists_orig_view = raft::make_device_matrix_view<EvalT, uint32_t>(
+      distances_original.data(), ps.num_queries, ps.k);
+    auto inds_flat_view =
+      raft::make_device_matrix_view<IdxT, uint32_t>(indices_from_flat.data(), ps.num_queries, ps.k);
+    auto dists_flat_view = raft::make_device_matrix_view<EvalT, uint32_t>(
+      distances_from_flat.data(), ps.num_queries, ps.k);
+
+    cuvs::neighbors::ivf_pq::search(
+      handle_, ps.search_params, index_interleaved, query_view, inds_orig_view, dists_orig_view);
+    cuvs::neighbors::ivf_pq::search(
+      handle_, ps.search_params, index_from_flat, query_view, inds_flat_view, dists_flat_view);
+
+    // Copy results to host for eval_neighbours
+    std::vector<IdxT> indices_orig_host(queries_size);
+    std::vector<IdxT> indices_flat_host(queries_size);
+    std::vector<EvalT> distances_orig_host(queries_size);
+    std::vector<EvalT> distances_flat_host(queries_size);
+    raft::update_host(indices_orig_host.data(), indices_original.data(), queries_size, stream_);
+    raft::update_host(indices_flat_host.data(), indices_from_flat.data(), queries_size, stream_);
+    raft::update_host(distances_orig_host.data(), distances_original.data(), queries_size, stream_);
+    raft::update_host(
+      distances_flat_host.data(), distances_from_flat.data(), queries_size, stream_);
+    raft::resource::sync_stream(handle_);
+
+    // Evaluate recall - should be very high since codes are identical
+    ASSERT_TRUE(eval_neighbours(indices_orig_host,
+                                indices_flat_host,
+                                distances_orig_host,
+                                distances_flat_host,
+                                ps.num_queries,
+                                ps.k,
+                                0.001,  // eps
+                                0.99))  // min_recall
+      << "Recall mismatch between original interleaved and repacked-from-flat indexes";
+  }
+
   template <typename BuildIndex>
   void run(BuildIndex build_index)
   {
@@ -589,6 +713,14 @@ class ivf_pq_test : public ::testing::TestWithParam<ivf_pq_inputs> {
   std::vector<EvalT> distances_ref;           // NOLINT
 };
 
+/**
+ * Separate test fixture for flat layout tests.
+ * This ensures flat_layout_tests() parameters only run check_flat_layout_codes(),
+ * not all the other tests defined in ivf_pq_test.
+ */
+template <typename EvalT, typename DataT, typename IdxT>
+class ivf_pq_flat_layout_test : public ivf_pq_test<EvalT, DataT, IdxT> {};
+
 template <typename EvalT, typename DataT, typename IdxT>
 class ivf_pq_filter_test : public ::testing::TestWithParam<ivf_pq_inputs> {
  public:
@@ -1147,7 +1279,60 @@ inline auto special_cases() -> test_cases_t
 #define TEST_BUILD_PRECOMPUTED(type) \
   TEST_P(type, build_precomputed) /* NOLINT */ { this->build_precomputed(); }
 
+#define TEST_FLAT_LAYOUT_CODES(type) \
+  TEST_P(type, flat_layout_codes) /* NOLINT */ { this->check_flat_layout_codes(); }
+
 #define INSTANTIATE(type, vals) \
   INSTANTIATE_TEST_SUITE_P(IvfPq, type, ::testing::ValuesIn(vals)); /* NOLINT */
 
+/**
+ * Test cases for flat layout comparison.
+ * Tests that FLAT layout produces the same PQ codes as INTERLEAVED layout
+ * by packing flat codes into a new interleaved index and comparing search results.
+ */
+inline auto flat_layout_tests() -> test_cases_t
+{
+  test_cases_t xs;
+
+  // Test with pq_bits = 8 (byte-aligned, 1 code per byte)
+  add_test_case(xs, [](ivf_pq_inputs& x) {
+    x.num_db_vecs          = 1000;
+    x.dim                  = 64;
+    x.index_params.n_lists = 10;
+    x.index_params.pq_bits = 8;
+    x.index_params.pq_dim  = 16;
+  });
+
+  // Test with pq_bits = 6 (packed codes)
+  add_test_case(xs, [](ivf_pq_inputs& x) {
+    x.num_db_vecs          = 1000;
+    x.dim                  = 64;
+    x.index_params.n_lists = 10;
+    x.index_params.pq_bits = 6;
+    x.index_params.pq_dim  = 16;
+  });
+
+  // Test with PER_CLUSTER codebook and pq_bits = 8
+  add_test_case(xs, [](ivf_pq_inputs& x) {
+    x.num_db_vecs                = 1000;
+    x.dim                        = 64;
+    x.index_params.n_lists       = 10;
+    x.index_params.pq_bits       = 8;
+    x.index_params.pq_dim        = 16;
+    x.index_params.codebook_kind = codebook_gen::PER_CLUSTER;
+  });
+
+  // Test with PER_CLUSTER codebook and pq_bits = 6
+  add_test_case(xs, [](ivf_pq_inputs& x) {
+    x.num_db_vecs                = 1000;
+    x.dim                        = 64;
+    x.index_params.n_lists       = 10;
+    x.index_params.pq_bits       = 6;
+    x.index_params.pq_dim        = 16;
+    x.index_params.codebook_kind = codebook_gen::PER_CLUSTER;
+  });
+
+  return xs;
+}
+
 }  // namespace cuvs::neighbors::ivf_pq
diff --git a/cpp/tests/neighbors/ann_ivf_pq/test_float_int64_t.cu b/cpp/tests/neighbors/ann_ivf_pq/test_float_int64_t.cu
index aad80bf317..47e5f8d0dd 100644
--- a/cpp/tests/neighbors/ann_ivf_pq/test_float_int64_t.cu
+++ b/cpp/tests/neighbors/ann_ivf_pq/test_float_int64_t.cu
@@ -7,8 +7,9 @@
 
 namespace cuvs::neighbors::ivf_pq {
 
-using f32_f32_i64        = ivf_pq_test<float, float, int64_t>;
-using f32_f32_i64_filter = ivf_pq_filter_test<float, float, int64_t>;
+using f32_f32_i64             = ivf_pq_test<float, float, int64_t>;
+using f32_f32_i64_flat_layout = ivf_pq_flat_layout_test<float, float, int64_t>;
+using f32_f32_i64_filter      = ivf_pq_filter_test<float, float, int64_t>;
 
 TEST_BUILD_HOST_INPUT_SEARCH(f32_f32_i64)
 TEST_BUILD_HOST_INPUT_OVERLAP_SEARCH(f32_f32_i64)
@@ -19,6 +20,9 @@ INSTANTIATE(f32_f32_i64,
             defaults() + small_dims() + big_dims_moderate_lut() + enum_variety_l2() +
               enum_variety_l2sqrt() + enum_variety_ip() + enum_variety_cosine());
 
+TEST_FLAT_LAYOUT_CODES(f32_f32_i64_flat_layout)
+INSTANTIATE(f32_f32_i64_flat_layout, flat_layout_tests());
+
 TEST_BUILD_SEARCH(f32_f32_i64_filter)
 INSTANTIATE(f32_f32_i64_filter,
             defaults() + small_dims() + big_dims_moderate_lut() + enum_variety_l2() +
diff --git a/go/ivf_pq/index_params.go b/go/ivf_pq/index_params.go
index d04b44621e..b35486061e 100644
--- a/go/ivf_pq/index_params.go
+++ b/go/ivf_pq/index_params.go
@@ -21,8 +21,8 @@ const (
 )
 
 var cCodebookKinds = map[codebookKind]int{
-	Subspace: C.PER_SUBSPACE,
-	Cluster:  C.PER_CLUSTER,
+	Subspace: C.CUVS_IVF_PQ_CODEBOOK_GEN_PER_SUBSPACE,
+	Cluster:  C.CUVS_IVF_PQ_CODEBOOK_GEN_PER_CLUSTER,
 }
 
 // Creates a new IndexParams
diff --git a/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pxd b/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pxd
index 072b61d643..f32089e8d6 100644
--- a/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pxd
+++ b/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pxd
@@ -24,9 +24,13 @@ cdef extern from "library_types.h":
 
 cdef extern from "cuvs/neighbors/ivf_pq.h" nogil:
 
-    ctypedef enum codebook_gen:
-        PER_SUBSPACE
-        PER_CLUSTER
+    ctypedef enum cuvsIvfPqCodebookGen:
+        CUVS_IVF_PQ_CODEBOOK_GEN_PER_SUBSPACE
+        CUVS_IVF_PQ_CODEBOOK_GEN_PER_CLUSTER
+
+    ctypedef enum cuvsIvfPqListLayout:
+        CUVS_IVF_PQ_LIST_LAYOUT_FLAT
+        CUVS_IVF_PQ_LIST_LAYOUT_INTERLEAVED
 
     ctypedef struct cuvsIvfPqIndexParams:
         cuvsDistanceType metric
@@ -37,10 +41,11 @@ cdef extern from "cuvs/neighbors/ivf_pq.h" nogil:
         double kmeans_trainset_fraction
         uint32_t pq_bits
         uint32_t pq_dim
-        codebook_gen codebook_kind
+        cuvsIvfPqCodebookGen codebook_kind
         bool force_random_rotation
         bool conservative_memory_allocation
         uint32_t max_train_points_per_pq_code
+        cuvsIvfPqListLayout codes_layout
 
     ctypedef cuvsIvfPqIndexParams* cuvsIvfPqIndexParams_t
 
diff --git a/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pyx b/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pyx
index d2552c2285..65b634631f 100644
--- a/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pyx
+++ b/python/cuvs/cuvs/neighbors/ivf_pq/ivf_pq.pyx
@@ -113,6 +113,15 @@ cdef class IndexParams:
         PER_SUBSPACE and PER_CLUSTER. In both cases, we will use
         pq_book_size * max_train_points_per_pq_code training points to
         train each codebook.
+    codes_layout : string, default = "interleaved"
+        Memory layout of the IVF-PQ list data.
+        Valid values ["flat", "interleaved"]
+
+            - flat: Codes are stored contiguously, one vector's codes after
+              another.
+            - interleaved: Codes are interleaved for optimized search
+              performance. This is the default and recommended for search
+              workloads.
     """
 
     def __cinit__(self):
@@ -134,7 +143,8 @@ cdef class IndexParams:
                  force_random_rotation=False,
                  add_data_on_build=True,
                  conservative_memory_allocation=False,
-                 max_train_points_per_pq_code=256):
+                 max_train_points_per_pq_code=256,
+                 codes_layout="interleaved"):
         self.params.n_lists = n_lists
         self.params.metric = <cuvsDistanceType>DISTANCE_TYPES[metric]
         self.params.metric_arg = metric_arg
@@ -143,9 +153,9 @@ cdef class IndexParams:
         self.params.pq_bits = pq_bits
         self.params.pq_dim = pq_dim
         if codebook_kind == "subspace":
-            self.params.codebook_kind = codebook_gen.PER_SUBSPACE
+            self.params.codebook_kind = cuvsIvfPqCodebookGen.CUVS_IVF_PQ_CODEBOOK_GEN_PER_SUBSPACE
         elif codebook_kind == "cluster":
-            self.params.codebook_kind = codebook_gen.PER_CLUSTER
+            self.params.codebook_kind = cuvsIvfPqCodebookGen.CUVS_IVF_PQ_CODEBOOK_GEN_PER_CLUSTER
         else:
             raise ValueError("Incorrect codebook kind %s" % codebook_kind)
         self.params.force_random_rotation = force_random_rotation
@@ -154,6 +164,12 @@ cdef class IndexParams:
             conservative_memory_allocation
         self.params.max_train_points_per_pq_code = \
             max_train_points_per_pq_code
+        if codes_layout == "flat":
+            self.params.codes_layout = cuvsIvfPqListLayout.CUVS_IVF_PQ_LIST_LAYOUT_FLAT
+        elif codes_layout == "interleaved":
+            self.params.codes_layout = cuvsIvfPqListLayout.CUVS_IVF_PQ_LIST_LAYOUT_INTERLEAVED
+        else:
+            raise ValueError("Incorrect codes layout %s" % codes_layout)
 
     def get_handle(self):
         return <size_t> self.params
@@ -210,6 +226,13 @@ cdef class IndexParams:
     def max_train_points_per_pq_code(self):
         return self.params.max_train_points_per_pq_code
 
+    @property
+    def codes_layout(self):
+        if self.params.codes_layout == cuvsIvfPqListLayout.CUVS_IVF_PQ_LIST_LAYOUT_FLAT:
+            return "flat"
+        else:
+            return "interleaved"
+
     def get_handle(self):
         return <size_t>self.params
 
@@ -539,8 +562,8 @@ def build_precomputed(IndexParams index_params, uint32_t dim, pq_centers, center
         Dimensionality of the input data
     pq_centers : CUDA array interface compliant tensor
         PQ codebook on device memory with required shape:
-        - codebook_kind PER_SUBSPACE: [pq_dim, pq_len, pq_book_size]
-        - codebook_kind PER_CLUSTER:  [n_lists, pq_len, pq_book_size]
+        - codebook_kind "subspace": [pq_dim, pq_len, pq_book_size]
+        - codebook_kind "cluster":  [n_lists, pq_len, pq_book_size]
         Supported dtype: float32
     centers : CUDA array interface compliant matrix
         Cluster centers in the original space [n_lists, dim_ext]
diff --git a/rust/cuvs-sys/build.rs b/rust/cuvs-sys/build.rs
index 1e49564669..99e811eeb9 100644
--- a/rust/cuvs-sys/build.rs
+++ b/rust/cuvs-sys/build.rs
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -79,7 +79,7 @@ fn main() {
         // Only generate bindings for cuvs/cagra types and functions
         .allowlist_type("(cuvs|bruteForce|cagra|DL).*")
         .allowlist_function("(cuvs|bruteForce|cagra).*")
-        .rustified_enum("(cuvs|cagra|DL|DistanceType|codebook_gen|cudaDataType_t).*")
+        .rustified_enum("(cuvs|cagra|DL|DistanceType|cudaDataType_t).*")
         // also need some basic cuda mem functions for copying data
         .allowlist_function("(cudaMemcpyAsync|cudaMemcpy)")
         .rustified_enum("cudaError")
diff --git a/rust/cuvs/src/ivf_pq/index_params.rs b/rust/cuvs/src/ivf_pq/index_params.rs
index 561f7f664b..e1f2d53656 100644
--- a/rust/cuvs/src/ivf_pq/index_params.rs
+++ b/rust/cuvs/src/ivf_pq/index_params.rs
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -8,7 +8,8 @@ use crate::error::{check_cuvs, Result};
 use std::fmt;
 use std::io::{stderr, Write};
 
-pub use ffi::codebook_gen;
+pub use ffi::cuvsIvfPqCodebookGen;
+pub use ffi::cuvsIvfPqListLayout;
 
 pub struct IndexParams(pub ffi::cuvsIvfPqIndexParams_t);
 
@@ -88,13 +89,24 @@ impl IndexParams {
         self
     }
 
-    pub fn set_codebook_kind(self, codebook_kind: codebook_gen) -> IndexParams {
+    pub fn set_codebook_kind(self, codebook_kind: cuvsIvfPqCodebookGen) -> IndexParams {
         unsafe {
             (*self.0).codebook_kind = codebook_kind;
         }
         self
     }
 
+    /// Memory layout of the IVF-PQ list data.
+    /// - FLAT: Codes are stored contiguously, one vector's codes after another.
+    /// - INTERLEAVED: Codes are interleaved for optimized search performance.
+    ///   This is the default and recommended for search workloads.
+    pub fn set_codes_layout(self, codes_layout: cuvsIvfPqListLayout) -> IndexParams {
+        unsafe {
+            (*self.0).codes_layout = codes_layout;
+        }
+        self
+    }
+
     /// Apply a random rotation matrix on the input data and queries even
     /// if `dim % pq_dim == 0`. Note: if `dim` is not multiple of `pq_dim`,
     /// a random rotation is always applied to the input data and queries