CPU和GPU互动模式
GPU线程组织模型(不停强化)
GPU存储模型
基本的编程问题
CPU和GPU互动模式
GPU存储器层次架构 硬件
基于OpenCL的图
分层次 主要是为了容量-速度的trade off
访存速度
Global Memory 和 Local Memory 实际上已经有cache了
Constant Memory在DRAM显存中 Texture Memory在DRAM显存中,对杂乱的小数据的临近性好一些。 Instruction Memory(invisible) DRAM中
GPU线程组织模型
Grid <- Block <- Thread
In fact:
线程组织架构说明:
一个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内存和线程的关系
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
编程模型: 如图像处理中像素的操作类似于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、被任意数量的处理器以任意顺序调度 处理器的数量具有可扩展性
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
1000 * 1000的矩阵的乘法 为1000,000次点乘,每次1000个乘法、1000个加法