1. 引言

前序博客：

• CUDA简介——基本概念
• CUDA简介——编程模式
• CUDA简介——For循环并行化
• CUDA简介——Grid和Block内Thread索引
• CUDA简介——CUDA内存模式

本文重点关注Thread同步和Barriers。

Threads并行执行，可能存在如下问题：

• 1）Race condition条件竞争：Thread A会在Thread B写入结果之前，读取某地址的值。

为此，需要引入Thread同步机制：

• 强迫一部分Device代码顺序执行，以强迫同一Block内的Threads同步。
  具体可为：
  • 1）实现Explicit Barrier：Barrier为Kernel内某个point，在该point，Block内所有Threads会stop并相互等待。当Block内所有Threads都到达该Barrier时，会继续各自执行。
    具体实现方式为：__syncthreads();
    
    
    

以数组左移为例：

• 由于a[i]=a[i+1]为读写操作，需确保a[i]先读后写。为此，需引入名为temp的register。并__syncthreads;等待所有读取操作完成。
• 为确保所有位移操作均已结束，再返回位移后的结果，需__syncthreads;等待所有写操作完成。

除此之外，还可实现Kernel launches间的Implicit Barrier：

• Host代码并不会等待Device代码执行结束返回后，再继续执行后续Host代码。即Host代码和Device代码是异步执行的。
• 为让Host代码等待kernel执行完成，需使用关键字：cudaDeviceSynchronize()。这样，Host代码会暂停，直到前一kernel执行完成。
  

不过，若连续启动2个kernel，则确保第二个kernel无法分配grid到device中执行，其implicitly需等待第一个kernel执行结束后，才会执行第二个kernel。

参考资料

[1] Intro to CUDA (part 6): Synchronization