CUDA核函数中__syncthreads()最优调用位置的性能影响

在保证代码逻辑正确、不存在共享内存读写竞态的前提下，尽可能晚调用__syncthreads()（即示例中的Option #2位置）存在明确的性能收益，调用位置对性能并非无影响；反之无意义的提前调用（Option #1位置）只会带来额外的性能损耗。

示例代码如下：

__global__ void kernel(const float* data) {
    __shared__ float shared_data[64];

    if (threadIdx.x < 64) {
        shared_data[threadIdx.x] = data[threadIdx.x];
    }
    // Option #1: 提前放置__syncthreads()
    // 此处为大量不访问shared_data的无关代码
    // Option #2: 延后放置__syncthreads()
    // 此处为访问shared_data的业务代码
}

• __syncthreads()本质是块内线程的强制屏障：块内所有线程必须全部抵达屏障点后，才能继续执行屏障后的指令。早到屏障的线程会进入等待状态，直到最后一个线程到达才会继续往下跑。
• 如果选择Option #1提前同步：示例中threadIdx.x >= 64的线程根本不参与共享内存写入，会直接跑到屏障点空等，必须等前64个线程完成共享内存写入、抵达屏障后，所有线程才能开始执行中间的无关代码。这段空等完全是无意义的开销，平白拉长了内核执行时间。就算所有线程都参与共享内存写入，提前同步也会让执行速度快的线程提前停在屏障点等慢线程，浪费算力。
• 如果选择Option #2延后同步：共享内存写入完成后，所有线程不需要等待，立刻并行执行中间的无关代码，直到马上要读取共享内存前才做同步。这时候原本需要花在屏障等待上的时间，完全可以被中间的计算、访存操作掩盖——等所有线程跑完无关代码抵达屏障时，之前的共享内存写入早就完成了，几乎不会出现空等stall。

延后同步有不可突破的红线：__syncthreads()必须放在第一次读取对应共享内存的指令之前，绝对不能让任何线程在所有共享内存写入完成前就读取对应的值，否则会出现读到脏数据、计算结果错误等未定义行为。

不存在提前同步性能更好的场景。只有当中间的无关代码长度为0（也就是两段位置完全重合）时，两种放置方式的性能才会一致。所有正常业务场景下，在保证正确性的前提下尽可能晚放__syncthreads()，都是更优的选择。

内容的提问来源于stack exchange，提问作者StableGeneous