我还在学CUDA编程（四）——内存管理及访问模式

内存管理

内存分配和释放

CUDA 编程模型中存在 Host 和 Device 两个异构系统，每个系统都有相对独立的内存空间。

在 Host 代码中，可以使用下面的函数分配 Device 全局内存：

cudaError_t cudaMalloc(void **devPtr, size_t count);

需要注意第一个参数是指针的指针。函数执行失败时，可能会返回 cudaErrorMemoryAllocation。

分配完地址后，可以使用下面的函数初始化 Device 内存：

cudaError_t cudaMemset(void *devPtr, int value, size_t count);

这里需要注意，cudaMemset() 是按字节设置内存，适合置 0 或设置某种字节模式，不能简单理解成像 C++ 赋值那样给任意类型数组设置任意值。

当内存不再使用时，使用下面的函数释放内存：

cudaError_t cudaFree(void *devPtr);

这个函数中的参数必须是 cudaMalloc() 分配的 Device 内存地址，否则可能会返回 cudaErrorInvalidDevicePointer 错误。

内存传输

分配好全局内存后，在传统离散内存模型中，由于 Host 不能直接访问 Device 内存，Device 也不能直接访问普通 Host 内存，所以需要显式传输数据：

cudaError_t cudaMemcpy(
    void *dst,
    const void *src,
    size_t count,
    enum cudaMemcpyKind kind
);

注意这里的参数是指针，而不是指针的指针。第一个参数 dst 是目标地址，第二个参数 src 是原始地址，count 是拷贝内存的大小，最后的 kind 是复制方向。

复制方向包括以下几种：

1. cudaMemcpyHostToHost
2. cudaMemcpyHostToDevice
3. cudaMemcpyDeviceToHost
4. cudaMemcpyDeviceToDevice

后面会提到的 pinned memory、zero-copy、UVA 和 Unified Memory，都是在这个传统模型基础上的扩展或例外。

固定内存

在 CUDA 架构的 Host 端，内存分为两种：一种是可分页内存（pageable memory），另一种是页锁定内存（pinned memory）。可分页内存通常由操作系统 API 或 malloc() 在主机上分配，页锁定内存则可以由 CUDA 函数 cudaMallocHost() 和 cudaHostAlloc() 在 Host 内存中分配。

页锁定内存的重要属性是 Host 的 OS 不会对这块内存进行分页和交换操作，确保该内存始终保留在物理内存中。GPU 知道页锁定内存的物理地址，可以通过直接内存访问（Direct Memory Access，DMA）技术直接在 Host 和 Device 之间复制数据，传输吞吐通常更高。

不过，页锁定内存并不是越多越好。由于每块页锁定内存都需要占用物理内存，并且不能交换到磁盘上，过度使用会影响系统分页能力，甚至降低系统整体性能。

GPU 不能在可分页的 Host 内存上安全地直接访问数据，因为当 Host 在物理位置上移动数据时，Device 无法控制。可分页内存从 Host 到 Device 的过程参见上图左侧。

CUDA 运行时可以使用下面的函数直接分配固定主机内存：

cudaError_t cudaMallocHost(void **devPtr, size_t count);

该函数会分配 count 字节的固定内存，这些内存是页锁定的，可以直接用于更高效的 Host-Device 数据传输。

Host 端页锁定内存释放必须使用：

cudaError_t cudaFreeHost(void *ptr);

页锁定内存虽然分配和释放成本更高，但是为大规模 Host-Device 数据传输提供了更高吞吐量。

零拷贝内存

前面说的内容都是：Host 不可以直接访问 Device 内存，Device 也不可以直接访问普通 Host 内存。**零拷贝内存（Zero-Copy）**是一个特殊情况，Host 和 Device 都可以访问这块内存。

CUDA 的 Kernel 使用零拷贝内存常见于以下几种情况：

• 当设备内存不足时，临时利用主机内存。
• 避免 Host 和 Device 之间的显式内存传输。
• 减少一次中转和额外内存占用。

需要注意，零拷贝并不意味着访问成本为零。Device 访问零拷贝内存时仍然会经过 PCIe，因此带宽和延迟通常不如 Device DRAM。它更适合少量、一次性、稀疏访问的数据。如果需要频繁读写或大量连续传输，性能可能显著下降。

另外，使用零拷贝共享 Host 和 Device 间数据时，必须同步 Host 和 Device 之间的内存访问，避免内存竞争。

零拷贝内存是页锁定内存，并且该内存会映射到设备地址空间中。可以通过以下函数创建固定内存映射：

cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);

其中 flags 参数可以有以下值：

• cudaHostAllocDefault：和 cudaMallocHost() 函数一致。
• cudaHostAllocPortable：返回能被所有 CUDA 上下文使用的固定内存，而不是只对执行内存分配的上下文可用。
• cudaHostAllocWriteCombined：返回写结合内存，在特定系统配置下通过 PCIe 传输很快，对缓冲区来说是个不错的选择。
• cudaHostAllocMapped：产生零拷贝内存。

此时 Device 不一定能直接通过 pHost 访问对应内存地址。Device 要访问 Host 上的零拷贝内存时，通常需要先获得一个 Device 端可用地址：

cudaError_t cudaHostGetDevicePointer(
    void **pDevice,
    void *pHost,
    unsigned flags
);

这里的 flags 通常置为 0。

还得注意，因为每次访问都要经过 PCIe，在频繁读写操作时，零拷贝内存的效率会显著降低。当有大量数据传输时，性能也可能显著下降。在《CUDA C 编程权威指南》和谭升老师的 Blog中都对此进行了测试。

统一虚拟寻址

CUDA 4.0 之后的版本开始支持统一虚拟寻址（Unified Virtual Addressing，UVA）。在 UVA 出现之前，我们需要自己管理哪些指针指向 Host 内存，哪些指针指向 Device 内存。UVA 使得指针指向的内存空间对程序代码变得更透明。

通过 UVA，对于 mapped pinned memory，Host 指针和 Device 指针的值在很多情况下可能相同，因此不一定需要再单独调用：

cudaError_t cudaHostGetDevicePointer(
    void **pDevice,
    void *pHost,
    unsigned flags
);

但是这件事不能绝对化。是否需要 cudaHostGetDevicePointer()，仍然取决于内存分配方式、设备属性以及当前 CUDA 环境。

统一内存

CUDA 6.0 中引入了统一内存（Unified Memory，UM），也就是托管内存（Managed Memory）这个特性。不同于 UVA，统一内存的重点是简化内存管理。它会创建一个托管内存池，池中已分配的空间可以用相同的内存地址（指针）在 Host 和 Device 上进行访问。底层系统会在需要时自动进行 Device 和 Host 之间的数据迁移。

Unified Memory 依赖 UVA，但是二者又完全不同。UVA 为系统中的处理器提供了一个统一的虚拟内存地址空间，但 UVA 本身不会自动迁移数据。

托管内存指的是由系统底层自动管理迁移的统一内存，未托管内存就是我们自己显式分配和管理的内存。可以同时传递给 Kernel 两种类型的内存，即托管内存和未托管内存。

托管内存可以动态分配，也可以静态声明。动态托管内存通过 Host 端调用 cudaMallocManaged() 分配：

cudaError_t cudaMallocManaged(
    void **devPtr,
    size_t size,
    unsigned int flags = 0
);

静态或全局托管变量可以使用 __managed__ 修饰，并在全局或命名空间作用域声明。Device 代码不能调用 cudaMallocManaged()。

内存访问模式

我们所编写的 GPU 程序容易受到内存带宽的限制，因此，最大限度利用全局内存带宽是提高 Kernel 性能的重要手段。

CUDA 执行模型的特征之一就是指令必须以 Warp 为单位发布和执行，存储操作也是如此。

对齐与合并访问

全局内存（DRAM）通过缓存来实现加载和存储，所有对全局内存的访问都会通过 L2 Cache，有些访问也会通过 L1 Cache。

当 Kernel 从 DRAM 中读取数据时，常见缓存粒度有 128 Byte 和 32 Byte。哪怕只读取 1 个 Byte 的变量，也可能触发更大粒度的数据传输。CUDA 支持通过编译选项控制部分全局内存加载是否使用 L1 Cache。

需要注意，128 Byte / 32 Byte 粒度、L1 是否可配置、全局内存请求如何切分，都和 GPU 架构、计算能力以及编译选项有关。早期资料中的描述很有参考价值，但实际项目中仍应结合具体 GPU 架构理解。

如果启用 L1 Cache，那么从 DRAM 加载数据时可能以 128 Byte 为粒度；如果不用 L1 Cache，只走 L2 Cache，则可能以 32 Byte 为粒度。

SM 执行的基础是线程束。也就是说，当一个 SM 中正在执行的某个线程需要访问内存时，和它同线程束的其他 31 个线程也会一起发起内存访问。这个基础表示，即使每个线程只访问一个字节，只要有内存请求，也至少会形成一个 warp 级别的访问请求。因此，不使用一级缓存的内存加载，一次粒度通常是 32 Byte，而不是更小。

所以在优化程序时，需要注意设备内存访问的 2 个特性：

• 对齐内存访问。
• 合并内存访问。

当设备内存事务的第一个地址，是用于服务该事务的缓存粒度的整数倍（例如 32 Byte L2 Cache 或 128 Byte L1 Cache）时，会出现对齐内存访问。非对齐加载会造成带宽浪费。

当一个 Warp 中的 32 个 Thread 访问一个连续内存 Block 时，会出现合并内存访问。

下图描述了对齐与合并内存的加载操作，其中蓝色表示全局内存，橙色是 Warp 要的数据，绿色是对齐的地址段：

下图描述了非对齐和未合并的内存访问：

上图的情况下，一个 warp 的请求跨越了 3 个 cache line / segment，因此需要 3 个内存事务，内存利用率约为对齐和合并情况下的 1/3。在极端情况下，每个 Thread 的请求都在不同的段，利用率会进一步下降。

全局内存读取

在 SM 中，根据不同设备和访问类型，加载数据有三种途径：

• L1 / L2 Cache
• 常量缓存
• 只读缓存

默认情况下会通过 L1 / L2 Cache，如果想通过其他途径需要显式说明。禁用 L1 Cache 的编译器选项为 -Xptxas -dlcm=cg，启用 L1 的编译器选项为 -Xptxas -dlcm=ca。

如果 L1 Cache 被禁用，那么所有对全局内存的加载请求会直接进入 L2 Cache；如果 L2 Cache 缺失，则由 DRAM 完成请求。

内存加载可分为 2 类：

1. 缓存加载
2. 无缓存加载

内存加载的访问模式有几个特点：

• 有无 Cache：如启用 L1 Cache，则内存加载可以被缓存。
• 对齐 / 非对齐：如内存访问首地址为 32 Byte 的整数倍，则是对齐加载。
• 合并 / 非合并：如 Warp 访问一个连续的数据块，则加载更容易合并。

缓存加载

缓存加载操作会经过 L1 Cache。下图为理想情况，对齐与合并的内存访问，利用率 100%。

下图为对齐访问，但引用地址不是连续 ID，而是 128 Byte 内的随机值。由于 Thread 访问的数据都在一个 block 内，利用率也是 100%。

下图为 warp 请求连续但非对齐的 32 个 4 Byte 数据。数据横跨 2 个 block，需要 2 个 128 Byte 的事务，利用率 50%。

下图中 warp 中所有 Thread 都请求相同的地址，只需请求一个内存事务。如果请求的数据为 4 Byte，那么利用率为 4 / 128 = 3.125%。

下图是最差的情况，warp 中所有 Thread 请求分散的 32 个 4 Byte 地址，所有数据分布在 N 个缓存行上。假设 1 行 L1 Cache 为 128 Byte，请求 32 个 4 Byte 的数据，就需要 N 个内存事务，利用率为 1 / N。

GPU 的 L1 Cache 在部分架构和编译选项下可以进行控制，CPU 的 L1 Cache 通常不由程序员直接控制。而且 CPU 的 L1 Cache 更强调时间局部性和空间局部性，GPU 的缓存设计则更服务于大规模并行访问模式。

无缓存加载

相比 L1 Cache 加载，无缓存加载的内存段粒度可以从 128 Byte 变成 32 Byte，更细粒度的加载会带来更好的总线利用率。下图是一种理想情况，对齐与合并的内存访问，总线利用率 100%：

下图为一种随机访问，但是利用率也是 100%：

下图为连续但是不对齐的情况，请求的地址落在 5 个内存段内。连续访问的情况下通常不可能超过 5 个段，利用率至少为 80%：

下图为 Warp 中所有 Thread 都请求相同的数据，地址落在同一个内存段内。利用率为 4（请求 4 Byte）/ 32（加载 32 Byte）= 12.5%：

下图是最坏的情况，Warp 请求 32 个分散在全局内存中的 4 Byte 数据。由于请求的 128 Byte 数据最多落在 N 个 32 Byte 的内存分段内，而不是 N 个 128 Byte 的内存分段内，所以比 Figure 4-13 的情况还要好一些。

只读缓存

只读缓存早期常与纹理路径相关。在一些设备中，只读缓存也支持使用全局内存加载来代替 L1 Cache。现代 GPU 中，只读数据缓存、texture cache 和 L1 的关系会随架构变化，因此具体行为仍应参考对应架构文档。

只读缓存的加载粒度通常为 32 Byte。有两种方式可以提示内存通过只读路径进行读取：

• 使用函数 __ldg()，注意这里是字母 l，不是数字 1。
• 在间接引用的指针上使用合适的修饰符。

下面代码使用函数 __ldg() 通过只读路径直接对数组进行读取访问：

__global__ void copyKernel(float *in, float *out)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;
    out[idx] = __ldg(&in[idx]);
}

也可以使用 __restrict__ 修饰符应用到指针上。这些修饰符可以帮助 nvcc 编译器识别无别名指针。

__global__ void copyKernel_2(int *__restrict__ out, const int *__restrict__ in)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    out[idx] = in[idx];
}

全局内存写入

内存的写入和加载（读取）是不同的。全局内存 store 通常经过 L2 写入路径；是否经过或影响 L1，与 GPU 架构和 cache policy 有关。存储操作在 32 Byte 的粒度上进行，内存事务也会被分为 1 段、2 段或者 4 段。例如，如果 2 个地址属于一个 128 Byte 的区域，但是不属于一个对齐的 64 Byte 区域，则可能会执行一个 4 段事务。

下图是理想情况，对齐的内存访问，并且 Warp 中所有 Thread 访问一个连续的 128 Byte 范围。存储操作使用一个 4 段事务完成。

下图数据分散在一个 192 Byte 的范围内，不连续，使用三个 1 段事务完成。

下图为对齐的访问，在一个 64 Byte 的范围内，使用一个 2 段事务完成。

结构体数组（SoA）& 数组结构体（AoS）

AoS（Array of Structures）是数组元素为结构体；SoA（Structure of Arrays）是结构体中按字段拆成多个数组。

如果使用 AoS 方法存储数据，先定义一个结构体，再定义一个结构体的数组。存储的是空间中相邻的数据，比如 x 和 y，在 CPU 上有较好的缓存局部性。

// Array of Structures (AoS)
struct AoS {
    float x, y, z, w;
};

AoS myAoS[N];

如果使用 SoA 方式存储数据，原结构体中每个字段的值都被分发到各自的数组中。这种方式不仅能将相同字段的数据紧密存储，还可以跨数组保存数据。

// Structure of Arrays (SoA)
struct SoA {
    float x[N];
    float y[N];
};

SoA mySoA;

如何区分 AoS 和 SoA 呢？我区分不了……（笑

其实 SoA 和 AoS 的中文名称很容易让人困惑。SoA 全称是 Structure of Arrays，也就是数组的结构体，本质是一个结构体，再看它的实例化，也是声明为结构体。AoS 全称是 Array of Structures，也就是结构体的数组，本质是一个数组，再看它的实例化，也是实例化为一个数组。

下图说明了两种方法的内存布局。

AoS 模式在 GPU 上存储数据并执行一个只使用 x 字段的代码时，可能会导致带宽浪费，同时 AoS 也在不需要访问的 y 等变量上浪费了传输带宽。

SoA 模式可以充分利用 GPU 内存带宽，实现更高效的全局内存访问。

指针数组和数组指针

单纯复习一下：

• 指针数组：可以说成是指针的数组，首先是一个数组，数组的所有元素都是指针类型。
• 数组指针：可以说成是数组的指针，首先是一个指针，指针存放着一个数组的首地址，或者说这个指针指向一个数组的首地址。

AoS 进行简单数学运算

#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"

struct naiveStruct {
    float a;
    float b;
};

void sumArrays(float *a, float *b, float *res, const int size)
{
    for (int i = 0; i < size; i++) {
        res[i] = a[i] + b[i];
    }
}

__global__ void sumArraysGPU(float *a, float *b, struct naiveStruct *res, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        res[i].a = a[i] + b[i];
    }
}

void checkResult_struct(float *res_h, struct naiveStruct *res_from_gpu_h, int nElem)
{
    for (int i = 0; i < nElem; i++) {
        if (res_h[i] != res_from_gpu_h[i].a) {
            printf("check fail!\n");
            exit(0);
        }
    }
    printf("result check success!\n");
}

int main(int argc, char **argv)
{
    int dev = 0;
    cudaSetDevice(dev);

    int nElem = 1 << 18;
    int offset = 0;
    if (argc >= 2) {
        offset = atoi(argv[1]);
    }

    printf("Vector size:%d\n", nElem);

    int nByte = sizeof(float) * nElem;
    int nByte_struct = sizeof(struct naiveStruct) * nElem;

    float *a_h = (float *)malloc(nByte);
    float *b_h = (float *)malloc(nByte);
    float *res_h = (float *)malloc(nByte_struct);
    struct naiveStruct *res_from_gpu_h = (struct naiveStruct *)malloc(nByte_struct);

    memset(res_h, 0, nByte);
    memset(res_from_gpu_h, 0, nByte);

    float *a_d, *b_d;
    struct naiveStruct *res_d;

    CHECK(cudaMalloc((float **)&a_d, nByte));
    CHECK(cudaMalloc((float **)&b_d, nByte));
    CHECK(cudaMalloc((struct naiveStruct **)&res_d, nByte_struct));
    CHECK(cudaMemset(res_d, 0, nByte_struct));

    initialData(a_h, nElem);
    initialData(b_h, nElem);

    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    dim3 block(1024);
    dim3 grid(nElem / block.x);

    double iStart, iElaps;
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;

    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte_struct, cudaMemcpyDeviceToHost));
    printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec\n", grid.x, block.x, iElaps);

    sumArrays(a_h, b_h, res_h, nElem);
    checkResult_struct(res_h, res_from_gpu_h, nElem);

    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);

    free(a_h);
    free(b_h);
    free(res_h);
    free(res_from_gpu_h);

    return 0;
}

SoA 进行简单数学运算

#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"

void sumArrays(float *a, float *b, float *res, int offset, const int size)
{
    for (int i = 0, k = offset; k < size; i++, k++) {
        res[i] = a[k] + b[k];
    }
}

__global__ void sumArraysGPU(float *a, float *b, float *res, int offset, int n)
{
    int i = blockIdx.x * blockDim.x * 4 + threadIdx.x;
    int k = i + offset;

    if (k + 3 * blockDim.x < n) {
        res[i] = a[k] + b[k];
        res[i + blockDim.x] = a[k + blockDim.x] + b[k + blockDim.x];
        res[i + blockDim.x * 2] = a[k + blockDim.x * 2] + b[k + blockDim.x * 2];
        res[i + blockDim.x * 3] = a[k + blockDim.x * 3] + b[k + blockDim.x * 3];
    }
}

int main(int argc, char **argv)
{
    int dev = 0;
    cudaSetDevice(dev);

    int block_x = 512;
    int nElem = 1 << 18;
    int offset = 0;

    if (argc == 2) {
        offset = atoi(argv[1]);
    } else if (argc == 3) {
        offset = atoi(argv[1]);
        block_x = atoi(argv[2]);
    }

    printf("Vector size:%d\n", nElem);

    int nByte = sizeof(float) * nElem;
    float *a_h = (float *)malloc(nByte);
    float *b_h = (float *)malloc(nByte);
    float *res_h = (float *)malloc(nByte);
    float *res_from_gpu_h = (float *)malloc(nByte);

    memset(res_h, 0, nByte);
    memset(res_from_gpu_h, 0, nByte);

    float *a_d, *b_d, *res_d;

    CHECK(cudaMalloc((float **)&a_d, nByte));
    CHECK(cudaMalloc((float **)&b_d, nByte));
    CHECK(cudaMalloc((float **)&res_d, nByte));
    CHECK(cudaMemset(res_d, 0, nByte));

    initialData(a_h, nElem);
    initialData(b_h, nElem);

    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    dim3 block(block_x);
    dim3 grid(nElem / block.x);

    double iStart, iElaps;

    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, offset, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;
    printf("warmup Time elapsed %f sec\n", iElaps);

    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, offset, nElem);
    cudaDeviceSynchronize();
    iElaps = cpuSecond() - iStart;

    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
    printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec --offset:%d \n", grid.x, block.x, iElaps, offset);

    sumArrays(a_h, b_h, res_h, offset, nElem);
    checkResult(res_h, res_from_gpu_h, nElem - 4 * block_x);

    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);

    free(a_h);
    free(b_h);
    free(res_h);
    free(res_from_gpu_h);

    return 0;
}

性能调整

优化设备内存带宽利用率的目标有二：

1. 对齐与合并内存访问，减少带宽浪费。
2. 使用足够并发的内存操作，隐藏内存延迟。

之前优化吞吐量的内容中提到，实现并发内存访问最大化可以通过：

1. 增加每个线程中执行独立内存操作的数量。
2. 对 Kernel 启动配置进行试验。

展开技术

包含内存操作的循环展开，可以增加更多独立的内存操作，从而提高并行度和隐藏内存延迟的能力。

增大并行性

可以通过调整 Block 的大小来实现并行性调整。不同 Block 大小会影响 occupancy、寄存器压力、共享内存使用量和内存访问并发度，因此通常需要结合具体 Kernel 做测试。

参考文章

[1] CUDA：页锁定内存（pinned memory）和按页分配内存（pageable memory） - 博客园

[2] CUDA C++ / 内存管理以及优化 - 博客园

[3] 【CUDA 基础】4.2 内存管理 | 谭升的博客

[4] CUDA：cudaMalloc vs cudaMallocHost - CSDN 博客

[5] 【CUDA 基础】4.3 内存访问模式 | 谭升的博客

[6] 优化数据排布，让你的程序加速 4 倍！ - 知乎

[7] 指针数组和数组指针的区别 - 菜鸟教程