Replace multiply_add with homogeneous_multiply_add to Address Clang Template Parameter Issue
#20142
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
👋 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 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 🚀 |
|
This pull request was exported from Phabricator. Differential Revision: D76555276 |
There was a problem hiding this comment.
Summary of Changes
Hello @wenxin0319, I'm Gemini Code Assist1! I'm currently reviewing this pull request and will post my feedback shortly. In the meantime, here's a summary to help you and other reviewers quickly get up to speed!
This pull request addresses a critical compilation issue with the Clang compiler by updating the template function used in specific CUTLASS epilogue components. The change ensures that the multiply_add operation is correctly instantiated, resolving build failures and maintaining compatibility across different compiler versions for CUDA kernels.
Highlights
- Clang Compiler Compatibility Fix: Replaced
cutlass::multiply_addwithcutlass::homogeneous_multiply_addwithinVisitorComputeandSm90Computetemplate instantiations. This change specifically addresses a known Clang compiler issue where template template parameters with one template parameter do not correctly match classes that take multiple template parameters with defaults, leading to compilation errors.
Using Gemini Code Assist
The full guide for Gemini Code Assist can be found on our documentation page, here are some quick tips.
Invoking Gemini
You can request assistance from Gemini at any point in your pull request via creating an issue comment (i.e. comment on the pull request page) using either /gemini <command> or @gemini-code-assist <command>. Below is a summary of the supported commands.
| Feature | Command | Description |
|---|---|---|
| Code Review | /gemini review |
Performs a code review for the current pull request in its current state. |
| Pull Request Summary | /gemini summary |
Provides a summary of the current pull request in its current state. |
| Comment | @gemini-code-assist | Responds in comments when explicitly tagged, both in issue comments and review comments. |
| Help | /gemini help |
Displays a list of available commands. |
Customization
To customize Gemini Code Assist for GitHub experience, repository maintainers can create a configuration file and/or provide a custom code review style guide (such as PEP-8 for Python) by creating and adding files to a .gemini/ folder in the base of the repository. Detailed instructions can be found here.
Limitations & Feedback
Gemini Code Assist is currently in preview and may make mistakes. Please leave feedback on any instances where its feedback is incorrect or counter productive. You can react with 👍 and 👎 on @gemini-code-assist comments to provide feedback.
You can also get AI-powered code generation, chat, as well as code reviews directly in the IDE at no cost with the Gemini Code Assist IDE Extension.
Footnotes
-
Review the Privacy Notices, Generative AI Prohibited Use Policy, Terms of Service, and learn how to configure Gemini Code Assist in GitHub here. Gemini can make mistakes, so double check it and use code with caution. ↩
Comment on lines
155
to
157
| using Compute1 = cutlass::epilogue::threadblock::VisitorCompute< | ||
| cutlass::multiply_add, ElementD, float, | ||
| cutlass::homogeneous_multiply_add, ElementD, float, | ||
| cutlass::FloatRoundStyle::round_to_nearest>; |
There was a problem hiding this comment.
The original code used cutlass::multiply_add, which caused a Clang template parameter issue. Replacing it with cutlass::homogeneous_multiply_add resolves this issue by ensuring that the template parameters match the class definition. This change is crucial for code correctness and compatibility with Clang.
Suggested change
| using Compute1 = cutlass::epilogue::threadblock::VisitorCompute< | |
| cutlass::multiply_add, ElementD, float, | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; |
Comment on lines
212
to
214
| using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute< | ||
| cutlass::multiply_add, ElementD, float, | ||
| cutlass::homogeneous_multiply_add, ElementD, float, | ||
| cutlass::FloatRoundStyle::round_to_nearest>; |
There was a problem hiding this comment.
The original code used cutlass::multiply_add, which caused a Clang template parameter issue. Replacing it with cutlass::homogeneous_multiply_add resolves this issue by ensuring that the template parameters match the class definition. This change is crucial for code correctness and compatibility with Clang.
Suggested change
| using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute< | |
| cutlass::multiply_add, ElementD, float, | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; |
Comment on lines
290
to
292
| using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute< | ||
| cutlass::multiply_add, ElementD, float, | ||
| cutlass::homogeneous_multiply_add, ElementD, float, | ||
| cutlass::FloatRoundStyle::round_to_nearest>; |
There was a problem hiding this comment.
The original code used cutlass::multiply_add, which caused a Clang template parameter issue. Replacing it with cutlass::homogeneous_multiply_add resolves this issue by ensuring that the template parameters match the class definition. This change is crucial for code correctness and compatibility with Clang.
Suggested change
| using ComputeScaleBiasA = cutlass::epilogue::threadblock::VisitorCompute< | |
| cutlass::multiply_add, ElementD, float, | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; |
Comment on lines
197
to
199
| using Compute1 = cutlass::epilogue::fusion::Sm90Compute< | ||
| cutlass::multiply_add, ElementD, float, | ||
| cutlass::homogeneous_multiply_add, ElementD, float, | ||
| cutlass::FloatRoundStyle::round_to_nearest>; |
There was a problem hiding this comment.
The original code used cutlass::multiply_add, which caused a Clang template parameter issue. Replacing it with cutlass::homogeneous_multiply_add resolves this issue by ensuring that the template parameters match the class definition. This change is crucial for code correctness and compatibility with Clang.
Suggested change
| using Compute1 = cutlass::epilogue::fusion::Sm90Compute< | |
| cutlass::multiply_add, ElementD, float, | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; |
Comment on lines
240
to
242
| using Compute1 = cutlass::epilogue::fusion::Sm90Compute< | ||
| cutlass::multiply_add, ElementD, float, | ||
| cutlass::homogeneous_multiply_add, ElementD, float, | ||
| cutlass::FloatRoundStyle::round_to_nearest>; |
There was a problem hiding this comment.
The original code used cutlass::multiply_add, which caused a Clang template parameter issue. Replacing it with cutlass::homogeneous_multiply_add resolves this issue by ensuring that the template parameters match the class definition. This change is crucial for code correctness and compatibility with Clang.
Suggested change
| using Compute1 = cutlass::epilogue::fusion::Sm90Compute< | |
| cutlass::multiply_add, ElementD, float, | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; |
Comment on lines
297
to
299
| using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute< | ||
| cutlass::multiply_add, ElementD, float, | ||
| cutlass::homogeneous_multiply_add, ElementD, float, | ||
| cutlass::FloatRoundStyle::round_to_nearest>; |
There was a problem hiding this comment.
The original code used cutlass::multiply_add, which caused a Clang template parameter issue. Replacing it with cutlass::homogeneous_multiply_add resolves this issue by ensuring that the template parameters match the class definition. This change is crucial for code correctness and compatibility with Clang.
Suggested change
| using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute< | |
| cutlass::multiply_add, ElementD, float, | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; |
Comment on lines
373
to
375
| using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute< | ||
| cutlass::multiply_add, ElementD, float, | ||
| cutlass::homogeneous_multiply_add, ElementD, float, | ||
| cutlass::FloatRoundStyle::round_to_nearest>; |
There was a problem hiding this comment.
The original code used cutlass::multiply_add, which caused a Clang template parameter issue. Replacing it with cutlass::homogeneous_multiply_add resolves this issue by ensuring that the template parameters match the class definition. This change is crucial for code correctness and compatibility with Clang.
Suggested change
| using ComputeScaleBiasA = cutlass::epilogue::fusion::Sm90Compute< | |
| cutlass::multiply_add, ElementD, float, | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; | |
| cutlass::homogeneous_multiply_add, ElementD, float, | |
| cutlass::FloatRoundStyle::round_to_nearest>; |
wenxin0319
added a commit
to wenxin0319/vllm
that referenced
this pull request
Jun 26, 2025
…ng Template Parameter Issue (vllm-project#20142) Summary: Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276
421741b to
fa0a92e
Compare
|
This pull request was exported from Phabricator. Differential Revision: D76555276 |
wenxin0319
added a commit
to wenxin0319/vllm
that referenced
this pull request
Jun 26, 2025
…ng Template Parameter Issue (vllm-project#20142) Summary: Pull Request resolved: vllm-project#20142 Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276
fa0a92e to
d4bd575
Compare
wenxin0319
added a commit
to wenxin0319/vllm
that referenced
this pull request
Jun 26, 2025
…ng Template Parameter Issue (vllm-project#20142) Summary: Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276
d4bd575 to
e92866e
Compare
|
This pull request was exported from Phabricator. Differential Revision: D76555276 |
wenxin0319
added a commit
to wenxin0319/vllm
that referenced
this pull request
Jun 26, 2025
…ng Template Parameter Issue (vllm-project#20142) Summary: Pull Request resolved: vllm-project#20142 Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276
e92866e to
f75aadd
Compare
wenxin0319
added a commit
to wenxin0319/vllm
that referenced
this pull request
Jun 27, 2025
…ng Template Parameter Issue (vllm-project#20142) Summary: Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276
f75aadd to
bd371c0
Compare
|
This pull request was exported from Phabricator. Differential Revision: D76555276 |
wenxin0319
added a commit
to wenxin0319/vllm
that referenced
this pull request
Jun 27, 2025
…ng Template Parameter Issue (vllm-project#20142) Summary: Pull Request resolved: vllm-project#20142 Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276
bd371c0 to
c7c4342
Compare
houseroad
approved these changes
Jun 27, 2025
ProExpertProg
approved these changes
Jun 27, 2025
zou3519
approved these changes
Jun 27, 2025
|
Can you merge from main and see if the CI failures are resolved? |
c7c4342 to
9055b7f
Compare
houseroad
pushed a commit
to wenxin0319/vllm
that referenced
this pull request
Jul 8, 2025
…ng Template Parameter Issue (vllm-project#20142) Summary: Pull Request resolved: vllm-project#20142 Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276
…ng Template Parameter Issue (vllm-project#20142) Summary: Pull Request resolved: vllm-project#20142 Replaced the `multiply_add` template with `homogeneous_multiply_add` in the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults. Error: ``` vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:155:84: error: template template argument has different template parameters than its corresponding template template parameter 155 | using Compute1 = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:160:69: error: unknown type name 'Compute1'; did you mean 'Compute0'? 160 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:162:72: error: unknown type name 'Compute1'; did you mean 'Compute0'? 162 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< Compute1, ScaleA, EVTCompute0, Bias> ::Arguments; | ^~~~~~~~ | Compute0 vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:148:7: note: 'Compute0' declared here 148 | using Compute0 = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:212:93: error: template template argument has different template parameters than its corresponding template template parameter 212 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:217:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 217 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:221:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 221 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:204:7: note: 'ComputeScaleB' declared here 204 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:290:93: error: template template argument has different template parameters than its corresponding template template parameter 290 | using ComputeScaleBiasA = cutlass::epilogue::threadblock::template VisitorCompute< cutlass::multiply_add, ElementD, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:295:69: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 295 | public: using EVTCompute = cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:299:72: error: unknown type name 'ComputeScaleBiasA'; did you mean 'ComputeScaleB'? 299 | using ArgumentType = typename cutlass::epilogue::threadblock::Sm80EVT< ComputeScaleBiasA, ScaleA, EVTComputeScaleB, Bias> ::Arguments; | ^~~~~~~~~~~~~~~~~ | ComputeScaleB vllm/__vllm_cpp_lib__/buck-headers/cutlass_extensions/epilogue/scaled_mm_epilogues_c2x.hpp:282:7: note: 'ComputeScaleB' declared here 282 | using ComputeScaleB = cutlass::epilogue::threadblock::VisitorCompute< cutlass::multiplies, float, float, cutlass::FloatRoundStyle::round_to_nearest> ; | ^ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:5668: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:11:6055: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:5680: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:12:6068: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:5691: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:13:6079: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:5697: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:14:6086: error: expected unqualified-id vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:5664: error: template template argument has different template parameters than its corresponding template template parameter third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/functional.h:538:1: note: too many template parameters in template template argument 538 | template< class A, class B = A, class C = A> | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ third-party/cutlass-3/__cutlass-3-headers__/buck-headers/cutlass/epilogue/threadblock/fusion/visitor_compute.hpp:56:1: note: previous template template parameter is here 56 | template< class > class ComputeFn, class | ^~~~~~~~~~~~~~~~~ In file included from scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:1: vllm/__vllm_cpp_lib__/__action___33__/csrc/quantization/cutlass_w8a8/scaled_mm_c2x.cu.pic.o.compute_90a.cudafe1.stub.c:15:6047: error: expected unqualified-id fatal error: too many errors emitted, stopping now [-ferror-limit=] 20 errors generated. ``` Test Plan: D76353588 all testing passed. Rollback Plan: Reviewed By: zhuhan0 Differential Revision: D76555276 Signed-off-by: Lu Fang <lufang@fb.com>
9055b7f to
a7fd612
Compare
Chen-zexi
pushed a commit
to Chen-zexi/vllm
that referenced
this pull request
Jul 13, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com>
patrickvonplaten
pushed a commit
to patrickvonplaten/vllm
that referenced
this pull request
Jul 15, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com> Signed-off-by: Patrick von Platen <patrick.v.platen@gmail.com>
LyrisZhong
pushed a commit
to LyrisZhong/vllm
that referenced
this pull request
Jul 23, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com>
avigny
pushed a commit
to avigny/vllm
that referenced
this pull request
Jul 31, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com> Signed-off-by: avigny <47987522+avigny@users.noreply.github.com>
Pradyun92
pushed a commit
to Pradyun92/vllm
that referenced
this pull request
Aug 6, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com>
npanpaliya
pushed a commit
to odh-on-pz/vllm-upstream
that referenced
this pull request
Aug 6, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com>
jinzhen-lin
pushed a commit
to jinzhen-lin/vllm
that referenced
this pull request
Aug 9, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com> Signed-off-by: Jinzhen Lin <linjinzhen@hotmail.com>
epwalsh
pushed a commit
to epwalsh/vllm
that referenced
this pull request
Aug 27, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com>
googlercolin
pushed a commit
to googlercolin/vllm
that referenced
this pull request
Aug 29, 2025
…ng Template Parameter Issue (vllm-project#20142) Signed-off-by: Lu Fang <lufang@fb.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Summary:
Replaced the
multiply_addtemplate withhomogeneous_multiply_addin the codebase to address a known Clang issue where template template parameters with one template parameter do not match classes that take multiple template parameters with defaults.Error:
Test Plan:
D76353588 all testing passed.
Rollback Plan:
Differential Revision: D76555276