[DO NOT MERGE] Use librtcx for jit-lto

cpp/CMakeLists.txt

@@ -353,8 +353,9 @@ 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)
+  include(cmake/thirdparty/get_rtcx.cmake)
 
   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,34 @@
+#==============================================================================
+# 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)
+
+  rapids_cpm_find(
+    rtcx ${VERSION}
+    GLOBAL_TARGETS rtcx::rtcx
+    CPM_ARGS
+    GIT_REPOSITORY https://github.com/arhag23/librtcx.git
+    GIT_TAG 75f90a8503766ce37d62f4418449315874bf2d87
+    GIT_SHALLOW FALSE
+    EXCLUDE_FROM_ALL TRUE
+  )
+
+  # 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})