From 1a75aad0e826a3dbad3f21debcd38d81d37fb054 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 20 Nov 2024 19:12:40 +0000 Subject: [PATCH 01/10] Upgrade to latest cutlass version --- cpp/cmake/thirdparty/get_cutlass.cmake | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/cpp/cmake/thirdparty/get_cutlass.cmake b/cpp/cmake/thirdparty/get_cutlass.cmake index 0123c4b07a..f09bf19b14 100644 --- a/cpp/cmake/thirdparty/get_cutlass.cmake +++ b/cpp/cmake/thirdparty/get_cutlass.cmake @@ -79,8 +79,9 @@ function(find_and_configure_cutlass) ) endfunction() +set(_cutlass_version 3.5.1) if(NOT RAFT_CUTLASS_GIT_TAG) - set(RAFT_CUTLASS_GIT_TAG v2.10.0) + set(RAFT_CUTLASS_GIT_TAG "v${_cutlass_version}") endif() if(NOT RAFT_CUTLASS_GIT_REPOSITORY) @@ -88,5 +89,5 @@ if(NOT RAFT_CUTLASS_GIT_REPOSITORY) endif() find_and_configure_cutlass( - VERSION 2.10.0 REPOSITORY ${RAFT_CUTLASS_GIT_REPOSITORY} PINNED_TAG ${RAFT_CUTLASS_GIT_TAG} + VERSION ${_cutlass_version} REPOSITORY ${RAFT_CUTLASS_GIT_REPOSITORY} PINNED_TAG ${RAFT_CUTLASS_GIT_TAG} ) From c4f1fcae38c54b1933f7a43b8adc500877b60b68 Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 14:47:41 -0500 Subject: [PATCH 02/10] kIsSingleSource --- cpp/include/raft/distance/detail/distance_ops/cosine.cuh | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/include/raft/distance/detail/distance_ops/cosine.cuh b/cpp/include/raft/distance/detail/distance_ops/cosine.cuh index 0883136c9f..71391f3f8d 100644 --- a/cpp/include/raft/distance/detail/distance_ops/cosine.cuh +++ b/cpp/include/raft/distance/detail/distance_ops/cosine.cuh @@ -29,6 +29,7 @@ struct cosine_cutlass_op { return static_cast(1.0) - static_cast(accVal / (aNorm * bNorm)); } __device__ AccT operator()(DataT aData) const noexcept { return aData; } + static const bool kIsSingleSource = false; }; /** From 5a6cf2615b09f7183c78e1dc900dea82e75499c0 Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 14:56:57 -0500 Subject: [PATCH 03/10] Move to correct class --- cpp/include/raft/distance/detail/distance_ops/cosine.cuh | 1 - .../raft/distance/detail/pairwise_distance_cutlass_base.cuh | 1 + .../distance/detail/pairwise_distance_epilogue_elementwise.h | 2 ++ 3 files changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/include/raft/distance/detail/distance_ops/cosine.cuh b/cpp/include/raft/distance/detail/distance_ops/cosine.cuh index 71391f3f8d..0883136c9f 100644 --- a/cpp/include/raft/distance/detail/distance_ops/cosine.cuh +++ b/cpp/include/raft/distance/detail/distance_ops/cosine.cuh @@ -29,7 +29,6 @@ struct cosine_cutlass_op { return static_cast(1.0) - static_cast(accVal / (aNorm * bNorm)); } __device__ AccT operator()(DataT aData) const noexcept { return aData; } - static const bool kIsSingleSource = false; }; /** diff --git a/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh b/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh index 1cc272f74e..aa3889bb75 100644 --- a/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh +++ b/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh @@ -83,6 +83,7 @@ std::enable_if_t::value> cutlassDistanceKernel(const Da AccT, // ElementZ_ OutT, // ElementT_ 1, // Elements per access 1 + false, // Single source DistanceFn, FinalLambda>; constexpr int batch_count = 1; diff --git a/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h b/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h index 2b2c04b9d3..960f28d644 100644 --- a/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h +++ b/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h @@ -49,6 +49,7 @@ template class PairwiseDistanceEpilogueElementwise { @@ -61,6 +62,7 @@ class PairwiseDistanceEpilogueElementwise { using ElementT = ElementT_; static int const kElementsPerAccess = ElementsPerAccess; static int const kCount = kElementsPerAccess; + static bool const kIsSingleSource = IsSingleSource; using DistanceOp = DistanceOp_; using FinalOp = FinalOp_; From 00f0c58795100cecb524aa5598779ecaa99425e0 Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 15:08:49 -0500 Subject: [PATCH 04/10] Single source --- .../raft/distance/detail/pairwise_distance_cutlass_base.cuh | 1 - .../distance/detail/pairwise_distance_epilogue_elementwise.h | 3 +-- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh b/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh index aa3889bb75..1cc272f74e 100644 --- a/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh +++ b/cpp/include/raft/distance/detail/pairwise_distance_cutlass_base.cuh @@ -83,7 +83,6 @@ std::enable_if_t::value> cutlassDistanceKernel(const Da AccT, // ElementZ_ OutT, // ElementT_ 1, // Elements per access 1 - false, // Single source DistanceFn, FinalLambda>; constexpr int batch_count = 1; diff --git a/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h b/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h index 960f28d644..f6dea987e5 100644 --- a/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h +++ b/cpp/include/raft/distance/detail/pairwise_distance_epilogue_elementwise.h @@ -49,7 +49,6 @@ template class PairwiseDistanceEpilogueElementwise { @@ -62,7 +61,7 @@ class PairwiseDistanceEpilogueElementwise { using ElementT = ElementT_; static int const kElementsPerAccess = ElementsPerAccess; static int const kCount = kElementsPerAccess; - static bool const kIsSingleSource = IsSingleSource; + static bool const kIsSingleSource = true; using DistanceOp = DistanceOp_; using FinalOp = FinalOp_; From f6460acdfd887caa1dbdd1ec81490c239e794250 Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 16:16:38 -0500 Subject: [PATCH 05/10] Debugging --- cpp/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 4ed9529a36..122c03814a 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -175,6 +175,7 @@ endif() add_library(raft INTERFACE) add_library(raft::raft ALIAS raft) +message(WARNING "\${RAFT_SOURCE_DIR}/include: ${RAFT_SOURCE_DIR}/include") target_include_directories( raft INTERFACE "$" "$" ) From 96e18fae4819a3b921996c4d318b5167c8d4546a Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 16:42:27 -0500 Subject: [PATCH 06/10] Patch cutlass --- .pre-commit-config.yaml | 5 +++- cpp/cmake/thirdparty/get_cutlass.cmake | 3 +++ .../patches/cutlass/build-export.patch | 27 +++++++++++++++++++ 3 files changed, 34 insertions(+), 1 deletion(-) create mode 100644 cpp/cmake/thirdparty/patches/cutlass/build-export.patch diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index d8ccf92ce5..679878607c 100644 --- a/.pre-commit-config.yaml +++ b/.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/thirdparty/patches/cutlass/build-export[.]patch$ - repo: https://github.com/pre-commit/pre-commit-hooks rev: v4.5.0 hooks: diff --git a/cpp/cmake/thirdparty/get_cutlass.cmake b/cpp/cmake/thirdparty/get_cutlass.cmake index f09bf19b14..712a5e1164 100644 --- a/cpp/cmake/thirdparty/get_cutlass.cmake +++ b/cpp/cmake/thirdparty/get_cutlass.cmake @@ -16,6 +16,8 @@ function(find_and_configure_cutlass) set(oneValueArgs VERSION REPOSITORY PINNED_TAG) cmake_parse_arguments(PKG "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN}) + find_package(Git REQUIRED) + # if(RAFT_ENABLE_DIST_DEPENDENCIES OR RAFT_COMPILE_LIBRARIES) set(CUTLASS_ENABLE_HEADERS_ONLY ON @@ -42,6 +44,7 @@ function(find_and_configure_cutlass) GIT_TAG ${PKG_PINNED_TAG} GIT_SHALLOW TRUE OPTIONS "CUDAToolkit_ROOT ${CUDAToolkit_LIBRARY_DIR}" + PATCH_COMMAND ${CMAKE_COMMAND} -E env GIT_COMMITTER_NAME=rapids-cmake GIT_COMMITTER_EMAIL=rapids.cmake@rapids.ai ${GIT_EXECUTABLE} am -3 ${CMAKE_CURRENT_SOURCE_DIR}/patches/cutlass/build-export.patch ) if(TARGET CUTLASS AND NOT TARGET nvidia::cutlass::cutlass) diff --git a/cpp/cmake/thirdparty/patches/cutlass/build-export.patch b/cpp/cmake/thirdparty/patches/cutlass/build-export.patch new file mode 100644 index 0000000000..a6423e9c08 --- /dev/null +++ b/cpp/cmake/thirdparty/patches/cutlass/build-export.patch @@ -0,0 +1,27 @@ +From e0a9597946257a01ae8444200f836ee51d5597ba Mon Sep 17 00:00:00 2001 +From: Kyle Edwards +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( + $ + $ + $ +- $ +- $ + ) + + # Mark CTK headers as system to supress warnings from them +-- +2.34.1 + From b8e8a4a6dac12c9f35bbe73f703b3a061aa30d6c Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 16:49:47 -0500 Subject: [PATCH 07/10] Remove debugging --- cpp/CMakeLists.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 122c03814a..4ed9529a36 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -175,7 +175,6 @@ endif() add_library(raft INTERFACE) add_library(raft::raft ALIAS raft) -message(WARNING "\${RAFT_SOURCE_DIR}/include: ${RAFT_SOURCE_DIR}/include") target_include_directories( raft INTERFACE "$" "$" ) From 283d35d4417dc2d1f88c649ed654cd2c24081c86 Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 18:19:17 -0500 Subject: [PATCH 08/10] Fix patch path --- cpp/cmake/thirdparty/get_cutlass.cmake | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/cmake/thirdparty/get_cutlass.cmake b/cpp/cmake/thirdparty/get_cutlass.cmake index 712a5e1164..84169627ce 100644 --- a/cpp/cmake/thirdparty/get_cutlass.cmake +++ b/cpp/cmake/thirdparty/get_cutlass.cmake @@ -44,7 +44,7 @@ function(find_and_configure_cutlass) GIT_TAG ${PKG_PINNED_TAG} GIT_SHALLOW TRUE OPTIONS "CUDAToolkit_ROOT ${CUDAToolkit_LIBRARY_DIR}" - PATCH_COMMAND ${CMAKE_COMMAND} -E env GIT_COMMITTER_NAME=rapids-cmake GIT_COMMITTER_EMAIL=rapids.cmake@rapids.ai ${GIT_EXECUTABLE} am -3 ${CMAKE_CURRENT_SOURCE_DIR}/patches/cutlass/build-export.patch + PATCH_COMMAND ${CMAKE_COMMAND} -E env GIT_COMMITTER_NAME=rapids-cmake GIT_COMMITTER_EMAIL=rapids.cmake@rapids.ai ${GIT_EXECUTABLE} am -3 ${CMAKE_CURRENT_FUNCTION_LIST_DIR}/patches/cutlass/build-export.patch ) if(TARGET CUTLASS AND NOT TARGET nvidia::cutlass::cutlass) From f08c2741409898fdc9d8997562440d0fc6733ff0 Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Wed, 20 Nov 2024 18:54:05 -0500 Subject: [PATCH 09/10] Use rapids_cpm_generate_patch_command() --- .pre-commit-config.yaml | 2 +- .../patches/cutlass/build-export.patch | 0 cpp/cmake/patches/cutlass_override.json | 16 ++++++++ cpp/cmake/thirdparty/get_cutlass.cmake | 37 +++++++++---------- 4 files changed, 34 insertions(+), 21 deletions(-) rename cpp/cmake/{thirdparty => }/patches/cutlass/build-export.patch (100%) create mode 100644 cpp/cmake/patches/cutlass_override.json diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 679878607c..e3b3c8c440 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -96,7 +96,7 @@ repos: exclude: | (?x) ^CHANGELOG[.]md$| - ^cpp/cmake/thirdparty/patches/cutlass/build-export[.]patch$ + ^cpp/cmake/patches/cutlass/build-export[.]patch$ - repo: https://github.com/pre-commit/pre-commit-hooks rev: v4.5.0 hooks: diff --git a/cpp/cmake/thirdparty/patches/cutlass/build-export.patch b/cpp/cmake/patches/cutlass/build-export.patch similarity index 100% rename from cpp/cmake/thirdparty/patches/cutlass/build-export.patch rename to cpp/cmake/patches/cutlass/build-export.patch diff --git a/cpp/cmake/patches/cutlass_override.json b/cpp/cmake/patches/cutlass_override.json new file mode 100644 index 0000000000..7bf818987f --- /dev/null +++ b/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" : "" + } + ] + } + } +} diff --git a/cpp/cmake/thirdparty/get_cutlass.cmake b/cpp/cmake/thirdparty/get_cutlass.cmake index 84169627ce..d5bdd4632f 100644 --- a/cpp/cmake/thirdparty/get_cutlass.cmake +++ b/cpp/cmake/thirdparty/get_cutlass.cmake @@ -13,11 +13,11 @@ # ============================================================================= 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}) - find_package(Git REQUIRED) - # if(RAFT_ENABLE_DIST_DEPENDENCIES OR RAFT_COMPILE_LIBRARIES) set(CUTLASS_ENABLE_HEADERS_ONLY ON @@ -36,15 +36,23 @@ 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}" - PATCH_COMMAND ${CMAKE_COMMAND} -E env GIT_COMMITTER_NAME=rapids-cmake GIT_COMMITTER_EMAIL=rapids.cmake@rapids.ai ${GIT_EXECUTABLE} am -3 ${CMAKE_CURRENT_FUNCTION_LIST_DIR}/patches/cutlass/build-export.patch ) if(TARGET CUTLASS AND NOT TARGET nvidia::cutlass::cutlass) @@ -82,15 +90,4 @@ function(find_and_configure_cutlass) ) endfunction() -set(_cutlass_version 3.5.1) -if(NOT RAFT_CUTLASS_GIT_TAG) - set(RAFT_CUTLASS_GIT_TAG "v${_cutlass_version}") -endif() - -if(NOT RAFT_CUTLASS_GIT_REPOSITORY) - set(RAFT_CUTLASS_GIT_REPOSITORY https://github.com/NVIDIA/cutlass.git) -endif() - -find_and_configure_cutlass( - VERSION ${_cutlass_version} REPOSITORY ${RAFT_CUTLASS_GIT_REPOSITORY} PINNED_TAG ${RAFT_CUTLASS_GIT_TAG} -) +find_and_configure_cutlass() From 4de0dd41d60cfebb7dcc64f4d91b36cf877b689a Mon Sep 17 00:00:00 2001 From: Kyle Edwards Date: Thu, 21 Nov 2024 09:08:13 -0500 Subject: [PATCH 10/10] Remove symbol exclusions --- .github/workflows/pr.yaml | 1 - .github/workflows/test.yaml | 1 - 2 files changed, 2 deletions(-) diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index 47951783ba..9c22edf74c 100644 --- a/.github/workflows/pr.yaml +++ b/.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 diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index 2bee8a3d1d..92020f6a76 100644 --- a/.github/workflows/test.yaml +++ b/.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/conda-cpp-tests.yaml@branch-24.12