Always build with JIT+LTO

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

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

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

@@ -30,6 +30,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-129_arch-aarch64.yaml

@@ -29,6 +29,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-129_arch-x86_64.yaml

@@ -31,6 +31,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-129_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-129_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-129_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-129_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

@@ -72,9 +72,7 @@ cache:
       - ninja
       - ${{ stdlib("c") }}
     host:
-      - if: cuda_major == "13"
-        then:
-          - libnvjitlink-dev
+      - libnvjitlink-dev
       - librmm =${{ minor_version }}
       - libraft-headers =${{ minor_version }}
       - nccl ${{ nccl_version }}
@@ -121,9 +119,7 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink-dev
+        - libnvjitlink-dev
       run:
         - ${{ pin_compatible("cuda-version", upper_bound="x", lower_bound="x") }}
         - libraft-headers =${{ minor_version }}
@@ -133,9 +129,7 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink
+        - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -149,9 +143,7 @@ outputs:
           - librmm
           - mkl
           - nccl
-          - if: cuda_major == "13"
-            then:
-              - libnvjitlink
+          - libnvjitlink
     about:
       homepage: ${{ load_from_file("python/libcuvs/pyproject.toml").project.urls.Homepage }}
       license: ${{ load_from_file("python/libcuvs/pyproject.toml").project.license }}
@@ -188,9 +180,7 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink-dev
+        - libnvjitlink-dev
       run:
         - ${{ pin_compatible("cuda-version", upper_bound="x", lower_bound="x") }}
         - ${{ pin_subpackage("libcuvs-headers", exact=True) }}
@@ -201,9 +191,7 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink
+        - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -217,9 +205,7 @@ outputs:
           - librmm
           - mkl
           - nccl
-          - if: cuda_major == "13"
-            then:
-              - libnvjitlink
+          - libnvjitlink
     about:
       homepage: ${{ load_from_file("python/libcuvs/pyproject.toml").project.urls.Homepage }}
       license: ${{ load_from_file("python/libcuvs/pyproject.toml").project.license }}
@@ -254,9 +240,7 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink-dev
+        - libnvjitlink-dev
       run:
         - ${{ pin_compatible("cuda-version", upper_bound="x", lower_bound="x") }}
         - ${{ pin_subpackage("libcuvs-headers", exact=True) }}
@@ -267,9 +251,7 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink
+        - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -280,9 +262,7 @@ outputs:
           - libcurand
           - libcusolver
           - libcusparse
-          - if: cuda_major == "13"
-            then:
-              - libnvjitlink
+          - libnvjitlink
           - librmm
           - mkl
           - nccl
@@ -419,9 +399,7 @@ outputs:
         - libcurand-dev
         - libcusolver-dev
         - libcusparse-dev
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink-dev
+        - libnvjitlink-dev
       run:
         - ${{ pin_subpackage("libcuvs-headers", exact=True) }}
         - ${{ pin_subpackage("libcuvs", exact=True) }}
@@ -431,9 +409,7 @@ outputs:
         - libcurand
         - libcusolver
         - libcusparse
-        - if: cuda_major == "13"
-          then:
-            - libnvjitlink
+        - libnvjitlink
       ignore_run_exports:
         by_name:
           - cuda-cudart
@@ -444,9 +420,7 @@ outputs:
           - libcurand
           - libcusolver
           - libcusparse
-          - if: cuda_major == "13"
-            then:
-              - libnvjitlink
+          - libnvjitlink
           - librmm
           - mkl
           - nccl

cpp/CMakeLists.txt

@@ -355,12 +355,8 @@ if(NOT BUILD_CPU_ONLY)
     )
   endif()
 
-  set(JIT_LTO_TARGET_ARCHITECTURE "")
-  set(JIT_LTO_COMPILATION OFF)
-  if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
-    set(JIT_LTO_TARGET_ARCHITECTURE "75-real")
-    set(JIT_LTO_COMPILATION ON)
-  endif()
+  set(JIT_LTO_TARGET_ARCHITECTURE "75-real")
+  set(JIT_LTO_COMPILATION ON)
 
   if(JIT_LTO_COMPILATION)
     # Generate interleaved scan kernel files at build time
@@ -780,6 +776,7 @@ 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>
+             CUDA::cuda_driver
       PRIVATE
         $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
         $<COMPILE_ONLY:nvidia::cutlass::cutlass>
@@ -842,6 +839,7 @@ 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
+             CUDA::cuda_driver
       PRIVATE
         $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
         $<$<BOOL:${JIT_LTO_COMPILATION}>:CUDA::nvJitLink>

cpp/include/cuvs/detail/jit_lto/AlgorithmLauncher.hpp

@@ -11,13 +11,13 @@
 #include <unordered_map>
 #include <vector_types.h>
 
-#include <cuda_runtime.h>
+#include <cuda.h>
 #include <memory>
 
 struct AlgorithmLauncher {
-  AlgorithmLauncher() : kernel{nullptr}, library{nullptr} {}
+  AlgorithmLauncher() : function{nullptr}, library{nullptr} {}
 
-  AlgorithmLauncher(cudaKernel_t k, cudaLibrary_t lib);
+  AlgorithmLauncher(CUfunction f, CUlibrary lib);
 
   ~AlgorithmLauncher();
 
@@ -28,18 +28,18 @@ struct AlgorithmLauncher {
   AlgorithmLauncher& operator=(AlgorithmLauncher&& other) noexcept;
 
   template <typename... Args>
-  void dispatch(cudaStream_t stream, dim3 grid, dim3 block, std::size_t shared_mem, Args&&... args)
+  void dispatch(CUstream stream, dim3 grid, dim3 block, std::size_t shared_mem, Args&&... args)
   {
     void* kernel_args[] = {const_cast<void*>(static_cast<void const*>(&args))...};
     this->call(stream, grid, block, shared_mem, kernel_args);
   }
 
-  cudaKernel_t get_kernel() { return this->kernel; }
+  CUfunction get_function() { return this->function; }
 
  private:
-  void call(cudaStream_t stream, dim3 grid, dim3 block, std::size_t shared_mem, void** args);
-  cudaKernel_t kernel;
-  cudaLibrary_t library;
+  void call(CUstream stream, dim3 grid, dim3 block, std::size_t shared_mem, void** args);
+  CUfunction function;
+  CUlibrary library;
 };
 
 std::unordered_map<std::string, std::shared_ptr<AlgorithmLauncher>>& get_cached_launchers();

cpp/include/cuvs/detail/jit_lto/cu_try.hpp

@@ -0,0 +1,25 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include <raft/util/cuda_rt_essentials.hpp>
+
+#include <cuda.h>
+
+#define CU_TRY(call)                                                                      \
+  do {                                                                                    \
+    CUresult const status = call;                                                         \
+    if (status != CUDA_SUCCESS) {                                                         \
+      std::string msg{};                                                                  \
+      const char* name;                                                                   \
+      cuGetErrorName(status, &name);                                                      \
+      const char* str;                                                                    \
+      cuGetErrorString(status, &str);                                                     \
+      SET_ERROR_MSG(                                                                      \
+        msg, "CUDA error encountered at: ", "call='%s', Reason=%s:%s", #call, name, str); \
+      throw raft::cuda_error(msg);                                                        \
+    }                                                                                     \
+  } while (0)

cpp/src/detail/jit_lto/AlgorithmLauncher.cpp

@@ -4,52 +4,55 @@
  */
 
 #include <cuvs/detail/jit_lto/AlgorithmLauncher.hpp>
+#include <cuvs/detail/jit_lto/cu_try.hpp>
 
-#include <raft/util/cuda_rt_essentials.hpp>
-
-AlgorithmLauncher::AlgorithmLauncher(cudaKernel_t k, cudaLibrary_t lib) : kernel{k}, library{lib} {}
+AlgorithmLauncher::AlgorithmLauncher(CUfunction f, CUlibrary lib) : function{f}, library{lib} {}
 
 AlgorithmLauncher::~AlgorithmLauncher()
 {
-  if (library != nullptr) { (void)cudaLibraryUnload(library); }
+  if (library != nullptr) { (void)cuLibraryUnload(library); }
 }
 
 AlgorithmLauncher::AlgorithmLauncher(AlgorithmLauncher&& other) noexcept
-  : kernel{other.kernel}, library{other.library}
+  : function{other.function}, library{other.library}
 {
-  other.kernel  = nullptr;
-  other.library = nullptr;
+  other.function = nullptr;
+  other.library  = nullptr;
 }
 
 AlgorithmLauncher& AlgorithmLauncher::operator=(AlgorithmLauncher&& other) noexcept
 {
   if (this != &other) {
     // Unload current library if it exists
-    if (library != nullptr) { cudaLibraryUnload(library); }
-    kernel        = other.kernel;
-    library       = other.library;
-    other.kernel  = nullptr;
-    other.library = nullptr;
+    if (library != nullptr) { cuLibraryUnload(library); }
+    function       = other.function;
+    library        = other.library;
+    other.function = nullptr;
+    other.library  = nullptr;
   }
   return *this;
 }
 
 void AlgorithmLauncher::call(
-  cudaStream_t stream, dim3 grid, dim3 block, std::size_t shared_mem, void** kernel_args)
+  CUstream stream, dim3 grid, dim3 block, std::size_t shared_mem, void** kernel_args)
 {
-  cudaLaunchAttribute attribute[1];
-  attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-  attribute[0].val.programmaticStreamSerializationAllowed = 1;
-
-  cudaLaunchConfig_t config;
-  config.gridDim          = grid;
-  config.blockDim         = block;
-  config.stream           = stream;
-  config.attrs            = attribute;
-  config.numAttrs         = 1;
-  config.dynamicSmemBytes = shared_mem;
-
-  RAFT_CUDA_TRY(cudaLaunchKernelExC(&config, kernel, kernel_args));
+  CUlaunchAttribute attribute[1];
+  attribute[0].id = CU_LAUNCH_ATTRIBUTE_PROGRAMMATIC_STREAM_SERIALIZATION;
+  attribute[0].value.programmaticStreamSerializationAllowed = 1;
+
+  CUlaunchConfig config;
+  config.gridDimX       = grid.x;
+  config.gridDimY       = grid.y;
+  config.gridDimZ       = grid.z;
+  config.blockDimX      = block.x;
+  config.blockDimY      = block.y;
+  config.blockDimZ      = block.z;
+  config.hStream        = stream;
+  config.attrs          = attribute;
+  config.numAttrs       = 1;
+  config.sharedMemBytes = shared_mem;
+
+  CU_TRY(cuLaunchKernelEx(&config, function, kernel_args, nullptr));
 }
 
 std::unordered_map<std::string, std::shared_ptr<AlgorithmLauncher>>& get_cached_launchers()