《并行程序设计》实验指导书之五

《并行程序设计》实验指导书之五实验5.1运行CUDA向量加法的样例代码实验目的1．掌握cuda开发环境的使用与配置；2．掌握利用cuda调试和运行；3．掌握cuda并行计算的原理。实验要求1．熟练掌握C++语言；2．掌握VisualStudio*.NET*集成开发环境的使用；3．掌握cuda开发环境。实验原理CUDA(Compute

Unified

Device

Architecture)是由NVIDIA公司创立的基于他们公司生产的图形处理器GPUs（Graphics

Processing

Units，可以通俗的理解为显卡）的一个并行计算平台和编程模型。通过CUDA，GPUs可以很方便地被用来进行通用计算（有点像在CPU中进行的数值计算等等）。在没有CUDA之前，GPUs一般只用来进行图形渲染（如通过OpenGL，DirectX）。有了CUDA之后，开发人员可以通过调用CUDA的API，来进行并行编程，达到高性能计算的目的。CUDA中常见的概念及名称：主机：CPU及系统的内存（内存条）；设备：

GPU及GPU本身的显示内存；线程(Thread)：资源调度单元，一般交给GPU的一个核去处理（有一维，二维，三维）。线程块(Block)：多个线程组合在一起就是线程块；各block并行执行，互相之间不能通信，执行时无法指定顺序；线程块有数量限制，最多可以有65535个线程块。线程格(Grid)：由多个线程块组成。线程、线程块、线程格的逻辑结构如下所示：线程束：线程束是一个集合，其中包含32个线程。这个集合里的线程被“组合在一起”并且“步调一致”地进行执行。对程序的每一行，线程束里的线程都将在不同数据上分别执行。GPU内存分类1.全局内存：就是指设备内存。2.共享内存：存储在全局内存中，使用时要添加关键字__shared__到变量声明中。CUDA对GPU上启动的每个线程块，都会保存一个共享变量的副本。在同一个线程块内的所有线程都要共享这块内存，但线程却不能看到因此也不能修改其他线程块的副本。这就实现了一个线程块中的多个线程能在计算上进行通信以及协作。3.常量内存：存储在全局内存中，使用时要添加关键字__constant__到变量声明中。常量内存用来保存在核函数执行期间不会发生变化的数据，变量是只读的，通过特殊的处理方式，有时候使用常量内存替代全局内存可以减少内存带宽，对性能提升有帮助；当需要拷贝数据到常量内存中必须使用cudaMemcpyToSymbol，如果使用cudaMemcpy会将数据复制到全局内存。4.纹理内存：存储在全局内存中。面向访问内存具有空间聚簇性的程序(例如图像处理方面的计算程序)设计，互相临近的线程所读取的数据在物理存储上也是临近的，可以减少访存次数，节约带宽，以此来提升效率。纹理内存有一维与二维两种：一维纹理内存的声明方式是texture<类型>，使用cudaBindTexture()函数绑定纹理内存，cudaUnbindTexture()函数解除绑定，读取内存数据时要使用tex1D()函数；二维纹理内存的声明方式是texture<类型,数字>，使用cudaBindTexture2D()函数绑定纹理内存，cudaUnbindTexture()函数解除绑定，读取内存数据时要使用tex2D()函数。5.固定内存：存储在主机内存中，又称为不可分页或页锁定内存。对于固定内存，操作系统不会对其分页，也不会交换到磁盘上，可以确保它始终驻留在物理内存上。在编写程序时可以直接访问这块物理地址，因为它不会被破坏或迁移。固定内存是为了提高访问速度而被设计出来的。GPU如果知道主机中的物理地址，就可通过DMA方式来复制主机与GPU之间的数据。当然，用户编写程序时要注意不可一味使用固定内存，这样将导致物理内存迅速消耗完。在使用固定内存时，一般将调用cudaMemcpy()函数时使用的源内存或目的内存设置为固定内存，在调用完后不再需要时立即释放掉。分配固定内存需要使用cudaHostAlloc()函数；释放固定内存需要使用cudaFreeHost()函数。注意复制固定内存是异步的方式。核函数（Kernel）核函数在GPU上面运行，在GPU上运行的函数都可视作是核函数；核函数可以使用标识符来修饰，通常使用的是__global__标识符。调用的方法与C语言有所区别，通常是通过<<<参数n1,参数n2>>>，调用时必须声明内核函数的执行参数；核函数需要通过线程格（Grid）来组织，线程格下包含若干线程块（block），而线程块下又包含若干个线程（thread）；核函数的执行单位是线程块（block）；编写程序时要注意的一点是，对kernel函数中所需要使用的数组或变量，一定要在调用前提前分配好空间，否则在GPU进行计算的时候会发生错误，例如越界错误，也有可能导致蓝屏。在编写程序需要知道如下这些CUDA对C语言的扩展：1.指定函数执行器件的扩展：定义了函数类型限定符，可用来确定函数是在CPU上执行或是在GPU上执行，以及是从CPU上调用还是从GPU上调用；有如下几种函数类型限定符：__device__，由它所修饰的函数从GPU上调用，且在GPU上执行，这样的函数能够由__device__或__global__修饰的函数调用。由它所修饰的函数的使用有限制，比如不允许使用函数指针；__global__，由它所修饰的函数从CPU上调用，且在GPU上执行，也就是前面提到过的内核(kernel)函数；它只能由主机调用，不是一个完整的程序，只是表示数据并行的步骤，其指令流由多个线程执行；__host__，由它所修饰的函数从CPU上调用，且在CPU上执行，这与传统的C函数没什么区别；2.变量存储位置的扩展定义了变量类型限定符，在传统的CPU程序中，变量的存储位置是由编译器来负责完成的，编写程序的人并不需要指定存储位置；但在CUDA架构中，需要用变量类型限定符来规定变量的存储位置，有以下几种可供选择：GPU里的缓存、共享存储器、寄存器；设备的显存；主机的内存等等。基于这些，CUDA中抽象出8种不同的存储器，下面列出的是常用的变量类型限定符以及其含义。__device__：这个修饰符声明的数据的存放位置是显存，主机通过运行时库可以对其进行访问，所有的线程都可以访问这样声明的数据；__shared__：这个修饰符声明的数据的存放位置是共享存储器，不是所有线程都能访问这样声明的数据，只有它所在的块里的线程可以对其进行访问；__constant__：这个修饰符声明的数据的存放位置是常量存储器，主机通过运行时库可以对其进行访问，所有的线程都可以访问这样声明的数据；3.执行配置扩展定义了执行配置运算符<<<>>>，调用内核函数时要将执行配置传递过去，<<<>>>就是用来传递执行配置的。执行配置的组成是4个参数：1.网格大小2.块大小，3.共享存储器大小（可以忽略，默认为0），4.执行的流（可以忽略，默认为0）。4.内建变量扩展内建变量用来在运行时获取线程索引以及块和网格的尺寸等信息。CUDA里一共有5个内建变量：1.gridDim，包含grid的维度。这是一个结构体，它包含三个元素x，y，z，用来表示网格在这几个方向上的尺寸，虽然CUDA设计的是三维，但目前只能使用二维；2.blockDim，包含block的维度。和gridDim一样也是一个由x，y，z三个元素组成的结构体，它表示的是块在这几个方向上的尺寸；3.blockIdx，包含网格的纬度。和gridDim一样也是一个由x，y，z三个元素组成的结构体，这几个元素分别表示当前线程所在的块在网格中x，y，z三个方向上的索引；4.threadIdx，和gridDim一样也是一个由x，y，z三个元素组成的结构体，分别表示当前线程在其所在的块中x，y，z三个方向上的索引；5.warpSize，它表明warp的尺寸。常用的GPU内存函数cudaMalloc()(1)原型：cudaError_tcudaMalloc(void**devPtr,size_tsize)(2)参数：devPtr——指向分配的设备内存；size——需要分配的内存大小；(3)作用：与malloc()函数类似，只是此函数在GPU的内存中进行内存分配。从设备上分配size比特的连续内存并返回一个指向此内存空间的指针*devPtr。分配的内存可用于存储任何类型的变量；(4)返回值：分配成功返回cudaSuccess，万一分配失败的话返回cudaErrorMemoryAllocationcudaMemcpy()原型：cudaError_tcudaMemcpy(void*dst,constvoid*src,size_tcount, enumcudaMemcpyKindkind)参数：dst

——目标内存地址，src——源内存地址，count——拷贝的字节数目，kind——数据拷贝的方向(3)作用：与memcpy()函数类似，从src指针指向的内存空间中拷贝count字节数据到dst指针指向的空间中。kind有以下几种取值可能：cudaMemcpyHostToHost，cudaMemcpyHostToDevice，cudaMemcpyDeviceToHost，或cudaMemcpyDeviceToDevice，指明了数据拷贝的方向。如果dst和src指向的空间不符合kind的方向，由此产生的行为是没有定义的。以同步方式执行，即当函数返回时，复制操作就已经完成了，并且在输出缓冲区中包含了复制进去的内容。(4)返回值：cudaSuccess，cudaErrorInvalidValue，cudaErrorInvalidDevicePointer，cudaErrorInvalidMemcpyDirection。cudaFree()(1)原型：cudaError_tcudaFree(void*devPtr)(2)参数：devPtr——指向要释放的内存空间；(3)作用：与free()函数类似，只是此函数释放的是cudaMalloc()分配的内存；(4)返回值：cudaSuccess，cudaErrorInvalidDevicePointer，cudaErrorInitializationError实验内容实验步骤：打开VisualStudio，新建一个项目，模版要选择NVIDIA下的CUDA；在生成项目中例子代码就是向量加法，直接编译运行，记录实验结果；实验4.2基于CUDA优化矩阵乘法实验目的1．掌握基于cuda优化科学计算的方法;2．比较各种cuda优化方式的性能提升效果；3．掌握利用加速比、运行时间、效率等测度分析并行程序性能；实验要求1．熟练掌握C++语言；2、掌握VisualStudio*.NET*集成开发环境的使用；3．掌握cuda开发环境；实验原理设A为

的矩阵，B为

的矩阵，那么称

的矩阵C为矩阵A与B的乘积，记作

，其中矩阵C中的第

i行第

j列元素可以表示为：本次实验中将A矩阵和B矩阵假设成大小相同的方形矩阵，这不影响性能分析。同时不做算法上层面的优化，直接使用三重for循环。程序逻辑图matgen函数的输入是矩阵首地址以及行数、列数，产生介于0与1之间的浮点随机数填充矩阵中的每个元素；matmult函数的输入是需要相乘的两个矩阵，对它们运行矩阵乘法程序。它是传统的CPU程序。既可以用来与其它方式进行矩阵乘法所得到的结果进行对比，判断正确与否，也可以用来作为性能对比的基准，为了提高精确度，中间结果使用双精度浮点类型来存储；compare_mat函数的输入是需要进行比较的两个矩阵，用来计算两个矩阵之间的平均相对误差以及最大相对误差，并打印出比较的结果；matmultCUDA函数的输入是需要相乘的两个矩阵，使用GPU实现矩阵乘法。它首先从显卡内存里申请用于存放矩阵的空间，之后将矩阵数据从主内存拷贝到显卡内存中。在进行内存拷贝时，如果用cudaMemcpy函数的话，将造成每个row分开进行拷贝，这样的话就需要多次调用cudaMemcpy函数，这会降低程序的效率。所以，在这里使用cudaMemcpy2D

函数来拷贝二维数组，这样就能够只执行一次函数调用就完成拷贝的工作。然后调用核函数matMultCUDA，它先通过bid以及tid计算出某个thread需要计算的行和列，将计算所得写入结果中。下面对程序做一些解释。1.在使用cuda进行计算时，用到了下面的代码：for(j=tid;j<n;j+=blockDim.x){ floatt=0; floaty=0; for(i=0;i<n;i++){ floatr; y-=data[i]*b[i*ldb+j]; r=t-y; y=(r-t)+y; t=r; } c[row*ldc+j]=t;}之所以要这样进行累加是基于下面的考虑：CPU上面的计算使用了64位浮点数来累加中间结果，而在GPU上却只能用32位浮点数进行累加，当累加的数据很大时，就会产生舍入误差。针对这个问题，CUDA在进行加、减、乘法的浮点运算时符合IEEE754规定的精确度，所以使用Kahan'sSummationFormula来提高精确度（思想是把每次截断的误差在一个较小的数字里进行累加）。2.在调用核函数时使用了这样的方式：函数名称<<<block数目,thread数目,sharedmemory大小>>>(参数...)；因为只有在同一个block中的线程可以共享内存，因此一行只能由同一个block里的线程来进行计算。另外需要共享内存存放整个row的数据。下面的代码可以把每行的数据放到共享内存中，便于后面使用共享内存进行计算：extern__shared__floatdata[];constinttid=threadIdx.x;constintrow=blockIdx.x;inti,j;for(i=tid;i<n;i+=blockDim.x){ data[i]=a[row*lda+i];}3.GPU在读取内存时会选择一个固定倍数的地址开始，例如64bytes的整数倍，这样可以在最大程度提升效率。但是实验中使用的矩阵未必是64的整数倍，这影响了效率。为了消除这个影响，代码在内存分配时使用了CUDA提供的cudaMallocPitch函数，它在配置内存时会自动使用最佳倍数。使用下面的代码来优化cudaMalloc部分：size_tpitch_a,pitch_b,pitch_c;

cudaMallocPitch((void**)&ac,&pitch_a,sizeof(float)*n,n);

cudaMallocPitch((void**)&bc,&pitch_b,sizeof(float)*n,n);

cudaMallocPitch((void**)&cc,&pitch_c,sizeof(float)*n,n);cudaMallocPitch函数会选择合适的倍数来配置内存，并传回配置宽度。将矩阵拷贝到显卡内存时要使用传回的宽度，代码如下：cudaMemcpy2D(ac,sizeof(float)*n,a,sizeof(float)*lda,sizeof(float)*n,n,cudaMemcpyHostToDevice);cudaMemcpy2D(bc,sizeof(float)*n,b,izeof(float)*ldb,sizeof(float)*n,n,cudaMemcpyHostToDevice);呼叫kernel的部分如下：matMultCUDA<<<n,NUM_THREADS,sizeof(float)*n>>>(ac,pitch_a/sizeof(float),bc,pitch_b/sizeof(float),cc,pitch_c/sizeof(float),n);把结果拷贝回主内存时，也要用上传回的宽度：cudaMemcpy2D(c,sizeof(float)*ldc,cc,pitch_c,sizeof(float)*n,n,cudaMemcpyDeviceToHost);实验内容1.打开VisualStudio，新建一个项目，模版要选择NVIDIA下的CUDA；2.将实验代码拷贝进去，编译运行，进行如下实验并记录数据（实验报告中给出数据并绘图）实验一：固定矩阵大小为250x250，调整块大小NUM_THREADS，记录程序运算时间及运算速度。基准运行时间(s)：基准运算速度(GFLOPS)：NUM_THREADS1632641282565