我想使用自适应时间步长和用户固定的 Courant 数为 2D 浅水方程实现 CUDA 求解器。算法伪代码如下所示：

while (t < end_time)
{
  evolve_flow_variables();
  dt = global_minimum_dt_over_all_elements_in_mesh();
  t += dt;
}

我已经看到了具有固定时间步长的示例 CUDA 代码（例如，波士顿大学的第 2 课幻灯片和代码），它从 CPU 上的时间循环进行多个 CUDA 内核调用，并且仅在最后同步：

const dt;
while (t < end_time)
{
  evolve_flow_variables<<<grid_size, block_size>>>();
  t += dt;
}
cudaDeviceSynchronize();

但我认为这对于自适应时间步长是不可行的。在每个时间步之后调用cudaDeviceSynchronize()似乎是个坏主意，而且根据我自己对波士顿大学示例代码的测量，代码可能会慢两倍。

我知道 CUDA 9 中引入的内核内网格同步可以让我将整个时间循环放在内核中，但我宁愿能够支持稍旧的设备。