[DO NOT MERGE] Use librtcx for jit-lto

conda/recipes/libcuvs/recipe.yaml

@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: Apache-2.0
 schema_version: 1
 
@@ -166,6 +166,7 @@ outputs:
           cmake --install cpp/build --component cuvs_shared
           cmake --install cpp/build --component c_api
           cmake --install cpp/build --component hnswlib
+          cmake --install cpp/build --component rtcx
     requirements:
       build:
         - cmake ${{ cmake_version }}
@@ -227,6 +228,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 }}

cpp/CMakeLists.txt

@@ -192,6 +192,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)
@@ -353,8 +354,8 @@ if(NOT BUILD_CPU_ONLY)
     set(JIT_LTO_TARGET_ARCHITECTURE "75-real")
   endif()
 
-  # Generate interleaved scan kernel files at build time
-  include(cmake/modules/generate_jit_lto_kernels.cmake)
+  # Generate interleaved scan kernel files at build time add_subdirectory(../../librtcx
+  # src/detail/librtcx)
 
   add_library(jit_lto_kernel_usage_requirements INTERFACE)
   target_include_directories(
@@ -1331,11 +1332,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
@@ -1460,6 +1456,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>
@@ -1540,7 +1537,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
@@ -1597,6 +1594,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-exports
+    INSTALL_EXPORT_SET  cuvs-exports
+    CPM_ARGS
+    GIT_REPOSITORY https://github.com/arhag23/librtcx.git
+    GIT_TAG edba294a388b69cb0f6006669fb7449d84116ac5
+    GIT_SHALLOW FALSE
+  )
+
+  # When CPM fetches from source (add_subdirectory), embed.cmake is not auto-included. Include it
+  # explicitly so add_embed/embed_includes/embed 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})