CUDA介绍与案例.ppt

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编译成二进制代码才能在设备端执行。,编程模型,变量和函数,案例,介绍,内核函数,编程模型,变量和函数,案例,介绍,串行代码在主机上执行，而并行代码在设备上执行 当调用kernel的时候，则CUDA的每个线程并行地执行一段指令，这与通常C函数只执行一次不一样。,执行模型

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

6、模型（体系结构） 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维组织，共享全局内存； Grid之间通过

7、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 ,情况类似： 对于一维块来说，两者是相同的； 对于大小为 (D

8、x，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 (A, B, C); 参见例：matrix_src.doc,编程模

10、型,变量和函数,案例,介绍,方法二： 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 (Ad, Bd, C

11、d); . ,编程模型,变量和函数,案例,介绍,说明内存和线程管理的基本特性 本地存储器、寄存器的用法 线程ID的用法 主机和设备之间数据传输的API 为了方便，以方形矩阵说明,4.2. 矩阵乘法,编程模型,变量和函数,案例,介绍,矩阵大小为 WIDTH x WIDTH 在没有采用分片优化算法的情况下： 一个线程计算P矩阵中的一个 元素 需要从全局存储器载入WIDTH次,方块矩阵乘法,编程模型,变量和函数,案例,介绍,CPU上的矩阵乘法 void MatrixMulOnHost(float* M, float* N, float* P, int Width) for (int i = 0; i

12、 Width; +i) for (int j = 0; j (Md, Nd, Pd，Width); / 从设备中读取P矩阵的数据 cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost); / 释放设备存储器中的空间 cudaFree(Md); cudaFree(Nd); cudaFree (Pd); ,编程模型,变量和函数,案例,介绍,/ 矩阵乘法的内核函数每个线程都要执行的代码 _global_ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) / 2维的线程ID号 int i = threadIdx.x; int j = threadIdx.y; /Pvalue用来保存被每个线程计算完成后的矩阵的元素 float Pvalue = 0;,编程模型,变量和函数,案例,介绍,/每个线程计算一个元素 for (int k = 0; k Width; +k) float Melement = Mdj* Width+k; float Nelement = Ndk * Width+i; Pvalue += Melement *