炼数成金 门户 CUDA 查看内容

CUDA线程通信

2015-9-9 12:51| 发布者: 炼数成金_小数| 查看: 1342| 评论: 0|原作者: 卜居|来自: csdn

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

存储 架构 GPU 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中的线程不能通过共享内存、同步进行通信,而应采用原子操作或主机介入。

鲜花

握手

雷人

路过

鸡蛋

最新评论

热门频道

  • 大数据
  • 商业智能
  • 量化投资
  • 科学探索
  • 创业

热门文章

     

    GMT+8, 2020-10-30 13:17 , Processed in 0.164336 second(s), 23 queries .