Use librtcx for JIT-LTO

conda/recipes/libcuvs/recipe.yaml

@@ -227,6 +227,7 @@ outputs:
         content: |
           cmake --install cpp/build --component cuvs_static
           cmake --install cpp/build --component hnswlib
+          cmake --install cpp/build --component rtcx
     requirements:
       build:
         - cmake ${{ cmake_version }}
@@ -256,6 +257,7 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
+        - zstd-static =1.5.7
       ignore_run_exports:
         by_name:
           - cuda-cudart

cpp/CMakeLists.txt

@@ -190,6 +190,7 @@ if(NOT BUILD_CPU_ONLY)
   include(cmake/thirdparty/get_cutlass.cmake)
   include(${rapids-cmake-dir}/cpm/cuco.cmake)
   rapids_cpm_cuco()
+  include(cmake/thirdparty/get_rtcx.cmake)
 endif()
 
 if(BUILD_TESTS OR BUILD_C_TESTS)
@@ -352,7 +353,6 @@ if(NOT BUILD_CPU_ONLY)
   endif()
 
   # Generate interleaved scan kernel files at build time
-  include(cmake/modules/generate_jit_lto_kernels.cmake)
 
   add_library(jit_lto_kernel_usage_requirements INTERFACE)
   target_include_directories(
@@ -1361,11 +1361,6 @@ if(NOT BUILD_CPU_ONLY)
     src/core/omp_wrapper.cpp
     src/util/file_io.cpp
     src/util/host_memory.cpp
-    src/detail/jit_lto/AlgorithmLauncher.cpp
-    src/detail/jit_lto/AlgorithmPlanner.cpp
-    src/detail/jit_lto/FragmentEntry.cpp
-    src/detail/jit_lto/nvjitlink_checker.cpp
-    src/detail/jit_lto/NVRTCLTOFragmentCompiler.cpp
     src/distance/detail/kernels/gram_matrix.cu
     src/distance/detail/kernels/kernel_factory.cu
     src/distance/detail/kernels/kernel_matrices.cu
@@ -1491,6 +1486,7 @@ if(NOT BUILD_CPU_ONLY)
     PRIVATE cuvs::cuvs_cpp_headers
             cuco::cuco
             nvidia::cutlass::cutlass
+            rtcx::rtcx
             ${CUVS_CTK_MATH_DEPENDENCIES}
             $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
             $<TARGET_NAME_IF_EXISTS:hnswlib::hnswlib>
@@ -1570,7 +1566,7 @@ if(NOT BUILD_CPU_ONLY)
              $<BUILD_LOCAL_INTERFACE:$<TARGET_NAME_IF_EXISTS:hnswlib::hnswlib>>
              $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3>
       PRIVATE $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX> $<COMPILE_ONLY:nvidia::cutlass::cutlass>
-              $<COMPILE_ONLY:cuco::cuco> CUDA::nvJitLink CUDA::nvrtc
+              $<COMPILE_ONLY:cuco::cuco> CUDA::nvJitLink CUDA::nvrtc rtcx::rtcx
     )
 
     # ensure CUDA symbols aren't relocated to the middle of the debug build binaries
@@ -1628,6 +1624,7 @@ SECTIONS
       PRIVATE $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
               CUDA::nvJitLink
               CUDA::nvrtc
+              rtcx::rtcx
               $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3>
               $<COMPILE_ONLY:nvidia::cutlass::cutlass>
               $<COMPILE_ONLY:cuco::cuco>

cpp/cmake/thirdparty/get_rtcx.cmake

@@ -0,0 +1,38 @@
+#==============================================================================
+# cmake-format: off
+# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+# cmake-format: on
+#=============================================================================-
+
+# This function finds rtcx
+function(find_and_configure_rtcx VERSION)
+
+  # Ensure rtcx installs its targets unconditionally. In shared builds the static library is
+  # absorbed into libcuvs.so, but CMake still requires the target to be in an export set for
+  # install(EXPORT) validation. In static builds consumers need to link librtcx.a directly.
+  set(RTCX_INSTALL ON)
+
+  include("${rapids-cmake-dir}/cmake/default_install_component.cmake")
+  rapids_cmake_default_install_component(DEFAULT_USE_PROJECT_NAME)
+
+  rapids_cpm_find(
+    rtcx ${VERSION}
+    GLOBAL_TARGETS rtcx::rtcx
+    BUILD_EXPORT_SET    cuvs-static-exports
+    INSTALL_EXPORT_SET  cuvs-static-exports
+    CPM_ARGS
+    GIT_REPOSITORY https://github.com/rapidsai/librtcx
+    GIT_TAG b5e753e9427e331c02f8c370143cb0d5f40da0fd
+    GIT_SHALLOW FALSE
+  )
+
+  # When CPM fetches from source (add_subdirectory), generate_jit_lto_kernels.cmake is not
+  # auto-included. Include it explicitly so necessary functions are available.
+  if(rtcx_ADDED OR DEFINED CPM_rtcx_SOURCE)
+    include("${rtcx_SOURCE_DIR}/generate_jit_lto_kernels.cmake")
+  endif()
+endfunction()
+
+set(RTCX_MIN_VERSION_cuvs "0.1")
+find_and_configure_rtcx(${RTCX_MIN_VERSION_cuvs})