Question about kernel synchronization in a warp group

[sleepwalker2017]

I see the code snippet here:

    if (warp_group_idx == 0) {
           // do computation..
           __syncthreads();
           cutlass::arch::NamedBarrier::arrive(kNThreads, static_cast<int>(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.

[beginlner]

The other warp group also goes through a path with the same barrier.

[sleepwalker2017]

The other warp group also goes through a path with the same barrier.

So the same goes for the __syncthreads function. It's a totally new understanding for me. Thank you!