CUDA GPU编程模型

26 阅读2分钟

CPU和GPU互动模式

GPU线程组织模型(不停强化)

GPU存储模型

基本的编程问题

CPU和GPU互动模式

tmp2789.png

GPU存储器层次架构 硬件 基于OpenCL的图 tmpBDB6.png

分层次 主要是为了容量-速度的trade off 访存速度 tmp50BB.png

Global Memory 和 Local Memory 实际上已经有cache了

Constant Memory在DRAM显存中 Texture Memory在DRAM显存中,对杂乱的小数据的临近性好一些。 Instruction Memory(invisible) DRAM中

GPU线程组织模型

Grid <- Block <- Thread

In fact:

tmpB939.png

线程组织架构说明:

一个kernnel启动一个grid,包含若干线程块。

在GPU中: Thread 对应 一个 线程处理器(ALU 所谓的小核) Thread Blocks对应一个 Multi-processor(SM 大核) Thread Grid对应一个Device 在CPU中: Thread对应一个scalar SSE Thread Blocks对应一个 Vector SSE Thread Grid 对应一个 Multi-core

GPU内存和线程的关系

tmp35BD.png

Local Memory -> Thread 独有 Shared Memory -> Block 独有 Global Memory -> grid拥有

设备GPU0和设备GPU1都可以和主机端存储器交互

GPU内存组织结构

瓶颈:访存

线程:私有的寄存器和Local Memory 读写均可 Block:Shared Memory 对Block内部的各个线程开放读写 Grid:Global Memroy(可读可写) Constant Memory(对GPU只读,对Host可读可写) Texture Memory(?)

Host通过PCIE总线R/W Global Memory和Constant Memory

tmp36C9.png

编程模型: 如图像处理中像素的操作类似于SIMD

SIMT:GPU版的SIMD 大量线程模型获得高度并行 线程切换获得延迟掩藏 多个线程执行相同的指令流 GPU上大量线程承载和调度

CUDA执行过程

CPU serial code -> GPU parallel kernel -> CPU serial code -> GPU parallel kernel

线程层次:

GRID(1维或2维的BLOCKs组成) -> BLOCK(1维、2维或3维) -> THREAD

BLOCKs内部的线程可以同步(synchronize)、访问共享存储器(shared memory)

线程块之间彼此独立执行: 1、任意顺序:并行或串行 2、被任意数量的处理器以任意顺序调度 处理器的数量具有可扩展性

tmp4DC4.png

8个block,在含有两个SM的Device上,经过4次调度完成。

一个块内部的线程: 共享容量有限的低延迟存储器 同步执行 合并访存 __syncThreads() -> Barrier:块内线程一起等待所有线程执行到某处语句 -> 轻量级

Host操作内存的具体接口: cudaMalloc() 在设备端申请内存 cudaFree() 在设备端释放内存 cudaMemcpy() Host 2 Host Device 2 Device Host 2 Device Device 2 Host

Matrix Multiply 算法提示

假定 A = B * C

tmp177.png

1000 * 1000的矩阵的乘法 为1000,000次点乘,每次1000个乘法、1000个加法