         d_vars->fin[x_dst][y_dst][f] = d_vars->fout[x][y][f];
     }
 
+#ifndef COMPUTE_ON_CPU
+    }
+#endif
 }
 
 void print_variables(lbm_vars *d_vars, lbm_vars *h_vars, double var[NX][NY][9]) {

        }
    }

}

#ifndef COMPUTE_ON_CPU
__global__ void lbm_streaming(lbm_vars *d_vars)
#else
void lbm_streaming(lbm_vars *d_vars, int x, int y)
#endif
{
#ifndef COMPUTE_ON_CPU
    int y = threadIdx.x / NX;
    int x = threadIdx.x % NX;
#endif

    // Streaming step

    HANDLE_ERROR(cudaMemcpy(d_vars, h_vars, sizeof(lbm_vars), cudaMemcpyHostToDevice));

    pgm_image* pgm = pgm_create(NX, NY);

#ifndef COMPUTE_ON_CPU
    dim3 dimBlock(1);
    dim3 dimGrid(getThreads(NX, NY));
#endif

    for (int time = 0; time < max_iter; time++) {
        RUN_KERNEL_1D(lbm_right_wall,             NY, d_vars);
