diff --git a/cmake/deps.txt b/cmake/deps.txt index f4ffa9fb2ec46..477383a418782 100644 --- a/cmake/deps.txt +++ b/cmake/deps.txt @@ -54,10 +54,10 @@ pytorch_cpuinfo;https://github.com/pytorch/cpuinfo/archive/4628dc060ce4e82345dc1 re2;https://github.com/google/re2/archive/refs/tags/2024-07-02.zip;646e1728269cde7fcef990bf4a8e87b047882e88 safeint;https://github.com/dcleblanc/SafeInt/archive/refs/tags/3.0.28.zip;23f252040ff6cb9f1fd18575b32fa8fb5928daac tensorboard;https://github.com/tensorflow/tensorboard/archive/373eb09e4c5d2b3cc2493f0949dc4be6b6a45e81.zip;67b833913605a4f3f499894ab11528a702c2b381 -cutlass;https://github.com/NVIDIA/cutlass/archive/refs/tags/v4.4.2.zip;4b0bae4428b84370407c0a71778b13dc2eee5be1 +cutlass;https://github.com/NVIDIA/cutlass/archive/refs/tags/v4.7.0.zip;51d4f1ba4b9977b94606e23bbc307463a6741383 extensions;https://github.com/microsoft/onnxruntime-extensions/archive/c24b7bab0c12f53da76d0c31b03b9f0f8ec8f3b4.zip;239063aee4946a9af147b473a4c3da78ba7413b4 directx_headers;https://github.com/microsoft/DirectX-Headers/archive/refs/tags/v1.613.1.zip;47653509a3371eabb156360f42faf582f314bf2e -cudnn_frontend;https://github.com/NVIDIA/cudnn-frontend/archive/refs/tags/v1.24.0.zip;a55a1980bf5c57692d66ae7bc3b39798f5535e1f +cudnn_frontend;https://github.com/NVIDIA/cudnn-frontend/archive/refs/tags/v1.27.0.zip;1e4c9a464d3437e388ab0163f3be068dba783c08 dawn;https://github.com/google/dawn/archive/refs/tags/v20260714.215939.zip;3056ed22d1606258ab43221b8c85b55b88614137 kleidiai;https://github.com/ARM-software/kleidiai/archive/refs/tags/v1.20.0.tar.gz;6895e72b3d5cf1173358164cb3d64c9d7d33cc84 # kleidiai-qmx is pinned to a specific commit as there are no tagged releases. When an appropriate tagged release becomes available, diff --git a/cmake/external/cutlass.cmake b/cmake/external/cutlass.cmake index cd9a9c5179615..b0a67deaeb963 100644 --- a/cmake/external/cutlass.cmake +++ b/cmake/external/cutlass.cmake @@ -4,7 +4,7 @@ onnxruntime_fetchcontent_declare( URL ${DEP_URL_cutlass} URL_HASH SHA1=${DEP_SHA1_cutlass} EXCLUDE_FROM_ALL - PATCH_COMMAND ${Patch_EXECUTABLE} --ignore-whitespace -p1 < ${PROJECT_SOURCE_DIR}/patches/cutlass/cutlass_4.4.2.patch + PATCH_COMMAND ${Patch_EXECUTABLE} --ignore-whitespace -p1 < ${PROJECT_SOURCE_DIR}/patches/cutlass/cutlass_4.7.0.patch ) # We only consume CUTLASS as a header-only dependency. Avoid FetchContent_MakeAvailable here diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index e500645f22377..ec2ef145d31db 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1063,6 +1063,7 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_common.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_loader.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cufft_loader.cc" + "${ONNXRUNTIME_ROOT}/core/providers/cuda/fpgeneric.cu" "${ONNXRUNTIME_ROOT}/core/providers/cuda/reduction/reduction_functions.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/reduction/reduction_functions.cu" "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_plugin_test_shims.cc" diff --git a/cmake/patches/cudnn_frontend/cudnn_frontend_win_dynamic_loading.patch b/cmake/patches/cudnn_frontend/cudnn_frontend_win_dynamic_loading.patch index 96242982a7468..c815d6305ddf8 100644 --- a/cmake/patches/cudnn_frontend/cudnn_frontend_win_dynamic_loading.patch +++ b/cmake/patches/cudnn_frontend/cudnn_frontend_win_dynamic_loading.patch @@ -15,7 +15,16 @@ const char *error = reinterpret_cast(dlerror()); if (!handle || error) { // If opening the library fails, throw an exception with the error message -@@ -85,13 +92,22 @@ +@@ -82,5 +89,8 @@ + if (user_lib) { + if (user_lib[0] != '\0') { ++#ifdef _WIN32 ++ SetLastError(0); ++#endif + HMODULE handle = dlopen(user_lib, RTLD_NOW); + const char *error = reinterpret_cast(dlerror()); + if (!handle || error) { +@@ -85,13 +95,22 @@ dlerror(); // List of potential libcudart libraries (Adding major version to support python package) diff --git a/cmake/patches/cutlass/cutlass_4.4.2.patch b/cmake/patches/cutlass/cutlass_4.7.0.patch similarity index 72% rename from cmake/patches/cutlass/cutlass_4.4.2.patch rename to cmake/patches/cutlass/cutlass_4.7.0.patch index 6776eded10640..afb2d21472483 100644 --- a/cmake/patches/cutlass/cutlass_4.4.2.patch +++ b/cmake/patches/cutlass/cutlass_4.7.0.patch @@ -26,11 +26,17 @@ index a8af62be..22e7332d 100644 virtual Status memsetDeviceImpl( void* destination, ///< Device memory pointer to be filled void const* fill_value, ///< Value to be filled in the buffer +diff --git a/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp b/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp +index ae31620a..fd559cbb 100644 +--- a/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp ++++ b/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp +@@ -53 +52,0 @@ public: +- CUTLASS_HOST_DEVICE diff --git a/include/cutlass/exmy_base.h b/include/cutlass/exmy_base.h index be207a49..6028e01d 100644 --- a/include/cutlass/exmy_base.h +++ b/include/cutlass/exmy_base.h -@@ -1021,18 +1021,18 @@ struct float_exmy_base +@@ -1021,17 +1021,17 @@ struct float_exmy_base /// Floating point conversion CUTLASS_HOST_DEVICE @@ -51,4 +57,15 @@ index be207a49..6028e01d 100644 + explicit float_exmy_base(unsigned x) { storage = static_cast(this)->convert_from_float(float(x)).storage; } - +diff --git a/include/cutlass/subbyte_reference.h b/include/cutlass/subbyte_reference.h +--- a/include/cutlass/subbyte_reference.h ++++ b/include/cutlass/subbyte_reference.h +@@ -456,7 +456,7 @@ public: + // + Storage assumed; + #if (__CUDACC_VER_MAJOR__ > 12) || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ >= 8) +- Storage original = __nv_atomic_load_n(ptr_, __NV_ATOMIC_RELAXED); ++ Storage original = __nv_atomic_load_n(ptr_, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_DEVICE); + #else + Storage original = *const_cast(ptr_); + #endif diff --git a/onnxruntime/contrib_ops/cuda/llm/cutlass_extensions/epilogue/collective/epilogue_moe_finalize.hpp b/onnxruntime/contrib_ops/cuda/llm/cutlass_extensions/epilogue/collective/epilogue_moe_finalize.hpp index 8ba877aa21a68..4aca09018377d 100644 --- a/onnxruntime/contrib_ops/cuda/llm/cutlass_extensions/epilogue/collective/epilogue_moe_finalize.hpp +++ b/onnxruntime/contrib_ops/cuda/llm/cutlass_extensions/epilogue/collective/epilogue_moe_finalize.hpp @@ -449,20 +449,20 @@ struct EpilogueMoeFusedFinalizeBuilder { // Dummy methods to perform different parts of TMA/Tensormap modifications - template + template CUTLASS_DEVICE void tensormaps_perform_update([[maybe_unused]] TensorMapStorage& shared_tensormaps, [[maybe_unused]] typename EpilogueOp::Params const& params, - [[maybe_unused]] cute::TmaDescriptor const* tensormap, [[maybe_unused]] ProblemShapeMNKL problem_shape, + [[maybe_unused]] TensorMaps const& tensormaps, [[maybe_unused]] ProblemShapeMNKL problem_shape, [[maybe_unused]] int32_t next_batch, [[maybe_unused]] int32_t warp_group_idx) { } - template + template CUTLASS_DEVICE void tensormaps_cp_fence_release([[maybe_unused]] TensorMapStorage& shared_tensormaps, - [[maybe_unused]] cute::TmaDescriptor const* tensormap, [[maybe_unused]] int32_t warp_group_idx) { + [[maybe_unused]] TensorMaps const& tensormaps, [[maybe_unused]] int32_t warp_group_idx) { } - template - CUTLASS_DEVICE void tensormaps_fence_acquire([[maybe_unused]] cute::TmaDescriptor const* tensormap) { + template + CUTLASS_DEVICE void tensormaps_fence_acquire([[maybe_unused]] TensorMaps const& tensormaps) { } };