CUDA简介——CUDA内存模式

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的通信吞吐量。