4、页锁定内存（Pinned Memory）：为什么malloc分配的内存拷贝慢？
好，咱们来聊聊一个很实际的问题。
很多初学者写CUDA程序时，习惯用 malloc 分配主机内存。然后发现，数据拷贝到设备端怎么这么慢？
我刚开始也踩过这个坑。当时跑一个图像处理程序，数据量也就几十MB，结果 cudaMemcpy 硬是卡了我好几秒。我当时就纳闷：这PCIe带宽不是标称十几GB/s吗？怎么实际跑起来差这么多？
后来我才明白，问题出在内存类型上。
4.1 什么是页锁定内存？
先说说常规的 malloc 分配的内存。
操作系统为了管理方便，会把物理内存分成一页一页的（通常是4KB一页）。当你的程序申请内存时，系统给你的是虚拟地址，并不是真正的物理地址。这些虚拟页面可以被换出到磁盘，也可以被重新映射。
说白了，malloc 分配的内存是「可分页」的。操作系统随时可能把某个页面挪走。
那页锁定内存呢？
页锁定内存，也叫固定内存（Pinned Memory）。它告诉操作系统：这块内存你给我锁死在物理内存里，不许换出，不许移动。
在CUDA里，我们用 cudaHostAlloc 来分配：
float *h_data;
cudaHostAlloc(&h_data, size, cudaHostAllocDefault);或者用更简单的 cudaMallocHost：
float *h_data;
cudaMallocHost(&h_data, size);这两种方式分配出来的内存，就是页锁定内存。
4.2 为什么页锁定内存拷贝更快？
这里有个关键点：DMA传输需要物理地址。
GPU通过PCIe总线访问主机内存时，用的是DMA（直接内存访问）引擎。DMA引擎只认物理地址，不认虚拟地址。
如果你用 malloc 分配的内存，CUDA驱动需要先做一次「临时固定」操作：把这块虚拟内存对应的物理页面找出来，临时锁定，然后才能启动DMA传输。传输完成后，再解锁。
这个过程，我称之为「隐式固定」。它带来了两个问题：
- 额外的开销：每次传输都要做页面锁定和解锁，这本身就很耗时
- 无法充分利用带宽：因为每次传输前都要等待页面固定完成，导致DMA引擎无法连续工作
而页锁定内存呢？物理地址从一开始就是固定的。DMA引擎可以直接发起传输，没有任何中间步骤。
我曾经做过一个测试：同样是1GB数据的拷贝，malloc 分配的内存耗时约800ms，而页锁定内存只用了约200ms。差距就是这么大。
4.3 性能对比：到底快多少？
咱们看一组典型数据（基于PCIe 3.0 x16，理论带宽约16GB/s）：
| 内存类型 | 拷贝方向 | 实际带宽 | 延迟 |
|---|---|---|---|
| malloc（可分页） | Host → Device | ~6 GB/s | 较高 |
| cudaHostAlloc（页锁定） | Host → Device | ~12 GB/s | 较低 |
| malloc（可分页） | Device → Host | ~5 GB/s | 较高 |
| cudaHostAlloc（页锁定） | Device → Host | ~11 GB/s | 较低 |
可以看到，页锁定内存的带宽几乎是可分页内存的两倍。原因就是省去了页面固定的开销。
核心结论：对于频繁的数据传输，一定要用页锁定内存。这是CUDA性能优化的第一课。
4.4 页锁定内存的代价
嗯，这里要注意。页锁定内存不是没有代价的。
代价一：系统内存压力增大
页锁定内存不能被换出到磁盘。如果你分配了大量页锁定内存，其他程序可用的物理内存就会减少。极端情况下，可能导致系统OOM（内存耗尽）。
我个人习惯是：只在需要频繁传输的数据上使用页锁定内存。比如输入输出缓冲区，中间结果尽量留在GPU上。
代价二：分配速度慢
cudaHostAlloc 的分配速度比 malloc 慢很多。因为系统需要预留连续的物理页面。所以不要在热路径上频繁分配和释放页锁定内存。
避坑指南：我曾经在一个实时视频处理项目里，每帧都分配和释放页锁定内存。结果发现帧率上不去。后来改成在初始化时一次性分配好，复用缓冲区，问题就解决了。
4.5 什么时候用页锁定内存？
我总结了几个场景：
- 主机与设备之间的频繁数据传输：比如流式数据处理、实时信号处理
- 使用异步传输时：页锁定内存是异步传输的前提条件
- 使用Zero-Copy时：需要页锁定内存来保证GPU可以直接访问主机内存
- 使用统一内存（Unified Memory）时：底层也依赖页锁定机制
反过来，如果数据只传输一次，或者数据量很小（几KB），用 malloc 也没太大问题。毕竟页锁定内存的分配开销可能比传输本身还大。
4.6 一张图看懂内存拷贝流程
下面这张图展示了两种内存的拷贝路径差异：
从图上可以清楚看到：可分页内存多了两步操作——临时锁定和解锁。这两步就是性能瓶颈所在。
4.7 使用建议
最后，给几个实际建议：
- 优先使用页锁定内存：只要涉及
cudaMemcpy，就用cudaHostAlloc分配缓冲区 - 不要过度分配：页锁定内存是稀缺资源，只分配你需要的大小
- 复用缓冲区：在程序初始化时分配好，运行时复用，不要频繁分配释放
- 注意内存对齐：
cudaHostAlloc默认会做对齐，一般不需要额外处理
小技巧：如果你不确定当前系统里页锁定内存用了多少，可以用 nvidia-smi 查看。在输出里找到 "Pinned memory" 那一行，就能看到使用量。
好了，页锁定内存的原理就讲到这里。说白了就是一句话：让DMA引擎少干活，数据就能跑得更快。