CUDA - 页锁内存

页锁内存（Page-locked Memory）简介

页锁内存，又称 固定内存（Pinned Memory）或 页固定内存，是指一块物理上固定分配的主机内存，操作系统不会对其进行分页（swapping）操作。这种内存的主要特性是可直接通过 DMA（Direct Memory Access） 方式与设备（如 GPU）进行数据传输，从而提高数据传输效率。

---

页锁内存的特点

1. 物理连续性：
   页锁内存的物理地址是连续的，便于 DMA 控制器直接访问。普通分页内存的虚拟地址是连续的，但物理地址可能是分散的，导致额外的开销。

2. 更高的数据传输性能：
   使用页锁内存进行主机和设备之间的数据传输，性能比普通分页内存高，因为避免了内存分页和映射操作。

3. 支持异步操作：
   页锁内存可与 GPU 的流（stream）配合使用，支持异步内存拷贝（如 cudaMemcpyAsync 或 hipMemcpyAsync），实现数据传输与计算的并行。

4. 访问直接性：
   页锁内存为 GPU 提供了直接访问主机内存的可能，减少不必要的数据拷贝。

---

页锁内存的优缺点

优点

1. 高效的数据传输：
   页锁内存避免了操作系统的分页机制导致的额外开销，能够以更高的带宽传输数据。

2. 异步拷贝的必要条件：
   异步内存拷贝需要源或目标内存为页锁定内存，从而实现数据传输与计算并行化。

3. 零拷贝（Zero-copy）支持：
   GPU 在某些场景下（如统一内存架构）可以直接使用页锁定内存中的数据，无需显式拷贝到设备端。

缺点

1. 内存资源有限：
   页锁内存分配后会减少主机可分页内存的容量。分配过多可能导致系统性能下降。

2. 分配成本较高：
   页锁内存的分配过程比普通内存慢，且操作系统需要对其进行特殊管理。

3. 增加程序复杂性：
   开发者需要显式管理页锁内存的分配和释放，同时考虑其影响范围。

---

页锁内存的使用

1. 在 CUDA 中使用页锁内存

CUDA 提供了 cudaHostAlloc 函数用于分配页锁内存：

float *host_data;
cudaHostAlloc((void**)&host_data, size * sizeof(float), cudaHostAllocDefault);

页锁内存释放：

cudaFreeHost(host_data);

异步拷贝示例：

cudaMemcpyAsync(device_ptr, host_data, size * sizeof(float), cudaMemcpyHostToDevice, stream);

2. 在 HIP 中使用页锁内存

HIP 提供了 hipHostMalloc 函数：

double *host_data;
hipHostMalloc((void**)&host_data, size * sizeof(double), 0);

释放页锁内存：

hipHostFree(host_data);

HIP 异步拷贝示例：

hipMemcpyAsync(device_ptr, host_data, size * sizeof(double), hipMemcpyHostToDevice, stream);

---

页锁内存的应用场景

1. 异步数据传输：
   在主机和设备之间传输大量数据时，配合流实现计算与数据传输并行。

2. 高性能计算（HPC）：
   在科学计算、机器学习等需要频繁主机-设备交互的场景中，利用页锁内存提高数据传输效率。

3. 零拷贝操作：
   在共享虚拟内存（Unified Memory）环境下，页锁内存可用作主机和设备的共享内存，减少拷贝操作。

---

页锁内存与普通内存的比较

属性	页锁内存	普通内存
分配方式	使用 cudaHostAlloc 或 hipHostMalloc	使用 malloc 或 new
是否分页	不分页	分页
传输性能	高性能	性能较低
异步拷贝支持	支持	不支持
开发复杂度	需要显式分配和释放	简单

---

页锁内存的注意事项

1. 控制分配规模：
   分配页锁内存时应合理控制其大小，避免因系统内存不足导致性能下降。

2. 释放内存：
   程序结束时应显式释放页锁内存，否则可能导致内存泄漏。

3. 性能测试：
   不同硬件平台对页锁内存的优化程度可能不同，建议在实际使用前测试性能提升是否显著。

4. 多线程程序中使用：
   页锁内存分配和使用需要注意线程安全问题，在多线程环境下需特别小心。

---

页锁内存实际测试

#include <iostream>
#include <cuda_runtime.h>

using namespace std;

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

    unsigned int isize = 1 << 22;
    unsigned int nbytes = isize * sizeof(float);

    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, dev);
    printf("%s starting at ", argv[0]);
    printf("device %d: %s memory size %d nbyte %5.2fMB\n", dev, deviceProp.name, isize, nbytes / (1024.0f * 1024.0f));

    float *h_a = (float *)malloc(nbytes);
    // float *h_a;
    // cudaError_t status = cudaMallocHost((float **)&h_a, nbytes);
    // if (status != cudaSuccess)
    // {
    //     fprintf(stderr, "Error returnd from pinned host memory allocation\n");
    //     exit(1);
    // }

    float *d_a;
    cudaMalloc((float **)&d_a, nbytes);

    for (size_t i = 0; i < isize; i++)
    {
        h_a[i] = 0.5f;
    }

    for (int i = 0; i < 100; i++)
    {
        cudaMemcpy(d_a, h_a, nbytes, cudaMemcpyHostToDevice);

        cudaMemcpy(h_a, d_a, nbytes, cudaMemcpyDeviceToHost);
    }

    free(h_a);
    // cudaFreeHost(h_a);
    cudaFree(d_a);

    cudaDeviceReset();
    return EXIT_SUCCESS;
}

采用cudaMalloc

采用cudaMallocHost

---

总结

页锁内存是主机与设备间高效数据传输的关键技术，尤其适用于需要异步拷贝或零拷贝的 GPU 应用场景。通过合理分配和使用页锁内存，可以显著提升程序的性能，但也需要权衡其对内存资源的影响以及编程复杂性。