同步

来源：互联网发布：python qt gui快速编程编辑：程序博客网时间：2024/06/11 23:22

设备和主机之间的同步

//显式同步，阻塞直到该语句前得所有kernel调用执行完毕cudaError_t cudaDeviceSynchronize(void);//隐式同步cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

线程块内的同步

同一线程块内的同步方式有两种：
Barriers
Memory Fences

__device__ void __syncthreads();__device__ void __threadfence_block();

Weakly-Ordered Memory Model指的是gpu对内存的写入操作未必是立即完成的，线程可能等到需要读取该段内存时，才将这之前的写入操作完成
当线程同步语句包含在执行分支中时，要确保同一线程束内的所有线程都执行相同的分支，否则程序可能僵死
```
if(threadIdx.x % 2 == 0) {    __syncthreads();}else {    __syncthreads();}
```
实际上，memory fence支持块内同步，网格内同步，主机-设备，设备-设备同步。调用memory fence的线程会阻塞直到对应的线程集合中所有对全局内存，共享内存，页锁定内存，其他设备的全局内存的写入都完成
```
void __threadfence_block();         //块同步void __threadfence();               //网格同步void __threadfence_system();        //系统同步
```

线程块之间同步

cuda不直接支持块间barrier，但由于同一stream中的多个kernel的执行是串行的，可以把需要块间同步的kernel分割成多个kernel，每个分割点对应一个barrier

0 0