1 / 76

CUDA 程序设计

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 ：更多资源用于缓存及流控制

joyce
Télécharger la présentation

CUDA 程序设计

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. CUDA程序设计

  2. 主要内容 • GPGPU及CUDA介绍 • CUDA编程模型 • 多线程及存储器硬件

  3. GPGPU及CUDA介绍

  4. 多核时代 • 多个适当复杂度、相对低功耗内核并行工作 • 配置并行硬件资源提高处理能力 • 核心时钟频率基本不变 Quad-core Opteron IBM Cell Broadband Engine nVidia GT200

  5. Control ALU ALU ALU ALU DRAM Cache DRAM GPU与CPU硬件架构的对比 • CPU：更多资源用于缓存及流控制 • GPU：更多资源用于数据计算 • 适合具备可预测、针对数组的计算模式 GPU CPU

  6. 应用范围 CPU: control processor • 不规则数据结构 • 不可预测存取模式 • 递归算法 • 分支密集型算法 • 单线程程序 GPU: data processor • 规则数据结构 • 可预测存取模式 • 油气勘探、金融分析、医疗成像、有限元、基因分析、地理信息系统、…

  7. GPGPU (General Purpose Computing on GPU)

  8. GPGPU • 核心思想 • 用图形语言描述通用计算问题 • 把数据映射到vertex或者fragment处理器 • 缺点 • 硬件资源使用不充分 • 存储器访问方式严重受限 • 难以调试和查错 • 高度图形处理和编程技巧

  9. . . . . . . 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);

  10. 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

  11. 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

  12. 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.

  13. 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; }

  14. 性能提升情况 S.S. Stone, et al, “Accelerating Advanced MRI Reconstruction using GPUs,” ACM Computing Frontier Conference 2008, Italy, May 2008.

  15. 计算结果对比

  16. CUDA成功案例 广泛应用于生命科学、机械、石油、金融、数学、天文和通信等行业

  17. 医疗成像 • MRI (磁共振成像) • GRAPPA 自动校准 • 加速网格化 • 快速重建 • Computed Tomography (CT) • GE • Digisens SnapCT Stone, UIUC Batenburg, Sijbers et al

  18. 量子化学 双电子积分 RI-MP2 correlation energy in Q-Chem 3.1 K Yasuda, Nagoya U, Japan Leslie Vogt, Harvard

  19. 分子动力学 • 现有的分子动力学软件 • NAMD / VMD (alpha release) • GROMACS (alpha release) • HOOMD • OpenMM: 分子建模 • https://simtk.org/home/openmm

  20. 金融 • Monte Calo模拟 • 投资组合优化 • 期权及衍生品定价 • 对冲基金 • 风险分析 CUDA中的 随机数发生器 SciFinance的Monte Calo定价模型 SciComp Co.

  21. 生物信息学和生命科学 • 序列对比 • 蛋白质对接 • 生物系统的随机仿真(SSA) • 人体视觉皮层的自组织计算模型 • 分析基因表达的DNA微阵列工具 Schatz et al, U Maryland

  22. 流体动力学 • 3D Lattice-Boltzman解算器 • 基于Lattice-Boltzman的PDE解算器 • 用于照明的Lattice Boltzman • Navier-Stokes解算器 • 等离子体湍流建模 Thibault and Senocak Tolke and Krafczy

  23. 电磁学和电磁力学 • GPMAD: 离子束动力学模拟 • FDTD法进行的光散射模拟 • Acceleware的解算器 FDTD加速 Accelerware

  24. 天气, 大气, 海洋科学与空间建模 • 天气研究与预测模型 (WRF) • 25% ~ 30%的性能提升 • 海啸模拟

  25. 加密编码

  26. 模式匹配

  27. CUDA编程模型

  28. CUDA设备与线程 • 计算设备（device） • 作为CPU（host）的协处理器 • 有独立的存储设备（device memory） • 同时启动大量线程 • 计算密集部分使用大量线程并行的kernel • GPU与CPU线程的区别 • GPU的线程非常轻量，线程切换~1 cycle，而CPU需要~1000 cycle • GPU上的线程数>1000时才能有效利用GPU的计算能力

  29. 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

  30. 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

  31. . . . . . . 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);

  32. 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);

  33. CUDA程序的编译 • 使用nvcc编译工具 nvcc <filename>.cu [-o excutable] • 调试选项：-g（debug）、-deviceemu（CPU模拟GPU）

  34. 并行线程组织 并行性的维度 • 一维 y = a + b • 二维 P = M  N • 三维 CT or MRI

  35. 并行线程组织结构 • Thread: 并行的基本单位 • Thread block: 互相合作的线程组 • Cooperative Thread Array (CTA) • 允许彼此同步 • 通过快速共享内存交换数据 • 以1维、2维或3维组织 • 最多包含512个线程 • Grid: 一组thread block • 以1维、2维或3维组织 • 共享全局内存 • Kernel: 在GPU上执行的核心程序 • One kernel ↔ one grid

  36. 线程层次

  37. Block and Thread IDs Blocks 和 Threads 具有IDs • threadIdx, blockIdx • Block ID: 1D or 2D • Thread ID: 1D, 2D or 3D • 由此决定相应处理数据

  38. 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 …

  39. 一个简单的例子——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); }

  40. CUDA线程的同步 • void __syncthreads(); • Barrier synchronization • 同步thread block之内的所有线程 • 避免访问共享内存时发生RAW/WAR/WAW 冒险 __shared__ float scratch[256]; scratch[threadID] = begin[threadID]; __syncthreads(); int left = scratch[threadID -1]; 在此等待，直至所有线程到达才开始执行下面的代码

  41. 存储器模型与内存分配 • 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

  42. GPU Global Memory分配 • cudaMalloc() • 分配显存中的global memory • 两个参数 • 对象数组指针 • 数组尺寸 • cudaFree() • 释放显存中的global memory • 一个参数 • 对象数组指针

  43. 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);

  44. Host - Device数据交换 • cudaMemcpy() • 在存储器直接传输数据 • 四个参数 • 目的对象数组指针 • 源对象数组指针 • 数组尺寸 • 传输方向 • Host到Host • Host到Device • Device到Host • Device到Device

  45. Host - Device数据交换 代码实例 • M.elements: CPU主存 • Md: GPU显存 • 符号常数: cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost cudaMemcpy(Md, M.elements, size, cudaMemcpyHostToDevice); cudaMemcpy(M.elements, Md, size, cudaMemcpyDeviceToHost);

  46. 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私有

  47. Built-in dim3 Type • 定义grid和thread block的组织 • dim3 dimGrid(2, 2); • dim3 dimBlock(4, 2, 2); • kernelFunction<<< dimGrid, dimBlock>>>(…);

  48. CUDA函数定义 • __global__定义kernel函数 • 必须返回void • __device__和__host__ 可以组合使用 • 则被定义的函数在CPU和GPU上都被编译

  49. CUDA函数定义 • __device__ 函数不能用&运算符取地址 • 限制 • 不支持递归调用 • 不支持静态变量(static variable) • 不支持可变长度参数函数调用 • type va_list(stdarg.h) • double average(int count, ...)

  50. 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 >>>(...);

More Related