CUDA并行架构及编程
摘 要 CUDA是一种由NVIDIA推出的并行计算架构，非常适合大规模数据密集型计算。CUDA使GPU的超高计算性能在数据处理和并行计算等通用计算领域发挥优势，本文讨论了CUDA的计算架构和基于GPU的CUDA C编程语言，CUDA使GPU流处理器阵列的性能得到充分发挥。极大地提高了并行计算程序的效率。
关键词 并行计算，GPU通用计算，CUDA
Abstract CUDA is a parallel computing architecture introduced by NVIDIA , it mainly used for large scale data-intensive computing. CUDA makes GPU a high performance in parallel computing ,data processing and other general computing. this paper discusses the CUDA computing architecture and CUDA C programming language based on GPU, CUDA makes GPU stream processor arrays fully used and Greatly improved the efficiency of parallel computing program.
Keywords parallel computing ，GPU general purpose computation，CUDA
1 引言
并行计算是指同时使用多种计算资源解决计算问题的过程。并行计算科学中主要研究的是空间上的并行问题。从程序和算法设计的角度来看，并行计算又可分为数据并行和任务并行。一般来说，GPU更注重于数据并行计算，主要是将一个大任务化解成相同的各个子任务。早期的GPU研究是通过可编程计算单元为屏幕上的每个像素计算出一个颜色值即渲染问题。自CUDA C出现后，基于GPU的通用计算已经成为一个新的研究领域。通常，像素着色器对各种颜色值进行合成并计算出最终的颜色值。实际上，输入值可以为任意数据，这样不一定非要使用GPU来处理图形，还可以实现某些通用计算。由于GPU有着很高的计算吞吐量，从而给大规模的数据计算应用提供了一种比CPU更加强大的计算能力。
CUDA是一种由NVIDIA推出的并行计算架构，该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构以及GPU内部的并行计算引擎[1]。随着显卡的发展，GPU越来越强大，在计算上已经超越了通用的CPU。如此强大的芯片如果只是作为显卡会造成计算能力的浪费，因此NVIDIA推出CUDA，让显卡可以用于图像渲染以外的目的。CUDA的GPU编程语言基于标准的C语言，通过在标准C语言的基础上增加一小部分关键字，任何有C语言基础的用户都很容易地开发CUDA的应用程序。数以千计的软件开发人员正在使用免费的CUDA软件开发工具来解决各种专业中的问题[2]。这些解决方案涵盖了石油天然气勘探、产品设计、医学成像以及科学研究等领域。
2 CUDA架构
CUDA程序架构分为两部分：主机和设备。一般而言，主机指的是CPU及
其内存，设备指的是GPU[1]。在CUDA程序架构中，主程序由CPU来执行，而当遇到数据并行处理的部分，CUDA就会将程序编译成GPU能执行的程序，并传送到GPU。这种函数在CUDA里叫做核函数。在GPU中要执行的线程，根据最有效的数据共享来创建线程块，其类型不止一维。在同一个块里的线程，使用同一个共享内存。每个线程由线程ID标识。这是线程块中的线程号。为了帮助基于线程ID的复杂寻址，还可以将线程块指定为任意大小的一维、二维或高维线程阵列，并使用多个索引分量来标识每个线程[2]。线程块的大小是也有限制的，所以不能把所有的线程都放到同一个块里。可以用同样维度和大小的块来组成一个网格做批处理。执行内核的线程被组织成线程块。而线程块又组成了线程格，如图1。线程块内一起协作的线程通过一些快速的共享内存有效地共享数据并同步执行，以协调内存访问。用户可以在核函数中指定同步点。线程块中的线程在到达此同步点时挂起。
图1 线程被组织为线程格与线程块 图2 SM单元的组成
在 GPU 内部， SM 代表流处理器，即计算核心。每个 SM 中又包含 8 个标量流处理器（SP）以及少量的其他计算单元[3]，如图2。实际上SP只是执行单元，并不是完整的处理核心。处理核心必须包含取指、解码、分发逻辑和执行单元。隶属同一 SM 的 8 个 SP 共用同一套取指和发射单元，也共用一块共享存储器。一个线程块必须被分配到一个 SM 中，但是一个 SM 中同一时刻可以有多个活动线程块等待执行。这可以更好地利用执行单元的资源，当一个线程块进行同步或者访问显存等高延迟操作时，另一个线程块就可以占用 GPU 执行资源。每个线程有一个私有的本地存储器，即每个标量处理器私有的缓存。每个线程块有一个共享存储器，块内所有的线程都可以访问。每个线程网格都可以访问设备内存，即GPU片上的全局存储器。每个线程均可以读写全局存储器。本地存储器用于储存设备代码中无法放进寄存器的变量[3]。
CUDA架构的所有这些功能都是为了使GPU不仅能执行传统的图形计算，还能高效的执行通用计算。
3 CUDA C编程
3.1 CUDA并行编程
运行在GPU上的CUDA并行计算函数称为核函数（kernel）。一个 kernel 函数并不是一个完整的程序，而是整个程序中的一个可以被并行执行的步骤[2]。内核函数必须通过 _global_ 函数类型限定符定义，如_global_ void kernel（void）。并且只能在主机端代码中调用。在调用时，必须声明内核函数的执行参数。现考虑给一个N位矢量的每一位加上一个常数。为完成此操作设计一个add函数，分别用C代码和CUDA C表示为：
#define N 10 #define N 10
void add（int *a,int b）{ _global_void add(int*a,int b) int index=0; {
while(index a[index]=a[index]+b; if(index 在CPU运行的程序中，通过while循环串行的给N位矢量每一位加上常数。值得注意的是右边的CUDA核函数，在实际运行时，CUDA会产生许多在GPU上执行的线程，每一个线程都会去执行内核这个程序，虽然程序是同一份，但是因为设置了变量blockIdx，这是一个内置变量，在CUDA运行中已经预先定义了这个变量，变量的值是当前执行设备代码线程块的索引，从而取得不同的数据来进行并行计算。这里考虑的实例是10位矢量的增量运算，当矢量的位数N远远大于10且进行的并不是简单的增量运算时，使用CUDA C就会有明显的差别。 int main(void){ int host_a[N]; int *device_a; cudaMalloc((void**)&device_a,N * sizeof(int)); //在设备上分配内存 for (int i=0; i cudaMemcpy(device_a, host_a, N * sizeof(int),cudaMemcpyHostToDevice); add<< cudaMemcpy( host_a,device_a,N*sizeof(int),cudaMemcpyDeviceToHost); cudaFree(device_a); //释放内存 return 0; } 在定义了核函数后，可以在主函数中调用它。add<< 3.2 线程的同步与通信 上一节矢量增值运算的例子，通过创建N个含有一个线程的进程快给矢量的各位各加上一个常数。上述运算并不需要考虑线程块内线程的同步与通信，但实际中的大多数应用都需要线程之间传递数据。这也是CUDA所提供的最重要的创新，它使得工作在GPU上的线程可以协作解决问题，允许应用程序更加高效的执行[2]。 在同一个 block 中的线程通过同步函数_syncthreads（）来保证块内的线程同步即线程块中的所有线程都执行到同一位置。为了在硬件上用很小的代价就能实现 _syncthreads() 函数，一个 block 中所有线程的数据都必须交由同一处理核心进行处理[3]。所以，这导致每个线程块中的线程数量受到处理核心硬件资源的限制。目前规定每个块里最多只能有 512 个线程[2]。 在同一个块中的线程通过共享存储器交换数据，为了保证线程块中的各个线程能够有效协作，访问共享存储器的延迟必须很小。所以在 GPU 中，共享存储器与执行单元的物理距离必须很小。在CUDA C中同一个块中的线程通过共享变量来实现通信，使用__shared__来定义一个共享变量。现利用一个大规模矩阵相乘的例子来说明线程块内线程同步与通信及线程之间的互斥。矩阵A（MxN）与矩阵B（NxK）的乘法公式为[5]： Cij??ait?bti,i?0,1,...m?1;j?0,1,...k?1 （1） t?0n?1为简单起见，这里仅考虑1xN矩阵与Nx1矩阵之积，于是公式简化为： a1a2N b1b2..bN?.??aibi （2） i?1.aN 可以先让每个进程计算aibi，由于要计算累加和每个进程需要保存一个临时变量，每计算完一个积，索引平移blockDim.x * gridDim.x个单位。这样线程块内的每个线程都保存着若干个aibi，还需要在每个线程块内设置一个共享数组变量__shared__ int var[threadsperblock]，这个数组的大小是线程块里线程的数量。这样线程块里每个线程都将自己保存的累加值赋给var[threadsperblock]。在给var赋值完毕后将var数组累加起来储存在var[0]中，这时需要调用同步函数_syncthreads（）确保进程块内所有进程都已完成计算方可进行累加。 至此，进程格中的每个进程块都有了一个累加和var[0]，需要再将这些属于各进程块的var[0]再累加起来储存在一个变量c中，这个操作可由线程块中任意一个线程来完成，不妨设为线程0。最终将结果c送回CPU。在第二步的累加中，为保证变量c被互斥的访问，需要一个原子锁lock（）。由上面的讨论可得核函数的主要部分： __global__ void matrix( Lock lock, float *a,float *b, float *c ) { __shared__ float var[threadsperblock]; int index = threadIdx.x + blockIdx.x * blockDim.x; int varnum = threadIdx.x; //初始化索引 float temp = 0; while (index < N) {temp += a[index] * b[index]; index += blockDim.x * gridDim.x; //索引偏移 } var[varnum] = temp; __syncthreads(); //同步,等待其他进程 int i = blockDim.x/2; while (i != 0) { if (varnum < i)var[varnum] += var[varnum + i]; __syncthreads(); i /= 2; } if (varnum == 0) { //互斥访问 lock.lock(); *c += var[0]; lock.unlock(); } } 在上面的核函数最后一部分使用了一个Lock结构，这个结构中包含了一个atomicCAS函数，它将判断和赋值的过程原子化[2]。使用Lock结构的好处是连加的过程可以在GPU上进行，CPU得到的是最后结果。核函数定义好后就可在主函数中调用，main函数的流程与上节相似，主要是同时启动的核函数数量不同，matrix<< matrix<< 4 结语 GPU在特定的科学计算方面比CPU有更强大的运算能力，如气象和医学图 像方面的相关算法，数据量都是非常的大，运用CUDA的计算能力，可以很显著的提高计算速度，具有广阔的应用前景。CUDA新技术为科学领域进行大规模运算提供了新的研究方法，由于GPU的特点是处理密集型数据和并行数据计算，因此CUDA非常适合需要大规模并行计算的领域。目前CUDA除了可以用C语言开发，也已经提供FORTRAN的应用接口，未来可以预计CUDA会支持C++、Java、Python等各类语言。但是想要充分利用GPU的计算能力，需要将计算问题进行合理的分解以适应GPU的计算模式，发挥多级内存和大规模并行计算的优势。 参考文献 [1]Shane Ryoo，Christopher I. Rodrigues，Sara S. Baghsorkhi，Sam S. Stone，David B. Kirk，Wen-mei．Optimization principles and application performance evaluation of a multithreaded GPU using CUDA［J］．Proceedings of the 13th ACMSIGPLAN Symposium on Principles and practice of parallel programming，2008． [2]Jason Sanders，Edward Kandrot．CUDA by Example[M].Addison WesleyProfes- -sional ，2010 ． [3]钱悦，图形处理器CUDA编程模型的应用研究[J]． 计算机与数字工程V36 No 12 ，2008． [4] John Nickolls ，Ian Buck，Michael Garland，Kevin Skadron．Scalable Parallel Programming with CUDA[J]．GPU Computing Queue Homepage archive Volume 6 Issue 2，2008． [5]肖江，胡柯良，邓元勇．矩阵乘法和FFT性能测试[J]．计算机工程35卷第 10期 2009． [6]刘勇，使用GPU加速通用科学计算[J]，科技信息，2008年24期．