diff --git a/tools/fregrid/conserve_interp.c b/tools/fregrid/conserve_interp.c index 50808d39..45028fdc 100644 --- a/tools/fregrid/conserve_interp.c +++ b/tools/fregrid/conserve_interp.c @@ -175,11 +175,12 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles lat_out_min_list[0:nx2*ny2], lat_out_max_list[0:nx2*ny2], \ lon_out_min_list[0:nx2*ny2], lon_out_max_list[0:nx2*ny2], \ lon_out_avg[0:nx2*ny2], n2_list[0:nx2*ny2],area_out2[0:nx2*ny2]) - +#ifdef _OPENACC get_list_inout(nx_out, ny_out, grid_out[n].lonc, grid_out[n].latc, lon_out_min_list, lon_out_max_list, lat_out_min_list, lat_out_max_list, lon_out_avg, lon_out_list, lat_out_list, area_out2, n2_list); +#endif nx_in=grid_in[0].nx; ny_in=grid_in[0].ny; @@ -229,11 +230,13 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles #pragma acc enter data copyin(mask[0:nx_in*ny_in]) #pragma acc enter data copyin(grid_in[m].lonc[0:(nx_in+1)*(ny_in+1)], grid_in[m].latc[0:(nx_in+1)*(ny_in+1)]) +#ifdef _OPENACC nxgrid_in = pre_create_xgrid_2dx2d_order2(&nx_in, &ny_now, &nx_out, &ny_out, grid_in[m].lonc+jstart*(nx_in+1), grid_in[m].latc+jstart*(nx_in+1), grid_out[n].lonc, grid_out[n].latc, mask, lon_out_list, lat_out_list, lat_out_min_list, lat_out_max_list, lon_out_min_list, lon_out_max_list, lon_out_avg, n2_list, area_out2, counts_per_ij, ij_start, ij_end); +#endif printf("%d ", nxgrid_in); @@ -253,13 +256,19 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles time_nxgrid = (time_end - time_start)/CLOCKS_PER_SEC; printf("time=%f ", time_nxgrid); - nxgrid = create_xgrid_2dx2d_order2(&nx_in, &ny_now, &nx_out, &ny_out, &nxgrid_in, +#ifdef _OPENACC + nxgrid = create_xgrid_2dx2d_order2_gpu(&nx_in, &ny_now, &nx_out, &ny_out, &nxgrid_in, grid_in[m].lonc+jstart*(nx_in+1), grid_in[m].latc+jstart*(nx_in+1), grid_out[n].lonc, grid_out[n].latc, mask, lon_out_list, lat_out_list, lat_out_min_list, lat_out_max_list, lon_out_min_list, lon_out_max_list, lon_out_avg, n2_list, area_out2, counts_per_ij, ij_start, ij_end, i_in, j_in, i_out, j_out, xgrid_area, xgrid_clon, xgrid_clat); +#else + nxgrid = create_xgrid_2dx2d_order2_cpu(&nx_in, &ny_now, &nx_out, &ny_out, grid_in[m].lonc+jstart*(nx_in+1), + grid_in[m].latc+jstart*(nx_in+1), grid_out[n].lonc, grid_out[n].latc, + mask, i_in, j_in, i_out, j_out, xgrid_area, xgrid_clon, xgrid_clat); +#endif #pragma acc exit data copyout(i_in[0:nxgrid_in], j_in[0:nxgrid_in], i_out[0:nxgrid_in], j_out[0:nxgrid_in], \ xgrid_area[0:nxgrid_in], xgrid_clon[0:nxgrid_in], xgrid_clat[0:nxgrid_in]) diff --git a/tools/libfrencutils/create_xgrid.c b/tools/libfrencutils/create_xgrid.c index 1c9568b2..4862804a 100644 --- a/tools/libfrencutils/create_xgrid.c +++ b/tools/libfrencutils/create_xgrid.c @@ -873,13 +873,6 @@ nxgrid = 0; };/* get_xgrid_2Dx2D_order1 */ -/******************************************************************************** - void create_xgrid_2dx1d_order2 - This routine generate exchange grids between two grids for the second order - conservative interpolation. nlon_in,nlat_in,nlon_out,nlat_out are the size of the grid cell - and lon_in,lat_in, lon_out,lat_out are geographic grid location of grid cell bounds. - mask is on grid lon_in/lat_in. -********************************************************************************/ #ifdef _OPENACC //MKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKL @@ -1098,7 +1091,7 @@ void get_list_inout( const int nx, const int ny, const double *lon_inout, const } } - +/// Set up for openacc(gpu) variant of create_xgrid_2dx2d_order2 int pre_create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const int *nlon_out, const int *nlat_out, const double *lon_in, const double *lat_in, const double *lon_out, const double *lat_out, const double *mask_in, @@ -1228,7 +1221,15 @@ int pre_create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const //MKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKLMKL -int create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const int *nlon_out, const int *nlat_out, const int *nxgrid_in, +/******************************************************************************** + int create_xgrid_2dx1d_order2 + This routine generate exchange grids between two grids for the second order + conservative interpolation. nlon_in,nlat_in,nlon_out,nlat_out are the size of the grid cell + and lon_in,lat_in, lon_out,lat_out are geographic grid location of grid cell bounds. + mask is on grid lon_in/lat_in. + Uses OpenACC gpu-accelerated algorithm +********************************************************************************/ +int create_xgrid_2dx2d_order2_gpu(const int *nlon_in, const int *nlat_in, const int *nlon_out, const int *nlat_out, const int *nxgrid_in, const double *lon_in, const double *lat_in, const double *lon_out, const double *lat_out, const double *mask_in, const double *lon_out_list, const double *lat_out_list, @@ -1364,7 +1365,8 @@ int create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const int };/* get_xgrid_2Dx2D_order2 */ #else -int create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const int *nlon_out, const int *nlat_out, +/// Original openmp accelerated algorithm for creating 2dx2d exchange grid +int create_xgrid_2dx2d_order2_cpu(const int *nlon_in, const int *nlat_in, const int *nlon_out, const int *nlat_out, const double *lon_in, const double *lat_in, const double *lon_out, const double *lat_out, const double *mask_in, int *i_in, int *j_in, int *i_out, int *j_out, double *xgrid_area, double *xgrid_clon, double *xgrid_clat) diff --git a/tools/libfrencutils/create_xgrid.h b/tools/libfrencutils/create_xgrid.h index f0c750d0..e10440ee 100644 --- a/tools/libfrencutils/create_xgrid.h +++ b/tools/libfrencutils/create_xgrid.h @@ -83,6 +83,7 @@ int pre_create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const const double *lon_out_avg, const int *n2_list, const double *area_out, int *counts_per_ij, int *ij_start, int *ij_end); +#ifdef _OPENACC int create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const int *nlon_out, const int *nlat_out, const int *nxgrid_in, const double *lon_in, const double *lat_in, const double *lon_out, const double *lat_out, const double *mask_in, @@ -92,6 +93,12 @@ int create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const int const double *lon_out_avg, const int *n2_list, const double *area_out, const int *counts_per_ij, const int *ij_start, const int *ij_end, int *i_in, int *j_in, int *i_out, int *j_out, double *xgrid_area, double *xgrid_clon, double *xgrid_clat); +#else +int create_xgrid_2dx2d_order2(const int *nlon_in, const int *nlat_in, const int *nlon_out, const int *nlat_out, + const double *lon_in, const double *lat_in, const double *lon_out, const double *lat_out, + const double *mask_in, int *i_in, int *j_in, int *i_out, int *j_out, + double *xgrid_area, double *xgrid_clon, double *xgrid_clat); +#endif int clip_2dx2d_great_circle(const double x1_in[], const double y1_in[], const double z1_in[], int n1_in, const double x2_in[], const double y2_in[], const double z2_in [], int n2_in,