JIT compile interleaved_scan_kernel for CUDA 13

conda/environments/all_cuda-131_arch-aarch64.yaml

@@ -31,6 +31,7 @@ dependencies:
 - libcurand-dev
 - libcusolver-dev
 - libcusparse-dev
+- libnvjitlink-dev
 - librmm==26.4.*,>=0.0.0a0
 - make
 - nccl>=2.19

conda/environments/all_cuda-131_arch-x86_64.yaml

@@ -31,6 +31,7 @@ dependencies:
 - libcurand-dev
 - libcusolver-dev
 - libcusparse-dev
+- libnvjitlink-dev
 - librmm==26.4.*,>=0.0.0a0
 - make
 - nccl>=2.19

conda/environments/bench_ann_cuda-131_arch-aarch64.yaml

@@ -30,6 +30,7 @@ dependencies:
 - libcusolver-dev
 - libcusparse-dev
 - libcuvs==26.4.*,>=0.0.0a0
+- libnvjitlink-dev
 - librmm==26.4.*,>=0.0.0a0
 - matplotlib-base>=3.9
 - nccl>=2.19

conda/environments/bench_ann_cuda-131_arch-x86_64.yaml

@@ -32,6 +32,7 @@ dependencies:
 - libcusolver-dev
 - libcusparse-dev
 - libcuvs==26.4.*,>=0.0.0a0
+- libnvjitlink-dev
 - librmm==26.4.*,>=0.0.0a0
 - matplotlib-base>=3.9
 - mkl-devel=2023

conda/environments/go_cuda-131_arch-aarch64.yaml

@@ -25,6 +25,7 @@ dependencies:
 - libcusolver-dev
 - libcusparse-dev
 - libcuvs==26.4.*,>=0.0.0a0
+- libnvjitlink-dev
 - libraft==26.4.*,>=0.0.0a0
 - nccl>=2.19
 - ninja

conda/environments/go_cuda-131_arch-x86_64.yaml

@@ -25,6 +25,7 @@ dependencies:
 - libcusolver-dev
 - libcusparse-dev
 - libcuvs==26.4.*,>=0.0.0a0
+- libnvjitlink-dev
 - libraft==26.4.*,>=0.0.0a0
 - nccl>=2.19
 - ninja

conda/environments/rust_cuda-131_arch-aarch64.yaml

@@ -22,6 +22,7 @@ dependencies:
 - libcusolver-dev
 - libcusparse-dev
 - libcuvs==26.4.*,>=0.0.0a0
+- libnvjitlink-dev
 - libraft==26.4.*,>=0.0.0a0
 - make
 - nccl>=2.19

conda/environments/rust_cuda-131_arch-x86_64.yaml

@@ -22,6 +22,7 @@ dependencies:
 - libcusolver-dev
 - libcusparse-dev
 - libcuvs==26.4.*,>=0.0.0a0
+- libnvjitlink-dev
 - libraft==26.4.*,>=0.0.0a0
 - make
 - nccl>=2.19

conda/recipes/libcuvs/recipe.yaml

@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
 # SPDX-License-Identifier: Apache-2.0
 schema_version: 1
 
@@ -72,6 +72,9 @@ cache:
       - ninja
       - ${{ stdlib("c") }}
     host:
+      - if: cuda_major == "13"
+        then:
+          - libnvjitlink-dev
       - librmm =${{ minor_version }}
       - libraft-headers =${{ minor_version }}
       - nccl ${{ nccl_version }}
@@ -118,6 +121,9 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink-dev
       run:
         - ${{ pin_compatible("cuda-version", upper_bound="x", lower_bound="x") }}
         - libraft-headers =${{ minor_version }}
@@ -128,6 +134,9 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -141,6 +150,9 @@ outputs:
           - librmm
           - mkl
           - nccl
+          - if: cuda_major == "13"
+            then:
+              - libnvjitlink
     about:
       homepage: ${{ load_from_file("python/libcuvs/pyproject.toml").project.urls.Homepage }}
       license: ${{ load_from_file("python/libcuvs/pyproject.toml").project.license }}
@@ -177,6 +189,9 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink-dev
       run:
         - ${{ pin_compatible("cuda-version", upper_bound="x", lower_bound="x") }}
         - ${{ pin_subpackage("libcuvs-headers", exact=True) }}
@@ -188,6 +203,9 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -201,6 +219,9 @@ outputs:
           - librmm
           - mkl
           - nccl
+          - if: cuda_major == "13"
+            then:
+              - libnvjitlink
     about:
       homepage: ${{ load_from_file("python/libcuvs/pyproject.toml").project.urls.Homepage }}
       license: ${{ load_from_file("python/libcuvs/pyproject.toml").project.license }}
@@ -235,6 +256,9 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink-dev
       run:
         - ${{ pin_compatible("cuda-version", upper_bound="x", lower_bound="x") }}
         - ${{ pin_subpackage("libcuvs-headers", exact=True) }}
@@ -246,6 +270,9 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -256,6 +283,9 @@ outputs:
           - libcurand
           - libcusolver
           - libcusparse
+          - if: cuda_major == "13"
+            then:
+              - libnvjitlink
           - librmm
           - mkl
           - nccl
@@ -393,6 +423,9 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink-dev
       run:
         - ${{ pin_subpackage("libcuvs-headers", exact=True) }}
         - ${{ pin_subpackage("libcuvs", exact=True) }}
@@ -403,6 +436,9 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
+        - if: cuda_major == "13"
+          then:
+            - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -413,6 +449,9 @@ outputs:
           - libcurand
           - libcusolver
           - libcusparse
+          - if: cuda_major == "13"
+            then:
+              - libnvjitlink
           - librmm
           - mkl
           - nccl

cpp/CMakeLists.txt

@@ -343,6 +343,29 @@ if(NOT BUILD_CPU_ONLY)
     )
   endif()
 
+  set(JIT_LTO_TARGET_ARCHITECTURE "")
+  set(JIT_LTO_COMPILATION OFF)
+  set(JIT_LTO_FILES "")
+  if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
+    set(JIT_LTO_TARGET_ARCHITECTURE "75-real")
+    set(JIT_LTO_COMPILATION ON)
+  endif()
+
+  if(JIT_LTO_COMPILATION)
+    # Generate interleaved scan kernel files at build time
+    include(cmake/modules/generate_jit_lto_kernels.cmake)
+    generate_jit_lto_kernels(cuvs_jit_lto_kernels)
+    add_library(cuvs::cuvs_jit_lto_kernels ALIAS cuvs_jit_lto_kernels)
+
+    set(JIT_LTO_FILES
+        src/detail/jit_lto/AlgorithmLauncher.cu
+        src/detail/jit_lto/AlgorithmPlanner.cu
+        src/detail/jit_lto/FragmentDatabase.cu
+        src/detail/jit_lto/FragmentEntry.cu
+        src/detail/jit_lto/nvjitlink_checker.cpp
+    )
+  endif()
+
   add_library(
     cuvs_objs OBJECT
     src/cluster/detail/minClusterDistanceCompute.cu
@@ -556,6 +579,7 @@ if(NOT BUILD_CPU_ONLY)
     src/stats/silhouette_score.cu
     src/stats/trustworthiness_score.cu
     ${CUVS_MG_ALGOS}
+    $<$<BOOL:${JIT_LTO_COMPILATION}>:${JIT_LTO_FILES}>
   )
 
   set_target_properties(
@@ -572,8 +596,10 @@ if(NOT BUILD_CPU_ONLY)
   )
 
   target_compile_definitions(
-    cuvs_objs PRIVATE $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:CUVS_BUILD_CAGRA_HNSWLIB>
-                      $<$<BOOL:${CUVS_NVTX}>:NVTX_ENABLED>
+    cuvs_objs
+    PRIVATE $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:CUVS_BUILD_CAGRA_HNSWLIB>
+            $<$<BOOL:${CUVS_NVTX}>:NVTX_ENABLED>
+            $<$<BOOL:${JIT_LTO_COMPILATION}>:CUVS_ENABLE_JIT_LTO>
   )
 
   target_link_libraries(
@@ -586,6 +612,13 @@ if(NOT BUILD_CPU_ONLY)
             $<TARGET_NAME_IF_EXISTS:hnswlib::hnswlib>
   )
 
+  target_include_directories(
+    cuvs_objs
+    PUBLIC "$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>"
+           "$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/../c/include>"
+    INTERFACE "$<INSTALL_INTERFACE:include>"
+  )
+
   # Endian detection
   include(TestBigEndian)
   test_big_endian(BIG_ENDIAN)
@@ -640,8 +673,10 @@ if(NOT BUILD_CPU_ONLY)
               "$<$<COMPILE_LANGUAGE:CUDA>:${CUVS_CUDA_FLAGS}>"
     )
     target_compile_definitions(
-      cuvs PUBLIC $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:CUVS_BUILD_CAGRA_HNSWLIB>
-                  $<$<BOOL:${CUVS_NVTX}>:NVTX_ENABLED>
+      cuvs
+      PUBLIC $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:CUVS_BUILD_CAGRA_HNSWLIB>
+             $<$<BOOL:${CUVS_NVTX}>:NVTX_ENABLED>
+             $<$<BOOL:${JIT_LTO_COMPILATION}>:CUVS_ENABLE_JIT_LTO>
     )
 
     target_link_libraries(
@@ -653,8 +688,12 @@ if(NOT BUILD_CPU_ONLY)
              $<BUILD_LOCAL_INTERFACE:$<TARGET_NAME_IF_EXISTS:NCCL::NCCL>>
              $<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>
+      PRIVATE
+        $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
+        $<COMPILE_ONLY:nvidia::cutlass::cutlass>
+        $<COMPILE_ONLY:cuco::cuco>
+        $<$<BOOL:${JIT_LTO_COMPILATION}>:CUDA::nvJitLink>
+        $<$<BOOL:${JIT_LTO_COMPILATION}>:$<LINK_LIBRARY:WHOLE_ARCHIVE,cuvs::cuvs_jit_lto_kernels>>
     )
 
     # ensure CUDA symbols aren't relocated to the middle of the debug build binaries
@@ -692,8 +731,10 @@ SECTIONS
 
     target_compile_options(cuvs_static PRIVATE "$<$<COMPILE_LANGUAGE:CXX>:${CUVS_CXX_FLAGS}>")
     target_compile_definitions(
-      cuvs_static PUBLIC $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:CUVS_BUILD_CAGRA_HNSWLIB>
-                         $<$<BOOL:${CUVS_NVTX}>:NVTX_ENABLED>
+      cuvs_static
+      PUBLIC $<$<BOOL:${BUILD_CAGRA_HNSWLIB}>:CUVS_BUILD_CAGRA_HNSWLIB>
+             $<$<BOOL:${CUVS_NVTX}>:NVTX_ENABLED>
+             $<$<BOOL:${JIT_LTO_COMPILATION}>:CUVS_ENABLE_JIT_LTO>
     )
 
     target_include_directories(cuvs_static INTERFACE "$<INSTALL_INTERFACE:include>")
@@ -709,8 +750,13 @@ SECTIONS
              ${CUVS_CTK_MATH_DEPENDENCIES}
              $<TARGET_NAME_IF_EXISTS:NCCL::NCCL> # needs to be public for DT_NEEDED
              $<BUILD_LOCAL_INTERFACE:$<TARGET_NAME_IF_EXISTS:hnswlib::hnswlib>> # header only
-      PRIVATE $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX> $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3>
-              $<COMPILE_ONLY:nvidia::cutlass::cutlass> $<COMPILE_ONLY:cuco::cuco>
+      PRIVATE
+        $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
+        $<$<BOOL:${JIT_LTO_COMPILATION}>:CUDA::nvJitLink>
+        $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3>
+        $<COMPILE_ONLY:nvidia::cutlass::cutlass>
+        $<COMPILE_ONLY:cuco::cuco>
+        $<$<BOOL:${JIT_LTO_COMPILATION}>:$<LINK_LIBRARY:WHOLE_ARCHIVE,cuvs::cuvs_jit_lto_kernels>>
     )
   endif()
 
@@ -751,9 +797,11 @@ target_compile_definitions(cuvs::cuvs INTERFACE $<$<BOOL:${CUVS_NVTX}>:NVTX_ENAB
   include(GNUInstallDirs)
   include(CPack)
 
-  set(target_names cuvs cuvs_static cuvs_cpp_headers cuvs_c)
-  set(component_names cuvs_shared cuvs_static cuvs_cpp_headers c_api)
-  set(export_names cuvs-shared-exports cuvs-static-exports cuvs-cpp-headers-exports cuvs-c-exports)
+  set(target_names cuvs cuvs_static cuvs_jit_lto_kernels cuvs_cpp_headers cuvs_c)
+  set(component_names cuvs_shared cuvs_static cuvs_static cuvs_cpp_headers c_api)
+  set(export_names cuvs-shared-exports cuvs-static-exports cuvs-static-exports
+                   cuvs-cpp-headers-exports cuvs-c-exports
+  )
   foreach(target component export IN ZIP_LISTS target_names component_names export_names)
     if(TARGET ${target})
       install(
@@ -794,6 +842,8 @@ target_compile_definitions(cuvs::cuvs INTERFACE $<$<BOOL:${CUVS_NVTX}>:NVTX_ENAB
     )
   endif()
 
+  list(REMOVE_DUPLICATES cuvs_components)
+  list(REMOVE_DUPLICATES cuvs_export_sets)
   include(cmake/modules/generate_cuvs_export.cmake)
   generate_cuvs_export(COMPONENTS ${cuvs_components} EXPORT_SETS ${cuvs_export_sets})
 

cpp/cmake/config.json

@@ -10,6 +10,15 @@
           "ADDITIONAL_DEP": "?",
           "PATH": "*"
         }
+      },
+      "embed_jit_lto_fatbin": {
+        "kwargs": {
+          "FATBIN_TARGET": 1,
+          "FATBIN_SOURCE": 1,
+          "EMBEDDED_TARGET": 1,
+          "EMBEDDED_HEADER": 1,
+          "EMBEDDED_ARRAY": 1
+        }
       }
     }
   },