[CUDA] upgrade cutlass 4.7 and cudnn-frontend 1.27

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,

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

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"

cmake/patches/cutlass/cutlass_4.4.2.patch → 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
@@ -50,5 +56,4 @@ index be207a49..6028e01d 100644
 -  explicit float_exmy_base<T, Derived>(unsigned x) {
 +  explicit float_exmy_base(unsigned x) {
      storage = static_cast<Derived*>(this)->convert_from_float(float(x)).storage;
-   }
-
+   }

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 <bool IsLoad, class ProblemShapeMNKL>
+    template <bool IsLoad, bool WaitForInflightTmaRequests = true, class ProblemShapeMNKL, class TensorMaps>
     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 <bool IsLoad>
+    template <bool IsLoad, bool WaitForInflightTmaRequests = true, class TensorMaps>
     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 <bool IsLoad>
-    CUTLASS_DEVICE void tensormaps_fence_acquire([[maybe_unused]] cute::TmaDescriptor const* tensormap) {
+    template <bool IsLoad, class TensorMaps>
+    CUTLASS_DEVICE void tensormaps_fence_acquire([[maybe_unused]] TensorMaps const& tensormaps) {
     }
   };