[Kernel][Build/CI] Bump CUTLASS to 3.8 and add initializers for cutlass epilogues

[henrylhtsang]

The original code was added in #9855. I am seeing errors like

cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp:149:194: error: missing initializer for member ‘cutlass::epilogue::fusion::detail::Sm90VisitorImplBase<cutlass::epilogue::fusion::Sm90RowOrScalarBroadcast<0, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, float, cute::tuple<cute::C<0>, cute::C<1>, cute::C<0> >, 4>, cutlass::epilogue::fusion::Sm90AccFetch, cutlass::epilogue::fusion::Sm90Compute<cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest, void> >::Arguments::op_2’ [-Werror=missing-field-initializers]
cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp:150:38: error: missing initializer for member ‘cutlass::epilogue::fusion::detail::Sm90VisitorImplBase<cutlass::epilogue::fusion::Sm90ColOrScalarBroadcast<0, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, float, cute::tuple<cute::C<1>, cute::C<0>, cute::C<0> >, 4>, cutlass::epilogue::fusion::Sm90TreeVisitor<cutlass::epilogue::fusion::Sm90Compute<cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest, void>, cutlass::epilogue::fusion::Sm90RowOrScalarBroadcast<0, cute::tuple<cute::C<128>, cute::C<128>, cute::C<128> >, float, cute::tuple<cute::C<0>, cute::C<1>, cute::C<0> >, 4>, cutlass::epilogue::fusion::Sm90AccFetch>, cutlass::epilogue::fusion::Sm90Compute<cutlass::multiplies, cutlass::half_t, float, cutlass::FloatRoundStyle::round_to_nearest, void> >::Arguments::op_2’ [-Werror=missing-field-initializers]
  150 |     return ArgumentType{a_args, evt0_args};
      |                                      ^

when I moved to cutlass 3.8. I think it is due to https://github.com/NVIDIA/cutlass/blob/main/include/cute/atom/mma_traits_sm100.hpp#L49-L52, which are newly added for 3.8.

From my setup, this change works. Feel free to suggest other changes that can help resolve the errors and/or tests to do before merging.

[github-actions]

👋 Hi! Thank you for contributing to the vLLM project.

💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in #pr-reviews, coordinate on features in #feat- channels, or join special interest groups in #sig- channels.

Just a reminder: PRs would not trigger full CI run by default. Instead, it would only run fastcheck CI which starts running only a small and essential subset of CI tests to quickly catch errors. You can run other CI tests on top of those by going to your fastcheck build on Buildkite UI (linked in the PR checks section) and unblock them. If you do not have permission to unblock, ping simon-mo or khluu to add you in our Buildkite org.

Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging.

To run CI, PR reviewers can either: Add ready label to the PR or enable auto-merge.

🚀

[henrylhtsang]

@LucasWilkinson any idea?

[tlrmchlsmth]

The fix looks good. Exact same change is in #13798, but maybe this one lands first as it's smaller.

Go ahead and update CUTLASS to 3.8 in this PR?

[henrylhtsang]

The fix looks good. Exact same change is in #13798, but maybe this one lands first as it's smaller.

Go ahead and update CUTLASS to 3.8 in this PR?

Thanks, it is this one right?

vllm/CMakeLists.txt

Line 250 in 4a8cfc7

GIT_TAG v3.7.0

btw noob question: why does the order of arguments seem to not matter? For example, for

  using EVTCompute0 =
      cutlass::epilogue::threadblock::Sm80EVT<Compute0, ScaleB, Accum>;

shouldn't it be

    typename EVTCompute0::Arguments evt0_args{{}, b_args, {}};

?

[jwfromm]

I think the argument order should matter and needs to be:

{argument 0, argument 1, operator}

In this case argument 0 is ScaleB, argument 1 is the accumulator (represented by {}) and argument 2 is multiplies (also represented by {}). so {b_args, {}, {}} is correct.

[henrylhtsang]

I think the argument order should matter and needs to be:

{argument 0, argument 1, operator}

In this case argument 0 is ScaleB, argument 1 is the accumulator (represented by {}) and argument 2 is multiplies (also represented by {}). so {b_args, {}, {}} is correct.

Thanks Josh. also sharing the source https://github.com/NVIDIA/cutlass/blob/eefa171318b79cbe2e78514d4cce5cd0fe919d0c/include/cutlass/epilogue/fusion/sm90_callbacks_tma_warpspecialized.hpp#L57-L58

template <class NodeOp, class... ChildOps>
struct Sm90TreeVisitor : Sm90VisitorImpl<ChildOps..., NodeOp> {

[henrylhtsang]

So I think this should be good to land