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 几种锁页内存总结

标志	特点与说明	使用建议
cudaHostAllocDefault	普通锁页内存，等价于 cudaMallocHost()	最常用，适合常规 H↔D 拷贝
cudaHostAllocPortable	多 context 多 GPU 共享主机内存	多 GPU / 多线程环境
cudaHostAllocWriteCombined	主机只写，优化 CPU→GPU 传输性能，读很慢	图像、音频、传感器流式写缓冲区
cudaHostAllocMapped	支持 ZeroCopy，GPU 可访问主机内存，需要配合 cudaHostGetDevicePointer() 使用	小数据共享、无需频繁 memcpy 场景

4.5 cudaHostAllocDefault补充说明

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

特性	说明
✅ 分配主机锁页（Pinned）内存	比普通 malloc 分配的 pageable memory 更适合 cudaMemcpy
✅ 提高 H2D / D2H 的数据传输速率	DMA 传输，绕过页交换机制，避免内核拷贝中断
✅ 适用于大多数单 GPU、单 context 应用	也是最不容易踩坑的分配方式
✅ 行为与 cudaMallocHost() 完全一致	所以也可以用它来替代后者

当我们在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() 对比

方法	是否锁页内存	GPU是否可直接访问	性能
malloc()	❌ 否	❌ 否	普通 CPU 内存，传输慢
cudaHostAllocDefault()	✅ 是	❌ 否	高效 H2D/D2H 拷贝
cudaHostAllocMapped()	✅ 是	✅ 是（需映射）	可 ZeroCopy，但访问慢

主机内存类型	拷贝方式	带宽性能（相对）
malloc()	cudaMemcpy()	1.0x
cudaHostAlloc()	cudaMemcpy()	🔺 1.5x ~ 2.5x
cudaHostAllocMapped() + 直接访问	ZeroCopy	⚠️ 慢，适合小数据

5 统一内存（Unified Memory） cudaMallocManaged

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

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

特点：

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

6 汇总比较

类型	分配方式	是否锁页内存	是否需 memcpy	GPU 是否直接访问	性能表现	适合场景
普通主机内存	malloc()	❌ 否	✅ 需要	❌ 否	⚠️ 最慢，H2D需拷贝	最普通的内存，不推荐传输用
锁页主机内存	cudaHostAllocDefault() 或 cudaMallocHost()	✅ 是	✅ 需要	❌ 否	✅ 快速拷贝（H2D/D2H）	高效拷贝用，推荐用于传输
零拷贝内存（ZeroCopy）	cudaHostAllocMapped()	✅ 是	❌ 不需要	✅ 是（映射）	⚠️ 延迟高，带宽低	小数据共享、开发阶段调试
统一内存（Unified Memory）	cudaMallocManaged()	✅ 是（托管）	❌ 不需要	✅ 是（自动迁移）	✅ 自动迁移但性能波动	开发方便，复杂数据结构共享等

补充说明

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