问当“研磨”自身时，Cuda异步MemCpy行为

EN

主机MemCpyAsync在上一份副本完成之前会阻塞吗？

不，通常不是，假设“块”是指阻塞CPU线程。发出给GPU的异步工作项进入队列，控制在工作开始之前立即返回到CPU线程。队列可以容纳“许多”项目。工作的发布可以一直进行，直到队列已满，而没有任何其他障碍或依赖。

记住流语义的两个规则之一是很重要的：

• 在同一流中发出的项按发出顺序执行。B项，在A项之后印发，直到A项完成后才开始。

因此，假设有这样的情况(假设h_ibuff和h_obuff指向固定的主机内存)：

cudaStream_t stream;
cudaStreamCreate(&stream);
for (int i = 0; i < frames; i++){
  cudaMemcpyAsync(d_ibuff, h_ibuff, cudaMemcpyHostToDevice, stream);
  kernel<<<...,stream>>>(...);
  cudaMemcpyAsync(h_obuff, d_obuff, cudaMemcpyDeviceToHost, stream);
  }

在循环的第二次传递中，cudaMemcpyAsync操作将被插入到队列中，但在流语义表明它们可以开始之前，不会开始执行(或执行任何事情)。这个循环发出的每一个op都是这样的。

一个合理的问题可能是“如果在循环的每一次传递中，我都希望h_ibuff中有不同的内容呢?”(相当明智)。那么你就需要具体地解决这个问题。例如，插入一个简单的memcpy操作来“重新加载”h_ibuff是行不通的。你需要某种同步。例如，当内核和后续的h_ibuff D->H操作发生时，您可能会决定要“重新填充”cudaMemcpyAsync。你可以这样做：

cudaStream_t stream;
cudaEvent_t event;
cudaEventCreat(&event);
cudaStreamCreate(&stream);
for (int i = 0; i < frames; i++){
  cudaMemcpyAsync(d_ibuff, h_ibuff, cudaMemcpyHostToDevice, stream);
  cudaEventRecord(event, stream);
  kernel<<<...,stream>>>(...);
  cudaMemcpyAsync(h_obuff, d_obuff, cudaMemcpyDeviceToHost, stream);
  cudaEventSynchronize(event);
  memcpy(h_ibuff, databuff+i*chunksize, chunksize); // "refill"
  }

这种重构将允许发出异步工作，以保持GPU的忙碌，并将副本“重叠”为“填充”h_ibuff。它还将防止“填充”操作从开始到以前的缓冲区内容安全地传输到设备，还将防止下一个缓冲区副本从开始一直到新内容被“重新加载”。

这不是唯一的方法；这是一种可能的方法。

对于上面问/回答的最后一个问题，您可能会问一个类似的问题：“如何处理输出缓冲区端？”这一机制可能非常相似，留给读者。

关于这个主题的结构化学习，您可能希望学习本讲座系列的CUDA并发部分。