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