168 lines
7.1 KiB
Markdown
168 lines
7.1 KiB
Markdown
# CUDA 的内存组织
|
||
|
||
< CPU >
|
||
// ----------------------------------------------------------------------------------
|
||
// 内存
|
||
// ----------------------------------------------------------------------------------
|
||
||
|
||
||
|
||
< GPU > ||
|
||
// ----------------------------------------------------------------------------------
|
||
// 全局内存
|
||
// ----------------------------------------------------------------------------------
|
||
||
|
||
// ----------------------------------------------------------------------------------
|
||
// 纹理内存
|
||
// ----------------------------------------------------------------------------------
|
||
||
|
||
// ----------------------------------------------------------------------------------
|
||
// 常量内存
|
||
// ----------------------------------------------------------------------------------
|
||
|| ||
|
||
// --------------------------------------- // -------------------------------
|
||
// 共享内存 [block0] // 共享内存 [block1]
|
||
// --------------------------------------- // -------------------------------
|
||
|| ||
|
||
// ----------------------- // -----------------------
|
||
// 局部内存 [thread00] // 局部内存 [thread01] ......
|
||
// 寄存器 // 寄存器
|
||
// ----------------------- // -----------------------
|
||
|
||
------
|
||
|
||
## CUDA 中不同类型的内存
|
||
|
||
CUDA 中的内存类型有：全局内存、常量内存、纹理内存、寄存器、局部内存、共享内存。
|
||
CUDA 的内存，即设备内存，主机无法直接访问。
|
||
|
||
------
|
||
|
||
### 全局内存
|
||
|
||
**全局内存（global memory）**，即核函数中所有线程都可以访问的内存，可读可写，由主机端分配和释放；
|
||
如 cudaMalloc() 的设备内存 d_x, d_y, d_z。
|
||
|
||
全局内存由于没有放到 GPU 芯片上，所以具有较高的延迟和较低的访问速度，但是容量大（显存）。
|
||
全局内存主要为核函数提供数据，并在主机和设备、设备和设备之间传递数据。
|
||
|
||
全局内存的生命周期由主机端维护，期间不同的核函数可以多次访问全局内存。
|
||
|
||
除以上动态分配的全局内存变量外，还可以使用 **静态全局内存变量**，其所占内存数量在编译器确定；
|
||
这样的静态全局内存变量必须在 所有主机和设备函数外部定义，例如：
|
||
|
||
```cuda
|
||
__device__ real epsilon; // 单个静态全局内存变量， `__device` 表示是设备中的变量。
|
||
__device__ real arr[10]; // 固定长度的静态全局内存数组变量。
|
||
```
|
||
|
||
对于静态全局内存变量，其访问权限：
|
||
|
||
1. 核函数中可以直接访问静态全局内存变量，不必以参数形式传给核函数；
|
||
2. 主机中不可以直接访问静态全局内存变量，可以通过 `cudaMemcpyToSymbol()` 和 `cudaMemcpyFromSymbol()` 调用。
|
||
|
||
------
|
||
|
||
### 常量内存
|
||
|
||
**常量内存（constant memory）**，仅有 64 kb，可见范围和生命周期与全局内存一样；具有缓存，从而高速；
|
||
常量内存仅可读、不可写。
|
||
|
||
使用常量内存的方法：一是在核函数外定义常量内存变量；二是向核函数传递常量参数，默认存放在常量内存：
|
||
|
||
1. 核函数中可以直接访问常量全局内存变量，不必以参数形式传给核函数，但不可更改（只读）；
|
||
2. 主机中不可以直接访问常量全局内存变量，可以通过 `cudaMemcpyToSymbol()` 和 `cudaMemcpyFromSymbol()` 调用。
|
||
|
||
------
|
||
|
||
### 纹理内存
|
||
|
||
**纹理内存(texture memory)**，类似常量内存，也是一种具有缓存的全局内存，具有相同可见范围和生命周期。
|
||
|
||
可以将某些只读的全局内存数据用 `__ldg()` 函数通过只读数据缓存（read-only data cache）读取，
|
||
既可以达到使用纹理内存的加速效果，又可使代码简洁：
|
||
|
||
int __ldg(const int* ptr); // 函数原型。
|
||
|
||
全局内存的读取在默认情况下就利用了 `__ldg()` 函数，所以不需要显式地使用。
|
||
|
||
------
|
||
|
||
### 寄存器
|
||
|
||
在核函数中定义的、不加任何限定符的变量一般存放在寄存器（register）；核函数中不加任何限定符的数组可能放在
|
||
寄存器，也可能放在局部内存中。寄存器可读可写。
|
||
|
||
各种内建变量，如 gridDim、blockDim 等都保存在特殊的寄存器中。
|
||
|
||
寄存器变量仅被一个线程看见，寄存器的生命周期也和所属线程相同。
|
||
|
||
寄存器内存在芯片上，是所有内存中访问速度最高的。一个寄存器占 32b（4字节），一个双精度浮点数占 2个寄存器。
|
||
|
||
------
|
||
|
||
### 局部内存
|
||
|
||
局部内存(local memory)也是全局内存的一部分，每个线程最多可以使用 512 kb 的局部内存，但过多使用会降低性能。
|
||
局部内存的用法类似寄存器。
|
||
|
||
------
|
||
|
||
### 共享内存
|
||
|
||
共享内存（shared memory）与寄存器类似，都是位于芯片上，读写速度较快。
|
||
|
||
共享内存对整个线程块可见，一个线程块上的所有线程都可以访问共享内存上的数据；共享内存的生命
|
||
周期也与所属线程块一致。
|
||
|
||
共享内存的主要作用是减少对全局内存的访问，或者改善对全局内存的访问模式。
|
||
|
||
------
|
||
|
||
### L1 和 L2 缓存
|
||
|
||
SM 层次的 L1 缓存（一级缓存）和设备层次 L2 缓存（二级缓存）。它们主要用来缓存全局内存和设备内存的访问。
|
||
|
||
------
|
||
|
||
## SM 及其占有率
|
||
|
||
一个 GPU 由多个 SM（流多处理器）构成，一个 SM 包含如下资源：
|
||
|
||
1. 一定数量的寄存器；
|
||
2. 一定数量的共享内存；
|
||
3. 常量内存的缓存；
|
||
4. 纹理内存的缓存；
|
||
5. L1 缓存；
|
||
6. 两个或四个线程束调度器，用于在不同线程上下文间迅速切换，及为准备就绪的线程束发出执行指令；
|
||
7. 执行核心。
|
||
|
||
一般来说，要尽量让 SM 的占有率不小于某值（如 25%），才有可能获得较高的性能。
|
||
|
||
+ 一个 SM 中最多拥有的线程块个数 Nb=16（开普勒和图灵架构）或 Nb=32（麦克斯韦、帕斯卡和伏特架构）；
|
||
+ 一个 SM 中最多拥有的线程格式为 Nt=1028（图灵架构）或 Nt=2048（开普勒到伏特架构）。
|
||
|
||
在线程块中，每 32 个连续线程为一个 **线程束**。
|
||
SM 中线程的执行是以线程束为单位的，所以最好将线程块大小取为线程束大小（32个线程）的整数倍（如 128）.
|
||
|
||
------
|
||
|
||
## CUDA 运行时 API 函数查询设备
|
||
|
||
使用 CUDA 运行时 API 函数查询所用GPU 规格。
|
||
|
||
Device id: 0
|
||
Device name: GeForce MX450
|
||
Compute capability: 7.5
|
||
Amount of global memory: 2 GB
|
||
Amount of constant memory: 64 KB
|
||
Maximum grid size: 2147483647, 65535, 65535
|
||
Maximum block size: 1024, 1024, 64
|
||
Number of SMs: 14
|
||
Maximum amount of shared memory per block: 48 KB
|
||
Maximum amount of shared memory per SM: 64 KB
|
||
Maximum number of registers per block: 64 K
|
||
Maximum number of registers per SM: 64 K
|
||
Maximum number of threads per block: 1024
|
||
Maximum number of threads per SM: 1024
|
||
------
|