[Kernel] Make vllm compatible with cutlass 4.0 #20460
Conversation
|
👋 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: D77771042 |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
There was a problem hiding this comment.
Summary of Changes
Hello @henrylhtsang, 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 aims to update vllm to be compatible with Cutlass 4.0. The changes involve adjusting internal configurations and header dependencies to align with the new library version's requirements, ensuring continued functionality and potential performance benefits.
Highlights
- Cutlass 4.0 Compatibility: Updated various Cutlass-related configurations to ensure compatibility with Cutlass 4.0, specifically by adjusting memory alignment requirements for
CDelements from 4 to 8 in severalscaled_mmimplementations. - Header Clean-up: Removed no longer necessary
#include "cute/tensor_predicate.hpp"directives fromsm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hppandmachete_mainloop.cuh, streamlining dependencies.
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. ↩
There was a problem hiding this comment.
Code Review
This pull request updates the codebase to be compatible with CUTLASS 4.0. The changes consist of removing an apparently obsolete header include (cute/tensor_predicate.hpp) and increasing a memory alignment constant (AlignmentCD) from 4 to 8 across several kernel definition files.
These modifications are consistent and appear to be necessary for the library upgrade. The increased alignment requirement is already met by stricter checks in the existing code, so these changes should be safe. Overall, the pull request looks good and addresses the compatibility goal.
henrylhtsang
changed the title
Summary: Rollback Plan: Differential Revision: D77771042
henrylhtsang
force-pushed
the
export-D77771042
branch
from
9e9bfca to
7e26967
Compare
|
This pull request was exported from Phabricator. Differential Revision: D77771042 |
| static constexpr int AlignmentAB = | ||
| 128 / cutlass::sizeof_bits<ElementAB>::value; | ||
| static constexpr int AlignmentCD = 4; | ||
| static constexpr int AlignmentCD = 8; |
There was a problem hiding this comment.
Would like to understand the motivation of this change
There was a problem hiding this comment.
This is about min alignments. Is it related to #17280?
|
This pull request has merge conflicts that must be resolved before it can be |
|
Nevermind, #19757 have the changes merged. |
4 participants
Purpose
Related to issue #20456.
I think we need these two changes to prepare for cutlass 4.0 upgrade. Since they seem to build fine on current trunk, I want to see if we can merge this first.
Though as a heads-up, we would still need to patch flash attention to include Dao-AILab/flash-attention#1719 and Dao-AILab/flash-attention#1723. cc @seemethere
Test plan
Tested build (in a different system).
Test result
Build was successful.
Differential Revision: D77771042