CUDA介绍与案例.ppt

《CUDA介绍与案例.ppt》由会员分享，可在线阅读，更多相关《CUDA介绍与案例.ppt（61页珍藏版）》请在淘文阁 - 分享文档赚钱的网站上搜索。

1、CUDA 介绍与案例1介绍编程模型变量和函数案例介绍p应用领域：游戏、图形动画、科学计算可视化、地质、生物、物理模拟等；编程模型变量和函数案例介绍2008年SIGGRAPH年会上，NVIDIA公司推出CUDA(Compute Unified Device Architecture)NVIDIA2008;CUDA是NVIDIA为自己的GPU编写的一套编译器及相关的库文件；将GPU 视作数据并行计算设备,是一种新的处理和管理GPU 计算的硬件和软件架构；编程模型变量和函数案例介绍GPU到CUDAGPU多多“核核”：SM真正意义的变革真正意义的变革GPGPU通用计算通用计算重要突破重要突破GPU通用

2、计通用计算开始出现算开始出现世界上第一个世界上第一个GPUGeForce 8800GeForce 6800GeForce 3GeForce 256编程模型变量和函数案例介绍pCPU 和 GPU 的每秒浮点运算次数和存储器带宽比较具有强大浮点具有强大浮点运算能力运算能力nGPU采用大量的执行单元，这些执行单元可以轻松的加载并行处理线程，而不像CPU那样的单线程处理；n主机和设备均维护自己的主机和设备均维护自己的 DRAM，分别称为主机存储器和设备存，分别称为主机存储器和设备存储器储器 CPU 和 GPU 之间浮点功能的差异，原因在于 GPU 专为高计算密集型（数学运算与存储器运算的比率）、高度并

3、行化的计算而设计；GPU 的设计能使更多晶体管用于数据处理，而非数据缓存和流控制；编程模型变量和函数案例介绍高度并行化、多线程、多核处理器，操作 系统的多任务机制负责管理多个并发运行 的CUDA和应用程序对GPU的访问；基于标准C语言,可自由地调用GPU的并行 处理架构;同时适用于图形和通用并行计算应用程序，成为图形处理器的主要发展趋势。编程模型变量和函数案例介绍 依次安装 Cuda Driver Cuda Toolkit Cuda SDK 在安装目录中包括：bin工具程序及动态链接库 doc文件 include 头文件 lib 链接库 open64基于open64的CUDA compiler

4、 src 一些自带例子，如 C:CUDA_SDKCsrc 安装程序会设定一些环境变量：CUDA_BIN_PATH CUDA_INC_PATH CUDA_LIB_PATH 编程模型变量和函数案例介绍CUDA软件栈包含多个层，如图所示：设备驱动程序、应用程序编程接口（API）及其运行时两个较高级别的通用数学库，即CUFFT和CUBLASCUDA安装后编程模型变量和函数案例介绍 CUDA程序编译pCUDA的源文件被nvcc编译 NVCC编译器会分离源码中设备代码和主机代码，主机代码交由一般的C/C+编译器（gcc等）编译，设备代码由NVCC编译；内核必须使用nvcc编译成二进制代码才能在设备端执行。

5、编程模型变量和函数案例介绍内核函数编程模型变量和函数案例介绍编程模型变量和函数案例介绍n串行代码在主机上串行代码在主机上执行，而并行代码执行，而并行代码在设备上执行在设备上执行n当调用当调用kernelkernel的时的时候，则候，则CUDACUDA的每个的每个线程并行地执行一线程并行地执行一段指令，这与通常段指令，这与通常C C函数只执行一次不函数只执行一次不一样。一样。执行模型执行模型编程模型变量和函数案例介绍说明：n在CUDA架构下，一个程序分为两部分，即host 端和 device端；Host 端 是在 CPU上执行的部分；Device端 是在显卡芯片上执行的部分，该端程序又称为内核

6、函数（kernel function）n通常Host端程序复制内核函数到显卡内存中，再由显卡芯片执行device端程序，完成后再由host端程序将结果从显卡内存中取回；编程模型变量和函数案例介绍nCUDA 允许程序员定义内核函数，实现配置；n 块组织为一个一维或二维线程块网格，维度由 语法的第一个参数指定；n 执行内核的每个线程都会被分配一个独特的线程 ID，可通过内置的 threadIdx 变量在内核中访问编程模型22.1线程模型（体系结构）2.2存储器模型（体系结构）编程模型变量和函数案例介绍2.1 线程模型编程模型变量和函数案例介绍并行线程结构：Thread:并行的基本单位索引：thre

7、adIdx（内置变量）Block(Thread block):线程块允许彼此同步，通过快速共享内存交换数据每个块的线程数应是 warp（调度任务的最小单位）块大小的倍数，每个块的线程数受限索引：blockIdx（内置变量）Grid:一组thread block共享全局内存Kernel:在GPU上执行的核心程序One kernel One grid编程模型变量和函数案例介绍说明：Grid 是一组Block,以1维、2维或3维组织，共享全局内存；Grid之间通过global memory交换数据 Block 是互相协作的线程组，通过global memory共享数据，允许彼此同步；以1维、2维或3

8、维组织；同一block内的thread可以通过shared memory和同步实现通信;编程模型变量和函数案例介绍一个内核可能由多个大小相同的线程块执行，因而线程总数线程总数应等于每个块的线程数乘以块的数;线程索引线程索引(Index)(Index)及及ID(ID(IDentity)，索引为(x,y)的线程ID为(x+yDx);对于大小为(Dx,Dy,Dz)的三维块，索引为(x,y,z)的线程ID为(x+yDx+zDxDy)；示意图编程模型变量和函数案例介绍 线程块索引及线程块索引及ID,ID,情况类似：对于一维块来说，两者是相同的；对于大小为(Dx，Dy)的二维块来说，索引为(x,y)的线程

9、块的ID 是（x+yDx）；对于大小为（Dx，Dy，Dz）的三维块来说，索引 为（x,y,z）的线程的ID 是（x+yDx+zDxDy）;/+图示说明编程模型变量和函数案例介绍 例:向量的和 _global_ void vecAdd(float*A,float*B,float*C)/内核函数 int i=threadIdx.x;Ci=Ai+Bi；int main()vecAdd(dA,dB,dC);/+配置？哪个线程对哪个数组下标求和；!第i个线程对相应的数组下标求和；！用内置变量确定数组的下标编程模型变量和函数案例介绍或int main()int threadsPerBlock=256;in

10、t blocksPerGrid=(N+threadsPerBlock-1)/threadsPerBlock;VecAdd(d_A,d_B,d_C,N);编程模型变量和函数案例介绍 _global_ void VecAdd(const float*A,const float*B,float*C,int N)/内核函数 int i=blockDim.x*blockIdx.x+threadIdx.x;if(i N)Ci=Ai+Bi;/参见例：vectorAdd!？哪个线程对哪个数组下标求和；即第几个Block（blockIdx）的第几个线程，对相应的数组下标求和；编程模型变量和函数案例介绍2.2 编

11、程模型变量和函数案例介绍编程模型变量和函数案例介绍说明：p所有线程都可访问全局存储器；p所有的线程都可以访问只读存储区域：“常量存储区”和“纹理存储区”；纹理存储器提供不同的寻址模式，为某些特定的数据格式提供数据过滤的能力；pBlock中所有thread都在同一个处理核心上运行。因此每个block的thread数目受到处理核心所拥有存储器资源的限制，在NVIDIA Tesla架构中，一个线程block最多可包含512个线程;编程模型变量和函数案例介绍p每个线程块都有一个共享存储器共享存储器(Shared MemoryShared Memory)，该存储器对于块内的所有线程块内的所有线程都是可见

12、可见的，并且与块具有相同的生命周期。每个线程都有一个私有的本地存储器（本地存储器（LocalMemoryLocalMemory）。p在同一个blockblock中的中的threadthread通过共享存储器共享数据，相互协作，实现同步同步：_syncthreads();编程模型变量和函数案例介绍p registerregister 访问延迟极低；每个线程占有的register有限，编程时不要为其分配过多 私有变量；plocal memorylocal memory 寄存器被使用完毕，数据将被存储在局部存储器中；此外，存储大型结构体或者数组（无法确定大小的数组），线程的输入和中间变量；编程模型变

13、量和函数案例介绍pshared memoryshared memory 访问速度与寄存器相似；实现线程间通信的延迟最小；保存公用的计数器,或block的公用结果；声明时使用关键字 _ shared_;pconstant memoryconstant memory 只读地址空间，位于显存，用于存储需频繁访问的只读参数；由于其在设备上有片上缓存，比全局存储器读取效率高很多；使用_ constant_ 关键字;定义在所有函数之外；编程模型变量和函数案例介绍pglobalmemoryglobalmemory 存在存在于显存中，也称为线性内存;cudaMalloc（）函数分配;cudaFree（）函数释

14、放;cudaMemcpy（）进行主机端与设备端数据传输;对于二维和三维数组：cudaMallocPitch（）和cudaMalloc3D（）分配线性存储空间；cudaMemcpy2D（），cudaMemcpy3D（）向设备存储器进行拷贝；编程模型变量和函数案例介绍编程模型变量和函数案例介绍 对象数组指针对象数组指针 用于指向 global memory的存储区域,由Host端分配，例：float*d_A，d_B，d_C;int N=50000;size_t size=N*sizeof(float);cudaMalloc(void*)&d_A,size);cudaMalloc(void*)&d_

15、B,size);cudaMalloc(void*)&d_C,size);cudaMemcpy(d_A,h_A,size,cudaMemcpyHostToDevice)；cudaMemcpy(d_B,h_B,size,cudaMemcpyHostToDevice)；编程模型变量和函数案例介绍p texture memorytexture memory：只读；不是一块专门的存储器，涉及显存、纹理缓存、纹理拾取等纹理流水线操作；在kernel中访问纹理存储器的操作,称为纹理拾取纹理拾取（texture fetching）纹理拾取使用的坐标与数据在显存中的位置可以不同，通过纹理参考约定二者映射方式；将

16、显存中数据与纹理参考关联的操作，称为纹理绑定（texture binding）；3变量和函数3.1变量类型申明变量类型申明3.2 3.2 函数类型申明函数类型申明3.3 3.3 内置变量内置变量3.4 3.4 内置向量内置向量3.5 3.5 数学函数数学函数3.6 3.6 与与OpenglOpengl互操作互操作编程模型变量和函数案例介绍3.13.1编程模型变量和函数案例介绍编程模型变量和函数案例介绍3.2 函数函数编程模型变量和函数案例介绍ngridDimgridDim（有多少（有多少BlockBlock）变量的类型为 dim3，包含网格的维度。nblockDim(blockDim(有多少有

17、多少ThreadThread）变量的类型为 dim3，包含块的维度。nblockIdx(blockIdx(第几个第几个Block)Block)变量的类型为 uint3，包含网格内的块索引 nthreadIdx(threadIdx(第几个第几个Thread)Thread)变量的类型为 uint3，包含块内的线程索引。备注：dim3 是整型的向量类型，未指定的任何组件都将初始化为1,如 dim3 dimBlock(4,2,2);uint3 是无符号整型的向量类型；3.3 3.3 内置变量内置变量编程模型变量和函数案例介绍 Type name:char1、uchar1、char2、uchar2、ch

18、ar3、uchar3、char4、uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1、int2、uint2、int3、uint3、int4、uint4、long1、ulong1、long2、ulong2、long3、ulong3、long4、ulong4、float1、float2、float3、float4、double2 均为结构体，第 1、2、3、4 个组件分别可通过字段 x、y、z 和 w 访问；均附带形式为 make_ 的构造函数 例如：int2 v=make_int2(6,9);

19、3.4 3.4 内置向量内置向量编程模型变量和函数案例介绍CUDA支持标量运算的数学函数,包括：pow、sqrt、cbrt、hypot、exp、exp2、espm1、log、log2、log10、log1p、sin、cos、tan、asin、acos、atan、atan2、sinh、cosh、tanh、asinh、acosh、atanh、ceil、floor、trunc、round等;当函数有前缀“_”时，运算速度较快。3.5 3.5 数学函数数学函数编程模型变量和函数案例介绍3.6 3.6 与与OpenglOpengl的互操作的互操作将一个Opengl缓冲对象注册到 CUDA，通过通过 cu

20、daGLRegisterBufferObjeccudaGLRegisterBufferObject()；例如：GLuint bufferObj;cuGLRegisterBufferObject(bufferObj);编程模型变量和函数案例介绍 注册完成后，内核即可使用 cuGLMapBufferObject(cuGLMapBufferObject();/);/映射缓冲对象 返回设备存储器地址，读取或写入缓冲对象；即将CUDA内存的指针指向OpenGL的缓冲区；例如：CUdeviceptr devPtr;int size;cuGLMapBufferObject(&devPtr,&size,buf

21、ferObj);用完之后 cuGLUnmapBufferObject(cuGLUnmapBufferObject()；/解除映射 cuGLUnregisterBufferObject(cuGLUnregisterBufferObject()；/取消注册；编程模型变量和函数案例介绍4案例编程模型变量和函数案例介绍在设计并行程序时，需要找到被分解问题关键参数和需要找到被分解问题关键参数和线程下标的对应关系，并建立映射机制线程下标的对应关系，并建立映射机制，以使GPU的并行计算能力得到充分发挥；此外，考虑GPU的软硬件资源，CUDA编程模型有一些参数上的限制，主要包括对所分配线程和线程块数量分配线程

22、和线程块数量的规定；如GeForces系列GPU在每个线程块支持的最大线程数是512，而每个网格所具有最大线程块的数量是65535编程模型变量和函数案例介绍/方法一：_global_ void VecAdd(float*A,float*B,float*C)int i=threadIdx.x;Ci=Ai+Bi;int main()/Kernel invocationVecAdd(A,B,C);参见例：matrix_src.doc编程模型变量和函数案例介绍方法二：intmain()intthreadsPerBlock=256;intblocksPerGrid=(N+threadsPerBlock-

23、1)/threadsPerBlock;VecAdd(d_A,d_B,d_C,N);编程模型变量和函数案例介绍_global_ void VecAdd(const float*A,const float*B,float*C,int N)int i=blockDim.x*blockIdx.x+threadIdx.x;if(i N)Ci=Ai+Bi;编程模型变量和函数案例介绍扩展：矩阵加法 _global_ void matAdd(float*Ad,float*Bd,float*Cd)int i=threadIdx.x;int j=threadIdx.y;Cdij=Adij+Bdij;int mai

24、n()/Kernel invocation dim3 dimBlock(N,N);matAdd(Ad,Bd,Cd);.编程模型变量和函数案例介绍p说明内存和线程管理的基本特性说明内存和线程管理的基本特性本地存储器、寄存器的用法本地存储器、寄存器的用法线程线程IDID的用法的用法主机和设备之间数据传输的主机和设备之间数据传输的APIAPI为了方便，以方形矩阵说明为了方便，以方形矩阵说明编程模型变量和函数案例介绍矩阵大小为矩阵大小为 WIDTH x WIDTHWIDTH x WIDTH在没有采用分片优化算法的情况下：在没有采用分片优化算法的情况下：一个线程计算一个线程计算P P矩阵中的一个矩阵中的

25、一个 元素元素需要从全局存储器载入需要从全局存储器载入WIDTHWIDTH次次方块矩阵乘法方块矩阵乘法编程模型变量和函数案例介绍CPU上的矩阵乘法void MatrixMulOnHost(float*M,float*N,float*P,int Width)for(int i=0;i Width;+i)for(int j=0;j Width;+j)double sum=0;for(int k=0;k Width;+k)double a=Mi*width+k;double b=Nk*width+j;sum+=a*b;Pi*Width+j=sum;编程模型变量和函数案例介绍向向GPUGPU传输矩阵数

26、据传输矩阵数据void MatrixMulOnDevice(float*M,float*N,float*P,int Width)int size=Width*Width*sizeof(float);float*Md,Nd,Pd;/设置调用内核函数时的线程数目设置调用内核函数时的线程数目 dim3 dimBlock(Width,Width);dim3 dimGrid(1,1);/在设备存储器上给在设备存储器上给M M和和N N矩阵分配空间，并将数据复制到设备存储器中矩阵分配空间，并将数据复制到设备存储器中 cudaMalloc(&Md,size);/Md线性存储矩阵元素；线性存储矩阵元素；cud

27、aMemcpy(Md,M,size,cudaMemcpyHostToDevice);cudaMalloc(&Nd,size);/Nd线性存储矩阵元素；线性存储矩阵元素；cudaMemcpy(Nd,N,size,CudaMemcpyHostToDevice);/在设备存储器上给在设备存储器上给P P矩阵分配空间矩阵分配空间 cudaMalloc(&Pd,size);GPU上的矩上的矩阵乘法乘法编程模型变量和函数案例介绍计算结果向主机传输计算结果向主机传输/内核函数调用，将在后续部分说明内核函数调用，将在后续部分说明/只使用了一个线程块只使用了一个线程块(dimGrid)(dimGrid)，此线程

28、块中有，此线程块中有Width*WidthWidth*Width个线程个线程 MatrixMulKernel(Md,Nd,Pd，Width);/从设备中读取从设备中读取P P矩阵的数据矩阵的数据cudaMemcpy(P,Pd,size,cudaMemcpyDeviceToHost);/释放设备存储器中的空间释放设备存储器中的空间 cudaFree(Md);cudaFree(Nd);cudaFree(Pd);编程模型变量和函数案例介绍/矩阵乘法的内核函数矩阵乘法的内核函数每个线程都要执行的代码每个线程都要执行的代码_global_ void MatrixMulKernel(float*Md,fl

29、oat*Nd,float*Pd,int Width)/2/2维的线程维的线程IDID号号 int i=threadIdx.x;int j=threadIdx.y;/P/Pvaluevalue用来保存被每个线程计算完成后的矩阵的元素用来保存被每个线程计算完成后的矩阵的元素 float Pvalue=0;编程模型变量和函数案例介绍/每个线程计算一个元素每个线程计算一个元素for(int k=0;k Width;+k)float Melement=Mdj*Width+k;float Nelement=Ndk*Width+i;Pvalue+=Melement*Nelement;/将计算结果写入设备存储器中将计算结果写入设备存储器中Pdj*Width+i=Pvalue;编程模型变量和函数案例介绍注记：一个线程块中的每个线程计算Pd中的一个元素每个线程作的事情是：载入矩阵Md中的一行;载入矩阵Nd中的一列;为每对Md和Nd元素执行了一次乘法和加法;缺点：受存储器延迟影响很大；矩阵的大小受到线程块所能容纳最大线程数（512个线程）的限制