# Implementing adaptive timestepping in CUDA

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)
{
evolve_flow_variables();
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.