760 likes | 1.02k Views
CUDA 程序设计. 主要内容. GPGPU 及 CUDA 介绍 CUDA 编程模型 多线程及存储器硬件. GPGPU 及 CUDA 介绍. 多核时代. 多个适当复杂度、相对低功耗内核并行工作 配置并行硬件资源提高处理能力 核心时钟频率基本不变. Quad-core Opteron. IBM Cell Broadband Engine. nVidia GT200. Control. ALU. ALU. ALU. ALU. DRAM. Cache. DRAM. GPU 与 CPU 硬件架构的对比. CPU :更多资源用于缓存及流控制
E N D
主要内容 • GPGPU及CUDA介绍 • CUDA编程模型 • 多线程及存储器硬件
多核时代 • 多个适当复杂度、相对低功耗内核并行工作 • 配置并行硬件资源提高处理能力 • 核心时钟频率基本不变 Quad-core Opteron IBM Cell Broadband Engine nVidia GT200
Control ALU ALU ALU ALU DRAM Cache DRAM GPU与CPU硬件架构的对比 • CPU:更多资源用于缓存及流控制 • GPU:更多资源用于数据计算 • 适合具备可预测、针对数组的计算模式 GPU CPU
应用范围 CPU: control processor • 不规则数据结构 • 不可预测存取模式 • 递归算法 • 分支密集型算法 • 单线程程序 GPU: data processor • 规则数据结构 • 可预测存取模式 • 油气勘探、金融分析、医疗成像、有限元、基因分析、地理信息系统、…
GPGPU • 核心思想 • 用图形语言描述通用计算问题 • 把数据映射到vertex或者fragment处理器 • 缺点 • 硬件资源使用不充分 • 存储器访问方式严重受限 • 难以调试和查错 • 高度图形处理和编程技巧
. . . . . . CUDA (Compute Unified Device Architecture) CUDA有效结合CPU+GPU编程 • 串行部分在CPU上运行 • 并行部分在GPU上运行 CPU Serial Code Grid 0 GPU Parallel Kernel KernelA<<< nBlk, nTid >>>(args); CPU Serial Code Grid 1 GPU Parallel Kernel KernelB<<< nBlk, nTid >>>(args);
CUDA极大提高了现有应用的效果 MRI Reconstruction Cartesian Scan Data Spiral Scan Data Gridding1 (b) (b) (c) (a) FFT Iterative Reconstruction Spiral scan data + Gridding + FFT Reconstruction requires little computation Based on Fig 1 of Lustig et al, Fast Spiral Fourier Transform for Iterative MR Image Reconstruction, IEEE Int’l Symp. on Biomedical Imaging, 2004
Advanced MRI Reconstruction Cartesian Scan Data Spiral Scan Data Gridding (b) (c) (a) (b) FFT Iterative Reconstruction Spiral scan data + Iterative recon Reconstruction requires a lot of computation
Advanced MRI Reconstruction Compute Q • Q只和扫描参数有关 • FHd是数据相关的 • 使用线性求解器计算ρ More than 99.5% of time Acquire Data Compute FHd Find ρ Haldar, et al, “Anatomically-constrained reconstruction from noisy data,” MR in Medicine.
Code CPU GPU for (p = 0; p < numP; p++) { for (d = 0; d < numD; d++) { exp = 2*PI*(kx[d] * x[p] + ky[d] * y[p] + kz[d] * z[p]); cArg = cos(exp); sArg = sin(exp); rFhD[p] += rRho[d]*cArg – iRho[d]*sArg; iFhD[p] += iRho[d]*cArg + rRho[d]*sArg; } } __global__ void cmpFhD(float* gx, gy, gz, grFhD, giFhD) { int p = blockIdx.x * THREADS_PB + threadIdx.x; // register allocate image-space inputs & outputs x = gx[p]; y = gy[p]; z = gz[p]; rFhD = grFhD[p]; iFhD = giFhD[p]; for (int d = 0; d < SCAN_PTS_PER_TILE; d++) { // s (scan data) is held in constant memory float exp = 2 * PI * (s[d].kx * x + s[d].ky * y + s[d].kz * z); cArg = cos(exp); sArg = sin(exp); rFhD += s[d].rRho*cArg – s[d].iRho*sArg; iFhD += s[d].iRho*cArg + s[d].rRho*sArg; } grFhD[p] = rFhD; giFhD[p] = iFhD; }
性能提升情况 S.S. Stone, et al, “Accelerating Advanced MRI Reconstruction using GPUs,” ACM Computing Frontier Conference 2008, Italy, May 2008.
CUDA成功案例 广泛应用于生命科学、机械、石油、金融、数学、天文和通信等行业
医疗成像 • MRI (磁共振成像) • GRAPPA 自动校准 • 加速网格化 • 快速重建 • Computed Tomography (CT) • GE • Digisens SnapCT Stone, UIUC Batenburg, Sijbers et al
量子化学 双电子积分 RI-MP2 correlation energy in Q-Chem 3.1 K Yasuda, Nagoya U, Japan Leslie Vogt, Harvard
分子动力学 • 现有的分子动力学软件 • NAMD / VMD (alpha release) • GROMACS (alpha release) • HOOMD • OpenMM: 分子建模 • https://simtk.org/home/openmm
金融 • Monte Calo模拟 • 投资组合优化 • 期权及衍生品定价 • 对冲基金 • 风险分析 CUDA中的 随机数发生器 SciFinance的Monte Calo定价模型 SciComp Co.
生物信息学和生命科学 • 序列对比 • 蛋白质对接 • 生物系统的随机仿真(SSA) • 人体视觉皮层的自组织计算模型 • 分析基因表达的DNA微阵列工具 Schatz et al, U Maryland
流体动力学 • 3D Lattice-Boltzman解算器 • 基于Lattice-Boltzman的PDE解算器 • 用于照明的Lattice Boltzman • Navier-Stokes解算器 • 等离子体湍流建模 Thibault and Senocak Tolke and Krafczy
电磁学和电磁力学 • GPMAD: 离子束动力学模拟 • FDTD法进行的光散射模拟 • Acceleware的解算器 FDTD加速 Accelerware
天气, 大气, 海洋科学与空间建模 • 天气研究与预测模型 (WRF) • 25% ~ 30%的性能提升 • 海啸模拟
CUDA设备与线程 • 计算设备(device) • 作为CPU(host)的协处理器 • 有独立的存储设备(device memory) • 同时启动大量线程 • 计算密集部分使用大量线程并行的kernel • GPU与CPU线程的区别 • GPU的线程非常轻量,线程切换~1 cycle,而CPU需要~1000 cycle • GPU上的线程数>1000时才能有效利用GPU的计算能力
Streaming Processor(SP) A fully pipelined, single-issue, inordermicroprocessor • 2 ALUs and a FPU • Register file • 32-bit scalar processing • No instruction fetch andscheduling • No cache
Streaming Multiprocessor(SM) An array of SPs • 8 streaming processor • 2 Special Function Units (SFU) • A 16KB read/write shared memory • Not a cache • But a software-managed data store • Multithreading issuing unit • Instruction and constant cache
. . . . . . CUDA 程序基本结构 • 串行部分在CPU上运行(host) • 并行部分在GPU上运行(device) CPU Serial Code (host) Grid 0 GPU Parallel Kernel (device) KernelA<<< nBlk, nTid >>>(args); CPU Serial Code (host) Grid 1 GPU Parallel Kernel(device) KernelB<<< nBlk, nTid >>>(args);
C扩展 • Declspecs • global, device, shared, local, constant • Keywords • threadIdx, blockIdx • Intrinsics • __syncthreads • Runtime API • Memory, symbol, execution management • Function launch __device__ float filter[N]; __global__ void convolve (float *image) { __shared__ float region[M]; ... region[threadIdx] = image[i]; __syncthreads() ... image[j] = result; } // Allocate GPU memory void *myimage = cudaMalloc(bytes) // 100 blocks, 10 threads per block convolve<<<100, 10>>> (myimage);
CUDA程序的编译 • 使用nvcc编译工具 nvcc <filename>.cu [-o excutable] • 调试选项:-g(debug)、-deviceemu(CPU模拟GPU)
并行线程组织 并行性的维度 • 一维 y = a + b • 二维 P = M N • 三维 CT or MRI
并行线程组织结构 • Thread: 并行的基本单位 • Thread block: 互相合作的线程组 • Cooperative Thread Array (CTA) • 允许彼此同步 • 通过快速共享内存交换数据 • 以1维、2维或3维组织 • 最多包含512个线程 • Grid: 一组thread block • 以1维、2维或3维组织 • 共享全局内存 • Kernel: 在GPU上执行的核心程序 • One kernel ↔ one grid
Block and Thread IDs Blocks 和 Threads 具有IDs • threadIdx, blockIdx • Block ID: 1D or 2D • Thread ID: 1D, 2D or 3D • 由此决定相应处理数据
threadID … float x = input[threadID]; float y = func(x); output[threadID] = y; … … float x = input[threadID]; float y = func(x); output[threadID] = y; … … float x = input[threadID]; float y = func(x); output[threadID] = y; … CUDA线程组织 • CUDA kernel函数由一系列线程组成 • 单指令多数据流(SPMD) • 通过IDs确定处理的数据 • 线程可划分为不同的Block • 在同一个block中,可以通过share memory、atomic operation和barrier synchronization进行协同 Thread Block 1 Thread Block N - 1 Thread Block 0 …
一个简单的例子——Increment Array Elements //CPU program void inc_cpu(float *a, float b, int N) { for (intidx = 0; idx<N; idx++) a[idx] = a[idx] + b; } void main() { … inc_cpu(a, b, N); } //CUDA program __global__ void inc_gpu(float *a, float b, int N) { intidx =blockIdx.x* blockDim.x+ threadIdx.x; if (idx < N) a[idx] = a[idx] + b; } void main() { … dim3 dimBlock (blocksize); dim3 dimGrid( ceil( N / (float)blocksize) ); inc_gpu<<<dimGrid, dimBlock>>>(a, b, N); }
CUDA线程的同步 • void __syncthreads(); • Barrier synchronization • 同步thread block之内的所有线程 • 避免访问共享内存时发生RAW/WAR/WAW 冒险 __shared__ float scratch[256]; scratch[threadID] = begin[threadID]; __syncthreads(); int left = scratch[threadID -1]; 在此等待,直至所有线程到达才开始执行下面的代码
存储器模型与内存分配 • R/W per-thread registers • 1-cycle latency • R/W per-thread local memory • Slow – register spilling to global memory • R/W per-block shared memory • 1-cycle latency • But bank conflicts may drag down • R/W per-grid global memory • ~500-cycle latency • But coalescing accessing could hide latency • Read only per-grid constant and texturememories • ~500-cycle latency • But cached
GPU Global Memory分配 • cudaMalloc() • 分配显存中的global memory • 两个参数 • 对象数组指针 • 数组尺寸 • cudaFree() • 释放显存中的global memory • 一个参数 • 对象数组指针
GPU Global Memory分配 代码实例 • 分配64 64单精度浮点数组 • 数组指针Md • 建议用“d”表示GPU显存数据结构 int BLOCK_SIZE = 64; float* Md; int size = BLOCK_SIZE * BLOCK_SIZE * sizeof(float); cudaMalloc((void**)&Md, size); cudaFree(Md);
Host - Device数据交换 • cudaMemcpy() • 在存储器直接传输数据 • 四个参数 • 目的对象数组指针 • 源对象数组指针 • 数组尺寸 • 传输方向 • Host到Host • Host到Device • Device到Host • Device到Device
Host - Device数据交换 代码实例 • M.elements: CPU主存 • Md: GPU显存 • 符号常数: cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost cudaMemcpy(Md, M.elements, size, cudaMemcpyHostToDevice); cudaMemcpy(M.elements, Md, size, cudaMemcpyDeviceToHost);
CUDA变量与函数 CUDA引入的变量修饰词 • __device__ • 储存于GPU上的global memory空间 • 和应用程序具有相同的生命期(lifetime) • 可被grid中所有线程存取, CPU代码通过runtime函数存取 • __constant__ • 储存于GPU上的constant memory空间 • 和应用程序具有相同的生命期(lifetime) • 可被grid中所有线程存取, CPU代码通过runtime函数存取 • __shared__ • 储存于GPU上thread block内的共享存储器 • 和thread block具有相同的生命期(lifetime) • 只能被thread block内的线程存取 • 无修饰(Local变量) • 储存于SM内的寄存器和local memory • 和具有相同的生命期(lifetime) • Thread私有
Built-in dim3 Type • 定义grid和thread block的组织 • dim3 dimGrid(2, 2); • dim3 dimBlock(4, 2, 2); • kernelFunction<<< dimGrid, dimBlock>>>(…);
CUDA函数定义 • __global__定义kernel函数 • 必须返回void • __device__和__host__ 可以组合使用 • 则被定义的函数在CPU和GPU上都被编译
CUDA函数定义 • __device__ 函数不能用&运算符取地址 • 限制 • 不支持递归调用 • 不支持静态变量(static variable) • 不支持可变长度参数函数调用 • type va_list(stdarg.h) • double average(int count, ...)
Kernel函数调用 • 调用时必须给出线程配置方式 __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 >>>(...);