本文介绍了CUDA编程中的线程协作和同步，包括线程块、线程格的概念，硬件限制以及共享内存的使用。通过共享内存实现线程间的通信，使用`__syncthreads()`进行同步，防止竞态条件。文章通过内积运算的例子深入讲解了同步的重要性，并警告了在条件语句中使用`__syncthreads()`可能导致的线程发散问题。
接触CUDA的时间并不长，最开始是在cuda-convnet的代码中接触CUDA代码，当时确实看的比较痛苦。最近得空，在图书馆借了本《GPU高性能编程 CUDA实战》来看看，同时也整理一些博客来加强学习效果。
在上篇博文中，我们已经用CUDA C编写了一个程序，知道了如何编写在GPU上并行执行的代码。但是对于并行编程来说，最重要的一个方面就是，并行执行的各个部分如何通过相互协作来解决问题。只有在极少数情况下，各个处理器才不需要了解其他处理器的执行状态而彼此独立地计算出结果。即使对于一些成熟的算法，也仍然需要在代码的各个并行副本之间进行通信和协作。因此，下面我们来讲讲不同线程之间的通信机制和并行执行线程的同步机制。
首先，我们来看一个线程块的网格示意图：
我们将并行线程块的集合称为线程格(Grid)，在上图的Grid中总共有6个线程块(block)，每个线程块有12个线程(thread)。
硬件限制：
- 线程块的数量限制为不超过65 535;
- 每个线程块的线程数量限制为不超过512。
解决线程块数量的硬件限制的方法就是将线程块分解为线程。
共享内存
线程协作主要是通过共享内存实现的。CUDA C支持共享内存，我们可以将CUDA C的关键字__share__添加到变量声明中，这将使这个变量驻留在共享内存中。
- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
附加知识：
变量类型限定符
__device__
该限定符声明位于设备上的变量。在接下来介绍的其他类型限定符中，最多只能有一种可与__device__限定符一起使用，以更具体地指定变量属于哪个存储器空间。如果未出现其他限定符，则变量具有以下特征：
最低0.47元/天 解锁文章
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
