Upgrade to latest cutlass version (#2503)

This upgrade is important because cutlass 3.4 fixed issues with kernel visibility that would otherwise lead to global kernel symbols from raft kernels using cutlass.

Authors:
  - Vyas Ramasubramani (https://github.com/vyasr)
  - Kyle Edwards (https://github.com/KyleFromNVIDIA)

Approvers:
  - Robert Maynard (https://github.com/robertmaynard)
  - Corey J. Nolet (https://github.com/cjnolet)
  - Kyle Edwards (https://github.com/KyleFromNVIDIA)

URL: #2503

.github/workflows/pr.yaml

@@ -80,7 +80,6 @@ jobs:
     with:
       build_type: pull-request
       enable_check_symbols: true
-      symbol_exclusions: raft_cutlass
   conda-python-build:
     needs: conda-cpp-build
     secrets: inherit

.github/workflows/test.yaml

@@ -23,7 +23,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       enable_check_symbols: true
-      symbol_exclusions: raft_cutlass
   conda-cpp-tests:
     secrets: inherit
     uses: rapidsai/shared-workflows/.github/workflows/[email protected]

.pre-commit-config.yaml

@@ -93,7 +93,10 @@ repos:
               - id: codespell
                 additional_dependencies: [tomli]
                 args: ["--toml", "pyproject.toml"]
-                exclude: (?x)^(^CHANGELOG.md$)
+                exclude: |
+                  (?x)
+                    ^CHANGELOG[.]md$|
+                    ^cpp/cmake/patches/cutlass/build-export[.]patch$
       - repo: https://github.com/pre-commit/pre-commit-hooks
         rev: v4.5.0
         hooks:

cpp/cmake/patches/cutlass/build-export.patch

@@ -0,0 +1,27 @@
+From e0a9597946257a01ae8444200f836ee51d5597ba Mon Sep 17 00:00:00 2001
+From: Kyle Edwards <[email protected]>
+Date: Wed, 20 Nov 2024 16:37:38 -0500
+Subject: [PATCH] Remove erroneous include directories
+
+These directories are left over from when CuTe was a separate
+CMake project. Remove them.
+---
+ CMakeLists.txt | 2 --
+ 1 file changed, 2 deletions(-)
+
+diff --git a/CMakeLists.txt b/CMakeLists.txt
+index 7419bdf5e..545384d82 100755
+--- a/CMakeLists.txt
++++ b/CMakeLists.txt
+@@ -665,8 +665,6 @@ target_include_directories(
+   $<INSTALL_INTERFACE:include>
+   $<BUILD_INTERFACE:${CUTLASS_INCLUDE_DIR}>
+   $<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}/include>
+-  $<BUILD_INTERFACE:${cute_SOURCE_DIR}/include>
+-  $<BUILD_INTERFACE:${cute_SOURCE_DIR}/examples>
+   )
+
+ # Mark CTK headers as system to supress warnings from them
+-- 
+2.34.1
+

cpp/cmake/patches/cutlass_override.json

@@ -0,0 +1,16 @@
+{
+  "packages" : {
+    "cutlass" : {
+      "version": "3.5.1",
+      "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",
+          "fixed_in" : ""
+        }
+      ]
+    }
+  }
+}

cpp/cmake/thirdparty/get_cutlass.cmake

@@ -13,7 +13,9 @@
 # =============================================================================
 
 function(find_and_configure_cutlass)
-  set(oneValueArgs VERSION REPOSITORY PINNED_TAG)
+  set(options)
+  set(oneValueArgs)
+  set(multiValueArgs)
   cmake_parse_arguments(PKG "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
 
   # if(RAFT_ENABLE_DIST_DEPENDENCIES OR RAFT_COMPILE_LIBRARIES)
@@ -34,13 +36,22 @@ function(find_and_configure_cutlass)
     set(CUDART_LIBRARY "${CUDA_cudart_static_LIBRARY}" CACHE FILEPATH "fixing cutlass cmake code" FORCE)
   endif()
 
+  include("${rapids-cmake-dir}/cpm/package_override.cmake")
+  rapids_cpm_package_override("${CMAKE_CURRENT_FUNCTION_LIST_DIR}/../patches/cutlass_override.json")
+
+  include("${rapids-cmake-dir}/cpm/detail/package_details.cmake")
+  rapids_cpm_package_details(cutlass version repository tag shallow exclude)
+
+  include("${rapids-cmake-dir}/cpm/detail/generate_patch_command.cmake")
+  rapids_cpm_generate_patch_command(cutlass ${version} patch_command)
+
   rapids_cpm_find(
-    NvidiaCutlass ${PKG_VERSION}
+    NvidiaCutlass ${version}
     GLOBAL_TARGETS nvidia::cutlass::cutlass
     CPM_ARGS
-    GIT_REPOSITORY ${PKG_REPOSITORY}
-    GIT_TAG ${PKG_PINNED_TAG}
-    GIT_SHALLOW TRUE
+    GIT_REPOSITORY ${repository}
+    GIT_TAG ${tag}
+    GIT_SHALLOW ${shallow} ${patch_command}
     OPTIONS "CUDAToolkit_ROOT ${CUDAToolkit_LIBRARY_DIR}"
   )
 
@@ -79,14 +90,4 @@ function(find_and_configure_cutlass)
   )
 endfunction()
 
-if(NOT RAFT_CUTLASS_GIT_TAG)
-  set(RAFT_CUTLASS_GIT_TAG v2.10.0)
-endif()
-
-if(NOT RAFT_CUTLASS_GIT_REPOSITORY)
-  set(RAFT_CUTLASS_GIT_REPOSITORY https://github.com/NVIDIA/cutlass.git)
-endif()
-
-find_and_configure_cutlass(
-  VERSION 2.10.0 REPOSITORY ${RAFT_CUTLASS_GIT_REPOSITORY} PINNED_TAG ${RAFT_CUTLASS_GIT_TAG}
-)
+find_and_configure_cutlass()

cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h

@@ -61,6 +61,7 @@ class PairwiseDistanceEpilogueElementwise {
   using ElementT                      = ElementT_;
   static int const kElementsPerAccess = ElementsPerAccess;
   static int const kCount             = kElementsPerAccess;
+  static bool const kIsSingleSource   = true;
 
   using DistanceOp = DistanceOp_;
   using FinalOp    = FinalOp_;