CUDA 学习（六）、线程网格

一、概述

       一个线程网格是由若干线程块组成的，每个线程块是二维、三维的，拥有X轴、Y轴、Z轴。此时，每次最多能开启Y*X*Z*T 个线程。

       通常线程块中线程数量最好是一个线程束大小的整数倍，即32 的整数倍。由于设备是整个线程束为单位进行调度，如果我们不把线程块上的线程数目设成32的整数倍，则最后一个线程束中有一部分线程是没有用的，因此我们必须设置一个限制条件进行限制，防止处理的元素超出X轴方向上所规定的范围。

        在程序中，要尽量避免使用小的线程块，因为这样做无法充分利用硬件。比如，一张图片分辨率为1920*1080，我们可以分配 dim3 block_num(1080, 10) ;   dim3  thread_num(192, 1);

      

在当前费米架构的硬件上，一个SM可以处理8个线程块，所以上述程序从应用层的角度来说一共需要1350个（总共10800个线程块 / 每个SM 能调度的8个线程块）SM来完成实现并行。

二、线程与线程块

       CUDA 中线程与线程块到底有什么联系？CUDA 的设计是用来将数据分解到并行的线程与线程块中。它允许我们定义一维、二维、三维的索引（Y*X*T）来方便我们在程序中引用一些并行结构。这样就使得我们程序的结构和内存数据的分布建立一一映射，处理的数据能被分配到到单独的SM 中。不论是在GPU上还是在CPU上，让数据与处理器保存紧密联系能使性能得到很大的提升。

       不过，在对数组进行布局的时候，有一点需要我们特别注意，那就是数组的宽度值最好是线程束大小的整数倍。如果不是，填补数组，使它能充满最后一个线程束。但是这样做会增加数据集的大小。此外，我们还需要注意对填充单元处理，它和数组中其他单元的处理不同的。

三、X 与Y 方向的线程索引

       在一个线程块上分布一个二维数组也意味着需要两个线程索引，这样我们才可以用到二维的方式访问数组：

   

      注意blockDim.x 与 blockDim.y 的使用，这个结构体是由CUDA 运行时库提供的，分布表示X轴和Y轴这两个维度上线程块的数目。

      例如，计算一个32*16 维的数组，假设调度四个线程块，我们可以有下面的分布方式：

    

       在实际中长方形的布局比正方向的布局优越，这是为什么呢？主要有两个原因：第一，同一个线程块中的线程可以通过共享内存进行通信，这是线程协作中一种比较快的方式。第二，在同一个线程束的线程存储访问合并在一起了，而在当前费米架构的设备中，高速缓存存储器的大小是128个字节，一次直接访问连续的128个字节比两个分布访问64个字节要高效得多。在正方形的布局中，0-15号线程映射在一个线程块中，它们访问一块内存数据，但与这块内存相连的数据区则是由另一个线程块访问的，因此，这两块连续的内存数据通过两次存储访问才获得，而在长方形的布局中，这只需要一次存储访问的操作。

      这两种布局方式，线程块和网格配置如下：

       

       线程网格、线程块及线程的维度如下：

        

     其中：

    gridDim.X ----线程网格X维度上线程块的数量

    gridDim.Y ----线程网格Y维度上线程块的数量

    blockDim.x ---一个线程块X维度上的线程数量

    blockDim.y ---一个线程块Y维度上的线程数量

    threadIdx.x ---一个线程块X维度上线程索引

    threadIdx.x ---一个线程块Y维度上线程索引

通过找出当前的行索引，然后乘以每一行的线程总数，最后加上在X轴方向上的偏移，我们便可以计算出相对于整个线程网格的绝对线程索引。具体点如下：

   thread_idx = ((gridDim.x * blockDim.x) * idy) + idx;

  其中：

         idx = (blockIdx.x * blockDim.x) + threadIdx.x;

         idy = (blockIdx.y * blockDim.y) + threadIdx.y;