模拟循环 GPU 利用率

Question

我正在努力使用模拟循环。每个周期启动 3 个内核。下一个时间步长由第二个内核计算。

while (time < end)
{
  kernel_Flux<<<>>>(...);
  kernel_Timestep<<<>>>(d_timestep);
  memcpy(&h_timestep, d_timestep, sizeof(float), ...);  
  kernel_Integrate<<<>>>(d_timestep);
  time += h_timestep;
}

我只需要复制回一个浮点数。避免不必要的同步的最有效方法是什么？

提前谢谢你。 :-)

Answer 1

在 CUDA 中，来自默认流的所有操作运行ning 都是同步的。因此，在您发布的代码中，内核将运行一个接一个。据我所知，内核 kernel_integrate() 取决于内核的结果 kernel_Timestep()，所以无法避免同步。无论如何，如果内核 kernel_Flux() 和 kernel_Timestep() 处理独立的数据，你可以尝试在两个不同的流中并行执行它们。

Answer 2

如果您非常关心迭代时间，您可以设置一个专用于 h_timestep 输出的 memcpy 的新流（在这种情况下您需要使用 cudaMemcpyAsync）。然后使用类似推测执行的方法，在您计算出时间之前循环继续进行。为此，您必须为接下来的几次迭代设置 GPU 内存缓冲区。您可能可以通过使用循环缓冲区来做到这一点。您还需要使用 cudaEventRecord 和 cudaStreamWaitEvent 来同步不同的流，这样只有当时间对应于您将要覆盖的缓冲区时，才允许进行下一次迭代，已计算（ memcpy stream 已经完成了这项工作），因为否则你将失去那次迭代的状态。

另一种可能的解决方案是利用动态并行性，我没有尝试过但我怀疑它会起作用。如果您的显卡支持，您可以将整个循环放在 GPU 中。

编辑：

抱歉，我才知道你有第三个内核。你同步延迟可能是因为你没有做cudaMemcpyAsync？第三个内核很可能运行比 memcpy 长。您应该能够毫不拖延地继续进行。唯一需要的同步是在每次迭代之后。

Answer 3

理想的解决方案是将所有内容都移至 GPU。但是，我不能这样做，因为我需要在每隔几次迭代后启动 CUDPP compact，并且它不支持 CUDA 流或动态并行性。我知道 Thrust 1.8 库有 copy_if 方法，它做同样的事情，并且它正在使用动态并行性。问题是它没有单独编译。

综上所述，现在我使用以下代码：

while (time < end)
{
  kernel_Flux<<<gs,bs, 0, stream1>>>();
  kernel_Timestep<<<gs,bs, 0, stream1>>>(d_timestep);
  cudaEventRecord(event, stream1);
  cudaStreamWaitEvent(mStream2, event, 0);
  memcpyasync(&h_timestep, d_timestep, sizeof(float), ..., stream2);  
  kernel_Integrate<<<>>>(d_timestep);
  cudaStreamSynchronize(stream2);
  time += h_timestep;
}

模拟循环 GPU 利用率

Simulation loop GPU utilization

cuda

utilization