Skip zero-sized reverse-edge launches in Vamana build

cpp/src/neighbors/detail/vamana/vamana_build.cuh

@@ -459,10 +459,13 @@ void batched_insert_vamana(
     }
 
     DistPair<IdxT, accT>* edge_dist_pair_ptr = edge_dist_pair.data_handle();
-    raft::linalg::map_offset(
-      res,
-      raft::make_device_vector_view<IdxT, int64_t>(edge_dest.data_handle(), total_edges),
-      [edge_dist_pair_ptr] __device__(size_t i) { return edge_dist_pair_ptr[i].idx; });
+    // No edges if the first inserted point is the medoid; skip to avoid zero-sized launches
+    if (total_edges > 0) {
+      raft::linalg::map_offset(
+        res,
+        raft::make_device_vector_view<IdxT, int64_t>(edge_dest.data_handle(), total_edges),
+        [edge_dist_pair_ptr] __device__(size_t i) { return edge_dist_pair_ptr[i].idx; });
+    }
 
     // Sort to group reverse edges by destination
     cub::DeviceMergeSort::SortPairs(temp_sort_storage.data_handle(),
@@ -480,7 +483,9 @@ void batched_insert_vamana(
     // Find which node IDs have reverse edges and their indices in the reverse edge list
     raft::copy(edge_dest_vec.data_handle(), edge_dest.data_handle(), total_edges, stream);
     auto unique_indices = raft::make_device_vector<int>(res, total_edges);
-    raft::linalg::map_offset(res, unique_indices.view(), raft::identity_op{});
+    if (total_edges > 0) {
+      raft::linalg::map_offset(res, unique_indices.view(), raft::identity_op{});
+    }
 
     thrust::unique_by_key(rmm::exec_policy_nosync(stream, large_ws),
                           thrust::device_pointer_cast(edge_dest_vec.data_handle()),

cpp/tests/neighbors/ann_vamana/test_float_uint32_t.cu

@@ -1,17 +1,48 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
 #include <gtest/gtest.h>
 
 #include "../ann_vamana.cuh"
 
+#include <numeric>
+
 namespace cuvs::neighbors::vamana {
 
 typedef AnnVamanaTest<float, float, std::uint32_t> AnnVamanaTestF_U32;
 TEST_P(AnnVamanaTestF_U32, AnnVamana) { this->testVamana(); }
 
 INSTANTIATE_TEST_CASE_P(AnnVamanaTest, AnnVamanaTestF_U32, ::testing::ValuesIn(inputs));
 
+// Returns a srand() seed for which the build's first inserted point is also its medoid.
+// The build calls rand() n times to shuffle the insert order, then once for the medoid.
+static unsigned seed_where_first_insert_is_medoid(uint32_t n)
+{
+  for (unsigned seed = 1;; seed++) {
+    srand(seed);
+    std::vector<uint32_t> order(n);
+    std::iota(order.begin(), order.end(), 0);
+    for (uint32_t i = 0; i < n; i++) {
+      std::swap(order[i], order[rand() % n]);
+    }
+    if (order[0] == rand() % n) { return seed; }
+  }
+}
+
+// https://github.com/NVIDIA/cuvs/issues/2767: the first batch then has no reverse edges.
+TEST(AnnVamanaBuildTest, FirstInsertIsMedoid)
+{
+  raft::resources res;
+  constexpr int64_t n = 410, dim = 16;
+  auto dataset = raft::make_device_matrix<float, int64_t>(res, n, dim);
+  raft::random::RngState r(1234ULL);
+  raft::random::uniform(res, r, dataset.data_handle(), n * dim, -1.0f, 1.0f);
+
+  srand(seed_where_first_insert_is_medoid(n));
+  auto idx = build(res, index_params{}, raft::make_const_mdspan(dataset.view()));
+  EXPECT_EQ(idx.graph().extent(0), n);
+}
+
 }  // namespace cuvs::neighbors::vamana