CUDA多核程序设计

•专用存储器通道 •有效隐藏存储器延时
General Purpose Computing on GPU (GPGPU)
GPGPU
核心思想
用图形语言描述通用计 算问题
把数据映射到vertex或 者fragment处理器
但是
硬件资源使用不充分 存储器访问方式严重受
限 难以调试和查错 高度图形处理和编程技
(4, 2)
CUDA存储器分配
cudaMalloc()
在全局存储器中分配空间
两个参数：
地址指针 空间大小
cudaFree()
在全局存储器中回收空间
参数
Host
回收空间地址指针
Grid Block (0, 0)
Block (1, 0)
Shared Memory Registers Registers
Texture
Vertex Processor
Rasterizer
Fragment Processor
Framebuffer
传统GPU架构
Graphics program Vertex processors
Fragment processors
Pixel operations Output image
Compute 1.3: double precision
GT200
CUDA成功案例
提纲
从GPGPU到CUDA 并行程序组织 并行执行模型 CUDA基础 存储器 CUDA程序设计工具 新一代Fermi GPU
并行性的维度
1维
y = a + b //y, a, b vectors
a[0] a[1] … a[n]
CUDA: 集成CPU + GPU C应用程序 通用并行计算模型
单指令、多数据执行模式 (SIMD)
所有线程执行同一段代码(1000s threads on the fly) 大量并行计算资源处理不同数据 隐藏存储器延时 提升计算／通信比例 合并相邻地址的内存访问 快速线程切换1 cycle@GPU vs. ~1000 cycles@CPU
(0, 0)
(1, 0)
(2, 0)
(3, 0)
(4, 0)
Thread Thread Thread Thread Thread
(0, 1)
(1, 1)
(2, 1)
(3, 1)
(4, 1)
Thread Thread Thread Thread Thread
(0, 2)
(1, 2)
(2, 2)
(3, 2)
CUDA多核程序设计
来自百度文库陈华
提纲
从GPGPU到CUDA 并行程序组织 并行执行模型 CUDA基础 存储器 CUDA程序设计工具 CUDA4.2
Graphic Processing Unit (GPU)
用于个人计算机、工作站和游戏机的专用图像显示设 备
显示卡 nVidia和ATI (now AMD)是主要制造商
– Intel准备通过Larrabee进入这一市场
主板集成 Intel
3维图像流水线
一帧典型图像 30 frames/s
1M triangles
30M triangles/s
3M vertices
90M vertices/s
25M fragments
750M fragments/s
CPU GPU
Compute 1.2: dramatically improved memory coalescing rules, double the register count, intra-warp voting primitives, atomic shared memory operations
GT21X
Shared Memory Registers Registers
Thread (0, 0) Thread (1, 0) Thread (0, 0) Thread (1, 0)
Global Memory
主机和设备之间的数据传输
巧
G80 GPU
Thread Processor
Host Input Assembler Vtx Thread Issue
Geom Thread Issue
Setup / Rstr / ZCull Pixel Thread Issue
SP
SP
SP
SP
SP
SP
SP
SP
SP
SP
SP
SP
SP
SP
SP
SP
若干线程还可以组成块 (Block)。
线程块可以呈一维、二维 或者三维结构。
Device
Grid 1
Block (0, 0)
Block (0, 1)
Block (1, 0)
Block (1, 1)
Block (2, 0)
Block (2, 1)
Block (1, 1)
Thread Thread Thread Thread Thread
TF
TF
L1
L1
TF
TF
L1
L1
TF
TF
L1
L1
TF
TF
L1
L1
L2 FB
L2 FB
L2 FB
L2 FB
Streaming Processor (SP)
Streaming Multiprocessor (SM)
L2 FB
L2 FB
CUDA: Compute Unified Device
Architecture
++
+
b[0] b[1] … b[n]
==
=
y[0] y[1] … y[n]
2维
P = M N //P, M, N matrices
=
3维
CT or MRI imaging
CUDA中的线程
线程：
CUDA中的基本执行单元； 硬件支持，开销很小；
所有线程执行相同的代码 （STMD）
每个线程有一个唯一的标 识ID——threadIdx。
Evolution of CUDA-Enabled GPUs
Compute 1.0: basic CUDA compatibility
G80
Compute 1.1: asynchronous memory copies and atomic global operations
G84, G86, G92, G94, G96, and G98
GPU的强大运算能力
•数据级并行: 计算一致性
120 GPU
100
CPU
80
60
G80 Ultra G80 G71
Memory bandwidth (GB/s)
40
NV40
20
NV30
Northwood
0
2003
2004
Prescott EE
Hapertown Woodcrest
2005
2006
2007