基于CUDA的GPU高性能计算

基于CUDA的GPU高性能计算

邹佳林 1420501011主要内容GPU以及GPU高性能计算技术介绍CUDA编程简介基于CUDA的GPU高性能计算的一个实例2一、GPU及GPU高性能计算技术1.1GPU简介：GPU英文全称GraphicProcessingUnit，中文翻译为“图形处理器。GPU有非常多的厂商都生产，和CPU一样，生产的厂商比较多，我们大家熟悉的有3个，分别是AMD、

NVIDIA、Intel这3个生产厂商。

Intel集成GPU英伟达TEGRA4GPUAMDRadeonR5M23031.1GPU简介GPU的发展大致可以分成四个阶段：第一个时代是1991年以前，CPU作为系统内唯一的通用处理器，包揽了图形处理在内的所有计算任务。第二个时代是1991-2001年。微软的Windows操作系统极大刺激了图形硬件的发展。第三个时代是2001-2006年。各种硬件加速计算的出现使显卡的性能突飞猛进。标志性的事件就是可编程的图形处理器的出现。第四个时代是2006年至今，这一时期的CPU从硬件设计之初就开始考虑到了GPGPU的应用。2006年NVDIA公布了统一着色器模型（UnifiedShaderModel）和GPU发展阶段GeForce8系列GPU，GPU从此进入了通用计算时代。4时间GPU特点1991年以前显示功能在CPU上实现1991年~2001年多为二维图形运算，功能单一2001年~2006年可编程图形处理器2006年至今统一着色器模型，通用计算GPU1.2GPU与CPU的区别51.2GPU与CPU的区别CPU由专为顺序串行处理而优化的几个核心组成。而GPU则由数以千计的更小、更高效的核心组成，这些核心专为同时处理多任务而设计，可高效地处理并行任务。61.2GPU与CPU的区别现如今GPU已经不再局限于3D图形处理了，在浮点运算、并行计算等部分计算方面，GPU可以提供数十倍乃至于上百倍于CPU的性能。71.2GPU与CPU的区别1.CPU是计算机的运算和控制核心，GPU主要用来做图形处理。2.由于其设计目标的不同，CPU需要很强的通用性来处理各种不同的数据类型，同时逻辑判断又会引入大量的分支跳转和中断的处理。而GPU面对的则是类型高度统一的、相互无依赖的大规模数据和不需要被打断的纯净的计算环境。3.CPU与GPU的区别还存在于片内的缓存体系和数字逻辑运算单元的结构差异：CPU虽然有多核，但核心总数没有超过16，每个核都有足够大的缓存和足够多的数字和逻辑运算单元，并辅助有很多加速分支判断甚至更复杂的逻辑判断的硬件；GPU的核数远超CPU，如NVIDIAFermi就有512个核。81.3GPU高性能计算技术什么是GPU加速的计算?GPU加速的计算是利用一颗图形处理器(GPU)以及一颗CPU来加速科学、工程以及企业

级应用程序。应用程序如何利用GPU实现加速？

密集计算代码（约占5%的代码量）由GPU负责完成，剩余串行代码由CPU负责执行。91.3GPU高性能计算技术GPGPU(GeneralPurposecomputingongraphicsprocessingunits，基于GPU的通用计算)。GPGPU并不是单纯的使用GPU进行通用计算，而是一种利用异构计算资源的大规模并行计算。异构计算：CPU+GPU是一个强大的组合，因为CPU包含几个专为串行处理而优化的核心，而GPU则由数以千计更小、更节能的核心组成，这些核心专为提供强劲的并行性能而设计。程序的串行部分在CPU上运行，而并行部分则在GPU上运行。GPU已经发展到成熟阶段，可轻松执行现实生活中的各种应用程序，而且程序运行速度已远远超过使用多核系统时的情形。未来计算架构将是并行核心GPU与多核CPU共同运行的混合型系统。101.3GPU高性能计算技术计算模型的分类

1.单指令单数据流（SISD）是非并行计算模型。

2.单指令多数据流（SIMD）是GPU的计算模型。

3.多指令单数据流（MISD）指在同一个数据流上执行不同的指令。4.多指令多数据流（MIMD）是多核CPU的计算模型。图：并行处理的费林分类法11Single

InstructionMultipleInstructionSingleDataSISDMISDMultipleDataSIMDMIMD1.3GPU高性能计算技术并行计算模型：SIMD

并行计算指的是在同一时刻存在多于一个计算任务被执行。SIMD的并行思路是让不同的线程处理它所对应的那部分数据。当线程数大于或者等于数据个数时，理论计算时间相当于处理一个数据的时间；如果线程数少于数据个数，则某些线程处理的数量会增加。

CUDA提出的SIMT（SingleInstructionMultipleThreads）属于SIMD的范畴，因为它也是在多个数据上执行相同的指令，SIMT允许由用户来分配线程，具体来说就是CUDA为每个线程指定了标识符（编号）。121.3GPU高性能计算技术SIMD的两大特点（即要使用GPU做并行计算，算法必须满足以下两点）：1）每个线程的任务互不相关。2）每个线程执行相同的指令。具有以下特点的算法能够在GPU上达到最高的执行效率：1）每个数据（数据包）都需要经过相同的流程来处理。2）数据之间没有相干性，即某些数据的计算不依赖于另外一些数据的计算结果。3）数据量庞大。13二、基于CUDA的GPU编程2.1CUDA简介：CUDA(ComputeUnifiedDeviceArchitecture)，是显卡厂商NVIDIA推出的运算平台。CUDA是一种由NVIDIA推出的通用并行计算架构，该架构使GPU能够解决复杂的计算问题。它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。开发人员现在可以使用C语言来为CUDA架构编写程序，C语言是应用最广泛的一种高级编程语言。所编写出的程序于是就可以在支持CUDA的处理器上以超高性能运行。142.1CUDA简介计算行业正在从只使用CPU的“中央处理”向CPU与GPU并用的“协同处理”发展。为打造这一全新的计算典范，NVIDIA发明了CUDA（ComputeUnifiedDeviceArchitecture，统一计算设备架构）这一编程模型，是想在应用程序中充分利用CPU和GPU各自的优点。现在，该架构现已应用于GeForce（精视）、ION（翼扬）、Quadro以及TeslaGPU（图形处理器）上，对应用程序开发人员来说，这是一个巨大的市场。15CPU+GPU协同处理2.1CUDA简介CUDA体系结构CUDA体系结构开发库运行环境驱动16CUDA开发环境CUDA开发环境nvccC语言编译器分析器CUDA编程手册2.1CUDA简介CUDA开发环境：

nvccC语言编译器：

适用于GPU（图形处理器）的CUDAFFT和BLAS库。

分析器：

适用于GPU（图形处理器）的gdb调试器（在2008年3月推出alpha版）

CUDA运行时（CUDAruntime）驱动程序（目前在标准的NVIDIAGPU驱动中也提供）。CUDA编程手册：CUDA开发者软件开发包（SDK）提供了一些范例（附有源代码），以帮助使用者开始CUDA编程。172.2CUDA执行过程CUDA模型的计算流程大致可以分为四个部分1）将待处理数据从内存传送到显存中。2）将程序指令从CPU转移到GPU中。3）在设备端GPU上执行相关指令，完成对显存数据的操作。在计算过程中可能涉及到与内存数据的频繁交换，最后将结果暂存在显存中。4）将计算结果从显存重新送回内存中。182.2CUDA执行过程CUDA程序的构成一个CUDAC程序通常由两部分构成：一部分在主机（CPU）上顺序执行，另外一部分则在设备（GPU）上启动成千上万个线程并行执行，它们在编译过程中由NVIDIA公司的C编译器（NVCC）区分开。NVCC简化了C语言或PTX的编译流程：它提供了简单的命令行选项，调用一系列的编译工具来执行它们。NVCC可同时编译由主机代码（在CPU上执行的代码）和设备代码（在GPU上执行的代码）组成的源文件。192.2CUDA执行过程CUDA编译流程分离主机端代码和设备端代码。将主机端代码输出为C语言代码供其他工具编译，或者NVCC在编译的最后阶段调用主机编译器将主机端代码输出为目标代码。编译设备端代码得到其二进制形式（cubin对象）或/和汇编形式（PTX）。将PTX源码或cubin对象利用CUDA驱动API装载并执行，或者使用链接将PTX源码或cubin对象加载到主机代码中并将其作为已初始化的全局数据数组导入并执行。202.2CUDA执行过程一般来说，控制从主机端向设备端的转移是通过在主机端调用由关键字__global__标示的且只能从主机端被调用的kernel函数。CUDA程序用类kernelFun<<<M,N>>>(d_a,d_b,d_c);的语句来启动kernel函数，其中<<<>>>运算符中的N和M是主机端设置设备端要启动kernel函数的参数，M表示线程块的数量（维度），N表示线程块的大小，(d_a,d_b,d_c)则为kernel函数的形参，与一般C函数没有区别，该函数体定义的语句即为后续每个线程要执行的代码。212.2CUDA执行过程当kernel函数启动运行后，执行过程转移到设备端（GPU），然后生成指定的线程数目，这些线程被组织成块（block）不同架构的支持CUDA的GPU一个block所能容纳的最大线程数目也不同，有512个（Tesla架构）也有1024个（Fermi架构、Kepler架构），块最后被组织成一个线程网格（grid）。每调用一次kernel函数会生成一个线程网格，当kernel函数中的所有线程都完成他们的执行任务后，相应的网格也会终止，并且在调用下一个kernel函数前，程序会转到主机端继续执行，即kernel函数和主机端代码是异步执行的。222.3CUDA的线程组织一般而言，在启动kernel函数时我们会把网格中的线程块组织成二维数组形式，将线程块中的线程组织成三维数组的形式。通过CUDAC拓展定义的一个类似C结构的数据类型dim3，它有3个无符号整数型字段，分别是x——标示x方向上线程或线程块的索引、y——标示y方向上线程或线程块的索引和z——标示z方向上线程或线程块的索引。同一个kernel中所有的线程都会执行kernel函数体定义的语句，唯一不同的是每个线程都要操作数据，所以基于CUDA的GPU计算模型为单指令多数据流（SIMD）。232.3CUDA的线程组织CUDA通过内置的一些预初始化变量——threadIdx.x、threadIdx.y、threadIdx.z标示当前线程所处的线程块的位置，以及blockIdx.x、blockIdx.y、blockIdx.z标示当前线程所处的线程块在整个网格中所处的位置，gridDim.x、gridDim.y、gridDim.z标示网格的维度和blockDim.x、blockDim.y、blockDim.z标示每个块的维度；使用上面这些CUDA提供的内置变量可以在同一个kernel函数中将各个块中的线程彼此区分开来，然后决定哪个线程要处理哪些数据。各个线程块间布局是相互独立的，所以不同线程块中的相对应位置的线程具有相同的threadIdx.x、threadIdx.y和threadIdx.z，所有的线程块拥有相同的线程数目。242.3CUDA的线程组织对网格中任意点(i,j)，在CUDA代码中表征一个线程，该线程对应的网格中的索引可以使用下面公式进行表示：i=threadIdx.x+blockIdx.x*blockDim.xj=threadIdx.y+blockIdx.y*blockDim.y252.3CUDA的并行优化存储器访问优化因为主机内存和设备端显存间传递数据的过程中耗时更长，所以在进行CUDA运算过程中，应尽量减少CPU和GPU间的数据传输，充分发挥GPU多线程并行计算的优势将更多的操作放置在设备端完成；尽量减少对全局存储器的访问，更多的使用共享存储器。指令优化1.在kernel函数中如果不是必要的尽量不要用__syncthreads()语句进行同步，另外，同一个warp中的线程不需要同步。2.遇到迭代累加的时候，如果给定累加的次数，可以通过将迭代过程展开的方式消除循环计数更新指令和分支指令。3.任何控制指令（if、else、switch、do、while、for）都会致使同一个warp中的线程产生分支。所以我们应尽量减少一个warp中线程产生分支的情况。26三、基于CUDA的GPU高性能计算的一个实例从矩阵相乘开始初始矩阵M,N,我们从矩阵M中取一行，矩阵N中取一列进行点积运算，从而得到结果矩阵P中的每个元素。从右图可以看到，P中不同的元素的点积运算是可以同时进行的。也就是说，这些点积运算之间互相不影响。27在C语言中实现矩阵乘法矩阵乘法中一个简单的主函数intmain(void){ 1.//分配和初始化矩阵M、N、P //执行I/O读取输入矩阵M、N 2.//在设备上执行M*N 3.//执行I/O写入输出矩阵P //释放矩阵M、N、P的存储空间 ... return0；}我们怎么实现矩阵M、N相乘呢？for(introw=0;row<row1;row++) { for(intcol=0;col<col2;col++) { intnum=col1;// for(intk=0;k<num;k++) result[row][col]+=p1[row][k]*p2[k][col]; } }2829

在CUDA上实现矩阵乘法

对于大型矩阵的乘法，点积的个数可能会非常大，如两个1000*1000的矩阵乘法就有1000000个独立的点积，每个点积涉及1000次乘法和1000次累加运算。因此，高维的矩阵乘法拥有大量的数据并行性。因此，我们可以用GPU进行高性能计算。voidMatrixMulKernel(float*M,float*N,float*P,intwidth){1.//把M和N传递到设备存储器中 cudaMalloc((void**)&Md,size); cudaMemcpy(Md,M,size,cudaMemcpyHostToDevice); cudaMalloc((void**)&Nd,size); cudaMemcpy(Nd,N,size,cudaMemcpyHostToDevice);//在设备上分配P cudaMalloc((void**)&Pd,size)2.//调用kernel的代码...3.//把P从设备上传递到主机上 cudaMemcpy(P,Pd,size,cudaMemcpyDeviceToHost);//释放设备上的矩阵 cudaFree(Md); cudaFree(Nd); cudaFree(Pd);}30Kernel函数：__global__voidMatrixMulKernel(int*dev_M,int*dev_N,int*dev_P,introw){ //计算P和M中元素的行索引

intRow=blockIdx.y*16+threadIdx.y; //计算P和N中元素的列索引

intCol=blockIdx.x*16+threadIdx.x; intPvalue=0; //每个线程负责计算块子矩阵的一个元素

if(Row<row&&Col<row) { for(intk=