-
Notifications
You must be signed in to change notification settings - Fork 120
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] Shrink lse_accum distributed tensor size in fmha splitkv combine kernel #1577
base: develop
Are you sure you want to change the base?
Conversation
constexpr index_t MThreads = kBlockSize / NThreads; | ||
constexpr index_t MPerThread = kMPerBlock / MThreads; | ||
constexpr index_t MThreadPerWarp = get_warp_size() / NThreads; | ||
constexpr index_t MWarps = kMPerBlock / MThreadPerWarp; |
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.
MWarps = kMPerBlock / MThreadPerWarp
? Confusing, why not MWarps = kBlockSize / get_warp_size()
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.
because I'm using hard-coded kBlockSize
=256, thus kBlockSize/ get_warp_size()
(=4) will exceed the real # warp needed here if kMPerBlock
is less than 64.
…ROCm/composable_kernel into feature/optimize-splitkv-combine-kernel
…ROCm/composable_kernel into feature/optimize-splitkv-combine-kernel
No description provided.