prefer CUDA 13.1 devcontainers, react to some cutlass removals in RAFT

.github/workflows/pr.yaml

@@ -334,7 +334,7 @@ jobs:
     uses: rapidsai/shared-workflows/.github/workflows/[email protected]
     with:
       arch: '["amd64", "arm64"]'
-      cuda: '["13.0"]'
+      cuda: '["13.1"]'
       node_type: "cpu8"
       rapids-aux-secret-1: GIST_REPO_READ_ORG_GITHUB_TOKEN
       env: |

cpp/CMakeLists.txt

@@ -192,6 +192,8 @@ if(NOT BUILD_CPU_ONLY)
   rapids_cpm_cccl(BUILD_EXPORT_SET cuvs-exports INSTALL_EXPORT_SET cuvs-exports)
   include(cmake/thirdparty/get_raft.cmake)
   include(cmake/thirdparty/get_cutlass.cmake)
+  include(${rapids-cmake-dir}/cpm/cuco.cmake)
+  rapids_cpm_cuco(BUILD_EXPORT_SET cuvs-exports INSTALL_EXPORT_SET cuvs-exports)
 endif()
 
 if(BUILD_TESTS OR BUILD_C_TESTS)
@@ -220,7 +222,7 @@ if(NOT BUILD_CPU_ONLY)
               "$<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}/include>"
               "$<INSTALL_INTERFACE:include>"
   )
-  target_link_libraries(cuvs_cpp_headers INTERFACE raft::raft rmm::rmm)
+  target_link_libraries(cuvs_cpp_headers INTERFACE nvidia::cutlass::cutlass raft::raft rmm::rmm)
 
   add_library(
     cuvs-cagra-search OBJECT
@@ -566,8 +568,11 @@ if(NOT BUILD_CPU_ONLY)
 
   target_link_libraries(
     cuvs_objs
-    PRIVATE cuvs::cuvs_cpp_headers ${CUVS_CTK_MATH_DEPENDENCIES}
-            $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX> $<TARGET_NAME_IF_EXISTS:hnswlib::hnswlib>
+    PRIVATE cuvs::cuvs_cpp_headers
+            ${CUVS_CTK_MATH_DEPENDENCIES}
+            $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
+            $<TARGET_NAME_IF_EXISTS:hnswlib::hnswlib>
+            $<TARGET_NAME_IF_EXISTS:cuco::cuco>
   )
 
   # Endian detection
@@ -638,6 +643,7 @@ if(NOT BUILD_CPU_ONLY)
              $<BUILD_LOCAL_INTERFACE:$<TARGET_NAME_IF_EXISTS:hnswlib::hnswlib>>
              $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3>
       PRIVATE nvidia::cutlass::cutlass $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
+              $<TARGET_NAME_IF_EXISTS:cuco::cuco>
     )
 
     # ensure CUDA symbols aren't relocated to the middle of the debug build binaries
@@ -693,7 +699,7 @@ SECTIONS
              $<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 nvidia::cutlass::cutlass $<TARGET_NAME_IF_EXISTS:OpenMP::OpenMP_CXX>
-              $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3>
+              $<$<BOOL:${CUVS_NVTX}>:CUDA::nvtx3> $<TARGET_NAME_IF_EXISTS:cuco::cuco>
     )
   endif()
 

cpp/cmake/patches/cutlass/Support-both-CUDA-12-and-13-cccl-header-locations.patch

@@ -0,0 +1,31 @@
+From 661c7e679ac72926d619da46834d09f52a727f5e Mon Sep 17 00:00:00 2001
+From: Robert Maynard <[email protected]>
+Date: Tue, 5 Aug 2025 15:05:57 -0400
+Subject: [PATCH] Support both CUDA 12 and 13 cccl header locations
+
+---
+ CMakeLists.txt | 8 +++++++-
+ 1 file changed, 7 insertions(+), 1 deletion(-)
+
+diff --git a/CMakeLists.txt b/CMakeLists.txt
+index 38dcca9f..4088b71f 100755
+--- a/CMakeLists.txt
++++ b/CMakeLists.txt
+@@ -704,8 +704,14 @@ target_include_directories(
+   CUTLASS
+   SYSTEM INTERFACE
+   $<BUILD_INTERFACE:${CUDA_TOOLKIT_ROOT_DIR}/include>
+-  $<BUILD_INTERFACE:${CUDA_TOOLKIT_ROOT_DIR}/include/cccl>
+   )
++if(CUDA_VERSION VERSION_GREATER_EQUAL 13.0)
++  target_include_directories(
++    CUTLASS
++    SYSTEM INTERFACE
++    $<BUILD_INTERFACE:${CUDA_TOOLKIT_ROOT_DIR}/include/cccl>
++    )
++endif()
+
+ install(
+   DIRECTORY
+--
+2.39.5 (Apple Git-154)

cpp/cmake/patches/cutlass_override.json

@@ -1,13 +1,13 @@
 {
   "packages" : {
     "cutlass" : {
-      "version": "3.5.1",
+      "version": "4.1.0",
       "git_url": "https://github.com/NVIDIA/cutlass.git",
       "git_tag": "v${version}",
       "patches" : [
         {
-          "file" : "${current_json_dir}/cutlass/build-export.patch",
-          "issue" : "Fix build directory export",
+          "file" : "${current_json_dir}/cutlass/Support-both-CUDA-12-and-13-cccl-header-locations.patch",
+          "issue" : "Support CUDA 12 CTK layout[https://github.com/NVIDIA/cutlass/pull/2543]",
           "fixed_in" : ""
         }
       ]

cpp/cmake/thirdparty/get_cutlass.cmake

@@ -1,6 +1,6 @@
 # =============================================================================
 # cmake-format: off
-# SPDX-FileCopyrightText: Copyright (c) 2021-2025, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
 # SPDX-License-Identifier: Apache-2.0
 # cmake-format: on
 # =============================================================================
@@ -24,9 +24,7 @@ function(find_and_configure_cutlass)
       CACHE BOOL "Disable CUTLASS to build with cuBLAS library."
   )
 
-  if (CUDA_STATIC_RUNTIME)
-    set(CUDART_LIBRARY "${CUDA_cudart_static_LIBRARY}" CACHE FILEPATH "fixing cutlass cmake code" FORCE)
-  endif()
+  set(CUDART_LIBRARY "${CUDA_cudart_static_LIBRARY}" CACHE FILEPATH "fixing cutlass cmake code" FORCE)
 
   include("${rapids-cmake-dir}/cpm/package_override.cmake")
   rapids_cpm_package_override("${CMAKE_CURRENT_FUNCTION_LIST_DIR}/../patches/cutlass_override.json")
@@ -49,6 +47,34 @@ function(find_and_configure_cutlass)
     add_library(nvidia::cutlass::cutlass ALIAS CUTLASS)
   endif()
 
+  # export cutlass so projects that depend on cuVS can find it
+  if(NvidiaCutlass_ADDED)
+    rapids_export(
+      BUILD NvidiaCutlass
+      EXPORT_SET NvidiaCutlass
+      GLOBAL_TARGETS nvidia::cutlass::cutlass
+      NAMESPACE nvidia::cutlass::
+    )
+  endif()
+
+  rapids_export_package(
+      BUILD NvidiaCutlass cuvs-exports GLOBAL_TARGETS nvidia::cutlass::cutlass
+  )
+  rapids_export_package(
+      INSTALL NvidiaCutlass cuvs-exports GLOBAL_TARGETS nvidia::cutlass::cutlass
+  )
+
+  # Tell cmake where it can find the generated NvidiaCutlass-config.cmake we wrote.
+  include("${rapids-cmake-dir}/export/find_package_root.cmake")
+  rapids_export_find_package_root(
+          INSTALL NvidiaCutlass [=[${CMAKE_CURRENT_LIST_DIR}/../]=]
+          EXPORT_SET cuvs-exports
+  )
+  rapids_export_find_package_root(
+          BUILD NvidiaCutlass [=[${CMAKE_CURRENT_LIST_DIR}]=]
+          EXPORT_SET cuvs-exports
+  )
+
 endfunction()
 
 find_and_configure_cutlass()

cpp/include/cuvs/util/cutlass_utils.hpp

@@ -0,0 +1,43 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include <raft/core/error.hpp>
+
+#include <cutlass/cutlass.h>
+
+namespace cuvs {
+
+/**
+ * @brief Exception thrown when a CUTLASS error is encountered.
+ */
+struct cutlass_error : public raft::exception {
+  explicit cutlass_error(char const* const message) : raft::exception(message) {}
+  explicit cutlass_error(std::string const& message) : raft::exception(message) {}
+};
+
+}  // namespace cuvs
+
+/**
+ * @brief Error checking macro for CUTLASS functions.
+ *
+ * Invokes a CUTLASS function call, if the call does not return cutlass::Status::kSuccess,
+ * throws an exception detailing the CUTLASS error that occurred.
+ *
+ */
+#define CUVS_CUTLASS_TRY(call)                        \
+  do {                                                \
+    cutlass::Status const status = call;              \
+    if (status != cutlass::Status::kSuccess) {        \
+      std::string msg{};                              \
+      SET_ERROR_MSG(msg,                              \
+                    "CUTLASS error encountered at: ", \
+                    "call='%s', Reason=%s",           \
+                    #call,                            \
+                    cutlassGetStatusString(status));  \
+      throw cuvs::cutlass_error(msg);                 \
+    }                                                 \
+  } while (0)

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

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -17,8 +17,8 @@
 
 #include "epilogue_elementwise.cuh"     // FusedDistanceNNEpilogueElementwise
 #include "gemm.h"                       // FusedDistanceNNGemm
+#include <cuvs/util/cutlass_utils.hpp>  // CUVS_CUTLASS_TRY
 #include <raft/util/cudart_utils.hpp>   // getMultiProcessorCount
-#include <raft/util/cutlass_utils.cuh>  // RAFT_CUTLASS_TRY
 
 #include <rmm/device_uvector.hpp>
 
@@ -152,11 +152,11 @@ void cutlassFusedDistanceNN(const DataT* x,
   // Instantiate CUTLASS kernel depending on templates
   fusedDistanceNN fusedDistanceNN_op;
   // Check the problem size is supported or not
-  RAFT_CUTLASS_TRY(fusedDistanceNN_op.can_implement(arguments));
+  CUVS_CUTLASS_TRY(fusedDistanceNN_op.can_implement(arguments));
   // Initialize CUTLASS kernel with arguments and workspace pointer
-  RAFT_CUTLASS_TRY(fusedDistanceNN_op.initialize(arguments, workspace.data(), stream));
+  CUVS_CUTLASS_TRY(fusedDistanceNN_op.initialize(arguments, workspace.data(), stream));
   // Launch initialized CUTLASS kernel
-  RAFT_CUTLASS_TRY(fusedDistanceNN_op.run(stream));
+  CUVS_CUTLASS_TRY(fusedDistanceNN_op.run(stream));
 }
 
 };  // namespace detail

cpp/src/distance/detail/pairwise_distance_cutlass_base.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2018-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -19,7 +19,7 @@
 #include "pairwise_distance_gemm.h"
 
 #include "distance_ops/cutlass.cuh"
-#include <raft/util/cutlass_utils.cuh>
+#include <cuvs/util/cutlass_utils.hpp>
 
 #include <rmm/device_uvector.hpp>
 
@@ -157,13 +157,13 @@ std::enable_if_t<ops::has_cutlass_op<OpT>::value> cutlassDistanceKernel(const Da
     // Instantiate CUTLASS kernel depending on templates
     cutlassDist cutlassDist_op;
     // Check the problem size is supported or not
-    RAFT_CUTLASS_TRY(cutlassDist_op.can_implement(arguments));
+    CUVS_CUTLASS_TRY(cutlassDist_op.can_implement(arguments));
 
     // Initialize CUTLASS kernel with arguments and workspace pointer
-    RAFT_CUTLASS_TRY(cutlassDist_op.initialize(arguments, workspace.data(), stream));
+    CUVS_CUTLASS_TRY(cutlassDist_op.initialize(arguments, workspace.data(), stream));
 
     // Launch initialized CUTLASS kernel
-    RAFT_CUTLASS_TRY(cutlassDist_op(stream));
+    CUVS_CUTLASS_TRY(cutlassDist_op(stream));
   }
 }