华南理工大学 陈虎 博士 tommychen74@yahoo.com.cn CUDA编程模型 华南理工大学 陈虎 博士 tommychen74@yahoo.com.cn

GPU与CPU的差异 GPU(Graphics Process Unit) 通用CPU 面向计算密集型和大量数据并行化的计算 大量的晶体管用于计算单元 面向通用计算 大量的晶体管用于Cache和控制电路 DRAM Cache ALU Control DRAM CPU GPU

GPU与CPU的峰值速度比较 1 Based on slide 7 of S. Green, “GPU Physics,” SIGGRAPH 2007 GPGPU Course. http://www.gpgpu.org/s2007/slides/15-GPGPU-physics.pdf

第一代GPU结构(GeForce 6800)

第二代GPU(GeForece 8800)

GeForce 8800的主要技术参数 16个流多处理器(SM) 每个SM中包含了8个流处理器SP 每个SP包含一个乘加单元 晶体管数目（百万） 681 工艺 90nm 芯片面积(mm2) 470 工作主频（GHZ） 1.5 峰值速度（Gflops） 576 处理器数目 128 片上存储器容量（KB） 488 功耗（W） 150 16个流多处理器(SM) 每个SM中包含了8个流处理器SP 每个SP包含一个乘加单元 每个SM管理了24个线程簇(warp)，共有768个线程 采用单线程多数据(Single-thread, Multiple-data)技术 每个周期在8个SP上并行执行一个线程簇

GF8800的存储层次结构 GeForce 8800 层次 容量 (KB) 延迟 (ns) 局部存储器 16 26 L1 Cache 5 280 L2 Cache 32 370 DDR 510 GeForce 8800

GPU最适合做什么？ 对多个数据进行同一种运算（STMD适用） 一次存储器访问，多次运算（外部DDR访问开销高，局部存储器容量较小） 浮点计算比例高（特别是单精度浮点） 典型计算：物理模拟，线性代数计算 应用领域： 计算生物学 图像处理

CUDA编程模型 Nvidia公司开发的编程模型 可以和VC 8.0集成使用 下载地址： http://www.nvidia.com/object/cuda_get.html

CUDA工具链 Compiler Target code C/C++ CUDA Application NVCC CPU Code float4 me = gx[gtid]; me.x += me.y * me.z; NVCC CPU Code 链接时的动态库： The CUDA runtime library (cudart) The CUDA core library (cuda) PTX Code Virtual Physical PTX to Target Compiler ld.global.v4.f32 {$f1,$f3,$f5,$f7}, [$r9+0]; mad.f32 $f1, $f5, $f3, $f1; G80 … GPU Target code

CUDA中的线程 线程： 每个线程有一个唯一的标识ID——threadIdx 若干线程还可以组成块(Block) 硬件支持，开销很小； 所有线程执行相同的代码（STMD） 每个线程有一个唯一的标识ID——threadIdx 若干线程还可以组成块(Block) 线程块可以呈一维、二维或者三维结构 Device Grid 1 Block (0, 0) (1, 0) (2, 0) (0, 1) (1, 1) (2, 1) Block (1, 1) Thread (3, 1) (4, 1) (0, 2) (1, 2) (2, 2) (3, 2) (4, 2) (3, 0) (4, 0)

CUDA存储器分配 cudaMalloc() cudaFree() 在全局存储器中分配空间 两个参数： 在全局存储器中回收空间 参数 地址指针 空间大小 cudaFree() 在全局存储器中回收空间 参数 回收空间地址指针 Grid Block (0, 0) Block (1, 0) Shared Memory Shared Memory Registers Registers Registers Registers Thread (0, 0) Thread (1, 0) Thread (0, 0) Thread (1, 0) Host Global Memory

示范代码 分配64 * 64的单精度浮点数组 存储器地址为Md.elements TILE_WIDTH = 64; Matrix Md /*“d” is often used to indicate a device data structure*/ int size = TILE_WIDTH * TILE_WIDTH * sizeof(float); cudaMalloc((void**)&Md.elements, size); cudaFree(Md.elements);

主机和设备之间的数据传输 cudaMemcpy() 异步传输 存储器数据传输 参数： 目的地址 源地址 传输字节数 传输类型 主机<->主机 主机<->全局存储器 全局存储器<->全局存储器 异步传输 Grid Block (0, 0) Block (1, 0) Shared Memory Shared Memory Registers Registers Registers Registers Thread (0, 0) Thread (1, 0) Thread (0, 0) Thread (1, 0) Host Global Memory

示范代码 传输 64 * 64 单精度浮点数组 M在主机存储器中 Md在设备存储器中 传输方向常数： cudaMemcpyHostToDevice cudaMemcpyDeviceToHost cudaMemcpy(Md.elements, M.elements, size, cudaMemcpyHostToDevice); cudaMemcpy(M.elements, Md.elements, size, cudaMemcpyDeviceToHost);

CUDA的函数声明 __global__定义一个内核函数 必须返回 void __device__ 和 __host__可以一起使用 在哪里执行: 只能从哪里调用: __device__ float DeviceFunc() 设备 __global__ void KernelFunc() 主机 __host__ float HostFunc() __global__定义一个内核函数 必须返回 void __device__ 和 __host__可以一起使用

CUDA 的函数声明(续) __device__ 函数不能使用函数指针； __device__函数在设备上执行，所以: 不能递归 不能在函数内定义静态变量 不能使用变量作为参数

调用一个内核函数– 线程生成 内核函数必须使用执行配置调用: 任何从CUDA 1.0调用的内核函数都是异步的, 阻塞需要显式的同步。 __global__ void KernelFunc(...); dim3 DimGrid(100, 50); // 5000 thread blocks dim3 DimBlock(4, 8, 8); // 256 threads per block size_t SharedMemBytes = 64; // 64 bytes of shared memory KernelFunc<<< DimGrid, DimBlock, SharedMemBytes >>> (...); //用<<<, >>>指明该函数调用的线程数 任何从CUDA 1.0调用的内核函数都是异步的, 阻塞需要显式的同步。

CUDA变量类型修饰符 当使用__local__, __shared__, 和 __constant__时， __device__ 是可选的 变量声明 存储器 作用域 生命期 __device__ __local__ int LocalVar; 本地存储器 thread __device__ __shared__ int SharedVar; 共享存储器 block __device__ int GlobalVar; 全局存储器 grid application __device__ __constant__ int ConstantVar; 常存储器 当使用__local__, __shared__, 和 __constant__时， __device__ 是可选的 没有任何修饰符的自动变量将默认放在寄存器内 数组除外，它们将会被摆放在本地存储器内

在哪里声明变量? 寄存器 (自动) 共享存储器 本地存储器 全局存储器 常存储器 20

G80 的CUDA存储器实现 每一个thread 可以: 可以读写寄存器 可以读写本地存储器 每一个block可以读写 共享存储器 每一个grid 可以读写 全局存储器 每一个grid 只可读 常数存储器 Grid Block (0, 0) Block (1, 0) Shared Memory Shared Memory Registers Registers Registers Registers Thread (0, 0) Thread (1, 0) Thread (0, 0) Thread (1, 0) Host Global Memory Constant Memory 21

通用的编程策略 全局存储器在设备内存中(DRAM) 它的访问速度比共享存储器要慢 在设备上执行运算的优化方法是： 把数据分割成子集，放入共享存储器中 以一个线程块操作一个数据子集： 从全局存储器读入数据子集到共享存储器，使用多个线程实现存储级并行 在共享存储器中对数据子集执行运算; 每一个线程可以有效率地多次访问任何数据元素 从共享存储器上拷贝结果到全局存储器 22

通用的编程策略(续) 常数存储器也在设备内存里(DRAM) –比访问共享存储器慢得多 通过访问模式小心划分数据 但是它可以高速缓存，对于只读数据的访问非常高效。 通过访问模式小心划分数据 只读 常数存储器(若在高速缓存内，非常快) 在块内读写共享 共享存储器(非常快) 线程可读写 寄存器 (非常快) 读写输入和结果 全局存储器(非常慢)

变量类型约束 指针只能指向在全局存储器已分配或声明的内存 : 在主机分配，然后传递给内核: __global__ void KernelFunc(float* ptr) 以全局变量的地址获得: float* ptr = &GlobalVar; 24

CUDA编程框架 //全局变量声明 __host__ ,…, __device__... __global__, __constant__, __texture__ //函数原型声明 __global__ void kernelOne(…) //内核函数 float handyFunction(…) //普通函数 main(){ cudaMalloc(&d_GlblVarPtr, bytes ) //在设备上分配空间 cudaMemCpy(d_GlblVarPtr, h_Gl…) //从主机传输数据到设备 执行内核函数时的配置 kernelOne<<<execution configuration>>>( args… );//调用内核函数 cudaMemCpy(h_GlblVarPtr,…); //从设备传输结果到主机 } __global__ void kernelOne(type args,…){ //内核函数 //局部变量声明 __local__, __shared__ //自动变量被默认分配到寄存器或本地存储器中 float handyFunction(int inVar…){ //普通函数