GPU并行计算的CUDA架构浅析

吴辉 罗清海 彭文武

摘要：本文阐述了GPU并行运算的一种主流架构——CUDA架构，包括CUDA编程模型、程序的运行模式、线程架构、存储器结构、指令结构等。

关键词：GPU；CUDA架构；并行计算

中图分类号：G642.41 文献标志码：A 文章编号：1674-9324（2019）06-0277-02

从20世纪90年代，第一款图形处理器诞生到现在，摩尔定律预测到由于制造工艺和设计的进步，单一芯片内部集成的晶体管数量大约每隔一年都会翻一翻，但是随着工艺进步的空间日趋缩小，制造工艺反而成了一种限制，CPU上的晶体管不能无限增加，所以怎么把任务更均衡地分配给GPU，实现CPU-GPU的负载均衡，是当代计算机技术的一个重要研究方向。

而CUDA是英伟达公司2006年推出的通用并行计算架构，其提供直接的硬件访问接口，使得程序开发人员不必再依赖应用程序编程接口（Application Programming Interface，API）来实现GPU的交互通信，完整地解决GPU设备的通用并行计算的程序开发及实现的问题，CUDA平台主要包含三个部分：开发库CUFFT（离散快速傅立叶变换）和CUBLAS（离散基本线性计算）等、运行期环境（宿主代码和设备代码）和驱动。

一、CUDA编程模型

CUDA的编程模型为中央处理器CPU与图形处理器GPU联合运行的异构编程平台，即也就是在该平台上，程序同时使用了CPU和GPU，其中CPU作为主机端（Host），GPU作为设备端（Device）或协处理器，CPU利用强大的逻辑门判断功能，能够快速解释计算机指令，分配数据处理事物和串行计算，并指挥设备端的运作，而GPU的逻辑门判断功能简单，但是算术逻辑单元数量众多，数据计算能力十分强大，负责处理可以高度并行化的任务。在Nvida显卡平台，CPU和GPU之间通过PCI-E数据总线进行数据传递，以保证数据传输拥有足够的带宽。具体执行模式是主机端（CPU）运行的程序段通过kernel函数来控制设备端（GPU）的程序的运行；此外，它们拥有各自互相独立的内存空间：显存和内存；另外CUDA对储存空间的调配方式也不同，CUDA平台使用调配函数CUDA API来调配设备端（GPU）上的显存进行空间开辟、空间释放、初始化显存等动作，对主机端（CPU）上的内存管理，则使用C语言集成的内存管理函数来读写。

二、CUDA程序的运行模式

a.开辟主机端存储器（主存）空间并进行初始化。

b.开辟设备端存储器（显存）空间。

c.将已初始化的主机端存储器的数据通过PCI-E数据总线复制并传递到设备端存储器上的开辟空间内。

d.GPU进行并行计算。

e.将GPU上已处理完成的数据拷贝回主机端。

f.主机端进行数据处理。

三、GPU的线程架构

GPU的线程架构是一种逻辑架构，这样使开发人员不必深入地学习每一个GPU的内部物理结构，从而从繁杂的计算机图形学工作中脱离出来，事实上，该线程架构包含了由浅入深的三个部分，由大到小依次为线程格（grid）、线程块（block）和线程（thread）。一个kernel函数映射到GPU上后，称为一个网格（grid）；一个线程格由若干个线程块构成，线程块与线程块之间通过共享存储器进行数据传输，并把数据发射到流多处理器（Streaming Multiprocessor，SM）上执行，这些线程块由数量和排列方式能够组成一维和二维结构，且每个线程块都用一个唯一的坐标（blockIdx）进行区别编号；同时，每个线程块都包含了若干个线程（thread），线程在组织和结构上与线程块相类似，但是线程能组成三维结构。线程块中的线程通过编号变量threadIdx.x、threadIdx.y、threadIdx.z中的x、y、z来使用线程，并构成在线程块中的一维、二维和三维中的索引标识。

四、CUDA存储器架构

CUDA存储器分为GPU片内内存：寄存器（register）、共享内存（shared memory，SM）、本地内存（local memory，LM）；板载内存：常量内存（constant memory，CM）、纹理内存（texture memory，TM）、全局内存（global memory，GM）。

共享内存位于图形处理器内部，每个流多处理器SM能使用的共享内存分为16KB、64K等几种，如果要获取较高的存储器带宽，则流多处理器上的共享内存一般被分割合并成一些16位或者32位的Bank，线程访问Bank的速度与寄存器的访问速度在理论上能达到一致，实际上由于几个线程读写相同的Bank时产生冲突，而造成存取延时，因此共享寄存器的延迟有30—40个时钟。本地内存用来存储寄存器溢出的数据，通常是可能消耗了大量寄存器空间的大数组或大结构、无法确定大小的数组和内核使用的寄存器溢出的任何变量。另外，由于本地内存存在于板载内存空间，所以所有的本地内存的读写访问延迟与全局内存一样搞，带宽和全局内存一样低。全局内存位于设备存储器中，储存空间最大，但距离运算执行单元流处理器最远，存取速度最慢的存储器，只有约177GB/s，访问延时高达400—600个时钟中期。设备存储器通过32，64或128字节的存储器通信访问，当一个束需要读取全局内存时，它会采用合并访问（coalesced access）的形式，将束内线程的存储器的访问一次或者多次完成，以提高访问效率。整合内存是指，若干线程的内存请求在硬件上被合并和整合成一次多数据的数据传输，从而提高全局内存的存取速度，这是实现高性能GPGPU编程极为重要的手段。常量内存是一段位于设备存储器上的只读地址空间，特点是对其访问可以获得缓存加速，如果从常量内存中获得所需要的常量时，只要一个时钟周期即可返回，大大提高了程序运行的速度，但是常量内存的容量一般只有64K大小，故只存储可能被频繁读取的只读参数。常量内存在函数体外使用_constant_变量类型限制符进行存放，能够被GPU的所有线程访问。本文的程序的权系数、流体的参数、系数矩阵等存储于常量内存空间。纹理内存的位置在设备存储器上，能够将部分全局内存的数据缓存到纹理缓存中，但是纹理内存也是一段只读的内存空间。

五、CUDA指令结构

可拓展的多线程流多处理器（SMs）架构是CUDA架构的核心，如果运行在主机端的CUDA程序调用kernel核函数后，线程格内块枚举并分发到处于空闲状态的流多处理器上执行，SM设计为能同时控制数百的线程进行运算，因此，流多处理器采用单指令多线程（SIMT）方式来同时操控如此多的线程。

六、结论

本文所述的GPU并行计算架构为NVIDIA的CUDA计算平台，这是当前GPU并行计算应用及研究的主流架构，本文通过介绍该架构的主要结构及运行模式，为其在各个领域的应用提供前瞻性的参考。

参考文献：

[1]雷德川，陈浩，王远，张成鑫，陈云斌，胡栋材.基于CUDA的多GPU加速SART迭代重建算法[J].强激光与粒子束，2013，（09）：2418-2422.

[2]岳俊，鄒进贵，何豫航.基于CPU与GPU/CUDA的数字图像处理程序的性能比较[J].地理空间信息，2012，（04）：45-47+180.