CUDA介绍与案例

1、CUDA 介绍与案例,1,介绍,应用领域： 游戏、图形动画、科学计算可视化、 地质、生物、物理模拟等；,编程模型,变量和函数,案例,介绍,2008年SIGGRAPH年会上，NVIDIA公司推出CUDA (Compute Unified Device Architecture) NVIDIA2008; CUDA是NVIDIA为自己的GPU编写的一套编译器及相关的库文件； 将GPU 视作数据并行计算设备,是一种新的处理和管理GPU 计算的硬件和软件架构；,编程模型,变量和函数,案例,介绍,GPU到CUDA,编程模型,变量和函数,案例,介绍,CPU 和 GPU 的 每秒浮点运算次数 和存储器带宽比较

2、,具有强大浮点 运算能力,GPU采用大量的执行单元，这些执行单元可以轻松的加载并行处理线程，而不像CPU那样的单线程处理； 主机和设备均维护自己的 DRAM，分别称为主机存储器和设备存储器,CPU 和 GPU 之间浮点功能的差异，原因在于 GPU 专为高计算密集型（数学运算与存储器运算的比率）、高度并行化的计算而设计； GPU 的设计能使更多晶体管用于数据处理，而非数据缓存和流控制 ；,编程模型,变量和函数,案例,介绍,高度并行化、多线程、多核处理器，操作 系统的多任务机制负责管理多个并发运行 的CUDA和应用程序对GPU的访问； 基于标准C语言, 可自由地调用GPU的并行 处理架构; 同时适

3、用于图形和通用并行计算应用程序， 成为图形处理器的主要发展趋势。,编程模型,变量和函数,案例,介绍,依次安装 Cuda Driver Cuda Toolkit Cuda SDK 在安装目录中包括： bin工具程序及动态链接库 doc文件 include 头文件 lib 链接库 open64基于open64的CUDA compiler src 一些自带例子，如 C: CUDA_SDKCsrc 安装程序会设定一些环境变量： CUDA_BIN_PATH CUDA_INC_PATH CUDA_LIB_PATH,编程模型,变量和函数,案例,介绍,CUDA 软件栈包含多个层，如图所示：设备驱动程序、应用程

4、序编程接口（API）及其运行时两个较高级别的通用数学库，即 CUFFT 和 CUBLAS,CUDA安装后,编程模型,变量和函数,案例,介绍,CUDA程序编译,CUDA的源文件被nvcc编译 NVCC编译器会分离源码中设备代码和主机代码，主机代码交由一般的C/C+编译器（gcc等）编译，设备代码由NVCC编译； 内核必须使用nvcc编译成二进制代码才能在设备端执行。,v_2011,编程模型,变量和函数,案例,介绍,内核函数,编程模型,变量和函数,案例,介绍,串行代码在主机上执行，而并行代码在设备上执行 当调用kernel的时候，则CUDA的每个线程并行地执行一段指令，这与通常C函数只执行一次不一

5、样。,执行模型,说明： 在CUDA架构下，一个程序分为两部分，即host 端和 device端； Host 端 是在 CPU上执行的部分； Device端 是在显卡芯片上执行的部分，该端程序又称为内核 函数（ kernel function） 通常Host端程序复制内核函数到显卡内存中，再由显卡芯片执行device端程序，完成后再由host端程序将结果从显卡内存中取回；,CUDA 允许程序员定义内核函数，实现配置； 块组织为一个一维或二维线程块网格，维度由 语法的第一个参数指定； 执行内核的每个线程都会被分配一个独特的线程 ID，可通过内置的 threadIdx 变量在内核中访问,编程模型,2

6、,2.1 线程模型（体系结构） 2.2 存储器模型（体系结构）,2.1 线程模型,并行线程结构： Thread: 并行的基本单位 索引：threadIdx（内置变量） Block (Thread block): 线程块 允许彼此同步，通过快速共享内存交换数据 每个块的线程数应是 warp （调度任务的最小单位）块大小的倍数，每个块的线程数受限 索引：blockIdx（内置变量） Grid: 一组thread block 共享全局内存 Kernel: 在GPU上执行的核心程序 One kernel One grid,说明： Grid 是一组Block, 以1维、2维或3维组织，共享全局内存； G

7、rid之间通过global memory交换数据 Block 是互相协作的线程组，通过global memory共享数据，允许彼此同步；以1维、2维或3维组织； 同一block内的thread可以通过shared memory和同步实现通信;,一个内核可能由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数; 线程索引(Index)及ID(IDentity)，索引为(x,y)的线程ID为(x+yDx);对于大小为(Dx,Dy,Dz)的三维块，索引为(x,y,z)的线程ID为(x+yDx+zDxDy)； 示意图,线程块索引及ID ,情况类似： 对于一维块来说，两者是相同的； 对

8、于大小为 (Dx，Dy) 的二维块来说，索引为 (x,y) 的线程块的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个线程对相应的数组下标求和； ！用

9、内置变量确定数组的下标,或 int main( ) int threadsPerBlock = 256; int 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; /参见例 ：vec

10、torAdd!,？哪个线程对哪个数组下标求和； 即第几个Block（blockIdx）的第几个线程，对相应的数组下标求和；,2.2,说明： 所有线程都可访问全局存储器； 所有的线程都可以访问只读存储区域：“常量存储区”和“纹理存储区”； 纹理存储器提供不同的寻址模式，为某些特定的数据格式提供数据过滤的能力； Block中所有thread都在同一个处理核心上运行。因此每个block的thread数目受到处理核心所拥有存储器资源的限制，在NVIDIA Tesla架构中，一个线程block最多可包含512个线程;,每个线程块都有一个共享存储器(Shared Memory)，该存储器对于块内的所有线程

11、都是可见的，并且与块具有相同的生命周期。 每个线程都有一个私有的本地存储器（LocalMemory）。 在同一个block中的thread通过共享存储器共享数据，相互协作，实现同步： _syncthreads( );,register 访问延迟极低； 每个线程占有的register有限，编程时不要为其分配过多 私有变量； local memory 寄存器被使用完毕，数据将被存储在局部存储器中； 此外，存储大型结构体或者数组（无法确定大小的数组），线程的输入和中间变量；,shared memory 访问速度与寄存器相似； 实现线程间通信的延迟最小； 保存公用的计数器,或block的公用结果； 声

12、明时使用关键字 _ shared_; constant memory 只读地址空间，位于显存，用于存储需频繁访问的只读参数； 由于其在设备上有片上缓存，比全局存储器读取效率高很多； 使用_ constant_ 关键字; 定义在所有函数之外；,globalmemory 存在于显存中，也称为线性内存; cudaMalloc（）函数分配; cudaFree（）函数释放; cudaMemcpy（）进行主机端与设备端数据传输; 对于二维和三维数组： cudaMallocPitch（）和cudaMalloc3D（）分配线性存储空间； cudaMemcpy2D（），cudaMemcpy3D（）向设备存储器进

13、行拷贝；,对象数组指针 用于指向 global memory的存储区域, 由Host端分配，例： float* d_A， d_B ，d_C; int N = 50000; size_t size = N * sizeof(float); cudaMalloc(void*) cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice) ； cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice) ；,texture memory：只读； 不是一块专门的存储器，涉及显存、纹理缓存、纹理拾取等纹理流水线操作； 在ke

14、rnel中访问纹理存储器的操作, 称为纹理拾取（texture fetching） 纹理拾取使用的坐标与数据在显存中的位置可以不同，通过纹理参考约定二者映射方式； 将显存中数据与纹理参考关联的操作，称为纹理绑定（texture binding）；,3,变量和函数,3.1 变量类型申明 3.2 函数类型申明 3.3 内置变量 3.4 内置向量 3.5 数学函数 3.6 与Opengl互操作,3.1,3.1,3.2 函数,gridDim（有多少Block） 变量的类型为 dim3，包含网格的维度。 blockDim (有多少Thread） 变量的类型为 dim3，包含块的维度。 blockIdx(

15、第几个Block) 变量的类型为 uint3，包含网格内的块索引 threadIdx(第几个Thread) 变量的类型为 uint3，包含块内的线程索引。 备注：dim3 是整型的向量类型，未指定的任何组件都将初始化为1,如 dim3 dimBlock(4, 2, 2); uint3 是无符号整型的向量类型；,3.3 内置变量,Type name : char1、uchar1、char2、uchar2、char3、uchar3、char4、 uchar4、short1、ushort1、short2、ushort2、short3、ushort3、short4、ushort4、int1、uint1

16、、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);,3.4 内置向量,CUDA支持标量运算的数学函数,包括： pow、sqrt、cbrt、hypot、exp、exp2、espm1、log、log2、log10、log1p、sin、c

17、os、tan、asin、acos、atan、atan2、sinh、cosh、tanh、asinh、acosh、atanh、ceil、floor、trunc、round等; 当函数有前缀“_”时，运算速度较快。,3.5 数学函数,3.6 与Opengl的互操作 将一个Opengl缓冲对象注册到 CUDA，通过 cudaGLRegisterBufferObject( ) ； 例如： GLuint bufferObj; cuGLRegisterBufferObject(bufferObj);,注册完成后，内核即可使用 cuGLMapBufferObject( ); /映射缓冲对象 返回设备存储器地址

18、，读取或写入缓冲对象； 即将CUDA内存的指针指向OpenGL的缓冲区； 例如： CUdeviceptr devPtr; int size; cuGLMapBufferObject(,用完之后 cuGLUnmapBufferObject( )； /解除映射 cuGLUnregisterBufferObject( )； /取消注册；,编程模型,变量和函数,案例,介绍,4,案例,4.1. 向量加法； 4.2. 矩阵乘积; 4.3. 矩阵分块乘积;,编程模型,变量和函数,案例,介绍,在设计并行程序时，需要找到被分解问题关键参数和线程下标的对应关系，并建立映射机制，以使GPU的并行计算能力得到充分发挥

19、； 此外，考虑GPU的软硬件资源，CUDA编程模型有一些参数上的限制，主要包括对所分配线程和线程块数量的规定； 如GeForces系列GPU在每个线程块支持的最大线程数是512，而每个网格所具有最大线程块的数量是65535,4.1.向量和矩阵加法,编程模型,变量和函数,案例,介绍,/方法一：_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,编程模

20、型,变量和函数,案例,介绍,方法二： int main( ) int threadsPerBlock = 256; int 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

21、 + 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 dimBlock(N, N); matAdd(Ad, Bd, Cd); . ,编程模型,变量和函数,案例,介绍,说明内存和线程管理的基本特性 本地存储器、寄存器的用法 线程ID的用法 主机和设备之间数据传输的API 为了方便，

22、以方形矩阵说明,4.2. 矩阵乘法,编程模型,变量和函数,案例,介绍,矩阵大小为 WIDTH x WIDTH 在没有采用分片优化算法的情况下： 一个线程计算P矩阵中的一个 元素 需要从全局存储器载入WIDTH次,方块矩阵乘法,编程模型,变量和函数,案例,介绍,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;

23、 +k) double a = Mi * width + k; double b = Nk * width + j; sum += a * b; Pi*Width + j = sum; ,编程模型,变量和函数,案例,介绍,向GPU传输矩阵数据 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和N矩阵分配空间，并将数据复制到设备存储器中 cudaMalloc(,GPU上的矩阵乘法,编程模型,变量和函数,案例,介绍,计算结果