设备和主机之间的同步
//显式同步,阻塞直到该语句前得所有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