Notes about CUDA Zero-Copy Memory

本文将mark下CUDA Zero-Copy Memory的相关notes。

原理

pin住host的dram内存
host CPU mmu建立VA到pinned memory的映射
GPU mmu建立VA到pinned memory的映射

通过将一段pin住的物理内存，同时映射到CPU和GPU的虚拟地址空间来实现。这样一来，GPU内核就能像访问自己的显存一样，直接读写这块主机内存，省去了显式的cudaMemcpy操作。

Advantages

GPU threads can directly access zero-copy memory. There are several advantages to using zero-copy memory in CUDA kernels, such as:

Leveraging host memory when there is insuffi cient device memory
Avoiding explicit data transfer between the host and device
Improving PCIe transfer rates

劣势与挑战

带宽瓶颈: 访问速度受限于PCIe总线带宽。对频繁访问的数据，效率远低于本地显存。
高延迟: 每次访问都需经过PCIe总线，延迟比本地显存高几个数量级。
非合并访问影响大: GPU对零拷贝内存的非合并访问将产生大量小的PCIe事务，严重降低性能。应尽量确保GPU对它的访问是合并的（即连续的线程访问连续的内存地址）。

实践

分配并映射新的锁页内存
这种方式直接分配并映射一块全新的主机锁页内存，是最常见的方法。

核心API: cudaHostAlloc
关键步骤:
- 分配：使用 cudaHostAlloc 并指定 cudaHostAllocMapped 标志。
- 获取设备指针：通过 cudaHostGetDevicePointer 获取该内存对应的GPU侧有效指针。
- 在核函数中使用：将设备指针传入核函数。

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

// GPU 核函数：直接对主机上的锁页内存进行操作
__global__ void vector_add_kernel(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        // 直接访问主机上的锁页内存，如同访问显存
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1024;
    size_t size = n * sizeof(float);

    // 1. 在主机上分配并映射锁页内存
    float *h_a, *h_b, *h_c;
    cudaHostAlloc((void **)&h_a, size, cudaHostAllocMapped);
    cudaHostAlloc((void **)&h_b, size, cudaHostAllocMapped);
    cudaHostAlloc((void **)&h_c, size, cudaHostAllocMapped);

    // 2. 在主机端初始化数据
    for (int i = 0; i < n; ++i) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    // 3. 获取该主机内存对应的GPU设备指针
    float *d_a, *d_b, *d_c;
    cudaHostGetDevicePointer((void **)&d_a, (void *)h_a, 0);
    cudaHostGetDevicePointer((void **)&d_b, (void *)h_b, 0);
    cudaHostGetDevicePointer((void **)&d_c, (void *)h_c, 0);

    // 4. 配置核函数并启动
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vector_add_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);
    cudaDeviceSynchronize();

    // 5. 检查结果，此时结果已直接写入主机的 h_c 中
    for (int i = 0; i < n; ++i) {
        if (h_c[i] != 3.0f) {
            printf("Error: h_c[%d] = %f\n", i, h_c[i]);
            break;
        }
    }
    printf("Test passed!\n");

    // 6. 清理资源
    cudaFreeHost(h_a);
    cudaFreeHost(h_b);
    cudaFreeHost(h_c);
    return 0;
}

参考资料:

cuda-samples
Page-Locked Host Memory
Professional CUDA C Programming
https://chat.deepseek.com/share/ran9srj0hi8t35xmsw

文章目录

原理

Advantages

劣势与挑战

实践