FunnyWii
FunnyWii
发布于 2024-07-15 / 147 阅读
0
0

我还在学CUDA编程(四)——内存的层次结构

CUDA 内存模型

存储器的类型有两种:

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

在CPU层次结构中,一级缓存(L1 Cache)和二级缓存(L2 Cache)都是不可编程的。可编程的存储器类型包括:

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

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

CUDA存储器类型.jpg

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

  • 一个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,cudaOccupancyMaxPotentialBlockSizecudaOccupancyMaxPotentialBlockSizeVariableSMem ,启发式地计算实现最大多处理器级占用率的执行配置。^[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变量,反之亦然。

总结

Table1 CUDA变量和类型修饰符
修饰符 变量名称 存储器 作用域 生命周期
float var 寄存器 线程 线程
float var[100] 本地 线程 线程
share float var* 共享
device float var* 全局 全局 应用程序
constant float var* 常量 全局 应用程序
Table2 设备存储器特征
存储器 片上/片外 缓存 存取 范围 生命周期
寄存器 片上 n/a R/W 一个线程 线程
本地 片外 2.x以上设备 R/W 一个线程 线程
共享 片上 n/a R/W 块内所有线程
全局 片外 2.x以上设备 R/W 所有线程+主机 主机配置
常量 片外 Yes R 所有线程+主机 主机配置
纹理 片外 Yes R 所有线程+主机 主机配置

参考文章

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


评论