Support pre-Volta (Maxwell/Pascal) GPU architectures

cpp/CMakeLists.txt

@@ -35,6 +35,52 @@ set(CMAKE_CXX_SCAN_FOR_MODULES OFF)
 # Write the version header
 rapids_cmake_write_version_file(include/cuvs/version_config.h)
 
+# ##################################################################################################
+# * target architecture capabilities -----------------------------------------
+#
+# Determine the lowest/highest SM version this build targets and expose it to the sources through a
+# generated header (rather than a global -D, which would invalidate every object file whenever the
+# architecture list changes). Sources use it to compile out code paths that simply cannot be built
+# for pre-Volta targets, most notably the CUTLASS-backed distance kernels.
+
+set(CUVS_MIN_CUDA_ARCH 0)
+set(CUVS_MAX_CUDA_ARCH 0)
+
+if(NOT BUILD_CPU_ONLY AND CMAKE_CUDA_ARCHITECTURES)
+  set(_cuvs_arch_values)
+  foreach(_cuvs_arch IN LISTS CMAKE_CUDA_ARCHITECTURES)
+    # Entries may look like "80", "80-real", "80-virtual", "90a", "100f"; keep the leading digits.
+    if(_cuvs_arch MATCHES "^([0-9]+)")
+      list(APPEND _cuvs_arch_values "${CMAKE_MATCH_1}")
+    endif()
+  endforeach()
+  if(_cuvs_arch_values)
+    list(SORT _cuvs_arch_values COMPARE NATURAL)
+    list(GET _cuvs_arch_values 0 _cuvs_arch_lo)
+    list(GET _cuvs_arch_values -1 _cuvs_arch_hi)
+    math(EXPR CUVS_MIN_CUDA_ARCH "${_cuvs_arch_lo} * 10")
+    math(EXPR CUVS_MAX_CUDA_ARCH "${_cuvs_arch_hi} * 10")
+  endif()
+  unset(_cuvs_arch_values)
+endif()
+
+# CUTLASS (and the cuVS epilogues layered on it) require sm_70 or newer to even compile. The kernels
+# are only ever dispatched to on sm_80+, so dropping them for older targets is behaviour-preserving.
+if(CUVS_MIN_CUDA_ARCH EQUAL 0 OR CUVS_MIN_CUDA_ARCH GREATER_EQUAL 700)
+  set(CUVS_CUTLASS_ENABLED 1)
+else()
+  set(CUVS_CUTLASS_ENABLED 0)
+  message(
+    STATUS
+      "cuVS: targeting SM ${CUVS_MIN_CUDA_ARCH}; CUTLASS distance kernels disabled (need sm_70+), falling back to SIMT/JIT kernels"
+  )
+endif()
+
+configure_file(
+  "${CMAKE_CURRENT_SOURCE_DIR}/cmake/arch_config.hpp.in"
+  "${CMAKE_CURRENT_BINARY_DIR}/include/cuvs/detail/arch_config.hpp" @ONLY
+)
+
 # ##################################################################################################
 # * build type ---------------------------------------------------------------
 
@@ -347,18 +393,37 @@ if(NOT BUILD_CPU_ONLY)
     )
   endif()
 
+  # The JIT-LTO fragments are compiled once, to LTO-IR, and retargeted to the running device by
+  # nvJitLink (see src/detail/jit_lto/AlgorithmPlanner.cpp). LTO-IR is forward compatible, so a
+  # single baseline architecture covers every newer device -- but it can *not* be linked down to an
+  # older one. The baseline therefore has to be no higher than the oldest architecture we build for.
   set(JIT_LTO_TARGET_ARCHITECTURE "70-real")
   if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
     set(JIT_LTO_TARGET_ARCHITECTURE "75-real")
   endif()
 
+  string(REGEX MATCH "^[0-9]+" _jit_lto_baseline "${JIT_LTO_TARGET_ARCHITECTURE}")
+  math(EXPR _jit_lto_baseline "${_jit_lto_baseline} * 10")
+  if(CUVS_MIN_CUDA_ARCH GREATER 0 AND CUVS_MIN_CUDA_ARCH LESS ${_jit_lto_baseline})
+    math(EXPR _jit_lto_arch "${CUVS_MIN_CUDA_ARCH} / 10")
+    set(JIT_LTO_TARGET_ARCHITECTURE "${_jit_lto_arch}-real")
+    message(
+      STATUS
+        "cuVS: lowering JIT-LTO baseline architecture to ${JIT_LTO_TARGET_ARCHITECTURE} to match CMAKE_CUDA_ARCHITECTURES"
+    )
+    unset(_jit_lto_arch)
+  endif()
+  unset(_jit_lto_baseline)
+
   # Generate interleaved scan kernel files at build time
 
   add_library(jit_lto_kernel_usage_requirements INTERFACE)
   target_include_directories(
     jit_lto_kernel_usage_requirements
     INTERFACE "${CMAKE_CURRENT_SOURCE_DIR}/include" "${CMAKE_CURRENT_SOURCE_DIR}/src"
               "${CMAKE_CURRENT_SOURCE_DIR}/../c/include"
+              # so the offline-compiled kernels can see the generated cuvs/detail/arch_config.hpp
+              "${CMAKE_CURRENT_BINARY_DIR}/include"
   )
   target_compile_options(
     jit_lto_kernel_usage_requirements INTERFACE "$<$<COMPILE_LANGUAGE:CXX>:${CUVS_CXX_FLAGS}>"
@@ -1532,6 +1597,19 @@ if(NOT BUILD_CPU_ONLY)
                                  ${CUVS_CUSPARSE_DEPENDENCY} ${CUVS_CURAND_DEPENDENCY}
   )
 
+  # Before CUDA 12.8 the runtime rejects a cudaKernel_t where a kernel entry point is expected, and
+  # cudaKernelSetAttributeForDevice does not exist at all, so the JIT-LTO kernel handles are queried
+  # and configured through the equivalent driver-API entry points instead (see
+  # src/util/jit_kernel_compat.hpp). That needs the driver stub at link time.
+  set(CUVS_CUDA_DRIVER_DEPENDENCY "")
+  if(CUDAToolkit_VERSION VERSION_LESS 12.8)
+    set(CUVS_CUDA_DRIVER_DEPENDENCY CUDA::cuda_driver)
+    message(
+      STATUS
+        "cuVS: CUDA ${CUDAToolkit_VERSION} does not accept kernel handles in the runtime API; using the driver API for JIT-LTO kernel attributes"
+    )
+  endif()
+
   if(NOT cuvs_compile_mode STREQUAL "static_only")
     add_library(cuvs SHARED $<TARGET_OBJECTS:cuvs_objs>)
     add_library(cuvs::cuvs ALIAS cuvs)
@@ -1571,6 +1649,7 @@ if(NOT BUILD_CPU_ONLY)
              $<$<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 rtcx::rtcx
+              ${CUVS_CUDA_DRIVER_DEPENDENCY}
     )
 
     # ensure CUDA symbols aren't relocated to the middle of the debug build binaries
@@ -1629,6 +1708,7 @@ SECTIONS
               CUDA::nvJitLink
               CUDA::nvrtc
               rtcx::rtcx
+              ${CUVS_CUDA_DRIVER_DEPENDENCY}
               $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3>
               $<COMPILE_ONLY:nvidia::cutlass::cutlass>
               $<COMPILE_ONLY:cuco::cuco>

cpp/cmake/arch_config.hpp.in

@@ -0,0 +1,59 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * AUTO-GENERATED by CMake from cpp/cmake/arch_config.hpp.in -- DO NOT EDIT.
+ *
+ * Exposes the range of CUDA compute capabilities that this build of cuVS is
+ * being compiled for. Unlike `__CUDA_ARCH__`, these macros are visible in both
+ * the host and the device compilation passes, which makes them usable to gate
+ * declarations, includes and template instantiations consistently.
+ */
+
+#pragma once
+
+/** Lowest SM version being compiled for, in `__CUDA_ARCH__` form (e.g. 500). */
+#define CUVS_MIN_CUDA_ARCH @CUVS_MIN_CUDA_ARCH@
+
+/** Highest SM version being compiled for, in `__CUDA_ARCH__` form (e.g. 900). */
+#define CUVS_MAX_CUDA_ARCH @CUVS_MAX_CUDA_ARCH@
+
+/**
+ * @brief Whether the CUTLASS-backed distance kernels can be compiled.
+ *
+ * CUTLASS itself, and the cuVS epilogues built on top of it, depend on
+ * facilities that are only available on Volta and newer:
+ *   - `cuda::binary_semaphore` (libcu++ hard-errors below sm_70),
+ *   - `cooperative_groups::binary_partition` (needs `__match_any_sync`).
+ *
+ * Those kernels are only ever *launched* on sm_80+ (see the runtime dispatch in
+ * `fused_l2_nn.cuh`, `fused_cosine_nn.cuh` and `pairwise_matrix/dispatch-inl.cuh`),
+ * so when the build targets an architecture older than sm_70 the whole CUTLASS
+ * path is dead code and is compiled out. The pre-existing SIMT / JIT fallbacks
+ * are used instead.
+ */
+#define CUVS_CUTLASS_ENABLED @CUVS_CUTLASS_ENABLED@
+
+/**
+ * @brief Whether native `__half` / `__half2` arithmetic is available.
+ *
+ * Half-precision arithmetic intrinsics require sm_53. On older devices the
+ * values have to be converted to `float` before doing any math.
+ */
+#define CUVS_NATIVE_HALF_ARITHMETIC (CUVS_MIN_CUDA_ARCH == 0 || CUVS_MIN_CUDA_ARCH >= 530)
+
+/**
+ * @brief Whether block-scoped atomics (`atomicAdd_block` & friends) are available.
+ *
+ * Scoped atomics require sm_60.
+ */
+#define CUVS_HAS_BLOCK_SCOPED_ATOMICS (CUVS_MIN_CUDA_ARCH == 0 || CUVS_MIN_CUDA_ARCH >= 600)
+
+/**
+ * @brief Whether Volta+ warp-synchronous primitives are available.
+ *
+ * `__match_any_sync`, `__match_all_sync`, `cuda::std::semaphore`,
+ * `cooperative_groups::binary_partition` and independent thread scheduling all
+ * require sm_70.
+ */
+#define CUVS_HAS_VOLTA_WARP_PRIMITIVES (CUVS_MIN_CUDA_ARCH == 0 || CUVS_MIN_CUDA_ARCH >= 700)

cpp/src/distance/detail/distance.cuh

@@ -5,6 +5,8 @@
 
 #pragma once
 
+#include <util/half_compat.cuh>  // cuvs::util::sqrt_op
+
 #include "distance_ops/all_ops.cuh"
 #include "pairwise_matrix/dispatch.cuh"
 #include <cuvs/distance/distance.hpp>
@@ -357,12 +359,12 @@ void distance_impl(raft::resources const& handle,
     raft::linalg::map(
       handle,
       raft::make_device_vector_view<DataT, IdxT>((DataT*)x, m * k),
-      raft::sqrt_op{},
+      cuvs::util::sqrt_op{},
       raft::make_const_mdspan(raft::make_device_vector_view<const DataT, IdxT>(x, m * k)));
     raft::linalg::map(
       handle,
       raft::make_device_vector_view<DataT, IdxT>((DataT*)y, n * k),
-      raft::sqrt_op{},
+      cuvs::util::sqrt_op{},
       raft::make_const_mdspan(raft::make_device_vector_view<const DataT, IdxT>(y, n * k)));
   } else {
     // Arrays overlap: sqrt the union of both arrays exactly once
@@ -373,7 +375,7 @@ void distance_impl(raft::resources const& handle,
     raft::linalg::map(
       handle,
       raft::make_device_vector_view<DataT, IdxT>((DataT*)start, union_size),
-      raft::sqrt_op{},
+      cuvs::util::sqrt_op{},
       raft::make_const_mdspan(raft::make_device_vector_view<const DataT, IdxT>(start, union_size)));
   }
 

cpp/src/distance/detail/distance_ops/l_inf.cuh

@@ -39,8 +39,13 @@ struct l_inf_distance_op {
 
   DI void core(AccT& acc, DataT& x, DataT& y) const
   {
-    const auto diff = raft::abs(x - y);
-    acc             = raft::max(acc, diff);
+    // Widen to the accumulator type before doing any arithmetic. `__hsub` / `__habs` are FP16 ALU
+    // instructions that need sm_53, and `raft::abs(__half)` static_asserts below that. The other
+    // element-wise ops (l1, lp_unexp, jensen_shannon, kl_divergence) already do this; the
+    // accumulation happens in `AccT` regardless, so the only change is that the subtraction is
+    // now performed at accumulator precision.
+    const auto diff = raft::abs(raft::to_float(x) - raft::to_float(y));
+    acc             = raft::max(acc, static_cast<AccT>(diff));
   };
 
   template <typename Policy>

cpp/src/distance/detail/fused_distance_nn/cutlass_base.cuh

@@ -5,6 +5,15 @@
 
 #pragma once
 
+#include <cuvs/detail/arch_config.hpp>  // CUVS_CUTLASS_ENABLED
+
+// The CUTLASS-backed fused distance-NN kernel is only ever dispatched to on sm_80+, but it cannot
+// even be *compiled* for pre-Volta targets: `cuda::binary_semaphore` hard-errors below sm_70 and
+// the reduced-vector output iterator relies on `cg::binary_partition`. When the build targets an
+// older architecture the whole file collapses to nothing and the callers fall back to the SIMT
+// kernel in `simt_kernel.cuh` (which is itself only compiled for `__CUDA_ARCH__ < 800`).
+#if CUVS_CUTLASS_ENABLED
+
 #pragma GCC diagnostic push
 #pragma GCC diagnostic ignored "-Wstrict-aliasing"
 #pragma GCC diagnostic ignored "-Wtautological-compare"
@@ -164,3 +173,5 @@ void cutlassFusedDistanceNN(const DataT* x,
 };  // namespace cuvs
 
 #pragma GCC diagnostic pop
+
+#endif  // CUVS_CUTLASS_ENABLED

cpp/src/distance/detail/fused_distance_nn/fused_cosine_nn.cuh

@@ -10,6 +10,7 @@
 #include "cutlass_base.cuh"
 #include "helper_structs.cuh"
 #include "simt_kernel.cuh"
+#include <cuvs/detail/arch_config.hpp>   // CUVS_CUTLASS_ENABLED
 #include <raft/core/kvp.hpp>             // raft::KeyValuePair
 #include <raft/core/operators.hpp>       // raft::identity_op
 #include <raft/linalg/contractions.cuh>  // Policy
@@ -66,6 +67,7 @@ void fusedCosineNN(OutT* min,
                                       decltype(distance_op),
                                       decltype(fin_op)>;
 
+#if CUVS_CUTLASS_ENABLED
   // Get pointer to fp32 SIMT kernel to determine the runtime architecture of the
   // current system. Other methods to determine the architecture (that do not
   // require a pointer) can be error prone. See:
@@ -109,7 +111,11 @@ void fusedCosineNN(OutT* min,
                                          redOp,
                                          pairRedOp,
                                          stream);
-  } else {
+    return;
+  }
+#endif
+
+  {
     // If device less than SM_80, use fp32 SIMT kernel.
     constexpr size_t shmemSize = P::SmemSize + ((P::Mblk + P::Nblk) * sizeof(DataT));
     dim3 grid                  = launchConfigGenerator<P>(m, n, shmemSize, kernel);

cpp/src/distance/detail/fused_distance_nn/fused_l2_nn.cuh

@@ -10,6 +10,7 @@
 #include "cutlass_base.cuh"
 #include "helper_structs.cuh"
 #include "simt_kernel.cuh"
+#include <cuvs/detail/arch_config.hpp>   // CUVS_CUTLASS_ENABLED
 #include <raft/core/kvp.hpp>             // raft::KeyValuePair
 #include <raft/core/operators.hpp>       // raft::identity_op
 #include <raft/linalg/contractions.cuh>  // Policy
@@ -74,6 +75,7 @@ void fusedL2NNImpl(OutT* min,
                                       decltype(distance_op),
                                       decltype(fin_op)>;
 
+#if CUVS_CUTLASS_ENABLED
   // Get pointer to fp32 SIMT kernel to determine the best compute architecture
   // out of all for which the kernel was compiled for that matches closely
   // to the current device. Other methods to determine the architecture (that do not
@@ -118,7 +120,11 @@ void fusedL2NNImpl(OutT* min,
                                          redOp,
                                          pairRedOp,
                                          stream);
-  } else {
+    return;
+  }
+#endif
+
+  {
     // If device less than SM_80, use fp32 SIMT kernel.
     constexpr size_t shmemSize = P::SmemSize + ((P::Mblk + P::Nblk) * sizeof(DataT));
     dim3 grid                  = launchConfigGenerator<P>(m, n, shmemSize, kernel);

cpp/src/distance/detail/kernels/kernel_matrices.cu

@@ -3,6 +3,10 @@
  * SPDX-License-Identifier: Apache-2.0
  */
 
+// Must precede any RAFT header: supplies atomicAdd(double*) / atomicAdd_block for
+// device passes below sm_60, where CUDA does not declare them.
+#include <util/atomic_compat.cuh>
+
 #include "../../../distance/distance.cuh"
 #include <cuvs/distance/grammian.hpp>
 

cpp/src/distance/detail/pairwise_distance_base.cuh

@@ -6,11 +6,13 @@
 #ifdef CUVS_DISTANCE_PAIRWISE_USE_JIT
 #include "pairwise_matrix/jit_lto_kernels/device_functions.cuh"
 #endif
+#include <util/jit_kernel_compat.hpp>  // kernel_max_active_blocks_per_multiprocessor
 #include <raft/linalg/contractions.cuh>       // raft::linalg::Contractions_NT
 #include <raft/util/cuda_dev_essentials.cuh>  // ceildiv
 #include <raft/util/cuda_rt_essentials.hpp>   // RAFT_CUDA_TRY
 
-#include <cstddef>  // size_t
+#include <cstddef>      // size_t
+#include <type_traits>  // std::is_same_v
 
 namespace cuvs {
 namespace distance {
@@ -310,8 +312,14 @@ dim3 launchConfigGenerator(IdxT m, IdxT n, std::size_t sMemSize, T func)
   int numBlocksPerSm = 0;
   dim3 grid;
 
-  RAFT_CUDA_TRY(
-    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm, func, P::Nthreads, sMemSize));
+  if constexpr (std::is_same_v<T, cudaKernel_t>) {
+    // A JIT-LTO kernel handle: the runtime occupancy API only accepts one on CUDA 12.8+.
+    RAFT_CUDA_TRY(cuvs::util::kernel_max_active_blocks_per_multiprocessor(
+      &numBlocksPerSm, func, P::Nthreads, sMemSize));
+  } else {
+    RAFT_CUDA_TRY(
+      cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm, func, P::Nthreads, sMemSize));
+  }
   std::size_t minGridSize = numSMs * numBlocksPerSm;
   std::size_t yChunks     = raft::ceildiv<int>(m, P::Mblk);
   std::size_t xChunks     = raft::ceildiv<int>(n, P::Nblk);

cpp/src/distance/detail/pairwise_distance_cutlass_base.cuh

@@ -5,6 +5,12 @@
 
 #pragma once
 
+#include <cuvs/detail/arch_config.hpp>  // CUVS_CUTLASS_ENABLED
+
+// See the note in fused_distance_nn/cutlass_base.cuh: CUTLASS requires sm_70+ to compile, and is
+// only ever dispatched to on sm_80+. Pre-Volta builds use the JIT pairwise-matrix kernels instead.
+#if CUVS_CUTLASS_ENABLED
+
 #pragma GCC diagnostic push
 #pragma GCC diagnostic ignored "-Wstrict-aliasing"
 #pragma GCC diagnostic ignored "-Wtautological-compare"
@@ -172,3 +178,5 @@ std::enable_if_t<ops::has_cutlass_op<OpT>::value> cutlassDistanceKernel(const Da
 };  // namespace cuvs
 
 #pragma GCC diagnostic pop
+
+#endif  // CUVS_CUTLASS_ENABLED