-
Notifications
You must be signed in to change notification settings - Fork 139
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.
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
[CK_TILE] Support moe with up gemm #1793
base: develop
Are you sure you want to change the base?
Conversation
bf07f14
to
b4a1c30
Compare
b4a1c30
to
2880f7a
Compare
78e7aae
to
af02d65
Compare
Please note that 2 function call to the |
typename Ts_::IndexDataType, | ||
// ck_tile::element_wise::FastGeluAsm, // | ||
// TODO: hardcoded | ||
ck_tile::element_wise::Silu, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
@huaiguxu
hardcode original activation to silu. If you need add new function into it, you really need to modify the template interface
cc @coderfeli
@@ -73,7 +73,7 @@ struct FusedMoeGemmPipeline_FlatmmUk | |||
constexpr index_t smem_0 = Policy::template GetUK_0<Problem>().GetSmemSize(); | |||
constexpr index_t smem_1 = Policy::template GetUK_1<Problem>().GetSmemSize(); | |||
constexpr index_t smem_bridge = | |||
BlockShape::Block_M0 * BlockShape::Block_N0 * sizeof(YDataType); | |||
BlockShape::Block_M0 * BlockShape::Block_N0 * (IsGateOnly ? 1 : 2); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
this is not correct, not multiply sizeop(dtype)
This reverts commit e674fb7.
Proposed changes
Support fused MoE with up gemm.
Checklist
Please put an
x
into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.clang-format
on all changed filesDiscussion
If this is a relatively large or complex change, feel free to start a discussion by explaining why you chose the solution you did and what alternatives you considered