CUDA架构

第二章 CUDA架构

2.1 CUDA的编程模型

CUDA(Compute Unified Device Architecture)，是一种由NVIDIA推出的并行计算架构，非常适合大规模数据密集型计算。CUDA使GPU的超高计算性能在数据处理和并行计算等通用计算领域发挥优势 。它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。随着显卡的发展，GPU越来越强大，在计算上已经超越了通用的CPU。如此强大的芯片如果只是作为显卡会造成计算能力的浪费，因此NVIDIA推出CUDA，让显卡可以用于图像渲染以外的目的。CUDA的GPU编程语言基于标准的C语言，通过在标准C语言的基础上增加一小部分关键字，任何有C语言基础的用户都很容易地开发CUDA的应用程序。CUDA3.0已经开始支持C++和FORTRAN。

2.1.1主机和设备

CUDA编程模型在设计结构上采用了异构编程的模式，将CPU作为宿主(Host)，GPU作为设备(Device)，在同一个系统中可以有同时存在多个设备，但是只能有一个宿主。

在CUDA程序架构中，主程序由CPU来执行，而当遇到数据并行处理的部分，CUDA就会将程序编译成GPU能执行的程序，并传送到GPU。CUDA使用被称为块（Block）的单元，每个块都由一些CUDA线程组成，线程是CUDA中最小的处理单元，将这些较小的子问题进一步划分为若干更小的细粒度的问题，我们便可以使用线程来解决这些问题了。对于一个普通的NVIDIA GPU，其CUDA线程数目通常能达到数千个甚至更多，因此，这样的问题划分模型便可以成倍地提升计算机的运算性能。GPU是由多个流水多处理器构成的，流水处理器以块（Block）为基本调度单元，因此，对于流水处理器较多的GPU，它一次可以处理的块（Block）更多，从而运算速度更快，时间更短。而反之对于流水处理器较少的GPU，其运算速度便会较慢。

CUDA C是C语言的一个扩展，它允许程序员定义一种被称为内核函数（Kernel

Functions）的C函数，内核函数运行在GPU上，一旦启动，CUDA中的每一个线程都将会同时并行地执行内核函数中的代码。内核函数使用关键字__global__来声明，运行该函数的CUDA线程数则通过<<<...>>>执行配置语法来设置。每一个执行内核函数的线程都由一个唯一的线程ID。当整个程序执行到kernel函数时，程序转入到GPU端处理程序的并行部分，当并行部分处理完成后，程序在转入到宿主端进行剩下的部分。一个完整的CUDA程序是由宿主端的串行程序和一系列的设备端kernel函数共同组成的。这些处理程序会按照程序中相应语句的顺序依次执行，满足顺序一致性。图2-1给出了CUDA的编程模型示意图。

图2-1 CUDA编程模型示意图

2.1.2 CUDA的线程结构

CUDA将线程分为了三个层次，以网格的形式组织，如图22所示[21]。CUDA的每一个线程都有其线程ID，线程的ID信息由变量threadIdx给出。threadIdx是CUDA C语言的内建变量，通常它用一个三维数组来表示。使用三维数组的方便之处在于可以很方便地表示一维、二维和三维线程索引，进而方便地表示一维、二维和三维线程块（thread block）。这样，无论是数组、矩阵还是体积的计算，都可以很容易地使用CUDA进行运算。

线程的索引与线程ID之间存在着直接的换算关系，对于一个索引为(x, y, z)的线程来说：

1、如果线程块（block）是一维的，则线程ID = x

2、如果线程块是二维的，假设块尺寸为（Dx，Dy），那么线程ID = x + y *

Dx

3、如果线程块是三维的，设其尺寸为（Dx，Dy，Dz），那么线程ID = x + y

* Dx + z * Dx * Dy

下面的例子展示了两个NxN矩阵相加的CUDA实现：

图2-2 线程层次结构图

每个线程块（block）中的线程数量是有限制的，因为依据前面所说，同一线程块（block）中的所有线程都会被分配到同一个处理器核上运行，共享有限的存储资源，因此对于当前的GPU，线程块所能包含的最大线程数目为1024。

上面的例子中numBlocks代表线程块的数量，这里的值为1。在一般的CUDA程序中，这个值通常大于1，也就是说将会有多个线程块被分配到多个处理器核中同时进行处理，这样就大大提高了程序的并行性。

在CUDA中，线程块包含在线程格（grid）当中，线程格可以是一维、二维或者三维的，线程格的尺寸一般根据待处理数据的规模或者处理器的数量来指定。线程格中所包含的线程块数目通常远远大于GPU处理器核心的数目。上图展示了线程格（grid）、线程块（block）以及线程（thread）之间的关系。

内核函数的调用可以简化为kernel<<>>(parameters)，在尖括号中，A代表线程格（grid）的尺寸，它可以是三维的，用类型dim3表示，也可以是一维的，用int类型表示。B代表线程块（block）的尺寸，它与A类似，也可分别用dim3或int类型表示。

在内核函数内部，CUDA为我们内建了一些变量用于访问线程格、线程块的尺寸和索引等信息，它们是：

1. gridDim：代表线程格（grid）的尺寸，gridDim.x为x轴尺寸，gridDim.y、gridDim.z类似。拿上图来说，它的gridDim.x = 3，gridDim.y = 2，gridDim.z

= 1。

2. blockIdx：代表线程块（block）在线程格（grid）中的索引值，拿上图来说，Block（1,1）的索引值为：blockIdx.x = 1，blockIdx.y = 1。

3. blockDim：代表线程块（block）的尺寸，blockDIm.x为x轴尺寸，其它依此类推。注意到Block（1,1）包含了4 * 3个线程，因此blockDim.x = 4,

blockDim.y = 3。

4. threadIdx：线程索引。

2.2 CUDA的存储器模型

2.2.1 寄存器

在CUDA的硬件结构中，每个SM含有一定数量的片上寄存器，每个寄存器的大小为32hit，执行单元可以以极低的代价访问寄存器。在计算能力为1.0/1.1的GPU中，每个SM共有为8192个寄存器，在计算能力为1.2/1.3的GPU中这个数目为16384。每当程序开始执行时，SM将这些寄存器平均地分配给每个线程。因此，当线程数目过多的时候，每个线程被分配的寄存器数目就极其有限。在线程内部，每个的非double类型私有变量都会占有一个寄存器，double类型占用两个，当线程定义的私有变量超出了其所拥有的寄存器数量，多余的数据将会被存放在局部存储器上(local memory)，local memory是每个线程所私有的，位于设备存储器上，其访问代价很高，因此通常编写程序的时候，要仔细的计算每个线程内部定义的变量数目是否超出了它所被分配的寄存器的数目。

CUDA也规定了存储器模型（如图2所示）和一系列用于主控CPU与GPU间通信的不同地址空间。图中红色的区域表示GPU片内的高速存储器，橙色区域表示DRAM中的的地址空间。 首先，是最底层的寄存器(register,REG)。对每个线程来说，寄存器都是私有的--这与CPU中一样。如果寄存器被消耗完，数据将被存储在本地存储器(local memory)。本地存储器对每个线程也是私有的，但是数据时被保存在帧缓冲区DRAM中。然后是用于线程间通信的共享存储器。共享存储器是一块可以被同一block中的所有thread（上节提到过，一个block最多可以有512个thread）访问的可读写存储器。访问共享存储器几乎和访问寄存器一样快，是实现线程间通信的延迟最小的方法。共享存储器可以实现许多不同的功能，如用于保存共用的计数器（例如计算循环迭代次数）或者block的公用结果（例如计算512个数的平均值）。除此以外，还有两种只读的地址空间—常数存储器和纹理存储器（constant memory and texture memory），它们是利用GPU用于图形计算的专用单元发展而来的。常数存储器空间较小(只有64k)，支持随机访问。纹理存储器尺寸则大得多，并且支持二维寻址（传统的缓存只支持一维寻址）。

这两种存储器实际存在于帧缓冲区DRAM中，但由于它们的只读性质，在GPU片内可以进行缓存，从而加快访问速度。这两种存储器并不要求缓存一致性—它们是只读的。但这也意味着如果CPU或者GPU要更改常数存储器或者纹理存储器的值，缓存中的值在更新完成之前也无法使用。CUDA程序中，常数存储器用于存储需要经常访问的只读参数，而是用插值或者滤波的纹理存储器访问对大尺寸的二维或者三维图象或者采样序列进行高带宽的流式访问。

最后是全局存储器（global memory），使用的是普通的显存。整个网格中的任意线程都能读写全局存储器的任意位置，并且既可以从CPU访问，也可以从CPU访问。由于全局存储器是可写的，GPU片内没有对其进行缓存。

图2-2 线程存储器模型

2.2.2 共享内存

共享内存是CUDA设备中非常重要的一个存储区域，有效地使用共享内存可以充分利用CUDA设备的潜能，极大提升程序性能。那么，共享内存有哪些特点呢？

1、共享内存（shared Memory）是集成在GPU处理器芯片上的（on-chip），因此相比于存在于显存颗粒中的全局内存（global Memory）和本地内存（local

Memory），它具有更高的传输带宽，一般情况下，共享内存的带宽大约是全局内存带宽的7-10倍。

2、共享内存的容量很小。根据NVIDIA官方文档的说法，在计算能力1.x的设备中，每一个流多处理器（Streaming Multiprocessor）上的共享内存容量为16KB。对于计算能力2.x、3.0及3.5的设备该参数为48KB。因此共享内存是稀有资源。

3、共享内存在物理上被划分为很多块，每一块被称为一个存储体（bank）。在同一时刻，CUDA设备可以同时访问多个存储体。因此，如果一次针对共享内存的访存操作需要读取n个地址，而这n个地址恰好分布在n个不同的存储体（bank）中，那么只需要一个存取周期就可以完成n个地址的访存任务了。对于计算能力1.x的设备，共享内存被平均划分为16个存储体。而对于计算能力2.x、3.0及3.5的设备此参数为32。在共享内存中，相邻两块32bit的数据分别属于相邻的两个存储体。存储体每两个时钟周期可以传输32位数据。

4、共享内存既可以静态分配，也可以动态分配。从共享内存的这些特点中我们可以看出，它实际上相当于一个程序员可以操控的缓存（cache），下面，我们使用矩阵乘法的例子来说明如何有效使用共享内存。 首先，我们使用最直观的方法来完成矩阵乘法C = A x B：读取A的每一行和B的每一列，顺次完成计算任务。矩阵乘法的示意图如下所示：

图2-3 使用共享内存的矩阵乘法

现在我们来使用共享内存（shared memory）实现矩阵乘法。假设矩阵C可以被划分为若干个较小的子方阵Csub，我们使用一个线程块（thread block）来负责某一子方阵的计算，线程块中的每一个线程（thread）正好负责子方阵Csub中一个元素的计算。这样划分后，任何一个结果子方阵Csub'（尺寸为block_size