deepseek-ai / deepseek-ai/FlashMLA

Question about kernel synchronization in a warp group

Open
#41 2 comments 0 reactions 0 assignees View on GitHub
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.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.