             }
         }
         // bottom-most process
-        MPI_Issend(ysbuf, nlayers*nx*ng*4, MPI_FLOAT, rank-1, tag, comm);
+        MPI_Issend(ysbuf, nlayers*nx*ng*4, MPI_FLOAT, rank-1, tag, comm, &request);
         MPI_Recv(yrbuf, nlayers*nx*ng*4, MPI_FLOAT, rank-1, tag, comm, &status);
         MPI_Wait(&request, &status);
 

       ky_offset += blocks[j * kernels[rank].x].y * threads[j * kernels[rank].x].y;
    }

    ky_offset = 0;

    for (int j = 0; j < kernels[rank].y; j++) {
       kx_offset = kernels[0].x*rank;
       for (int i = 0; i < kernels[rank].x; i++) {
           evolve_fv_fluxes<<<blocks[j * kernels[rank].x + i],
                              threads[j * kernels[rank].x + i]>>>(

    Evolve system through nt timesteps, saving data to filename every dprint timesteps.
    */

    // set up GPU stuff
    int count;
    cudaGetDeviceCount(&count);

    cudaError_t err = cudaGetLastError();

    if (err != cudaSuccess)
        printf("Error: %s\n", cudaGetErrorString(err));

    printf("Found %i CUDA devices\n", count);

    int maxThreads = 256;
    int maxBlocks = 256; //64;

