试图在 CUDA 中理解为什么零复制也通过 PCIe 传输速度更快？

Question

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

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

Answer 1

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

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

另一个区别在于数据的可访问性。固定/零复制数据可在主机和设备之间同时访问，这对于某些类型的通信模式很有用cudaMemcpyAsync，例如，否则会更复杂。

这里有两个相当简单的设计模式，其中使用零复制而不是cudaMemcpy.

1. 当您拥有大量数据并且不确定需要什么时。假设我们有一个很大的数据表，比如 1GB，并且 GPU 内核需要访问它。假设内核设计是这样的，每次内核调用只需要表中的一个或几个位置，我们不知道这些位置将是哪些位置。我们可以用cudaMemcpy将整个 1GB 传输到 GPU。这当然可以工作，但可能需要很长的时间（例如~0.1s）。还假设我们不知道更新了哪个位置，并且在内核调用之后我们需要访问主机上修改后的数据。将需要另一次转移。在这里使用固定/零复制方法将主要消除与移动数据相关的成本，并且由于我们的内核只访问少数位置，内核使用零复制这样做的成本远小于 0.1 秒。

2. 当您需要检查搜索或收敛算法的状态时。假设我们有一个算法，它包含一个循环，该循环在每次循环迭代中调用一个内核。内核正在执行某种搜索或收敛类型的算法，因此我们需要“停止条件”测试。这可能就像一个布尔值一样简单，我们从内核活动返回给主机，以指示我们是否已经到达停止点。如果到达停止点，则循环终止。否则循环继续下一次内核启动。这里甚至可能存在“双向”通信。例如，主机代码可能将布尔值设置为 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();
    }