From 5eaf57005065441af2c7223eec01f4526571e00c Mon Sep 17 00:00:00 2001 From: Wenxin Cheng <115043072+wenxin0319@users.noreply.github.com> Date: Tue, 8 Jul 2025 17:30:18 -0700 Subject: [PATCH] Replace `multiply_add` with `homogeneous_multiply_add` to Address Clang Template Parameter Issue (#20142) Signed-off-by: Lu Fang --- .../epilogue/scaled_mm_epilogues_c2x.hpp | 6 +++--- .../epilogue/scaled_mm_epilogues_c3x.hpp | 8 ++++---- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp b/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp index 64b7ddae3d2d7..ad8c0067d4a99 100644 --- a/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp +++ b/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp @@ -153,7 +153,7 @@ struct ScaledEpilogueBias cutlass::epilogue::threadblock::Sm80EVT; using Compute1 = cutlass::epilogue::threadblock::VisitorCompute< - cutlass::multiply_add, ElementD, float, + cutlass::homogeneous_multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest>; public: @@ -210,7 +210,7 @@ struct ScaledEpilogueBiasAzp EVTComputeAzp>; using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute< - cutlass::multiply_add, ElementD, float, + cutlass::homogeneous_multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest>; public: @@ -288,7 +288,7 @@ struct ScaledEpilogueBiasAzpToken EVTComputeAcc>; using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute< - cutlass::multiply_add, ElementD, float, + cutlass::homogeneous_multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest>; public: diff --git a/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp b/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp index 62b848a0a9635..cf79507e19973 100644 --- a/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp +++ b/csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp @@ -195,7 +195,7 @@ struct ScaledEpilogueBias cutlass::epilogue::fusion::Sm90EVT; using Compute1 = cutlass::epilogue::fusion::Sm90Compute< - cutlass::multiply_add, ElementD, float, + cutlass::homogeneous_multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest>; public: @@ -238,7 +238,7 @@ struct ScaledEpilogueColumnBias cutlass::epilogue::fusion::Sm90EVT; using Compute1 = cutlass::epilogue::fusion::Sm90Compute< - cutlass::multiply_add, ElementD, float, + cutlass::homogeneous_multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest>; public: @@ -295,7 +295,7 @@ struct ScaledEpilogueBiasAzp cutlass::epilogue::fusion::Sm90EVT; using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute< - cutlass::multiply_add, ElementD, float, + cutlass::homogeneous_multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest>; public: @@ -371,7 +371,7 @@ struct ScaledEpilogueBiasAzpToken cutlass::epilogue::fusion::Sm90EVT; using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute< - cutlass::multiply_add, ElementD, float, + cutlass::homogeneous_multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest>; public: