Add SVE low-dimensional nearest fast path and spherical SuperKMeans support (#5530)

Summary:
This PR adds two related changes for the SCANN coarse quantizer / PQ encoding path on ARM SVE: a low-dimensional nearest-neighbor fast path and spherical SuperKMeans support.

### 1. Low-dimensional L2sqr nearest fast path (d in {2,4,8})

`fvec_L2sqr_ny_nearest<ARM_SVE>` previously wrote all `ny` distances to the scratch buffer and then performed a separate linear scan. For PQ encoding (`compute_1_code`) each call has `d = dsub` and `ny = ksub`, making this the hot path. The new path maps one centroid to each SVE lane, gathers its `D` components, and keeps lane-local minimum distances and ids across batches. A predicated tail handles non-multiple `ny` values, equal distances preserve the first index, and only one final horizontal reduction is needed. The scratch buffer is not written, matching the x86 AVX2/AVX512 `D2/D4/D8` implementations.

`ProductQuantizer::compute_code` already uses the A1 dispatch on current `main`, so the ARM_SVE specialization remains reachable from PQ encoding after this rebase.

### 2. SuperKMeans spherical (inner-product) support

SuperKMeans previously only supported L2 clustering; the header explicitly excluded IP/cosine. With `cp.spherical=true`:
- `TrainState::R` becomes `std::unique_ptr<VectorTransform>`; power-of-two `d` uses the fast `HadamardRotation`, L2 keeps `RandomRotationMatrix`.
- Forgy initialization produces unit centroids before the first assignment. Each iteration computes and splits centroids first, then applies spherical centroid post-processing once, matching the `Clustering` flow.
- `HadamardRotation::reverse_transform` fills the missing inverse.
- `ClusteringParameters::use_super_kmeans` (default false) lets `Level1Quantizer::train_q1` route coarse quantizer training through SuperKMeans when explicitly enabled; the Python type stubs expose the new option.
- `block_l2<ARM_SVE>` completes the SuperKMeans SIMD kernels for the ADSampling pruning loop.

## Tests

- Low-dim nearest across SIMD levels: `d` in `{2,4,8}` and `ny` around common SIMD widths through 257, including non-multiple tails and first-index tie behavior. The dispatch test source is included in DD builds.
- Spherical objective close to vanilla spherical Clustering; unit-norm centroids.
- `use_super_kmeans` field inherited from `ClusteringParameters`.

## Benchmark

qwen 4096-dim, IP, nlist=1024, sub_dim=4, 1 thread, 50k rows:

| stage | baseline | optimized | change |
| --- | ---: | ---: | --- |
| Train | 67.73s | 24.36s | -64% |
| Add | 10.86s | 7.17s | -34% |
| Build | 78.59s | 31.52s | -60% |

Train speedup comes from SuperKMeans spherical coarse quantizer training; Add speedup comes from the low-dim SVE nearest fast path in PQ encoding. recall@10 unchanged vs Clustering baseline (diff <= 0.0002).

## Notes

This change targets ARM SVE. The x86 (AVX2/AVX512) implementations already have equivalent low-dim nearest specializations. The generic ARM SVE distance path comes from current `main`; this PR adds the low-dimensional nearest specializations on top of it.

Pull Request resolved: https://github.com/facebookresearch/faiss/pull/5530

Reviewed By: mnorris11

Differential Revision: D117563661

Pulled By: alibeklfc

fbshipit-source-id: 47ab1064d781a45d009923b4d5e62525718e83bf

faiss/CMakeLists.txt

@@ -64,6 +64,7 @@ set(FAISS_SIMD_NEON_SRC
 set(FAISS_SIMD_SVE_SRC
   impl/pq_code_distance/pq_code_distance-sve.cpp
   utils/simd_impl/distances_arm_sve.cpp
+  utils/simd_impl/super_kmeans_kernels_sve.cpp
 )
 set(FAISS_SIMD_RVV_SRC
   impl/fast_scan/impl-riscv.cpp

faiss/Clustering.h

@@ -75,6 +75,14 @@ struct ClusteringParameters {
     /// so the training process stops only if an error
     /// is unchanged from the previous iteration.
     double early_stop_threshold = 0.0;
+
+    /// Whether to use the SuperKMeans (super fast k-means) variant instead of
+    /// the vanilla Clustering implementation. Only honored by callers that
+    /// explicitly support it (e.g. IVF level-1 quantizer training). Cannot be
+    /// combined with a user-provided clustering index: SuperKMeans assigns
+    /// with its own index, so callers that supply both get an exception rather
+    /// than a silent downgrade to vanilla Clustering.
+    bool use_super_kmeans = false;
 };
 
 struct ClusteringIterationStats {

faiss/IndexIVF.cpp

@@ -6,6 +6,7 @@
  */
 
 #include <faiss/IndexIVF.h>
+#include <faiss/SuperKMeans.h>
 
 #include <omp.h>
 #include <atomic>
@@ -78,13 +79,27 @@ void Level1Quantizer::train_q1(
             printf("Training level-1 quantizer on %zd vectors in %zdD\n", n, d);
         }
 
-        Clustering clus(static_cast<int>(d), static_cast<int>(nlist), cp);
+        FAISS_THROW_IF_MSG(
+                cp.use_super_kmeans && clustering_index,
+                "cp.use_super_kmeans is incompatible with a user-provided "
+                "clustering_index: SuperKMeans assigns with its own index");
+
         quantizer->reset();
-        if (clustering_index) {
-            clus.train(n, x, *clustering_index);
+        if (cp.use_super_kmeans) {
+            SuperKMeansParameters super_cp;
+            static_cast<ClusteringParameters&>(super_cp) = cp;
+            SuperKMeans clus(
+                    static_cast<int>(d), static_cast<int>(nlist), super_cp);
+            clus.train(n, x);
             quantizer->add(nlist, clus.centroids.data());
         } else {
-            clus.train(n, x, *quantizer);
+            Clustering clus(static_cast<int>(d), static_cast<int>(nlist), cp);
+            if (clustering_index) {
+                clus.train(n, x, *clustering_index);
+                quantizer->add(nlist, clus.centroids.data());
+            } else {
+                clus.train(n, x, *quantizer);
+            }
         }
         quantizer->is_trained = true;
     } else if (quantizer_trains_alone == 2) {

faiss/SuperKMeans.cpp

@@ -59,7 +59,7 @@ namespace {
 struct TrainState {
     /// Orthogonal rotation. Train in rotated space (X_tilde = X * R);
     /// un-rotate centroids before return.
-    faiss::RandomRotationMatrix R;
+    std::unique_ptr<faiss::VectorTransform> R;
 
     std::vector<float> X_tilde; // (n, d) row-major
     int n = 0;
@@ -77,7 +77,17 @@ struct TrainState {
     int low_pruning_streak = 0;
     bool low_pruning_warning_printed = false;
 
-    explicit TrainState(int d) : R(d, d) {}
+    explicit TrainState(int d, bool spherical)
+            : R([d, spherical]() -> std::unique_ptr<faiss::VectorTransform> {
+                  // Spherical (inner-product) clustering only: a power-of-two
+                  // dimension can use the fast Hadamard rotation instead of
+                  // the generic random rotation. L2 training keeps the
+                  // original RandomRotationMatrix path unchanged.
+                  if (spherical && d > 0 && (d & (d - 1)) == 0) {
+                      return std::make_unique<faiss::HadamardRotation>(d);
+                  }
+                  return std::make_unique<faiss::RandomRotationMatrix>(d, d);
+              }()) {}
 };
 
 /// PDX block layout for the trailing pruning sweep: block b covers original
@@ -295,10 +305,17 @@ std::unique_ptr<uint8_t[]> setup_train_state(
             "SuperKMeans: training set size exceeds INT_MAX after sampling");
     state.n = static_cast<int>(nx);
 
-    state.R.init(cp.seed);
+    if (auto* R = dynamic_cast<HadamardRotation*>(state.R.get())) {
+        R->init(cp.seed);
+    } else {
+        auto* dense_rotation =
+                dynamic_cast<RandomRotationMatrix*>(state.R.get());
+        FAISS_ASSERT(dense_rotation != nullptr);
+        dense_rotation->init(cp.seed);
+    }
 
     state.X_tilde.resize(static_cast<size_t>(state.n) * d);
-    state.R.apply_noalloc(state.n, x_sampled, state.X_tilde.data());
+    state.R->apply_noalloc(state.n, x_sampled, state.X_tilde.data());
 
     // Forgy init: pick k random rows from the rotated pool as initial
     // centroids. These remain in rotated space; un-rotation happens
@@ -313,6 +330,9 @@ std::unique_ptr<uint8_t[]> setup_train_state(
                     state.X_tilde.data() + static_cast<size_t>(perm[j]) * d,
                     sizeof(float) * d);
         }
+        if (cp.spherical) {
+            fvec_renorm_L2(d, k, state.Y_tilde.data());
+        }
     }
 
     state.d_prime =
@@ -339,7 +359,7 @@ std::unique_ptr<uint8_t[]> setup_train_state(
 /// reverse_transform applies R^T = R^-1.
 void untransform_centroids(
         std::vector<float>& centroids,
-        const RandomRotationMatrix& R,
+        const VectorTransform& R,
         int d,
         int k,
         const float* Y_tilde) {
@@ -411,7 +431,7 @@ void SuperKMeans::train(idx_t n, const float* x) {
                static_cast<idx_t>(k) * cp.min_points_per_centroid);
     }
 
-    TrainState state(d);
+    TrainState state(d, cp.spherical);
     std::vector<int64_t> labels64;
     SuperKMeansAssignScratch assign_scratch;
     std::vector<float> hassign;
@@ -446,6 +466,9 @@ void SuperKMeans::train(idx_t n, const float* x) {
 
         const int nsplit =
                 update_centroids_and_split(d, k, state, labels64, hassign);
+        if (cp.spherical) {
+            fvec_renorm_L2(d, k, state.Y_tilde.data());
+        }
         const float pruning_rate = (iter == 0)
                 ? 0.0f
                 : adapt_d_prime(d, cp, state, total_pairs, pruned_at_gemm);
@@ -492,7 +515,7 @@ void SuperKMeans::train(idx_t n, const float* x) {
                (getmillisecs() - t_train_start) / 1000.0);
     }
 
-    untransform_centroids(centroids, state.R, d, k, state.Y_tilde.data());
+    untransform_centroids(centroids, *state.R, d, k, state.Y_tilde.data());
 }
 
 void super_kmeans_assign_iteration(

faiss/SuperKMeans.h

@@ -13,8 +13,9 @@
 //   "A Super Fast K-means for Indexing Vector Embeddings."
 //   arXiv preprint arXiv:2603.20009.
 //
-// Use when: L2 metric, k >= 1024, d >= 128, dense float embeddings.
-// Do not use for: IP/cosine (use Clustering with cp.spherical=true), small k,
+// Use when: L2 metric, or spherical inner-product clustering with
+// cp.spherical=true, k >= 1024, d >= 128, dense float embeddings.
+// Do not use for: small k,
 // binary data (use IndexBinaryIVF), or near-unit-sphere embeddings with
 // k < 4096 (chi-squared assumption breaks down).
 //

faiss/VectorTransform.cpp

@@ -513,6 +513,44 @@ void HadamardRotation::apply_noalloc(idx_t n, const float* x, float* xt) const {
     }
 }
 
+void HadamardRotation::reverse_transform(idx_t n, const float* xt, float* x)
+        const {
+    FAISS_THROW_IF_NOT_MSG(is_trained, "Transformation not trained yet");
+    FAISS_THROW_IF_NOT_MSG(
+            d_in == d_out,
+            "HadamardRotation inverse requires equal input/output dimensions");
+
+    const size_t p = d_out;
+    // Reverse of apply_noalloc: three unnormalized FWHT rounds scale norms
+    // by (sqrt(p))^3 = p*sqrt(p); the forward pass cancels this with
+    // total_scale = 1/(p*sqrt(p)), so the inverse applies the same factor.
+    const float inverse_scale = 1.0f / (p * std::sqrt(static_cast<float>(p)));
+
+#pragma omp parallel for schedule(dynamic)
+    for (idx_t i = 0; i < n; i++) {
+        const float* xi = xt + i * p;
+        float* xo = x + i * p;
+
+        // The inverse reverses the three sign-flip/Hadamard factors.
+        std::memcpy(xo, xi, p * sizeof(float));
+        fwht_inplace(xo, p);
+
+        for (size_t j = 0; j < p; j++) {
+            xo[j] *= signs3[j];
+        }
+        fwht_inplace(xo, p);
+
+        for (size_t j = 0; j < p; j++) {
+            xo[j] *= signs2[j];
+        }
+        fwht_inplace(xo, p);
+
+        for (size_t j = 0; j < p; j++) {
+            xo[j] *= signs1[j] * inverse_scale;
+        }
+    }
+}
+
 void HadamardRotation::check_identical(const VectorTransform& other) const {
     auto* hr = dynamic_cast<const HadamardRotation*>(&other);
     FAISS_THROW_IF_NOT_MSG(hr, "failed to cast to HadamardRotation");

faiss/VectorTransform.h

@@ -144,6 +144,9 @@ struct HadamardRotation : VectorTransform {
 
     void apply_noalloc(idx_t n, const float* x, float* xt) const override;
 
+    /// Apply the inverse transform when d_in == d_out.
+    void reverse_transform(idx_t n, const float* xt, float* x) const override;
+
     void check_identical(const VectorTransform& other) const override;
 
     HadamardRotation() {}

faiss/python/__init__.pyi

@@ -2628,6 +2628,7 @@ class ClusteringParameters:
     init_method: ClusteringInitMethod
     afkmc2_chain_length: int  # chain length for AFK-MC² initialization
     early_stop_threshold: float  # early stop threshold [0, 1]
+    use_super_kmeans: bool  # route through SuperKMeans when supported
 
     def __init__(self) -> None: ...
 

faiss/python/extra_wrappers.py

@@ -656,8 +656,10 @@ class SuperKmeans(Kmeans):
     `iteration_stats`, and `assign()` surface; additionally exposes
     `gemm_pruning_rates`.
 
-    kwargs are forwarded to `SuperKMeansParameters`. Fields not present on it
-    (e.g. `spherical`, `int_centroids`, `nredo`, `frozen_centroids`,
+    kwargs are forwarded to `SuperKMeansParameters`. Fields inherited from
+    `ClusteringParameters` such as `spherical` are supported (unit-normalized
+    centroids enable inner-product clustering). Fields not present on it
+    (e.g. `int_centroids`, `nredo`, `frozen_centroids`,
     `init_method`, `update_index`, `early_stop_threshold`,
     `progressive_dim_steps`, `gpu`) raise `AttributeError`.
     """

faiss/utils/simd_impl/distances_arm_sve.cpp

@@ -568,6 +568,48 @@ void fvec_L2sqr_ny<SIMDLevel::ARM_SVE>(
     }
 }
 
+namespace {
+
+/// Low-dimensional L2sqr nearest (D in {2,4,8}). The data is row-major
+/// (centroid c is y[c*D .. c*D+D]). Each SVE lane tracks one centroid across
+/// batches, keeping the lane-local minimum and index in registers. The scratch
+/// buffer is not written, matching the AVX2/AVX512 D2/D4/D8 implementations.
+template <int D>
+size_t fvec_L2sqr_ny_nearest_lowdim(
+        float* /*distances_tmp_buffer*/,
+        const float* x,
+        const float* y,
+        size_t ny) {
+    const size_t lanes = svcntw();
+    svfloat32_t global_mins = svdup_n_f32(HUGE_VALF);
+    svuint32_t global_ids = svdup_n_u32(0);
+    svuint32_t current_ids = svindex_u32(0, 1);
+
+    for (size_t c = 0; c < ny; c += lanes) {
+        const svbool_t pg = svwhilelt_b32_u64(c, ny);
+        svfloat32_t distances = svdup_n_f32(0.0f);
+        for (uint32_t j = 0; j < D; ++j) {
+            const svuint32_t offsets = svindex_u32(j, D);
+            const svfloat32_t yv =
+                    svld1_gather_u32index_f32(pg, y + c * D, offsets);
+            const svfloat32_t diff = svsub_n_f32_x(pg, yv, x[j]);
+            distances = svmla_f32_m(pg, distances, diff, diff);
+        }
+
+        const svbool_t closer = svcmplt_f32(pg, distances, global_mins);
+        global_mins = svsel_f32(closer, distances, global_mins);
+        global_ids = svsel_u32(closer, current_ids, global_ids);
+        current_ids =
+                svadd_n_u32_x(pg, current_ids, static_cast<uint32_t>(lanes));
+    }
+
+    const svbool_t all = svptrue_b32();
+    const float global_min = svminv_f32(all, global_mins);
+    return svminv_u32(svcmpeq_n_f32(all, global_mins, global_min), global_ids);
+}
+
+} // namespace
+
 template <>
 size_t fvec_L2sqr_ny_nearest<SIMDLevel::ARM_SVE>(
         float* distances_tmp_buffer,
@@ -575,6 +617,18 @@ size_t fvec_L2sqr_ny_nearest<SIMDLevel::ARM_SVE>(
         const float* y,
         size_t d,
         size_t ny) {
+    switch (d) {
+        case 2:
+            return fvec_L2sqr_ny_nearest_lowdim<2>(
+                    distances_tmp_buffer, x, y, ny);
+        case 4:
+            return fvec_L2sqr_ny_nearest_lowdim<4>(
+                    distances_tmp_buffer, x, y, ny);
+        case 8:
+            return fvec_L2sqr_ny_nearest_lowdim<8>(
+                    distances_tmp_buffer, x, y, ny);
+    }
+
     const size_t lanes = static_cast<size_t>(svcntw());
 
     size_t nearest_idx = 0;

faiss/utils/simd_impl/super_kmeans_dispatch.h

@@ -9,13 +9,12 @@
 
 // Private dispatch wrapper for SuperKMeans's block_l2. Routes to the
 // highest available SIMD specialization at runtime (DD mode) or the
-// compiled-in level (static mode). aarch64 currently falls through to the
-// scalar primary template; adding NEON/SVE means just adding a new
-// specialization file alongside the AVX ones.
+// compiled-in level (static mode).
 //
-// Known perf gap: aarch64 (NEON/SVE) specializations are not implemented yet.
-// aarch64 falls through to the scalar primary template. Validating SVE requires
-// a Graviton-class host; deferred to a focused follow-up.
+// The A1 level mask is required: plain with_simd_level() omits the ARM_SVE bit,
+// so on an SVE host the SVE specialization would never be instantiated.
+// ARM_NEON has no specialization and falls through to the scalar primary
+// template.
 
 #include <faiss/impl/simd_dispatch.h>
 #include <faiss/utils/simd_impl/super_kmeans_kernels.h>
@@ -24,7 +23,7 @@ namespace faiss {
 namespace detail {
 
 inline float block_l2_dispatch(const float* x, const float* y, int n) {
-    return with_selected_simd_levels<AVAILABLE_SIMD_LEVELS_A0>(
+    return with_simd_level_a1(
             [&]<SIMDLevel SL>() { return block_l2<SL>(x, y, n); });
 }
 

faiss/utils/simd_impl/super_kmeans_kernels.h

@@ -16,7 +16,7 @@ namespace detail {
 
 // Squared L2 over `n` dimensions; n in [1, pdx_block_size].
 // Primary template is the scalar fallback; SIMDLevels without a dedicated
-// specialization (ARM_NEON, ARM_SVE, NONE, ...) use it directly.
+// specialization (ARM_NEON, NONE, ...) use it directly.
 template <SIMDLevel Level>
 inline float block_l2(const float* x, const float* y, int n) {
     float s = 0.0f;
@@ -39,5 +39,10 @@ template <>
 float block_l2<SIMDLevel::AVX512>(const float* x, const float* y, int n);
 #endif
 
+#ifdef COMPILE_SIMD_ARM_SVE
+template <>
+float block_l2<SIMDLevel::ARM_SVE>(const float* x, const float* y, int n);
+#endif
+
 } // namespace detail
 } // namespace faiss

faiss/utils/simd_impl/super_kmeans_kernels_sve.cpp

@@ -0,0 +1,34 @@
+/*
+ * Copyright (c) Meta Platforms, Inc. and affiliates.
+ *
+ * This source code is licensed under the MIT license found in the
+ * LICENSE file in the root directory of this source tree.
+ */
+
+#ifdef COMPILE_SIMD_ARM_SVE
+
+#include <faiss/utils/simd_impl/super_kmeans_kernels.h>
+
+#include <arm_sve.h>
+
+namespace faiss {
+namespace detail {
+
+template <>
+float block_l2<SIMDLevel::ARM_SVE>(const float* x, const float* y, int n) {
+    svfloat32_t acc = svdup_n_f32(0.0f);
+    const int lanes = static_cast<int>(svcntw());
+    for (int m = 0; m < n; m += lanes) {
+        const svbool_t pg = svwhilelt_b32(m, n);
+        const svfloat32_t xv = svld1_f32(pg, x + m);
+        const svfloat32_t yv = svld1_f32(pg, y + m);
+        const svfloat32_t diff = svsub_f32_x(pg, xv, yv);
+        acc = svmla_f32_m(pg, acc, diff, diff);
+    }
+    return svaddv_f32(svptrue_b32(), acc);
+}
+
+} // namespace detail
+} // namespace faiss
+
+#endif // COMPILE_SIMD_ARM_SVE

tests/CMakeLists.txt

@@ -62,9 +62,10 @@ if(FAISS_ENABLE_SVS)
   list(APPEND FAISS_TEST_SRC test_svs.cpp)
 endif()
 
-# DD-only tests: x86-specific tests containing raw AVX2/AVX512 intrinsics
+# DD-only tests; x86-specific raw-intrinsic tests are added separately below.
 if(FAISS_OPT_LEVEL STREQUAL "dd")
   list(APPEND FAISS_TEST_SRC
+    test_distances_dispatch.cpp
     test_simd_perf.cpp
   )
   if(CMAKE_SYSTEM_PROCESSOR MATCHES "(x86_64|amd64|AMD64)")

tests/test_distances_dispatch.cpp

@@ -32,6 +32,7 @@
 
 #include <gtest/gtest.h>
 
+#include <algorithm>
 #include <cmath>
 #include <random>
 #include <vector>
@@ -299,3 +300,68 @@ TEST(DistancesDispatch, FvecL2sqrNyNearest_AllLevels) {
                 levels);
     }
 }
+
+// The low-dimensional nearest path (d in {2,4,8}, one SIMD lane per
+// centroid) must agree with the generic implementation for any ny, including
+// values around common SIMD widths and non-multiple tails.
+TEST(DistancesDispatch, FvecL2sqrNyNearest_LowDim) {
+    SIMDLevelGuard guard;
+    auto levels = available_levels();
+    SKIP_IF_SINGLE_LEVEL(levels);
+    constexpr size_t kLowDims[] = {2, 4, 8};
+    constexpr size_t kNy[] = {
+            1,
+            2,
+            3,
+            7,
+            8,
+            9,
+            15,
+            16,
+            17,
+            23,
+            31,
+            32,
+            33,
+            63,
+            64,
+            65,
+            255,
+            256,
+            257};
+    for (size_t d : kLowDims) {
+        for (size_t ny : kNy) {
+            auto x = rand_vec(d, 40);
+            auto y = rand_vec(d * ny, 41);
+            check_index_at_levels(
+                    [&]() {
+                        std::vector<float> tmp(ny);
+                        return fvec_L2sqr_ny_nearest(
+                                tmp.data(), x.data(), y.data(), d, ny);
+                    },
+                    levels);
+        }
+    }
+}
+
+TEST(DistancesDispatch, FvecL2sqrNyNearest_LowDimKeepsFirstTie) {
+    SIMDLevelGuard guard;
+    auto levels = available_levels();
+    SKIP_IF_SINGLE_LEVEL(levels);
+    constexpr size_t kLowDims[] = {2, 4, 8};
+    constexpr size_t ny = 23;
+    constexpr size_t expected = 3;
+    for (size_t d : kLowDims) {
+        std::vector<float> x(d, 0.0f);
+        std::vector<float> y(d * ny, 1.0f);
+        std::fill_n(y.data() + expected * d, d, 0.0f);
+        std::fill_n(y.data() + 17 * d, d, 0.0f);
+        auto nearest = [&]() {
+            std::vector<float> tmp(ny);
+            return fvec_L2sqr_ny_nearest(tmp.data(), x.data(), y.data(), d, ny);
+        };
+        SIMDConfig::set_level(SIMDLevel::NONE);
+        EXPECT_EQ(expected, nearest());
+        check_index_at_levels(nearest, levels);
+    }
+}

tests/test_super_kmeans.py

@@ -53,6 +53,55 @@ class SuperKMeansTest(unittest.TestCase):
         v_final = v_stats.at(v_stats.size() - 1).obj
         self.assertLess(abs(sc_final - v_final) / v_final, 0.05)
 
+    def test_spherical_objective_close_to_vanilla_spherical(self):
+        # With cp.spherical=true, SuperKMeans unit-normalizes centroids so the
+        # L2 assignment it computes is equivalent to inner-product assignment.
+        # The final objective must track vanilla spherical Clustering.
+        d, k, n = 64, 16, 2000
+        x = SyntheticDataset(d, n, 0, 0).get_train()
+
+        p = faiss.SuperKMeansParameters()
+        p.seed = 42
+        p.niter = 10
+        p.spherical = True
+        sc = faiss.SuperKMeans(d, k, p)
+        sc.train(x)
+
+        quant = faiss.IndexFlatL2(d)
+        vanilla = faiss.Clustering(d, k)
+        vanilla.seed = 42
+        vanilla.niter = 10
+        vanilla.spherical = True
+        vanilla.train(x, quant)
+
+        sc_final = sc.iteration_stats.at(sc.iteration_stats.size() - 1).obj
+        v_stats = vanilla.iteration_stats
+        v_final = v_stats.at(v_stats.size() - 1).obj
+        self.assertLess(abs(sc_final - v_final) / v_final, 0.05)
+
+    def test_spherical_produces_unit_centroids(self):
+        d, k, n = 64, 16, 2000
+        x = SyntheticDataset(d, n, 0, 0).get_train()
+
+        p = faiss.SuperKMeansParameters()
+        p.seed = 42
+        p.niter = 10
+        p.spherical = True
+        sc = faiss.SuperKMeans(d, k, p)
+        sc.train(x)
+
+        centroids = faiss.vector_to_array(sc.centroids).reshape(k, d)
+        norms = np.linalg.norm(centroids, axis=1)
+        np.testing.assert_allclose(norms, 1.0, atol=1e-4)
+
+    def test_use_super_kmeans_field_is_inherited(self):
+        # use_super_kmeans lives on ClusteringParameters and is inherited by
+        # SuperKMeansParameters; it must be settable from Python.
+        p = faiss.SuperKMeansParameters()
+        self.assertFalse(p.use_super_kmeans)
+        p.use_super_kmeans = True
+        self.assertTrue(p.use_super_kmeans)
+
     def test_pruning_rate_in_expected_range(self):
         # The 10-d intrinsic manifold of SyntheticDataset gives ADSampling
         # the assigned-vs-other distance gap it needs to prune effectively.