< CPU >
// ----------------------------------------------------------------------------------
// 内存
// ----------------------------------------------------------------------------------
||
||
< GPU > ||
// ----------------------------------------------------------------------------------
// 全局内存
// ----------------------------------------------------------------------------------
||
// ----------------------------------------------------------------------------------
// 纹理内存
// ----------------------------------------------------------------------------------
||
// ----------------------------------------------------------------------------------
// 常量内存
// ----------------------------------------------------------------------------------
|| ||
// --------------------------------------- // -------------------------------
// 共享内存 [block0] // 共享内存 [block1]
// --------------------------------------- // -------------------------------
|| ||
// ----------------------- // -----------------------
// 局部内存 [thread00] // 局部内存 [thread01] ......
// 寄存器 // 寄存器
// ----------------------- // -----------------------
CUDA 中的内存类型有:全局内存、常量内存、纹理内存、寄存器、局部内存、共享内存。 CUDA 的内存,即设备内存,主机无法直接访问。
全局内存(global memory),即核函数中所有线程都可以访问的内存,可读可写,由主机端分配和释放;
如 cudaMalloc() 的设备内存 d_x, d_y, d_z。
全局内存由于没有放到 GPU 芯片上,所以具有较高的延迟和较低的访问速度,但是容量大(显存)。
全局内存主要为核函数提供数据,并在主机和设备、设备和设备之间传递数据。
全局内存的生命周期由主机端维护,期间不同的核函数可以多次访问全局内存。
除以上动态分配的全局内存变量外,还可以使用 静态全局内存变量,其所占内存数量在编译器确定;
这样的静态全局内存变量必须在 所有主机和设备函数外部定义,例如:
__device__ real epsilon; // 单个静态全局内存变量, `__device` 表示是设备中的变量。
__device__ real arr[10]; // 固定长度的静态全局内存数组变量。
对于静态全局内存变量,其访问权限:
- 核函数中可以直接访问静态全局内存变量,不必以参数形式传给核函数;
- 主机中不可以直接访问静态全局内存变量,可以通过
cudaMemcpyToSymbol()
和cudaMemcpyFromSymbol()
调用。
常量内存(constant memory),仅有 64 kb,可见范围和生命周期与全局内存一样;具有缓存,从而高速; 常量内存仅可读、不可写。
使用常量内存的方法:一是在核函数外定义常量内存变量;二是向核函数传递常量参数,默认存放在常量内存:
- 核函数中可以直接访问常量全局内存变量,不必以参数形式传给核函数,但不可更改(只读);
- 主机中不可以直接访问常量全局内存变量,可以通过
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)与寄存器类似,都是位于芯片上,读写速度较快。
共享内存对整个线程块可见,一个线程块上的所有线程都可以访问共享内存上的数据;共享内存的生命
周期也与所属线程块一致。
共享内存的主要作用是减少对全局内存的访问,或者改善对全局内存的访问模式。
SM 层次的 L1 缓存(一级缓存)和设备层次 L2 缓存(二级缓存)。它们主要用来缓存全局内存和设备内存的访问。
一个 GPU 由多个 SM(流多处理器)构成,一个 SM 包含如下资源:
- 一定数量的寄存器;
- 一定数量的共享内存;
- 常量内存的缓存;
- 纹理内存的缓存;
- L1 缓存;
- 两个或四个线程束调度器,用于在不同线程上下文间迅速切换,及为准备就绪的线程束发出执行指令;
- 执行核心。
一般来说,要尽量让 SM 的占有率不小于某值(如 25%),才有可能获得较高的性能。
- 一个 SM 中最多拥有的线程块个数 Nb=16(开普勒和图灵架构)或 Nb=32(麦克斯韦、帕斯卡和伏特架构);
- 一个 SM 中最多拥有的线程格式为 Nt=1028(图灵架构)或 Nt=2048(开普勒到伏特架构)。
在线程块中,每 32 个连续线程为一个 线程束。 SM 中线程的执行是以线程束为单位的,所以最好将线程块大小取为线程束大小(32个线程)的整数倍(如 128).
使用 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