《CUDA编程》6.CUDA的内存组织

前面几章讲了一些编写高性能CUDA程序的要点，但还有很多其他需要注意的，其中最重要的就是合理的使用设备内存

1 CUDA的内存组织简介

现代计算机中的内存存在一种组织结构(hierachy)，即不同类型的内存具有不同的容量和访问延迟（可以理解为处理器等待内存的时间）。一般来说，延迟低的内存容量小，延迟高的内存容量大。

下表是CUDA设备（显卡）中的几种内存和主要特征

下图是组织示意图和数据移动发方向图

2 不同内存的简介

2.1 全局内存

全局内存（global memory）： 核函数中的所有线程都能够访问。

特点：

1. 容量大： 通常是GPU上最大的内存区域。
2. 访问延迟高： 相比于其他内存，访问延迟较高。
3. 全局可见： 对所有线程和线程块可见，可在核函数中自由读取、可读可写。
4. 持久性： 全局内存中的数据在内核函数执行期间保持不变，可以跨多个内核调用使用。

全局内存主要是为核函数提供数据，并在主机和设备、设备和设备之间传输数据

全局内存的生命周期由主机端决定，所以：

• cudaMalloc()函数是主机在设备的全局内存中分配一段指定大小的内存区域
• cudaFree()函数是主机把该内存释放

以上所说的全局内存称为线性内存（linear memory），还有一种不对用户透明的内存称为CUDA Array，专为纹理拾取服务。

和C++函数一样，cudaMalloc()函数是动态地分配内存，CUDA中也允许使用静态全局内存变量，定义方法如下：

• __device__ T x; 定义单个变量
• __device__ T y[N]; 定义固定长度地数组

其中，修饰符 __ device __ 说明该变量是设备中的变量，而不是主机中的变量；T 是变量的 类型；

在核函数中，可以直接访问静态全局内存变量，不需要以参数的形式传入；注意，主机函数无法直接访问，只能用cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()在主机内存和静态全局内存之间传输数据，下面是两个函数的结构：

①将主机数据复制到静态全局内存中

②将静态全局内存数据复制到主机中

之后会讨论一种利用静态全局内存加速程序的技巧，现在给出使用例子：

#include <cuda.h>
#include <cuda_runtime.h>
#include "error_check.cuh"

__device__ int d_x = 1;
__device__ int d_y[2];

__global__ void cudaOut(void) {
	d_y[0]+= d_x;
	d_y[1]+= d_x;
	printf("Device:  d_x = %d,d_y[0]=%d, d_y[1]=%d\n",d_x,d_y[0], d_y[1]);
	printf("\n");
}

int main(void) {
	int h_y[2] = { 10,20 };
	CHECK(cudaMemcpyToSymbol(d_y, h_y, sizeof(int) * 2));

	cudaOut<<<1,1>>>();
	CHECK(cudaDeviceSynchronize());

	CHECK(cudaMemcpyFromSymbol(h_y, d_y, sizeof(int) * 2));
	printf("Host:  h_y[0]=%d, h_y[1]=%d\n", h_y[0], h_y[1]);

	return 0;
}

输出结果是：

可以看到主机数据和静态全局内存变量中的数据交流成功。

2.2 常量内存

常量内存（constant memory）： 是有常量缓存的全局内存，数量有限，最多有 64 KB。

特点：

1. 容量小： 通常不超过64KB。
2. 访问延迟低： 由于常量内存是通过缓存来访问的，当缓存命中时，其访问速度可以非常快，前提是一个线程束中的线 程（一个线程块中相邻的 32 个线程）要读取相同的常量内存数据。
3. 全局可见性： 对整个设备上的所有线程都是可见的，即所有的线程块都可以访问。
4. 只读性： 仅可读、不可写。

常量内存非常适合存储那些在内核执行期间不需要更改且被频繁访问的数据，例如转换矩阵、查找表等。

全局内存不能直接通过cudaMalloc()进行分配，只能通过cudaMemcpyToSymbol() 或cudaMemcpyToSymbolAsync()函数将数据从主机内存复制到设备的常量内存中，当程序结束时，CUDA驱动会自动回收该资源。

定义一个常量内存是在核函数外面用__constant__修饰符，下面是定义单个变量、数组、结构体的代码：

注意：核函数外定义的常量内存，在核函数中可以直接调用，不需要通过参数传导进去

#include <iostream>
#include <cuda_runtime.h>
#include "error_check.cuh"

// 定义常量内存中的数组大小
const int N = 4;
// 定义结构体
struct ConstStruct {
    float array[N];
    float singleValue;
};

// 核函数原型声明
__global__ void kernelFunction(float* d_result, float* d_result_s);

// 常量内存变量声明
__constant__ float d_constArray[N];    // 数组
__constant__ float d_singleValue;      // 单个变量
__constant__ ConstStruct d_constStruct; // 结构体

int main() {
    // 在主机上准备数据
    float h_array[N] = { 0 };
    ConstStruct h_constStruct;

    for (int i = 0; i < N; ++i) {
        h_array[i] = 2.1;
        h_constStruct.array[i] = 3.1f;
    }
    float h_singleValue = 1.1;
    h_constStruct.singleValue = 2.1f;

    // 将数据从主机复制到常量内存
    CHECK(cudaMemcpyToSymbol(d_constArray, h_array, N * sizeof(float)));
    CHECK(cudaMemcpyToSymbol(d_singleValue, &h_singleValue, sizeof(float)));
    CHECK(cudaMemcpyToSymbol(d_constStruct, &h_constStruct, sizeof(ConstStruct)));

    // 准备设备上的输出数组
    float* d_result;
    float* d_result_s;
    cudaMalloc(&d_result, N * sizeof(float));
    cudaMalloc(&d_result_s, N * sizeof(float));

    // 启动核函数
    kernelFunction << <1, N >> > (d_result, d_result_s);

    // 检查是否有错误
    CHECK(cudaGetLastError());

    // 将结果从设备复制回主机
    float h_result[N];
    float h_result_s[N];
    cudaMemcpy(h_result, d_result, N * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy(h_result_s, d_result_s, N * sizeof(float), cudaMemcpyDeviceToHost);

    // 输出结果
    for (int i = 0; i < N; ++i) {
        std::cout << "Result[" << i << "] = " << h_result[i] << std::endl;
        std::cout << "S_Result[" << i << "] = " << h_result_s[i] << std::endl;
    }

    // 清理资源
    cudaFree(d_result);
    cudaFree(d_result_s);

    return 0;
}

// 定义核函数
__global__ void kernelFunction(float* d_result, float* d_result_s) {
    int idx = threadIdx.x;
    if (idx < N) {
        // 从常量内存中读取常量数组和单个值，将结果存入 d_result
        d_result[idx] = d_constArray[idx] + d_singleValue;

        // 从常量内存中的结构体读取数据，将数组中的值和单个值相加，并存入 d_result_s
        d_result_s[idx] = d_constStruct.array[idx] + d_constStruct.singleValue;
    }
}

2.3 纹理内存

纹理内存（Texture Memory）： 具有缓存的全局内存。

特点：

1. 容量较大： 因为实际上它们是全局内存的一部分，所以容量还是比较大。
2. 主要用于只读： 对写操作支持非常有限，更适合读操作
3. 全局可见
4. 访问速度： 有专门的缓存优化，比全局内存快，适用于图像处理、计算机图形学等需要插值和边界处理的场景。

注意：使用__ldg(const T* addr)函数可以把某些只读的全局内存数据通过缓存，到达使用纹理内存的加速效果，后面会讨论该函数

2.4 表面内存

表面内存（surface memory）： 具有缓存的全局内存。

特点：

1. 容量较大： 因为实际上它们是全局内存的一部分，所以容量还是比较大。
2. 支持读写： 可以在核函数中直接修改表面内存中的数据。
3. 全局可见
4. 访问速度： 有专门的缓存优化，比全局内存快，用于需要频繁读写多维数据的应用，如图像处理和数据流处理。

2.5 寄存器

寄存器内存（Register Memory） 是所有内存中访问速度最快延迟最低的，但内存有限。

特点

1. 容量有限： 寄存器的数量是有限的，如果一个线程使用的寄存器数量超过了限制，编译器会自动将部分数据溢出到局部内存（Local Memory），这会导致性能下降。
2. 高速访问： 寄存器内存的访问速度非常快，通常只需要一个时钟周期。
3. 私有性： 每个线程都有自己的寄存器集合，不能被其他线程访问，仅仅被一个线程可见。
4. 可读可写
5. 自动管理： 寄存器内存由编译器自动管理，不需要手动指定哪些变量存储在寄存器中。编译器会根据变量的使用频率和生命周期来决定哪些变量应该存储在寄存器中。
6. 生命周期： 寄存器的生命周期也与所属线程的生命周期 一致，从定义它开始，到线程消失时结束。

在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器中。

以前提到过的各种内建变量，如 gridDim、blockDim、blockIdx、threadIdx 及 warpSize 都保存在特殊的寄存器中

例如下列代码

const int n = blockDim.x * blockIdx.x + threadIdx.x;

这里的 n 就是一个寄存器变量。

寄存器的数量和GPU相关，下图是一些不同计算能力和寄存器的指标：

2.6 局部内存

局部内存（Local Memory）： 当线程使用的寄存器数量超过限制时，编译器会将部分数据溢出到局部内存中。局部内存实际上位于全局内存中，但具有特定的访问特性和优化。

特点：

1. 访问延迟： 由于局部内存位于全局内存中，访问局部内存的延迟较高。
2. 私有性： 局部内存是每个线程独有的，不能被其他线程访问。
3. 容量： 每个线程最多能使用高达 512 KB 的局部内存，但使用 过多会降低程序的性能。

2.7 共享内存

共享内存（Shared Memory）： 是一种非常重要的内存类型，它位于每个SM中.

特点：

1. 高速访问： 共享内存的访问速度非常快，通常只需要几个时钟周期。
2. 线程块内共享： 共享内存只能被同一个线程块中的线程访问，不能跨线程块共享。
3. 有限容量： 每个SM的共享内存容量是有限的，具体容量取决于GPU的架构。
4. 手动管理： 共享内存需要程序员手动分配和管理

共享内存的主要作用是减少对全局内存的访问，或 者改善对全局内存的访问模式。

2.8 L1、L2缓存

从费米架构开始，有了 SM 层次的 L1 缓存（一级缓存）和设备（一个设备有多个 SM） 层次的 L2 缓存（二级缓存）。它们主要用来缓存全局内存和局部内存的访问，减少延迟。

从编程的角度来看，共享内存是可编程的缓存（共享内存的使用完全由用户操控），而 L1 和 L2 缓存是不可编程的缓存（用户最多能引导编译器做一些选择）。

对某些架构来说，还可以针对单个核函数或者整个程序改变 L1 缓存和共享内存的比例：

3 SM及其占有率

SM（Streaming Multiprocessor，流多处理器）是GPU的基本计算单元。每个SM包含多个流处理器（Streaming Processors，SP），这些流处理器负责执行实际的计算任务。

3.1 SM的构成

一个SM包含以下内容：

• 一定数量的寄存器
• 一定数量的共享内存
• 常量内存的缓存
• 纹理和表面内存的缓存
• L1 缓存
• 两个（计算能力 6.0）或 4 个（其他计算能力）线程束调度器（warp scheduler），用于
  在不同线程的上下文之间迅速地切换，以及为准备就绪的线程束发出执行指令。
• 执行核心，包括：
  – 若干整型数运算的核心（INT32）。
  – 若干单精度浮点数运算的核心（FP32）。
  – 若干双精度浮点数运算的核心（FP64）。
  – 若干单精度浮点数超越函数（transcendental functions）的特殊函数单元（Special Function Units，SFUs）。
  – 若干混合精度的张量核心（tensor cores，由伏特架构引入，适用于机器学习中的 低精度矩阵计算）。

3.2 SM的占有率

SM的占有率： 在指定的时间内，SM中的流处理器（SP）被有效利用的程度。高SM占有率意味着更多的计算资源被充分利用，从而提高整体性能，一般来说，要尽量让 SM 的占有率不小于 某个值，比如 25%，才有可能获得较高的性能。

要分析 SM 的理论占有率（theoretical occupancy），还需要知道两个指标：

• 一个 SM 中最多能拥有的线程块个数为 Nb = 16（开普勒架构和图灵架构）或者 Nb = 32（麦克斯韦架构、帕斯卡架构和伏特架构）；
• 一个 SM 中最多能拥有的线程个数为 Nt = 2048（从开普勒架构到伏特架构）或者 Nt = 1024（图灵架构），所以之前强调一个线程块（无论几维的）中的线程数不能超 过 1024。

下面在并行规模足够大（即核函数执行配置中定义的总线程数足够多）的前提下分几 种情况来分析 SM 的理论占有率：

①寄存器和共享内存使用量很少的情况：

• SM 的占有率完全由执行配置中的线程块大小决定。关于线程块大小，之前总是用 128。这是因为SM中线程的执行是以线程束为单位的，所以最好将线程块大小取为线程束大小（32 个线 程）的整数倍。
  例如，假设将线程块大小定义为 100，那么一个线程块中将有 3 个完 整的线程束（一共 96 个线程）和一个不完整的线程束（只有 4 个线程）。在执行核函数中的指令时，不完整的线程束花的时间和完整的线程束花的时间一样，这就无形中浪费了计算资源，所以建议将线程块大小取为 32 的整数倍。

②有限的寄存器个数的情况：

• 假设每个SM有 $T$ 个寄存器
• 假设每个线程需要 $R$ 个寄存器
• 假设每个线程块有 $B$ 个线程。

所以

• 每个线程块的寄存器总数=$R\times B$
• 每个SM可以容纳的线程块数量=$\lfloor T/R\times B\rfloor$，(向下取整)

又已知

• 每个线程块的线程束数量=$\lceil B/32\rceil$

所以

SM占有率计算公式如下：

③有限的共享内存对占有率的约束情况： 略，后面会详细介绍。

以 上 单 独 分 析 了 线 程 块 大 小、 寄 存 器 数 量 及 共 享 内 存 数 量 对 SM 占 有 率 的 影 响。 一 般 情 况 下， 需 要 综 合 以 上 三 点 分 析。 在 CUDA 工 具 箱 中， 有 一 个 名 为 CUDA_Occupancy_Calculator.xls 的 Excel 文 档， 可 用 来 计 算 各 种 情 况 下 的 SM 占有率，感兴趣的读者可以去尝试使用。

注： 用编译器选项 --ptxas-options=-v 可以报道每个核函数的寄存器使 用数量。CUDA 还提供了核函数的 launch_bounds() 修饰符和 --maxrregcount= 编 译选项来让用户分别对一个核函数和所有核函数中寄存器的使用数量进行控制。

4 用CUDA运行时API函数查询设备

该段介绍用 CUDA 运行时 API 函数查询所用 GPU 的规格 ，可以通过以下代码查看显卡的信息：

#include <iostream>
#include <cuda_runtime.h>
#include "error_check.cuh"

int main(int argc, char* argv[]) {
	int device_id = 0;//如果你不止一个显卡，可以切换ID，输出不同显卡的信息
	if (argc > 1) device_id = atoi(argv[1]);
	CHECK(cudaSetDevice(device_id));

	cudaDeviceProp prop;
	CHECK(cudaGetDeviceProperties(&prop, device_id));

	printf("Device id: %d\n", device_id);
	printf("Device name: %s\n", prop.name);
	printf("Compute capability: %d.%d\n", prop.major, prop.minor);
	printf("Amount of global memory: %g GB\n", prop.totalGlobalMem / (1024.0 * 1024 * 1024));
	printf("Amount of constant memory: %g KB\n", prop.totalConstMem / 1024.0);
	printf("Maximum grid size: %d %d %d\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
	printf("Maximum block size: %d %d %d\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
	printf("Number of SMs: %d\n", prop.multiProcessorCount);
	printf("Maximum amount of shared memory per block: %g KB\n", prop.sharedMemPerBlock / 1024.0);//每个线程块可以使用的最大共享内存
	printf("Maximum amount of shared memory per SM: %g KB\n", prop.sharedMemPerMultiprocessor / 1024.0);//个SM可以分配的最大共享内存总量
	printf("Maximum number of registers per block: %d K\n", prop.regsPerBlock / 1024);//每个线程块可以使用的最大寄存器数量
	printf("Maximum number of registers per SM: %d K\n", prop.regsPerMultiprocessor / 1024);//每个SM可以分配的最大寄存器总量
	printf("Maximum number of threads per block: %d\n", prop.maxThreadsPerBlock);
	printf("Maximum number of threads per SM: %d\n", prop.maxThreadsPerMultiProcessor);//每个SM可以同时运行的最大线程数量
	return 0;
}

输出结果如下：