试图在 CUDA 中理解为什么零拷贝也通过 PCIe 传输会更快？答案

【问题标题】：Trying to understand in CUDA why is zero copy faster if it also travel through PCIe?试图在 CUDA 中理解为什么零拷贝也通过 PCIe 传输会更快？
【发布时间】：2021-01-09 00:20:47
【问题描述】：

据说在满足“只读取和/或写入一次”约束的情况下应该使用零拷贝。没关系。

我已经理解这一点，但我的问题是为什么零拷贝首先要快？毕竟，无论我们是通过 cudamemcpy 还是零拷贝使用显式传输，在这两种情况下，数据都必须通过 pci 快速总线传输。或者是否存在任何其他路径（即通过传递设备 RAM 直接在 GPU 寄存器中复制发生？

【问题讨论】：

降低延迟...
由于避免了 cudamemcpy 函数的开销而导致的低延迟？
就数据传输率而言，它并不快。所以这个问题的前提是可疑的。如果主机和设备之间存在传输，则无法避免 PCIE（对于 PCIE 连接的 GPU）。没有其他路径。
如果不是更快，那为什么要使用它？
与使用cudaMemcpy 相比，在许多情况下，零拷贝可能会提高整体应用程序性能。我敢打赌，如果您寻找示例，您可以找到示例。如果您愿意，我可以写一个描述可能 2 个示例的答案。

标签： cuda

【解决方案1】：

纯粹从数据传输速率的角度考虑，我不知道为什么在比较使用零复制方法与使用零复制方法移动数据时，通过 PCIE 在主机和设备之间移动数据的数据传输速率应该有什么不同。使用cudaMemcpy移动它。

但是，这两种操作都有与之相关的开销。我能想到的零拷贝的主要开销来自于主机内存的固定。这具有明显的时间开销（例如，与使用 malloc 或 new 分配相同数量的数据相比）。想到cudaMemcpy 的主要开销是至少几微秒的每次传输开销，这与使用进行传输的底层 DMA 引擎的设置成本相关。

另一个区别是数据的可访问性。固定/零拷贝数据可以在主机和设备之间同时访问，这对于some kinds of communication patterns 很有用，否则cudaMemcpyAsync 会更复杂。

这里有两个相当简单的设计模式，使用零拷贝而不是cudaMemcpy 可能更有意义。

当您拥有大量数据并且不确定需要什么时。假设我们有一个大数据表，比如 1GB，GPU 内核需要访问它。假设内核设计是这样的，每次内核调用只需要表中的一个或几个位置，并且我们不知道这些位置将是哪些位置。我们可以使用cudaMemcpy 将整个 1GB 传输到 GPU。这肯定会奏效，但可能需要很长的时间（例如~0.1s）。还假设我们不知道更新了哪个位置，并且在内核调用之后我们需要访问主机上修改后的数据。将需要另一次转移。在这里使用固定/零拷贝方法将主要消除与移动数据相关的成本，并且由于我们的内核只访问几个位置，因此内核使用零拷贝这样做的成本远低于 0.1 秒。
当您需要检查搜索或收敛算法的状态时。假设我们有一个算法，该算法由一个循环组成，该循环在每次循环迭代中调用一个内核。内核正在执行某种搜索或收敛类型的算法，因此我们需要一个“停止条件”测试。这可能就像一个布尔值一样简单，我们从内核活动与主机通信，以指示我们是否已经到达停止点。如果到达停止点，则循环终止。否则循环继续下一次内核启动。这里甚至可能存在“双向”通信。例如，主机代码可能将布尔值设置为 false。如果迭代需要继续，内核可能会将其设置为 true，但内核不会将该标志设置为 false。因此，如果需要继续，主机代码将标志设置为 false 并再次调用内核。我们可以通过cudaMemcpy 实现这一点：
```
 bool *d_continue;
 cudaMalloc(&d_continue, sizeof(bool));
 bool h_continue = true;
 while (h_continue){
   h_continue = false;
   cudaMemcpy(d_continue, &h_continue, sizeof(bool), cudaMemcpyHostToDevice); 
   my_search_kernel<<<...>>>(..., d_continue);
   cudaMemcpy(&h_continue, d_continue, sizeof(bool), cudaMemcpyDeviceToHost);
 }
```
上述模式应该是可行的，但即使我们只传输少量数据（1 个字节），cudaMemcpy 操作每个都需要大约 5 微秒。如果这是性能问题，我们几乎肯定可以通过以下方式降低时间成本：
```
 bool *z_continue;
 cudaHostAlloc(&z_continue, sizeof(bool), ...);
 *z_continue = true;
 while (*z_continue){
   *z_continue = false;
   my_search_kernel<<<...>>>(..., z_continue);
   cudaDeviceSynchronize();
 }
```

【讨论】：

【解决方案2】：

例如，假设您编写了一个 cuda 加速编辑器算法来修复书籍的拼写错误。如果一个 2MB 的文本数据只有 5 个字节的错误，它只需要编辑它的 5 个字节。因此不需要将整个阵列从 GPU VRAM 复制到系统 RAM。在这里，零拷贝版本将只访问拥有 5 字节字的页面。如果没有零复制，则需要复制整个 2MB 文本。复制 2MB 会比复制 5 个字节（或仅拥有这些字节的页面）花费更多时间，因此会降低书籍/秒的吞吐量。

另一个例子，可以有一个稀疏路径跟踪算法来为游戏场景中的几个小对象添加闪亮的表面。结果可能只需要更新 10-100 像素而不是 1920x1080 像素。零拷贝会更好。

也许稀疏矩阵乘法在零拷贝下效果更好。如果 8192x8192 矩阵相乘，但只有 3-5 个元素是非零的，那么在写入结果时，零拷贝仍然会产生影响。

【讨论】：