Fix cuVS build (#5107)

Summary:

  D96034911 adds filtered search to BinaryCuvsCagra.cu (calling cuvs::neighbors::cagra::search with a filter_ref), but does not update
  fbcode/faiss/gpu/CMakeLists.txt.

  In the cmake build, cuVS source files are compiled with -fvisibility=hidden to ensure that RAFT/cuVS function calls resolve locally within libfaiss.so rather than
  through the PLT to libcuvs.so. This is critical because RAFT resources (cublas handles, CUDA streams) must be created and used within the same dynamic library
  context.

  Look at fbcode/faiss/gpu/CMakeLists.txt:302-316:

  set_source_files_properties(
      GpuIndexCagra.cu              # ✓ float CAGRA
      GpuDistance.cu
      GpuIndexIVFFlat.cu
      GpuIndexIVFPQ.cu
      GpuIndexFlat.cu
      StandardGpuResources.cpp
      impl/CuvsCagra.cu             # ✓ float CAGRA impl
      impl/CuvsFlatIndex.cu
      impl/CuvsIVFFlat.cu
      impl/CuvsIVFPQ.cu
      utils/CuvsFilterConvert.cu
      utils/CuvsUtils.cu
      TARGET_DIRECTORY faiss
      PROPERTIES COMPILE_OPTIONS "-fvisibility=hidden")

  Missing:
  - GpuIndexBinaryCagra.cu
  - impl/BinaryCuvsCagra.cu

  Before D96034911, this didn't matter because GpuIndexBinaryCagra::search() threw on any params (FAISS_THROW_IF_NOT_MSG(!params, "params not implemented")). After
  D96034911, it processes params and calls cuvs::neighbors::cagra::search with a filter, making the code sensitive to cross-DSO symbol resolution. Without hidden
  visibility, cuVS template instantiations (particularly cagra::search<uint8_t, uint32_t> with the filter type) leak out of libfaiss.so and can conflict with/resolve
   to libcuvs.so symbols, causing RAFT context mismatches.

  The float CAGRA IDSelector tests (TestGpuIndexCagra) pass because GpuIndexCagra.cu and impl/CuvsCagra.cu are in the hidden visibility list. The binary CAGRA tests
  fail because they aren't.

  Buck doesn't have this issue because it links everything statically — there's no DSO boundary.

Differential Revision: D101072304

Author: Michael Norris

faiss/gpu/CMakeLists.txt

@@ -287,32 +287,24 @@ if(FAISS_ENABLE_CUVS)
   target_compile_definitions(faiss_avx512 PUBLIC USE_NVIDIA_CUVS=1)
   target_compile_definitions(faiss_avx512_spr PUBLIC USE_NVIDIA_CUVS=1)
 
-  # Mark all functions as hidden so that we don't generate
-  # global 'public' functions that also exist in libcuvs.so
+  # Mark cuVS implementation files as hidden so that cuVS/RAFT
+  # template instantiations resolve locally within libfaiss.so
+  # rather than through the PLT to libcuvs.so. This is needed to
+  # ensure that RAFT cublas resources are created and used within
+  # the same dynamic library + CUDA runtime context.
   #
-  # This ensures that faiss functions will call the local version
-  # inside libfaiss.so . This is needed to ensure that things
-  # like raft cublas resources are created and used within the same
-  # dynamic library + CUDA runtime context which are requirements
-  # for valid execution
-  #
-  # To still allow these classes to be used by consumers, the
-  # respective classes/types in the headers are explicitly marked
-  # as 'public' so they can be used by consumers
+  # Only impl/utils files are hidden here — they contain the cuVS
+  # template instantiations and are internal to libfaiss.so.
+  # Public-facing GpuIndex*.cu files are NOT hidden because their
+  # headers lack visibility("default") annotations and hiding them
+  # would break the shared library's public API.
   set_source_files_properties(
-    GpuIndexCagra.cu
-    GpuDistance.cu
-    GpuIndexIVFFlat.cu
-    GpuIndexIVFPQ.cu
-    GpuIndexFlat.cu
-    StandardGpuResources.cpp
+    impl/BinaryCuvsCagra.cu
     impl/CuvsCagra.cu
     impl/CuvsFlatIndex.cu
     impl/CuvsIVFFlat.cu
     impl/CuvsIVFPQ.cu
-    utils/CuvsFilterConvert.cu
     utils/CuvsUtils.cu
-    TARGET_DIRECTORY faiss
     PROPERTIES COMPILE_OPTIONS "-fvisibility=hidden")
   target_compile_definitions(faiss_gpu_objs PUBLIC USE_NVIDIA_CUVS=1)
 endif()

faiss/gpu/impl/BinaryCuvsCagra.cu

@@ -36,6 +36,7 @@
 
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
+#include <thrust/transform.h>
 
 namespace faiss {
 namespace gpu {
@@ -284,11 +285,20 @@ void BinaryCuvsCagra::search(
             distances_float_view,
             filter_ref);
 
-    thrust::copy(
+    // Copy uint32_t indices to idx_t, replacing sentinel values from cuVS
+    // filtered search with -1. cuVS CAGRA returns large sentinel values
+    // (e.g., INT32_MAX) for result slots where no valid neighbor was found
+    // due to filtering. Faiss convention is -1 for "no result".
+    auto n = n_;
+    thrust::transform(
             raft::resource::get_thrust_policy(raft_handle),
             indices_copy.data_handle(),
             indices_copy.data_handle() + indices_copy.size(),
-            indices_view.data_handle());
+            indices_view.data_handle(),
+            [n] __device__(uint32_t idx) -> idx_t {
+                return idx < static_cast<uint32_t>(n) ? static_cast<idx_t>(idx)
+                                                      : idx_t{-1};
+            });
     auto distances_view = raft::make_device_matrix_view(
             outDistances.data(),
             static_cast<int64_t>(numQueries),

faiss/gpu/impl/CuvsCagra.cu

@@ -35,6 +35,7 @@
 
 #include <thrust/copy.h>
 #include <thrust/device_ptr.h>
+#include <thrust/transform.h>
 #include <optional>
 
 namespace faiss {
@@ -315,11 +316,20 @@ void CuvsCagra<data_t>::search(
             indices_copy.view(),
             distances_view,
             filter_ref);
-    thrust::copy(
+    // Copy uint32_t indices to idx_t, replacing sentinel values from cuVS
+    // filtered search with -1. cuVS CAGRA returns large sentinel values
+    // (e.g., INT32_MAX) for result slots where no valid neighbor was found
+    // due to filtering. Faiss convention is -1 for "no result".
+    auto n = n_;
+    thrust::transform(
             raft::resource::get_thrust_policy(raft_handle),
             indices_copy.data_handle(),
             indices_copy.data_handle() + indices_copy.size(),
-            indices_view.data_handle());
+            indices_view.data_handle(),
+            [n] __device__(uint32_t idx) -> idx_t {
+                return idx < static_cast<uint32_t>(n) ? static_cast<idx_t>(idx)
+                                                      : idx_t{-1};
+            });
 }
 
 template <typename data_t>

faiss/gpu/utils/CuvsFilterConvert.h

@@ -39,3 +39,4 @@ void convert_to_bitset(
         cuvs::core::bitset_view<uint32_t, int64_t> bitset,
         int num_threads = 0);
 } // namespace faiss::gpu
+#pragma GCC visibility pop