CUDA中的内存管理、锁页内存、UVA统一虚拟地址、零拷贝、统一内存

0 前言

翻了下以前关于CUDA的UVA、零拷贝、统一内存的笔记，感觉顺序有些乱，而且里面有个描述还是错的，这次重新整理一下。

1 swap内存跟锁页内存

Swap 是操作系统提供的一种“虚拟内存扩展机制”。当物理内存（比如一根 4GB 内存条）不够用时，操作系统会将某些暂时不活跃的内存页（比如后台程序的数据）换出（swap out）到磁盘上的交换空间（Swap 分区或 Swap 文件），从而释放物理内存，给当前活跃的程序使用。

通俗地说：

“房子床位不够，就把出差的人行李先收拾塞进仓库（磁盘），腾出床位给新来的程序。老住户回来时，再从仓库拿回来换入（swap in）。”

与此相反，锁页内存就是告诉操作系统，这块内存是我“强占”的，不能随便给我换到磁盘上去。

2 UVA(Unified Virtual Addressing)统一虚拟地址

“以前 CPU 和 GPU 各自管理自己的虚拟地址空间，彼此之间的指针不能通用。而有了统一虚拟地址（UVA）后，CPU 内存和 GPU 显存共享同一个虚拟地址空间，指针在 CPU 和 GPU 间可以直接传递、访问更自然、管理更统一。”

3 先看最普通的cuda内存分配、释放、传输

先看内存分配和释放

float * devMem=NULL;
cudaError_t cudaMalloc((float**) devMem, count)
cudaError_t cudaMemset(void * devPtr,int value,size_t count)
cudaError_t cudaFree(void * devPtr)

然后是内存传输

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

然后最后一个参数有下面四个枚举值

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

4 申请锁页内存

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

这里的第三个参数flags可以有下面四个选项，这个 flags 可以是以下这些值之一或多个（可 OR 组合）

cudaHostAllocDefalt
cudaHostAllocPortable
cudaHostAllocWriteCombined
cudaHostAllocMapped

4.1 cudaHostAllocDefault

• 说明：这是默认行为，分配常规的锁页内存。
• 特性：
  • 分配的内存是锁页内存（pinned），不允许被操作系统换出（swap）。
  • 适合用于 CPU↔GPU 传输的缓冲区，支持高效的 DMA 传输。
  • 不保证能被所有 CUDA 上下文共享，也不保证是写结合（write-combined）内存。
• 使用场景：
  • 一般常规的主机缓冲区分配，兼顾传输性能和通用性

4.2 cudaHostAllocPortable

• 说明：分配的内存是可被所有 CUDA 上下文共享的锁页内存。
• 特性：
  • 与 cudaHostAllocDefault 类似，但确保这块内存在所有 CUDA 上下文中都可用。
  • 适合多 GPU 或多上下文应用程序。
• 使用场景：
  • 多 GPU 环境下，多个上下文都需要访问同一块主机内存。

4.3 cudaHostAllocWriteCombined

• 说明：分配写结合（Write-Combined, WC）内存。
• 特性：
  • 写结合内存不保证 CPU 读取效率高，但对 GPU 写入性能有利。
  • 适合CPU 主要写入、GPU 主要读取的场景。
  • CPU 读这块内存时速度可能较慢（因为写结合内存是针对写优化）。
  • 仍是锁页内存，支持高速传输。
• 使用场景：
  • CPU 向缓冲区写数据，GPU 读取数据的流式处理场景，如视频解码后处理。

4.3 cudaHostAllocMapped

• 说明：分配映射的锁页内存，允许 GPU 直接访问这块主机内存。
• 特性：
  • 这块内存同时映射到 CPU 和 GPU 地址空间。
  • GPU 可以通过特定设备指针直接访问主机内存，实现零拷贝（Zero-Copy）。
  • 减少了显存占用和显存间的显式数据拷贝，但访问速度受限于 PCIe 带宽。
• 使用场景：
  • 适合小数据量、对延迟敏感、不想显式拷贝的场景。
  • 需要调用 cudaHostGetDevicePointer() 获取对应的 GPU 设备指针。
  • 零拷贝场景。

4.4 几种锁页内存总结

4.5 cudaHostAllocDefault补充说明

这个cudaHostAllocDefault也是比较常用的一个flag.

当我们在host端申请内存，而我们可能需要再host和device相互之间memcpy这块内存的时候，用这个申请内存要比用malloc申请的内存更快。

因为 malloc() 分配的是 pageable memory（可分页内存）：

• 操作系统可以把它 swap 到磁盘；
• 在进行 cudaMemcpy() 时，驱动需要：
  1. 临时创建一块锁页缓冲区；
  2. 先从 malloc 的内存拷贝到临时锁页内存；
  3. 再拷贝到 GPU 显存；
• 整个过程是 双拷贝 + page fault 风险，速度较慢。

而 cudaHostAllocDefault() 分配的是 pinned memory（锁页内存）：

• 操作系统保证这块内存 不会被分页；
• 可以被 CUDA 驱动直接用于 DMA（直接内存访问）拷贝；
• 是真正的 单次、高速拷贝。

4.6 cudaMallocHost

cudaError_t cudaMallocHost(void ** devPtr,size_t count)
cudaError_t cudaFreeHost(void *ptr)

一句话总结，cudaMallocHost相当于是cudaHostAlloc的第三个参数选cudaHostAllocDefault。

4.7 零拷贝内存

    /**
    * Allocate ZeroCopy mapped memory, shared between CUDA and CPU.
    *
    * @note although two pointers are returned, one for CPU and GPU, they both resolve to the same physical memory.
    *
    * @param[out] cpuPtr Returned CPU pointer to the shared memory.
    * @param[out] gpuPtr Returned GPU pointer to the shared memory.
    * @param[in] size Size (in bytes) of the shared memory to allocate.
    *
    * @returns `0` if the allocation succeeded, otherwise faield.
    * @ingroup cudaMemory
    */
    int cudaAllocMapped(void** cpuPtr, void** gpuPtr, size_t size) {
        if (!cpuPtr || !gpuPtr || size == 0)
            return -1;
 
        CUDA_SAFECALL(cudaHostAlloc(cpuPtr, size, cudaHostAllocMapped), "cudaHostAlloc failed", -1);
        CUDA_SAFECALL(cudaHostGetDevicePointer(gpuPtr, *cpuPtr, 0), "cudaHostGetDevicePointer failed", -1);
 
        memset(*cpuPtr, 0, size);
        VLOG(3) << "[InferServer] cudaAllocMapped " << size << " bytes, CPU " << *cpuPtr << " GPU " << *gpuPtr;
        return 0;
    }
 
 
    /**
    * Allocate ZeroCopy mapped memory, shared between CUDA and CPU.
    *
    * @note this overload of cudaAllocMapped returns one pointer, assumes that the
    *       CPU and GPU addresses will match (as is the case with any recent CUDA version).
    *
    * @param[out] ptr Returned pointer to the shared CPU/GPU memory.
    * @param[in] size Size (in bytes) of the shared memory to allocate.
    *
    * @returns `0` if the allocation succeeded, otherwise failed.
    * @ingroup cudaMemory
    */
    int cudaAllocMapped(void** ptr, size_t size) {
        void* cpuPtr{};
        void* gpuPtr{};
 
        if (!ptr || size == 0)
            return cudaErrorInvalidValue;
 
        auto error = cudaAllocMapped(&cpuPtr, &gpuPtr, size);
        if (error != cudaSuccess)
            return error;
 
        CUDA_SAFECALL(cpuPtr != gpuPtr, "cudaAllocMapped() - addresses of CPU and GPU pointers don't match", cudaErrorMemoryAllocation);
 
        *ptr = gpuPtr;
        return cudaSuccess;
    }

这个零拷贝内存其实就是在申请锁页内存的基础上，用cudaHostGetDevicePointer获取了跟锁页内存对应的GPU设备内存指针。cpuPtr 和 gpuPtr 实际上指向的是同一块物理内存。这是通过CUDA的统一虚拟寻址（Unified Virtual Addressing, UVA）实现的。

4.7.1 补充说明：ZeroCopy 的注意事项

• cudaHostAllocMapped 产生的内存虽然 CPU 和 GPU 都能访问，但仍然是主机内存，GPU 访问时通过 PCIe 远程访问(不是cuda驱动自动背后memcpy)，性能远不如显存；
• 即使有了 UVA（统一虚拟寻址），也仍需调用 cudaHostGetDevicePointer() 获取 GPU 可访问的地址；
• 不适合大数据频繁访问，用在小数据共用、高效开发场景更好。

4.8 malloc、cudaHostAllocDefault()、cudaHostAllocMapped() 对比

5 统一内存（Unified Memory） cudaMallocManaged

统一内存是从 CUDA 6.0 引入的一项机制，其核心目标是：

✅ 简化内存管理 —— 让 CPU 和 GPU 使用同一个指针访问数据，CUDA 运行时自动在主机和设备之间迁移数据，无需手动调用 cudaMemcpy。

特点：

• 使用 cudaMallocManaged() 分配的托管内存，可以被 CPU 和 GPU 共同访问；
• 背后会在 CPU/GPU 之间 自动分页迁移（通过页错误机制），无需手动拷贝；
• 依赖于 UVA（统一虚拟地址）实现统一指针；
• 内存不再需要分别分配 host 和 device 内存再同步内容；
• 更适合新手开发、代码更简洁，但有时性能不如手动拷贝。

6 汇总比较

补充说明

• ✅ 锁页内存（Pinned Memory）：不能被操作系统 swap，提高了传输效率。
• ⚠️ ZeroCopy：虽然不需要显式拷贝，但实际通过 PCIe 总线远程访问，延迟和带宽都劣于显存。
• ✅ Unified Memory：托管内存在访问时由 CUDA 运行时系统自动分页迁移，适合开发快速验证，性能不易控制。
• cudaMemcpy()：适用于大数据高吞吐传输，配合显存使用效率最佳。