Fix cuVS build (#5107)

Summary:
Pull Request resolved: #5107

```
  The sentinel bug exists in both builds — the code is identical. The reason it manifests only
  on CMake/GHA:

  Different GPU hardware. Internal Buck tests run on RE with A100/H100 GPUs. GHA uses T4
  (compute capability 7.5). On the T4, cuVS CAGRA's filtered search doesn't explore enough graph
   to fill all k result slots, leaving sentinel values (INT32_MAX). On more powerful GPUs with
  higher throughput, the search finds all k valid neighbors before hitting the iteration limit.

  Possibly different cuVS versions. Buck uses Meta's vendored fbsource//third-party/cuvs:cuvs.
  GHA installs libcuvs=26.02 from the rapidsai conda channel. The sentinel behavior (what value,
   when it appears) could differ between versions.

  The underlying bug — blindly thrust::copying uint32_t indices to idx_t without sanitizing
  sentinels — exists in both builds. It just doesn't trigger internally because the GPU hardware
   is powerful enough to always fill all k slots. The thrust::transform fix is correct
  regardless of which GPU it runs on.
```

Reviewed By: junjieqi

Differential Revision: D101072304

fbshipit-source-id: fe191f61b2daba9f34f10bbd0f42c8530d51ad73

Author: Michael Norris

faiss/gpu/impl/BinaryCuvsCagra.cu

@@ -284,11 +284,12 @@ void BinaryCuvsCagra::search(
             distances_float_view,
             filter_ref);
 
-    thrust::copy(
-            raft::resource::get_thrust_policy(raft_handle),
+    faiss::gpu::sanitizeCuvsIndices(
+            resources_,
             indices_copy.data_handle(),
-            indices_copy.data_handle() + indices_copy.size(),
-            indices_view.data_handle());
+            indices_view.data_handle(),
+            indices_copy.size(),
+            n_);
     auto distances_view = raft::make_device_matrix_view(
             outDistances.data(),
             static_cast<int64_t>(numQueries),

faiss/gpu/impl/CuvsCagra.cu

@@ -315,11 +315,12 @@ void CuvsCagra<data_t>::search(
             indices_copy.view(),
             distances_view,
             filter_ref);
-    thrust::copy(
-            raft::resource::get_thrust_policy(raft_handle),
+    faiss::gpu::sanitizeCuvsIndices(
+            resources_,
             indices_copy.data_handle(),
-            indices_copy.data_handle() + indices_copy.size(),
-            indices_view.data_handle());
+            indices_view.data_handle(),
+            indices_copy.size(),
+            n_);
 }
 
 template <typename data_t>

faiss/gpu/utils/CuvsUtils.cu

@@ -33,6 +33,7 @@
 #include <thrust/gather.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/reduce.h>
+#include <thrust/transform.h>
 
 namespace faiss {
 namespace gpu {
@@ -115,5 +116,23 @@ idx_t inplaceGatherFilteredRows(
     return n_rows_valid;
 }
 
+void sanitizeCuvsIndices(
+        GpuResources* res,
+        uint32_t* src,
+        idx_t* dst,
+        size_t count,
+        idx_t n) {
+    raft::device_resources& raft_handle = res->getRaftHandleCurrentDevice();
+    thrust::transform(
+            raft_handle.get_thrust_policy(),
+            src,
+            src + count,
+            dst,
+            [n] __device__(uint32_t idx) -> idx_t {
+                return idx < static_cast<uint32_t>(n) ? static_cast<idx_t>(idx)
+                                                      : idx_t{-1};
+            });
+}
+
 } // namespace gpu
 } // namespace faiss

faiss/gpu/utils/CuvsUtils.h

@@ -71,6 +71,17 @@ idx_t inplaceGatherFilteredRows(
         GpuResources* res,
         Tensor<float, 2, true>& vecs,
         Tensor<idx_t, 1, true>& indices);
+
+/// Copy uint32_t indices to idx_t, replacing any index >= n with -1.
+/// cuVS CAGRA returns sentinel values (e.g. INT32_MAX) for result slots
+/// where filtered search couldn't find a valid neighbor.
+void sanitizeCuvsIndices(
+        GpuResources* res,
+        uint32_t* src,
+        idx_t* dst,
+        size_t count,
+        idx_t n);
+
 } // namespace gpu
 } // namespace faiss
 #pragma GCC visibility pop