cuda程序设计

模式匹配
Lenovo Confidential | 2010 Lenovo
CUDA编程模型 编程模型
Lenovo Confidential | 2010 Lenovo
CUDA设备与线程 设备与线程
计算设备（device） 计算设备（ ）
– 作为 作为CPU（host）的协处理器 （ ） – 有独立的存储设备（device memory） 有独立的存储设备（ ） – 同时启动大量线程
Lenovo Confidential | 2010 Lenovo
量子化学
双电子积分 K Yasuda, Nagoya U, Japan
RI-MP2 correlation energy in Q-Chem 3.1 Leslie Vogt, Harvard
Lenovo Confidential | 2010 Lenovo
Thibault and Senocak
Tolke and Krafczy
Lenovo Confidential | 2010 Lenovo
电磁学和电磁力学
GPMAD: 离子束动力学模拟 FDTD法进行的光散射模拟 法进行的光散射模拟 Acceleware的解算器 的解算器
FDTD加速 加速 Accelerware
Spiral scan data + Iterative recon Reconstruction requires a lot of computation
Lenovo Confidential | 2010 Lenovo
Advanced MRI Reconstruction
( F H F +λW H W ) ρ = F H d
金融
Monte Calo模拟 模拟
投资组合优化 期权及衍生品定价 对冲基金 风险分析
SciFinance的Monte Calo定价模型 的 定价模型 SciComp Co. CUDA中的 中的 随机数发生器
Lenovo Confidential | 2010 Lenovo
生物信息学和生命科学
CPU：更多资源用于缓存及流控制 ： GPU：更多资源用于数据计算 ：
– 适合具备可预测、针对数组的计算模式 适合具备可预测、
ALU Control ALU Cache
ALU ALU
DRAM
DRAM
CPU
GPU
Lenovo Confidential | 2010 Lenovo
应用范围
CPU: control processor 不规则数据结构 不可预测存取模式 递归算法 分支密集型算法 单线程程序 GPU: data processor 规则数据结构 可预测存取模式
Compute Q More than Acquire Data 99.5% of time Compute FHd
Q只和扫描参数有关 只和扫描参数有关 FHd是数据相关的 是数据相关的 使用线性求解器计算 使用线性求解器计算ρ
Find ρ
Haldar, et al, “Anatomically-constrained reconstruction from noisy data,” MR in Medicine.
20
10 1.0 0.3 0
.B as e
sc e
lo c
PU .C on st M em
un e
P
rig
P
G PU .R eg Al
G PU .C oa le
G PU
PU .F
PU .T
S.S. Stone, et al, “Accelerating Advanced MRI Reconstruction using GPUs,” ACM Computing Frontier Conference 2008, Italy, May 2008.
Grid 1 ...
Lenovo Confidential | 2010 Lenovo
CUDA极大提高了现有应用的效果 极大提高了现有应用的效果
MRI Reconstruction
Cartesian Scan Data Gridding1
(b)
Spiral Scan Data
(a)
(b)
(c)
FFT
– 油气勘探、金融分析、医疗成像、有限 油气勘探、金融分析、医疗成像、 元、基因分析、地理信息系统、… 基因分析、地理信息系统、
Lenovo Cwenku.baidu.comnfidential | 2010 Lenovo
GPGPU (General Purpose Computing on GPU)
Lenovo Confidential | 2010 Lenovo
GPGPU
核心思想
– 用图形语言描述通用计算问题 – 把数据映射到 把数据映射到vertex或者 或者 fragment处理器 处理器
缺点
– 硬件资源使用不充分 – 存储器访问方式严重受限 – 难以调试和查错 – 高度图形处理和编程技巧
Lenovo Confidential | 2010 Lenovo
多核时代
多个适当复杂度、相对低功耗内核并行工作 多个适当复杂度、
– 配置并行硬件资源提高处理能力
核心时钟频率基本不变
Quad-core Opteron
IBM Cell Broadband Engine
nVidia GT200
Lenovo Confidential | 2010 Lenovo
GPU与CPU硬件架构的对比 与 硬件架构的对比
计算密集部分使用大量线程并行的 计算密集部分使用大量线程并行的kernel GPU与CPU线程的区别 与 线程的区别
Lenovo Confidential | 2010 Lenovo
Code
CPU
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; } }
Lenovo Confidential | 2010 Lenovo
天气, 大气, 天气, 大气, 海洋科学与空间建模
天气研究与预测模型 (WRF) 25% ~ 30%的性能提升 的性能提升 海啸模拟
Lenovo Confidential | 2010 Lenovo
加密编码
Lenovo Confidential | 2010 Lenovo
Lenovo Confidential | 2010 Lenovo
性能提升情况
600 47.5 53.9 GFLOPS Time 495.2 60
500
50
34.1 300 22.5
28.8
30
200
127.5 100 8.0 0 16.8 7.0 11.1 13.1 85.8 4.4 1.2
152.5
Lenovo Confidential | 2010 Lenovo
Advanced MRI Reconstruction
Cartesian Scan Data Gridding Spiral Scan Data
ky
(b)
(a)
(b)
kx
(c)
FFT
Iterative Reconstruction
序列对比 蛋白质对接 生物系统的随机仿真(SSA) 生物系统的随机仿真 人体视觉皮层的自组织计算模型 分析基因表达的 分析基因表达的DNA微阵列工具 微阵列工具
Schatz et al, U Maryland
Lenovo Confidential | 2010 Lenovo
流体动力学
3D Lattice-Boltzman解算器 解算器 基于 基于Lattice-Boltzman的PDE解算器 的 解算器 用于照明的 用于照明的Lattice Boltzman Navier-Stokes解算器 解算器 等离子体湍流建模
CUDA (Compute Unified Device Architecture)
CUDA有效结合 有效结合CPU+GPU编程 有效结合 编程
串行部分在 串行部分在CPU上运行 上运行 并行部分在 并行部分在GPU上运行 上运行
CPU Serial Code GPU Parallel Kernel KernelA<<< nBlk, nTid >>>(args); CPU Serial Code GPU Parallel Kernel KernelB<<< nBlk, nTid >>>(args); Grid 0 ...
Iterative Spiral scan data + Gridding + FFT Reconstruction
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
Lenovo Confidential | 2010 Lenovo
医疗成像
MRI (磁共振成像 磁共振成像) 磁共振成像 – GRAPPA 自动校准
– 加速网格化 – 快速重建
Stone, UIUC
Computed Tomography (CT) – GE
– Digisens SnapCT
Batenburg, Sijbers et al
分子动力学
现有的分子动力学软件 – NAMD / VMD (alpha release)
– GROMACS (alpha release) – HOOMD
OpenMM: 分子建模 – https://simtk.org/home/openmm
Lenovo Confidential | 2010 Lenovo
GPU
__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; }
Lenovo Confidential | 2010 Lenovo
G
G
G
G
PU .M ul ti
PU .S
.D
C
C
as tT
PU
Time (min)
400 GFLOPS
40
计算结果对比
Lenovo Confidential | 2010 Lenovo
CUDA成功案例 成功案例
广泛应用于生命科学、机械、石油、金融、数学、 广泛应用于生命科学、机械、石油、金融、数学、天文和通信等行业
CUDA程序设计 程序设计
Lenovo Confidential | 2010 Lenovo
主要内容
GPGPU及CUDA介绍 及 介绍 CUDA编程模型 编程模型 多线程及存储器硬件
Lenovo Confidential | 2010 Lenovo
GPGPU及CUDA介绍 及 介绍
Lenovo Confidential | 2010 Lenovo