为什么CUDA固定内存如此快速？

Question

当我使用钉住内存的方式进行CUDA数据传输时，我观察到数据传输速度大幅提高。在Linux中，实现这一点的底层系统调用是mlock。从mlock的手册页中可以看到，锁定页面可防止它被交换出：

mlock()锁定从addr开始并持续len字节的地址范围内的页面。包含指定地址范围部分的所有页面在调用成功返回时都将保证驻留在RAM中;

在我的测试中，我的系统上有几个GB的空闲内存，因此从未存在过内存页面会被交换出的风险，但我仍然观察到速度提升。请问有人能解释这里的真正原因吗？非常感谢任何洞见或信息。

- Gearoid Murphy

---

你有测量mlock本身的时间吗？ - osgx

不，执行mlock调用所需的实际时间被认为是可以忽略不计的（如果这是你要问的）。真正的开销在于实际数据传输，而在我的算法中，这占总周期时间的相当大一部分。 - Gearoid Murphy

你的CPU是什么？也许，启用NUMA节点将无法从简单的mlock()中受益。 - osgx

AMD Phenom(tm) II X4 970 处理器 - Gearoid Murphy

Answer 1

CUDA驱动程序检查内存范围是否已锁定，然后使用不同的代码路径。锁定的内存存储在物理内存（RAM）中，因此设备可以在没有CPU帮助的情况下获取它（DMA，又称异步复制；设备只需要物理页面列表）。未锁定的内存可能会在访问时生成页面故障，并且它不仅存储在内存中（例如，它可以在交换中），因此驱动程序需要访问非锁定内存的每个页面，将其复制到固定缓冲区并将其传递给DMA（同步，逐页复制）。

如http://forums.nvidia.com/index.php?showtopic=164661所述

通过cudaMallocHost或cudaHostAlloc锁定异步mem复制调用使用的主机内存。我还建议在developer.download.nvidia.com上查看cudaMemcpyAsync和cudaHostAlloc手册。HostAlloc表示cuda驱动程序可以检测到固定内存：

该函数(cudaHostAlloc)分配的虚拟内存范围由驱动程序跟踪，并自动加速对诸如cudaMemcpy()之类的函数的调用。

- osgx

---

我想知道当另一个线程尝试在发出异步复制命令后解锁页面时，会造成多大的混乱？ - Zan Lynx

Zan Lynx，有趣的问题。你为什么想要解锁这个内存？即使在32位PC上，可能会有2-4GB的内存被锁定，当PCI-express卡访问64位（实际上是40或48位）寻址时，更多的内存会被锁定。购买更多的内存比支付经验丰富的（在SO上有18k声望！）程序员要便宜得多。在Linux中，我认为munlock将被阻止或返回错误，并且不会对系统造成任何损害。 - osgx

我可以将cudaHostRegister应用于内存映射文件的指针吗？ - Tomilov Anatoliy

Answer 2

CUDA使用DMA将固定内存传输到GPU。由于页面式主机内存可能驻留在磁盘上，因此无法与DMA一起使用可页面化的主机内存。 如果内存没有被固定（即页面锁定），则首先将其复制到一个页面锁定的“暂存”缓冲区，然后通过DMA复制到GPU。 因此，使用固定内存可以节省从可页面化主机内存复制到页面锁定主机内存的时间。

- Shen Yang

Answer 3

如果内存页尚未被访问过，它们很可能根本没有被交换进来。特别是，新分配的页面将是通用“零页”的虚拟副本，在写入之前不会有物理实例化。与此类似，磁盘上的文件映射也将纯粹保留在磁盘上，直到它们被读取或写入。

- R.. GitHub STOP HELPING ICE

---

我认为这不是一个问题（在我早期回答中写过），因为这是一个真正的程序，而且mlock()在程序中运行得很快（请查看问题的第二条评论）。 - osgx

Answer 4

关于将非锁定页面复制到锁定页面的详细说明。

如果在具有有限CPU RAM的繁忙系统上，操作系统将非锁定页面交换出去，那么代价可能非常昂贵。然后，将通过昂贵的磁盘IO操作触发页面故障以将页面加载到CPU RAM中。

在CPU RAM很珍贵的系统上，固定页面也可能导致虚拟内存抖动。如果发生抖动，CPU的吞吐量会大大降低。

- Izana