CUDA从入门到精通（八）：线程通信

我们前面几节主要介绍了三种利用GPU实现并行处理的方式：线程并行，块并行和流并行。在这些方法中，我们一再强调，各个线程所进行的处理是互不相关的，即两个线程不回产生交集，每个线程都只关注自己的一亩三分地，对其他线程毫无兴趣，就当不存在。。。。

 

当然，实际应用中，这样的例子太少了，也就是遇到向量相加、向量对应点乘这类才会有如此高的并行度，而其他一些应用，如一组数求和，求最大（小）值，各个线程不再是相互独立的，而是产生一定关联，线程2可能会用到线程1的结果，这时就需要利用本节的线程通信技术了。

 

线程通信在CUDA中有三种实现方式：

1. 共享存储器；

2. 线程 同步；

3. 原子操作；

 

最常用的是前两种方式，共享存储器，术语Shared Memory，是位于SM中的特殊存储器。还记得SM吗，就是流多处理器，大核是也。一个SM中不仅包含若干个SP（流处理器，小核），还包括一部分高速Cache，寄存器组，共享内存等，结构如图所示：

从图中可看出，一个SM内有M个SP，Shared Memory由这M个SP共同占有。另外指令单元也被这M个SP共享，即SIMT架构（单指令多线程架构），一个SM中所有SP在同一时间执行同一代码。

 

为了实现线程通信，仅仅靠共享内存还不够，需要有同步机制才能使线程之间实现有序处理。通常情况是这样：当线程A需要线程B计算的结果作为输入时，需要确保线程B已经将结果写入共享内存中，然后线程A再从共享内存中读出。同步必不可少，否则，线程A可能读到的是无效的结果，造成计算错误。同步机制可以用CUDA内置函数：__syncthreads()；当某个线程执行到该函数时，进入等待状态，直到同一线程块（Block）中所有线程都执行到这个函数为止，即一个__syncthreads()相当于一个线程同步点，确保一个Block中所有线程都达到同步，然后线程进入运行状态。

 

综上两点，我们可以写一段线程通信的伪代码如下：

 

//Begin

if this is thread B

     write something to Shared Memory;

end if

__syncthreads();

if this is thread A

    read something from Shared Memory;

end if

//End

 

上面代码在CUDA中实现时，由于SIMT特性，所有线程都执行同样的代码，所以在线程中需要判断自己的身份，以免误操作。

注意的是，位于同一个Block中的线程才能实现通信，不同Block中的线程不能通过共享内存、同步进行通信，而应采用原子操作或主机介入。

对于原子操作，如果感兴趣可以翻阅《GPU高性能编程CUDA实战》第九章“原子性”。

 

本节完。下节我们给出一个实例来看线程通信的代码怎么设计。