同步

设备和主机之间的同步

//显式同步，阻塞直到该语句前得所有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