A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 15, 02:15 UTC
Replace `multiply_add` with `homogeneous_multiply_add` to Address Clang Template Parameter Issue (#20142)
Signed-off-by: Lu Fang <lufang@fb.com>
Wenxin Cheng committed
Jul 8, 2025 at 17:30 UTC
5eaf57005065441af2c7223eec01f4526571e00c
2 files changed
+7
-7
csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp
+3
-3
@@ -153,7 +153,7 @@ struct ScaledEpilogueBias
153
cutlass::epilogue::threadblock::Sm80EVT<Compute0, ScaleB, Accum>;
154
155
using Compute1 = cutlass::epilogue::threadblock::VisitorCompute<
156
- cutlass::multiply_add, ElementD, float,
156
+ cutlass::homogeneous_multiply_add, ElementD, float,
157
cutlass::FloatRoundStyle::round_to_nearest>;
158
159
public:
@@ -210,7 +210,7 @@ struct ScaledEpilogueBiasAzp
210
EVTComputeAzp>;
211
212
using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute<
213
- cutlass::multiply_add, ElementD, float,
213
+ cutlass::homogeneous_multiply_add, ElementD, float,
214
cutlass::FloatRoundStyle::round_to_nearest>;
215
216
public:
@@ -288,7 +288,7 @@ struct ScaledEpilogueBiasAzpToken
288
EVTComputeAcc>;
289
290
using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute<
291
- cutlass::multiply_add, ElementD, float,
291
+ cutlass::homogeneous_multiply_add, ElementD, float,
292
cutlass::FloatRoundStyle::round_to_nearest>;
293
294
public:
csrc/cutlass_extensions/epilogue/scaled_mm_epilogues_c3x.hpp
+4
-4
@@ -195,7 +195,7 @@ struct ScaledEpilogueBias
195
cutlass::epilogue::fusion::Sm90EVT<Compute0, ScaleB, Accum>;
196
197
using Compute1 = cutlass::epilogue::fusion::Sm90Compute<
198
- cutlass::multiply_add, ElementD, float,
198
+ cutlass::homogeneous_multiply_add, ElementD, float,
199
cutlass::FloatRoundStyle::round_to_nearest>;
200
201
public:
@@ -238,7 +238,7 @@ struct ScaledEpilogueColumnBias
238
cutlass::epilogue::fusion::Sm90EVT<Compute0, ScaleB, Accum>;
239
240
using Compute1 = cutlass::epilogue::fusion::Sm90Compute<
241
- cutlass::multiply_add, ElementD, float,
241
+ cutlass::homogeneous_multiply_add, ElementD, float,
242
cutlass::FloatRoundStyle::round_to_nearest>;
243
244
public:
@@ -295,7 +295,7 @@ struct ScaledEpilogueBiasAzp
295
cutlass::epilogue::fusion::Sm90EVT<ComputeScaleB, ScaleB, EVTComputeAzp>;
296
297
using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute<
298
- cutlass::multiply_add, ElementD, float,
298
+ cutlass::homogeneous_multiply_add, ElementD, float,
299
cutlass::FloatRoundStyle::round_to_nearest>;
300
301
public:
@@ -371,7 +371,7 @@ struct ScaledEpilogueBiasAzpToken
371
cutlass::epilogue::fusion::Sm90EVT<ComputeScaleB, ScaleB, EVTComputeAcc>;
372
373
using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute<
374
- cutlass::multiply_add, ElementD, float,
374
+ cutlass::homogeneous_multiply_add, ElementD, float,
375
cutlass::FloatRoundStyle::round_to_nearest>;
376
377
public: