CUDA介绍与案例

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

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

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

4、e 头文件 lib 链接库 open64基于open64的CUDA compiler src 一些自带例子，如 C: CUDA_SDKCsrc 安装程序会设定一些环境变量： CUDA_BIN_PATH CUDA_INC_PATH CUDA_LIB_PATH 编程模型变量和函数案例介绍 CUDA 软件栈包含多个层，如图所示：设备驱动程序、应用程序编程接口（API）及其运行时两个较高级别的通用数学库，即 CUFFT 和 CUBLAS CUDA安装后编程模型变量和函数案例介绍 CUDA程序编译p CUDA的源文件被nvcc编译 NVCC编译器会分离源码中设备代码和主机代码，主机代码交由一般的C/C+

5、编译器（gcc等）编译，设备代码由NVCC编译； 内核必须使用nvcc编译成二进制代码才能在设备端执行。v_2011编程模型变量和函数案例介绍标注字体使用 行距背景图片出处声明英文 Century Gothic中文 微软雅黑正文 互联网是一个开放共享的平台OfficePLUS 部分设计灵感与元素来源于网络如有建议请联系 OfficePLUS内核函数编程模型变量和函数案例介绍编程模型变量和函数案例介绍n 串行代码在主机上串行代码在主机上执行，而并行代码执行，而并行代码在设备上执行在设备上执行n 当调用当调用kernelkernel的时的时候，则候，则CUDACUDA的每个的每个线程并行地执行一线

6、程并行地执行一段指令，这与通常段指令，这与通常C C函数只执行一次不函数只执行一次不一样。一样。执行模型执行模型编程模型变量和函数案例介绍说明：n在CUDA架构下，一个程序分为两部分，即host 端和 device端； Host 端 是在 CPU上执行的部分； Device端 是在显卡芯片上执行的部分，该端程序又称为内核 函数（ kernel function） n 通常Host端程序复制内核函数到显卡内存中，再由显卡芯片执行device端程序，完成后再由host端程序将结果从显卡内存中取回；编程模型变量和函数案例介绍n CUDA 允许程序员定义内核函数，实现配置；n 块组织为一个一维或二维线

7、程块网格，维度由 语法的第一个参数指定；n 执行内核的每个线程都会被分配一个独特的线程 ID，可通过内置的 threadIdx 变量在内核中访问编程模型22.1 线程模型（体系结构）2.2 存储器模型（体系结构）编程模型变量和函数案例介绍2.1 线程模型编程模型变量和函数案例介绍并行线程结构：Thread: 并行的基本单位索引：threadIdx（内置变量）Block (Thread block): 线程块允许彼此同步，通过快速共享内存交换数据每个块的线程数应是 warp （调度任务的最小单位）块大小的倍数，每个块的线程数受限索引：blockIdx（内置变量） Grid: 一组thread b

8、lock共享全局内存Kernel: 在GPU上执行的核心程序One kernel One grid编程模型变量和函数案例介绍说明： Grid 是一组Block, 以1维、2维或3维组织，共享全局内存； Grid之间通过global memory交换数据 Block 是互相协作的线程组，通过global memory共享数据，允许彼此同步；以1维、2维或3维组织； 同一block内的thread可以通过shared memory和同步实现通信;编程模型变量和函数案例介绍 一个内核可能由多个大小相同的线程块执行，因而线程总数线程总数应等于每个块的线程数乘以块的数; 线程索引线程索引(Index)(

9、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) 的线程块的ID 是 （x + yDx）； 对于大小为 （Dx，Dy， Dz） 的三维块来说，索引 为（x, y, z） 的线程的ID 是 （x + yDx +zDxDy）; /+图示说明编程模型变量和函数案例介绍 例: 向

10、量的和 _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; int blocksPerGrid = (N + threadsPerBlock - 1) /threadsPerBlock; Vec

11、Add(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 编程模型变量和函数案例介绍编程模型变量和函数案例介绍说明：p

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

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

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

15、udaMemcpy（）进行主机端与设备端数据传输; 对于二维和三维数组： 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) ;

16、 cudaMalloc(void*)&d_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 fe

17、tching） 纹理拾取使用的坐标与数据在显存中的位置可以不同，通过纹理参考约定二者映射方式； 将显存中数据与纹理参考关联的操作，称为纹理绑定（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 函数函数编程模型变量和函数案例介绍n gridDimgridDim（有多少（有多少Blo

18、ckBlock） 变量的类型为 dim3，包含网格的维度。n blockDimblockDim ( (有多少有多少ThreadThread） 变量的类型为 dim3，包含块的维度。n blockIdxblockIdx( (第几个第几个Block)Block) 变量的类型为 uint3，包含网格内的块索引 n threadIdxthreadIdx( (第几个第几个Thread)Thread) 变量的类型为 uint3，包含块内的线程索引。 备注：dim3 是整型的向量类型，未指定的任何组件都将初始化为1,如 dim3 dimBlock(4, 2, 2); uint3 是无符号整型的向量类型；3.

19、3 3.3 内置变量内置变量编程模型变量和函数案例介绍 Type name : char1、uchar1、char2、uchar2、char3、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

20、 个组件分别可通过字段 x、y、z 和 w 访问； 均附带形式为 make_ 的构造函数 例如： int2 v=make_int2(6, 9); 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 数学函数数

21、学函数编程模型变量和函数案例介绍3.6 3.6 与与OpenglOpengl的互的互操作操作将一个Opengl缓冲对象注册到 CUDA，通过通过 cudaGLRegisterBufferObjeccudaGLRegisterBufferObject( ) ；例如： GLuint bufferObj; cuGLRegisterBufferObject(bufferObj);编程模型变量和函数案例介绍 注册完成后，内核即可使用 cuGLMapBufferObjectcuGLMapBufferObject( ( ); / ); /映射缓冲对象 返回设备存储器地址，读取或写入缓冲对象； 即将CUDA内

22、存的指针指向OpenGL的缓冲区； 例如： CUdeviceptr devPtr; int size; cuGLMapBufferObject(&devPtr, &size, bufferObj); 用完之后 cuGLUnmapBufferObjectcuGLUnmapBufferObject( ( ) )； /解除映射 cuGLUnregisterBufferObjectcuGLUnregisterBufferObject( ( ) )； /取消注册；编程模型变量和函数案例介绍4案例 4.1. 4.1. 向量加法；向量加法； 4.2. 4.2. 矩阵乘积矩阵乘积; ; 4.3

23、.4.3. 矩阵分块乘积矩阵分块乘积; ; 编程模型变量和函数案例介绍在设计并行程序时，需要找到被分解问题关键参数和需要找到被分解问题关键参数和线程下标的对应关系，并建立映射机制线程下标的对应关系，并建立映射机制，以使GPU的并行计算能力得到充分发挥；此外，考虑GPU的软硬件资源，CUDA编程模型有一些参数上的限制，主要包括对所分配线程和线程块数量分配线程和线程块数量的规定； 如GeForces系列GPU在每个线程块支持的最大线程数是512，而每个网格所具有最大线程块的数量是655354.1.向量和矩阵加法编程模型变量和函数案例介绍/方法一：_global_ void VecAdd(float

24、* A, float* B, float* C) int i = threadIdx.x;Ci = Ai + Bi;int main() / Kernel invocationVecAdd(A, B, C); 参见例：matrix_src.doc编程模型变量和函数案例介绍方法二：int main( ) int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;VecAdd(d_A, d_B, d_C, N); 编程模型变量和函数案例介绍_global_ void VecA

25、dd(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 main( ) / Kernel invocation dim3 dimB

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

27、从全局存储器载入需要从全局存储器载入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*

28、Width + j = sum; 编程模型变量和函数案例介绍向向GPUGPU传输矩阵数据传输矩阵数据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线性存储矩阵元素；线性存储矩阵元素； cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); cudaMalloc(&Nd, size); / Nd线性存储矩阵元素；线性存储矩阵元素； cudaMemcpy(Nd, N, size, CudaMemcpyHostToDevice); /在设备存储器上给在设备存储器上给P P矩阵分配空间矩阵分配空间 cudaMalloc(&