using Flash Attention version 2.5.7, upgraded CUTLASS to version 3.5 then encountered the following compilation error.

[ccccjunkang]

Hello, I am currently using Flash Attention version 2.5.7, which depends on CUTLASS version 3.4. I directly upgraded CUTLASS to version 3.5 (because other models in the project depend on it), and encountered the following compilation error. Could you please help me understand the reason for this issue?

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(99): error: "prefetch" has already been declared in the current scope
  prefetch(Copy_Atom<Copy_Traits<CopyOp, CT_Args...>, CA_Args...> const& atom,
  ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(95): error: template parameter pack not at end of parameter list
  template <class CopyOp, class... CT_Args, class... CA_Args,
                                   ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(95): error: template parameter pack not at end of parameter list
  template <class CopyOp, class... CT_Args, class... CA_Args,
                                                     ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(97): warning #1835-D: attribute "always_inline" does not apply here
  __inline__ __attribute__((always_inline)) __attribute__((host)) __attribute__((device))
                            ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(97): warning #1835-D: attribute "__host__" does not apply here
  __inline__ __attribute__((always_inline)) __attribute__((host)) __attribute__((device))
                                                           ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(99): error: incomplete type is not allowed
  prefetch(Copy_Atom<Copy_Traits<CopyOp, CT_Args...>, CA_Args...> const& atom,
  ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(99): error: identifier "Copy_Atom" is undefined
  prefetch(Copy_Atom<Copy_Traits<CopyOp, CT_Args...>, CA_Args...> const& atom,
           ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(99): error: type name is not allowed
  prefetch(Copy_Atom<Copy_Traits<CopyOp, CT_Args...>, CA_Args...> const& atom,
                     ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(99): error: type name is not allowed
  prefetch(Copy_Atom<Copy_Traits<CopyOp, CT_Args...>, CA_Args...> const& atom,
                                                      ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(99): error: expected a ")"
  prefetch(Copy_Atom<Copy_Traits<CopyOp, CT_Args...>, CA_Args...> const& atom,
                                                                ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(101): error: expected a ";"
  {
  ^

/cutlass3.5-llm/3rdparty/cutlass/include/cute/algorithm/prefetch.hpp(99): error: An inline __device__/__constant__/__managed__ variable must have internal linkage when the program is compiled in whole program mode (-rdc=false)
  prefetch(Copy_Atom<Copy_Traits<CopyOp, CT_Args...>, CA_Args...> const& atom,
  ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(127): error: SM80_CP_ASYNC_CACHEGLOBAL is not a template
          SM80_CP_ASYNC_CACHEGLOBAL<cute::uint128_t>,
          ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(131): error: type name is not allowed
          make_tiled_copy(Copy_Atom<Gmem_copy_struct, Element>{},
                                    ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(131): error: type name is not allowed
          make_tiled_copy(Copy_Atom<Gmem_copy_struct, Element>{},
                                                      ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(131): error: expected an expression
          make_tiled_copy(Copy_Atom<Gmem_copy_struct, Element>{},
                                                              ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(135): error: type name is not allowed
          make_tiled_copy(Copy_Atom<DefaultCopy, Element>{},
                                    ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(135): error: type name is not allowed
          make_tiled_copy(Copy_Atom<DefaultCopy, Element>{},
                                                 ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(135): error: expected an expression
          make_tiled_copy(Copy_Atom<DefaultCopy, Element>{},
                                                         ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(147): error: type name is not allowed
          make_tiled_copy(Copy_Atom<DefaultCopy, ElementAccum>{},
                                    ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(147): error: type name is not allowed
          make_tiled_copy(Copy_Atom<DefaultCopy, ElementAccum>{},
                                                 ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(147): error: expected an expression
          make_tiled_copy(Copy_Atom<DefaultCopy, ElementAccum>{},
                                                              ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(152): error: type name is not allowed
          make_tiled_copy(Copy_Atom<UniversalCopy<uint64_t>, Element>{},
                                    ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(152): error: type name is not allowed
          make_tiled_copy(Copy_Atom<UniversalCopy<uint64_t>, Element>{},
                                                             ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(152): error: expected an expression
          make_tiled_copy(Copy_Atom<UniversalCopy<uint64_t>, Element>{},
                                                                     ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(156): error: type name is not allowed
          make_tiled_copy(Copy_Atom<DefaultCopy, Element>{},
                                    ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(156): error: type name is not allowed
          make_tiled_copy(Copy_Atom<DefaultCopy, Element>{},
                                                 ^

/cutlass3.5-llm/cpp/tensorrt_llm/kernels/FAFwd/src/kernel_traits.h(156): error: expected an expression
          make_tiled_copy(Copy_Atom<DefaultCopy, Element>{},
                                                         ^

/usr/include/c++/9.5.0/fenv.h(58): error: the global scope has no "fenv_t"
    using ::fenv_t;
            ^

/usr/include/c++/9.5.0/fenv.h(59): error: the global scope has no "fexcept_t"
    using ::fexcept_t;
            ^

/usr/include/c++/9.5.0/fenv.h(62): error: the global scope has no "feclearexcept"
    using ::feclearexcept;
            ^

/usr/include/c++/9.5.0/fenv.h(63): error: the global scope has no "fegetexceptflag"
    using ::fegetexceptflag;
            ^

/usr/include/c++/9.5.0/fenv.h(64): error: the global scope has no "feraiseexcept"
    using ::feraiseexcept;
            ^

/usr/include/c++/9.5.0/fenv.h(65): error: the global scope has no "fesetexceptflag"
    using ::fesetexceptflag;
            ^

/usr/include/c++/9.5.0/fenv.h(66): error: the global scope has no "fetestexcept"
    using ::fetestexcept;
            ^

/usr/include/c++/9.5.0/fenv.h(68): error: the global scope has no "fegetround"
    using ::fegetround;
            ^

/usr/include/c++/9.5.0/fenv.h(69): error: the global scope has no "fesetround"
    using ::fesetround;
            ^

/usr/include/c++/9.5.0/fenv.h(71): error: the global scope has no "fegetenv"
    using ::fegetenv;
            ^

/usr/include/c++/9.5.0/fenv.h(72): error: the global scope has no "feholdexcept"
    using ::feholdexcept;
            ^

/usr/include/c++/9.5.0/fenv.h(73): error: the global scope has no "fesetenv"
    using ::fesetenv;
            ^

/usr/include/c++/9.5.0/fenv.h(74): error: the global scope has no "feupdateenv"
    using ::feupdateenv;
            ^

/usr/include/c++/9.5.0/cfenv(61): error: the global scope has no "fenv_t"
    using ::fenv_t;
            ^

/usr/include/c++/9.5.0/cfenv(62): error: the global scope has no "fexcept_t"
    using ::fexcept_t;
            ^

/usr/include/c++/9.5.0/cfenv(65): error: the global scope has no "feclearexcept"
    using ::feclearexcept;
            ^

/usr/include/c++/9.5.0/cfenv(66): error: the global scope has no "fegetexceptflag"
    using ::fegetexceptflag;
            ^

/usr/include/c++/9.5.0/cfenv(67): error: the global scope has no "feraiseexcept"
    using ::feraiseexcept;
            ^

/usr/include/c++/9.5.0/cfenv(68): error: the global scope has no "fesetexceptflag"
    using ::fesetexceptflag;
            ^

/usr/include/c++/9.5.0/cfenv(69): error: the global scope has no "fetestexcept"
    using ::fetestexcept;
            ^

/usr/include/c++/9.5.0/cfenv(71): error: the global scope has no "fegetround"
    using ::fegetround;
            ^

/usr/include/c++/9.5.0/cfenv(72): error: the global scope has no "fesetround"
    using ::fesetround;
            ^

/usr/include/c++/9.5.0/cfenv(74): error: the global scope has no "fegetenv"
    using ::fegetenv;
            ^

/usr/include/c++/9.5.0/cfenv(75): error: the global scope has no "feholdexcept"
    using ::feholdexcept;
            ^

/usr/include/c++/9.5.0/cfenv(76): error: the global scope has no "fesetenv"
    using ::fesetenv;
            ^

/usr/include/c++/9.5.0/cfenv(77): error: the global scope has no "feupdateenv"
    using ::feupdateenv;
            ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(57): error: name followed by "::" must be a class or namespace name
              static_assert(platform::is_same<Transform, UnaryTransform::Identity>::value ||
                            ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(57): error: expected a string literal
              static_assert(platform::is_same<Transform, UnaryTransform::Identity>::value ||
                                                         ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(62): error: name followed by "::" must be a class or namespace name
              if (platform::is_same<Transform, UnaryTransform::Identity>::value )
                  ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(62): error: type name is not allowed
              if (platform::is_same<Transform, UnaryTransform::Identity>::value )
                                               ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(62): error: the global scope has no "value"
              if (platform::is_same<Transform, UnaryTransform::Identity>::value )
                                                                          ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(69): error: name followed by "::" must be a class or namespace name
              else if (platform::is_same<Transform, UnaryTransform::Conjugate>::value )
                       ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(69): error: type name is not allowed
              else if (platform::is_same<Transform, UnaryTransform::Conjugate>::value )
                                                    ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(69): error: the global scope has no "value"
              else if (platform::is_same<Transform, UnaryTransform::Conjugate>::value )
                                                                                ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(86): error: name followed by "::" must be a class or namespace name
              static_assert(platform::is_same<Transform, UnaryTransform::Identity>::value ||
                            ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(86): error: expected a string literal
              static_assert(platform::is_same<Transform, UnaryTransform::Identity>::value ||
                                                         ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(90): error: name followed by "::" must be a class or namespace name
              if (platform::is_same<Transform, UnaryTransform::Identity>::value )
                  ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(90): error: type name is not allowed
              if (platform::is_same<Transform, UnaryTransform::Identity>::value )
                                               ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(90): error: the global scope has no "value"
              if (platform::is_same<Transform, UnaryTransform::Identity>::value )
                                                                          ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(94): error: name followed by "::" must be a class or namespace name
              else if (platform::is_same<Transform, UnaryTransform::Conjugate>::value )
                       ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(94): error: type name is not allowed
              else if (platform::is_same<Transform, UnaryTransform::Conjugate>::value )
                                                    ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/transform/thread/unary_op.h(94): error: the global scope has no "value"
              else if (platform::is_same<Transform, UnaryTransform::Conjugate>::value )
                                                                                ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(145): error: namespace "cute::std" has no member "fesetround"
      std::fesetround(
           ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(146): error: namespace "cute::std" has no member "nearbyint"
      return (result_type)std::nearbyint(s);
                               ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(162): error: namespace "cute::std" has no member "fesetround"
      std::fesetround(
           ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(163): error: namespace "cute::std" has no member "nearbyint"
      return (result_type)std::nearbyint(s);
                               ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(277): error: namespace "cute::std" has no member "fesetround"
      std::fesetround(
           ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(278): error: namespace "cute::std" has no member "nearbyint"
      int32_t intermediate = (int32_t)std::nearbyint(s);
                                           ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(281): error: namespace "cute::std" has no member "max"
      intermediate = std::max(intermediate, (int32_t)std::numeric_limits<int8_t>::lowest());
                          ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(281): error: namespace "cute::std" has no member "numeric_limits"
      intermediate = std::max(intermediate, (int32_t)std::numeric_limits<int8_t>::lowest());
                                                          ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(281): error: type name is not allowed
      intermediate = std::max(intermediate, (int32_t)std::numeric_limits<int8_t>::lowest());
                                                                         ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(281): error: the global scope has no "lowest"
      intermediate = std::max(intermediate, (int32_t)std::numeric_limits<int8_t>::lowest());
                                                                                  ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(284): error: namespace "cute::std" has no member "min"
      intermediate = std::min(intermediate, (int32_t)std::numeric_limits<int8_t>::max());
                          ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(284): error: namespace "cute::std" has no member "numeric_limits"
      intermediate = std::min(intermediate, (int32_t)std::numeric_limits<int8_t>::max());
                                                          ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(284): error: type name is not allowed
      intermediate = std::min(intermediate, (int32_t)std::numeric_limits<int8_t>::max());
                                                                         ^

/cutlass3.5-llm/3rdparty/cutlass/include/cutlass/numeric_conversion.h(284): error: no instance of overloaded function "max" matches the argument list
      intermediate = std::min(intermediate, (int32_t)std::numeric_limits<int8_t>::max());
                                                                                ^

[tridao]

You'd want to use the latest version of flash-attention.