1. 引言
前序博客:
- CUDA简介——基本概念
- CUDA简介——编程模式
- CUDA简介——For循环并行化
- CUDA简介——Grid和Block内Thread索引
CUDA内存模式,采用分层设计,是CUDA程序与正常C程序的最大不同之处:
- Thread-Memory Correspondence
- Block-Memory Correspondence
- Grid-Memory Correspondence
总体的CUDA内存模式为:
- 1)Registers & Local Memory:声明在Kernel内的常规变量。小空间变量存储于Register中,大空间变量存储于Local Memory中。因Local Memory操作速度慢,应尽量避免使用Local Memory。
- 2)Shared Memory:允许同一Block内的Threads间相互通信。
- 3)Constant Memory:用于存储Kernel执行期间不变的数据。会缓存到On-Chip memory中,可大大降低Kernel执行期间Global Memory的通信吞吐量。
- 4)Global Memory:用于存储与Host交互的数据。
2. Thread-Memory Correspondence
Thread-Memory Correspondence,即Threads等价为Local Memory (and Registers):
- 其范围为:对相应Thread是private的,对其它Threads来说是不可访问的。
- 其生命周期为:Thread。当Thread执行完成,与该Thread相关的任何local memory (and Registers) 都将自动释放。
- 不过,local memory和registers存在完全不同的性能特性。
3. Block-Memory Correspondence
Block-Memory Correspondence,即Blocks等价为Shared Memory:
- 其范围为:同一Block内的每个Thread均可访问。
- 其生命周期为:Block。当Block执行完成,其shared memory内容都将自动释放。
4. Grid-Memory Correspondence
Grid-Memory Correspondence,即Grids等价为Global Memory:
- 其范围为:所有 Grids内的每个Thread均可访问。
- 其生命周期为:Host代码内的整个main()程序。或可通过在Host代码中手工调用
cudaFree(...)
来释放。
5. Device内存模式
- Host:由CPU及机器内存等组成。
- Device:由GPU及其DRAM等组成。【Device图上的绿色方格,表示的是一个CUDA core。】
Device的DRAM中,有:
- Global Memory物理空间
- Local Memory物理空间。需注意,此处的Local,并不是指其物理位置;此处的Local是指该内存空间的scope(范围)和 lifetime(生命周期)。
而Device的GPU中,有:
- Registers物理空间
- Shared Memory物理空间
Device图上的绿色方格,表示的是一个CUDA core。Device上的CUDA cores组合在一起,成为streaming Multiprocessor(简称为SM)。
Device图上的黄色方格,表示SM。黄色方格组合在一起为CUDA cores集合。
位于SM上的内存,称为:
- “On-Chip” Device memory。因此,Registers和Shared Memory均对应为 “On-Chip” Device memory。因Registers和Shared Memory均物理存在与GPU的streaming Multiprocessor中。
非SM上的内存,称为::
- “Off-Chip” Device memory。因其并不存在与GPU之上。对应,Global Memory和Local Memory均为“Off-Chip” Device memory。也即Device上的DRAM为 “Off-Chip” Device memory。
以NVIDIA GPU Geforce Titan 物理布局为例:
- 上图蓝色框所示,为Device的DRAM,均为 “Off-Chip” Device memory。
- 上图绿色框,即为实际的GPU,对应为“On-Chip” Device memory。
理解Blocks如何映射到SM,是设计kernel的基本要求,从而获得优化的GPU计算性能。
5.1 内存速度
不同的内存空间,其带宽和延迟各不相同:
- on-chip memory操作速度,要快于off-chip memory。
5.2 Global Memory访问
Global Memory访问方式有:
- cudaMalloc()
- cudaMemset()
- cudaMemCopy()
- cudaFree()
无法避免不使用Global Memory,因必须使用Global Memory空间,来将数据由host传递到device。不过,应尽量减少Global Memory通讯量,因为其速度很慢。
Global Memory的优势在于:
- 其通常很大。如计算机内存为8GB或16GB,而Titan和Tesla k40,均有6GB的global DRAM。
5.3 Registers and Local Memory
Kernel中声明的变量,存储于Register中:
- 对应为On-Chip Device memory。
- 为最快的内存形式。
太大不适于Register的数组,将存储在Local Memory中:
- 对应为Off-Chip Device memory。
- 由编译器控制。
- “Local”是指范围,而不是位置。此处“Local”,是指相对每个Thread的Local Memory。
- 每个Thread均有其自己的private local memory和registers,对其它Thread不可访问。
- 应尽量避免,因Local Memory为最慢的内存形式之一。Registers为最快的内存形式。
- 因此,设计目标为避免将更多信息存储于local变量中,以免超过register的存储空间。
- register space为稀缺硬件资源。应更好地规划充分利用。
5.4 Shared Memory
借助Shared Memory:
- 支持同一Block内的Threads之间相互通信:
- 同步方式通信
- Shared Memory为非常特殊的内存,对实现计算性能和正确性至关重要:
- Shared Memory处理速度快,仅次于Registers。因其为On-Chip device memory。
- 支持Block内的Threads相互通信,可将其看成是用户定义的L1 Cache,可用作“scratch-pad(高速暂存存储器)”内存。后续将介绍Shared Memory和L1 Cache关系密切。
使用__shared___
关键字来表示分配的为shared memory:
5.5 Constant Memory
Constant Memory为Device Memory的特殊区域:
- 用于存储Kernel执行过程中不变的数据。
- 对Kernel来说是只读的。
- Constant Memory为Off-Chip Device Memory。
- Constant Memory 积极缓存到 On-Chip Memory中。
Constant Memory的思想在于:
- GPU没有很大的cache。从而可使用constant memory来实现很简单的cache类型。
- Constant Memory可以很大,因其实际位于Device DRAM中。所有Threads均可访问Constant Memory,但其为只读内存。
- 对于需频繁访问,但Kernel执行过程中不变的数据,可使用Constant Memory。
- Constant Memory是在off-chain DRAM硬件中实现的,但实际其内容会积极缓存到On-Chip Memory中。因此,使用Constant Memory,可大大降低Kernel执行期间Global Memory的通信吞吐量。
参考资料
[1] Intro to CUDA (part 5): Memory Model