I want to implement a CUDA solver for the 2D shallow water equations using adaptive timestepping with a Courant number fixed by the user. The algorithm pseudocode looks something like this:

while (t < end_time)
  dt = global_minimum_dt_over_all_elements_in_mesh();
  t += dt;

I've seen example CUDA code with fixed timestepping (e.g. lecture 2 slides and code from Boston University) that makes multiple CUDA kernel calls from the time loops on the CPU, and synchronises only at the end:

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

But I don't think this is feasible with adaptive timestepping. Calling cudaDeviceSynchronize() after every timestep seems like a bad idea and, from my own measurements with Boston University's example code, the code could be up to twice as slow.

I'm aware of in-kernel grid synchronisation introduced in CUDA 9 that could allow me to put the entire time loop inside the kernel, but I'd rather be able to support slightly older devices.

  • 3
    $\begingroup$ Not sure if it is helpful but there was a paper published by colleagues that uses adaptivity in both space and time on GPUs that may be of interest: doi.org/10.1029/2019MS001635. Note that there's also an arXiv version as well along with code on github. $\endgroup$ Commented Sep 19, 2019 at 18:26
  • $\begingroup$ @KyleMandli I hadn't seen that paper, it has an excellent overview of how to implement a Godunov-type structured mesh solver in CUDA which will be very helpful. $\endgroup$ Commented Sep 20, 2019 at 10:05


Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Browse other questions tagged or ask your own question.