CUDA 内存模型

存储器的类型有两种：

1. 可编程：显式控制哪些数据存放
2. 不可编程：不能决定数据存储位置

在CPU层次结构中，一级缓存（L1 Cache）和二级缓存（L2 Cache）都是不可编程的。可编程的存储器类型包括：

1. 寄存器
2. 共享内存
3. 本地内存
4. 常量内存
5. 纹理内存
6. 全局内存

下图为上面提到的存储器类型：

根据这张图我们可以分析：

• 一个Kernel中的所有Threads都有私有的本地内存。
• 一个Block有自己的共享内存，对本Block内的所有Thread可见，且其内容持续Block的整个生命周期。
• 所有Thread都可以访问全局内存。
• 所有Thread也都可以访问只读的常量内存和纹理内存

对一个应用程序来说，全局内存、常量内存、纹理内存中的内容有相同的生命周期。

寄存器

是GPU上运行速度最快的内存空间。寄存器对于每个 Thread 来说是私有的，一个 Kernel 使用寄存器来保存需要频繁访问的 Thread 私有变量。寄存器变量和 Kernel 的生命周期相同。

如果 Kernel 使用超过硬件限制数量的寄存器，则会用本地内存替代多占用的寄存器，但是会降低性能。nvcc 编译器使用启发式策略（指编译器根据一套预设的规则和经验来决定如何编译代码，这些规则和经验旨在提高代码的运行效率，但并不总是能够达到全局最优解^[chatgpt]^）

• 占用计算器 API，cudaOccupancyMaxActiveBlocksPerMultiprocessor，可以根据Kernel的Block大小和共享内存使用情况提供占用预测。此函数根据每个多处理器的并发线程块数报告占用情况。
• • 请注意，此值可以转换为其他指标。乘以每个块的 Warp数得出每个多处理器的并发 Warp数；进一步将并发 Warp除以每个多处理器的最大 Warp得到占用率作为百分比。
• 基于占用率的启动配置器 API，cudaOccupancyMaxPotentialBlockSize 和 cudaOccupancyMaxPotentialBlockSizeVariableSMem ，启发式地计算实现最大多处理器级占用率的执行配置。^[1]^
• 代码中为 Kernel 显式地加入额外信息：

__global__ void
__launch_bounds__(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP)
fooKernel(int *inArr, int *outArr)
{
    // ... Computation of kernel
}

MAX_THREADS_PER_BLOCK 指定了每个 Block 的最大 Thread 数；MIN_BLOCKS_PER_MP 指出每个SM中预期的最小常驻 Block 数。

本地内存

本地内存的特点是高延迟、低带宽。

Kernel 中能够被存储在寄存器，但是不能进入该 Kernel 分配的寄存器空间的变量将会被溢出到本地内存中。这种变量包括：

1. 在编译时 使用位置索引引用的本地数组
2. 可能占用大量寄存器空间的较大本地结构体和数组
3. 其他不满足 Kernel 寄存器条件的变量

要注意：溢出到本地内存的变量，本质上和全局内存在同一块存储区域。

共享内存

Kernel 中使用 __shared__ 修饰的变量存放在共享内存中。共享内存相比全局内存和本地内存，有更高的带宽和更低的延迟。

每个SM都有一定数量由Block分配的共享内存，因此不能过度使用共享内存，否则会限制活跃Warp的数量。

共享内存的生命周期伴随着整个 Block 。当一个 Block 的执行结束后， 其分配的共享内存将被释放，并重新分配给其他 Block

线程间通信的主要方式就是共享内存，访问共享内存必须使用 void __syncthreads()以避免潜在的数据冲突。

SM中的 L1 Cache 和共享内存都使用片上内存，每个SM上有64KB的片上内存，默认通过静态划分，但在运行的时候也可以使用 cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig); 进行动态配置，其中 cacheConfig 支持如下缓存配置：

cudaFuncCachePreferNone: no preference(default)
cudaFuncCachePreferShared: prefer 48KB shared memory and 16 KB L1 cache
cudaFuncCachePreferL1: prefer 48KB L1 cache and 16 KB shared memory
cudaFuncCachePreferEqual: prefer 32KB L1 cache and 32 KB shared memory

常量内存

常量内存使用 __constant__进行修饰，保存在设备内存中，并在每个SM专用的常量Cache中缓存。

常量变量必须在全局空间内和Kernel函数外声明，而且只能声明64KB的常量内存。此外，常量内存是静态声明的。

Kernel 只能从常量内存读取数据，因此常量内存必须在 HOST 初始化： cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src,size_t count) 整个函数将 count 个字节，从 src复制到 symbol。

纹理内存

纹理内存驻留在设备内存中，并在每个SM的只读Cache中缓存。纹理内存是一种通过指定的只读Cache进行访问的全局内存。

此外，在Warp中使用纹理内存访问二维数据的Thread可以获得最优的性能。

全局内存

全局内存是GPU中最大，延迟最高，使用最多的内存类型。

要注意多个线程访问全局内存时，因为Thread的执行不能跨Block同步，不同Block的多个Thread并发修改同一位置的全局内存时，可能会导致未定义的行为。

可以通过 32/64/128 Byte 的内存事务访问全局内存，这些内存事务必须自然对其，也就是首地址必须是32/64/128 Byte 的整数倍。当一个Warp执行内存加载和存储时，需满足的传输数量取决于：

• 跨线程的内存地址分布
• 每个事务内存地址的对齐方式

GPU缓存

和CPU缓存一样都是不可编程的。但是CPU内存的加载和存储都可以被缓存；GPU上只有内存加载操作可以被缓存。

GPU上有4种缓存：

• 1级缓存：每个SM都有一个1级缓存，用来存储本地内存&全局内存的数据，包括溢出部分。
• 2级缓存：所有SM共享一个2级缓存，存储内容同1级缓存。
• 只读常量缓存：
• 只读纹理缓存：

静态全局内存

CPU内存有动态分配和静态分配两种类型，从内存位置来说，动态分配在堆上进行，静态分配在栈上进行，在代码上的表现是一个需要new，malloc等类似的函数动态分配空间，并用delete和free来释放。在CUDA中也有类似的动态静态之分，我们前面用的都是要cudaMalloc的，所以对比来说就是动态分配。与动态分配相同是，静态分配也需要显式的将内存copy到设备端，我们用下面代码来看一下程序的运行结果:

#include <cuda_runtime.h>
#include <stdio.h>
__device__ float devData;
__global__ void checkGlobalVariable(){
    printf("Device: The value of the global variable is %f\n",devData);
    devData+=2.0;
}
int main(){
    float value=3.14f;
    cudaMemcpyToSymbol(devData,&value,sizeof(float));
    printf("Host: copy %f to the global variable\n",value);
    checkGlobalVariable<<<1,1>>>();
    cudaMemcpyFromSymbol(&value,devData,sizeof(float));
    printf("Host: the value changed by the kernel to %f \n",value);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}

其中的 cudaMemcpyToSymbol(devData,&value,sizeof(float)); ，是由于__device__ float devData;是DEVICE上的变量定义，和HOST变量定义不同。在主函数，全局变量的值通过cudaMemcpyToSymbol初始化，在kernel中运行完后，新的值通过cudaMemcpyFromSymbol复制回HOST。

虽然上面DEVICE和HOST的代码在同一个文件中，也在同一个文件中可见，但是HOST代码不能直接访问DEVICE变量，反之亦然。

总结

参考文章

[1] CUDA 编程手册系列第五章: 性能指南
[2] 【CUDA 基础】4.1 内存模型概述 | 谭升的博客 (face2ai.com)