CUDA 学习（四）、线程

一、概述

       英伟达的硬件调度是采用SPMD(单程序多数据)，属于SIMD（单指令多数据）的一种变体。线程是并行程序的基本构建块。

二、CPU与GPU的差异

       GPU和CPU设备的架构是迥异的。CPU的设计是用来运行少量比较复杂的任务。GPU的设计则是用来运行大量比较简单的任务。CPU的设计主要是针对大量离散而不相关任务系统。而GPU的设计主要是针对解决那些可以解决成千上万个小块并可独立运行的问题。因此，CPU适合运行操作系统和应用程序软件。

        CPU与GPU支持线程的方式不同。CPU的每个核只有少量的寄存器，每个寄存器都将在执行任何已分配的任务中被用到。为了能执行不同的任务，CPU将在任务与任务之间进行上下文切换。从时间的角度来看，CPU上下文切换的代价是非常昂贵的，因为每一次上下文切换都要将寄存器里的数据保存到RAM中，等到重新执行这个任务时，又从RAM中恢复。相比之下，GPU同样用到上下文切换只需设置一个寄存器组调度者，用于将当前寄存器组里的内容换进、换出，它的速度比将数据保存到RAM中要快好几个数量级。

       CPU和GPU的一个主要差异就是每台设备上处理器数量的巨大差异。CPU是典型的双核或四核设备。也就是说它有一定数量的执行核可供程序运行。而目前费米架构的GPU拥有16个SM（流多处理器簇），每个SM可以看作是CPU的一个核。GPU默认就是并行的模式。GPU为每个SM提供了唯一并且高速的寄存器，即共享内存。

三、任务执行模式

GPU的指令队列中的每条指令都会分配到SM的每个SP中。每个SM就相当于嵌入了N个计算核心（SP）的处理器。GPU所用的SPMD模式是将同一条指令送到N个逻辑执行单元，也就是说GPU只需要相对于传统的处理器1/N的指令带宽。

                 

四、GPU 线程

        看下面一段代码，就是内核函数：

        

        __global__ 前缀是告诉编译器在编译这个函数的时候生成是GPU代码而不是CPU代码，并且这段GPU代码在CPU上是全局可见的。

        CPU和GPU有各自独立的内存空间，因此在GPU代码中，不可以直接访问CPU端的参数，反过来在CPU代码中，也不可以直接访问GPU端的参数。

        线程的信息是由一个结构体存储的。如，threadIdx.x

         

         线程0中的thread_idx值为0，线程1的为1，依次类推，线程127中的thread_idx值为127。每个线程都进行了两次读内存操作，一次乘法操作，一次存储操作，然后结束。我们注意到，每个线程执行的代码是一样的，但数据却不相同。这就是CUDA的核心---SPMD模型。

五、硬件执行线程

       GPU 的每个线程组被送到SM 中，然后N 个SP 开始执行代码。所有的线程不是立即发生的。实际上，当从存储子系统取得所需要的数后，已经过去了400-600个GPU时钟周期。这一期间，这一组中的N个线程都被挂起。

       实际上，线程都是以每32个一组，当所有32个线程都在等待诸如内存读取这样的操作时，他们就会被挂起。这些线程组叫线程束（32个线程）或半个线程束（16个线程）。

       因此，我们可以将这128个线程分成4组，每组32个线程。首先让所有的线程提取线程标号，计算得到数组地址，然后发出一条内存获取的指令。接着下一条指令是做乘法，但这必须是在从内存读取数据的时间很长，因此线程会挂起。当这组中的32个线程全部挂起，硬件就会切换到另外一个线程束。

               

        在图5-5 中我们可以看到，当线程束0 由于内存读取操作而挂起时，线程束1 就成为了整治执行的线程束。GPU一直以此种方式运行直到所有的线程束到成为挂起状态（如图5-6所示）。

        当连续的线程发出读取内存的指令时，读取操作会被合并或组合在一起执行。由于硬件在管理请求时会产生一定的开销，因此这样做将减少延迟（响应请求的时间）。由于合并，内存读取会返回数组线程所需要的数据，一般可以返回整个线程所需的数据。

        在完成内存读取之后，这些线程将再次置成就绪状态，当再次遇到阻塞操作时，例如另一个线程进行内存读取，GPU可能将这个线程束用作另一块内存读取。

        当所有的线程束（每组32个线程）都在等待内存读取操作完成时，GPU将会闲置。但到达某个时间点之后，GPU将从存储子系统返回一个内存块序列，并且这个序列的顺序通常与发出请求的顺序是一致的。

        假设数组下标为0-31 的元素在同一时间返回，线程束0进入就绪队列。如果当前没有如何线程束正在执行，则线程束0将自动进入执行状态（如图5-4）。渐渐地其他所有挂起的线程束也都完成了内存读取操作，紧接着它们也会返回到就绪队列。

       一旦线程束0 的乘法指令执行完毕，它就只剩下一条指令需要执行，即将计算得到的结果写入相同下标的数组a 中。由于再没有依赖该操作的其他指令，线程束0 全部执行完毕然后消失。其他线程束也像这样，最终发出一条写数据的请求，完成之后便消亡。当所有的线程束都消亡后，整个内核函数也就结束了，最终将控制返回到CPU端。