线程同步

线程同步

• 同一个 warp 内的 threads 天然同步。（有硬件支持）
• block 与 block 之间是异步的，不存在相互等待。当前 kernel 执行结束后， block 之间自然同步了。
• CUDA 中没有全局同步，全局同步的系统开销会很大。
• 一个 block 内的所有 threads 有时候是需要步的。当一个 thread 执行到 __syncthreadas() 时，这个 thread 会看它所在的 block 内的其他所有 threads 情况，如果发现还有其他 threads 没有执行到这个位置，则这个 thread 等待其他 threads 。直到 block 中所有 active 的 threads 都执行到此，接着向下执行。避免相互等待进入死锁。

同步异步

遇到 cudaMemcpy() 执行变成同步的，也就是说，所有指令必须等待其他指令执行到此，才可以一起向下继续执行。如果没有 cudaMemcpy()，可以使用 cudaDeviceSynchronize() 实现同步。

t1 = myCPUTimer();
saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);
cudaDeviceSynchronize();   // 没有只一句的话，得到的时间是 CPU 调用 kernel的时间，而不是GPU执行kernel的时间。
t2 = myCPUTimer();

CPU 与 GPU 是异步的，只有加上 cudaDeviceSynchronize()，告诉 CPU 等待 GPU 把 kernel 执行完毕。