        kx_offset = 0;
        for (int i = 0; i < kernels[rank].x; i++) {
             swe_from_compressible<<<blocks[k_offset + j * kernels[rank].x + i], threads[k_offset + j * kernels[rank].x + i]>>>(q_fd, qf_swe, nxf, nyf, nz, gamma_up, rho, gamma, kx_offset, ky_offset, p_floor);
+
             kx_offset += blocks[k_offset + j * kernels[rank].x + i].x *
                 threads[k_offset + j * kernels[rank].x + i].x - 2*ng;
        }

    }

    for (int i = k_offset; i < cumulative_kernels[rank]; i++) {
        printf("blocks: (%i, %i) , threads: (%i, %i)\n",
               blocks[i].x, blocks[i].y,
               threads[i].x, threads[i].y);
    }

    // gpu variables

            }

            cudaMemcpy(Uc_d, Uc_h, nx*ny*nlayers*3*sizeof(float), cudaMemcpyHostToDevice);

            // update old_phi
            for (int i = 0; i < nlayers*nx*ny; i++) {
                pphi[i] = Uc_h[i*3];
            }
            cudaMemcpy(old_phi_d, pphi, nx*ny*nlayers*sizeof(float), cudaMemcpyHostToDevice);


            cout << "\nCoarse grid after rk3\n\n";
            for (int y = 0; y < ny; y++) {
