deepseek-ai / deepseek-ai/FlashMLA
Question about kernel synchronization in a warp group
- Dominant language
- C++
- Stars
- 12.9k
- Forks
- 1.2k
- Avg merge
- 4h 20m
- Merged PRs (30d)
- 2
Description
I see the code snippet here:
```
if (warp_group_idx == 0) {
// do computation..
__syncthreads();
cutlass::arch::NamedBarrier::arrive(kNThreads, static_cast(NamedBarriers::SReady));
}
```
I print this value: kNThreads, it's 256 here.
And the cuda block size is also 256. It's divided into 2 warp groups, 128 threads for each group.
My question is:
Is it ok to wait for 256 threads in a path where only 128 threads will go through?
Doesn't it cause the cuda block hang?
Also, the same question for the usage of `__syncthreads` in a branch where only half threads goes to.
Contributor guide
No contributing guide indexed for this repository
Assessment
This issue has not been assessed yet.