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$ – Kyle Mandli Sep 19 '19 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$ – hertzsprung Sep 20 '19 at 10:05

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service, privacy policy and cookie policy

Browse other questions tagged or ask your own question.