矩阵乘法性能优化

多核时代
• 多个适当复杂度、相对低功耗内核并行工作
– 配置并行硬件资源提高处理能力
• 核心时钟频率基本不变
Quad-core Opteron
IBM Cell Broadband Engine
nVidia GT200
Lenovo Confidential | © 2010 Lenovo
GPU与CPU硬件架构的对比
CartesiLeabharlann Baidun Scan Data Spiral Scan Data Gridding1
ky kx
ky
(b)
ky
kx
(a)
(b)
kx
(c)
FFT
Iterative Spiral scan data + Gridding + FFT Reconstruction
Reconstruction requires little computation
• CPU：更多资源用于缓存及流控制 • GPU：更多资源用于数据计算
– 适合具备可预测、针对数组的计算模式
ALU Control
ALU
ALU Cache
ALU
DRAM
DRAM
CPU
GPU
Lenovo Confidential | © 2010 Lenovo
应用范围
CPU: control processor GPU: data processor
CUDA (Compute Unified Device Architecture)
CUDA有效结合CPU+GPU编程
• 串行部分在CPU上运行
• 并行部分在GPU上运行
CPU Serial Code GPU Parallel Kernel KernelA<<< nBlk, nTid >>>(args);
Lenovo Confidential | © 2010 Lenovo
GPGPU
• 核心思想
– 用图形语言描述通用计算问题 – 把数据映射到vertex或者 fragment处理器
• 缺点
– 硬件资源使用不充分 – 存储器访问方式严重受限 – 难以调试和查错 – 高度图形处理和编程技巧
Lenovo Confidential | © 2010 Lenovo
• 不规则数据结构
• 不可预测存取模式 • 递归算法 • 分支密集型算法 • 单线程程序
• 规则数据结构
• 可预测存取模式
– 油气勘探、金融分析、医疗成像、有限 元、基因分析、地理信息系统、…
Lenovo Confidential | © 2010 Lenovo
GPGPU (General Purpose Computing on GPU)
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; }
CPU Serial Code GPU Parallel Kernel KernelB<<< nBlk, nTid >>>(args);
Grid 0
...
Grid 1 ...
Lenovo Confidential | © 2010 Lenovo
CUDA极大提高了现有应用的效果
MRI Reconstruction
Based on Fig 1 of Lustig et al, Fast Spiral Fourier Transform for Iterative MR Image Reconstruction, IEEE Int’l Symp. on Biom edical Imaging, 2004
Lenovo Confidential | © 2010 Lenovo
Lenovo Confidential | © 2010 Lenovo
Advanced MRI Reconstruction
( F F W W ) F d
H H H
Compute Q More than Acquire Data 99.5% of time Compute FHd
• Q只和扫描参数有关
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; } }
Advanced MRI Reconstruction
Cartesian Scan Data Spiral Scan Data Gridding
ky kx
ky
(b)
ky
kx
(a)
(b)
kx
(c)
FFT
Iterative Reconstruction
Spiral scan data + Iterative recon Reconstruction requires a lot of computation
CUDA程序设计
Lenovo Confidential | © 2010 Lenovo
主要内容
• GPGPU及CUDA介绍 • CUDA编程模型 • 多线程及存储器硬件
Lenovo Confidential | © 2010 Lenovo
GPGPU及CUDA介绍
Lenovo Confidential | © 2010 Lenovo
• FHd是数据相关的
• 使用线性求解器计算ρ
Find ρ
Haldar, et al, “Anatomically-constrained reconstruction from noisy data,” MR in Medicine.
Lenovo Confidential | © 2010 Lenovo