Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

feat: modify group-gemm stage number #497

Merged
merged 1 commit into from
Sep 13, 2024
Merged

feat: modify group-gemm stage number #497

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion include/flashinfer/group_gemm/wrapper.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -85,7 +85,7 @@ cudaError_t CutlassSegmentGEMMWrapper(CutlassSegmentGEMMHandler* handler, DType*
cutlass::gemm::GemmShape<16, 8, 16>, // Instruction Shape
cutlass::epilogue::thread::LinearCombination<DType, 8, float, float>, // Epilogue
cutlass::gemm::threadblock::GemmBatchedIdentityThreadblockSwizzle, // Swizzling Operator
8 // Stages
4 // Stages
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Adjust the size based on different shared memory instead of directly changing 8 to 4. cc @yzh119

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The value 4 was chosen based on the configuration in cutlass, which should be compatible with most scenarios. 😄

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The example of the cutlass you provided is the SM80, not the SM90.

>::GemmKernel;

using EpilogueOutputOp = typename GemmKernel::Epilogue::OutputOp;
Expand Down