-1

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

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

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

先感谢您。:-)

4

3 回答 3

1

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

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

编辑：

对不起，我刚刚意识到你有第三个内核。您因同步而延迟可能是因为您没有在做什么cudaMemcpyAsync？第三个内核很可能会比 memcpy 运行更长的时间。您应该能够毫不拖延地继续进行。唯一需要的同步是在每次迭代之后。

于 2015-01-31T06:59:08.637 回答
1

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

于 2015-01-30T13:49:24.630 回答
0

理想的解决方案是将所有东西都移到 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;
}
于 2015-02-04T13:22:42.797 回答