Skip to content

Commit

Permalink
Merge pull request #249 from mlee03/struct_outlists
Browse files Browse the repository at this point in the history 
Create structures for list arrays and Add OpenACC data transfer directives
  • Loading branch information
@ngs333
ngs333 authored Oct 25, 2023
2 parents f6d1cde + bfd5ce0 commit af6c817
Show file tree
Hide file tree
Showing 8 changed files with 403 additions and 382 deletions.
50 changes: 27 additions & 23 deletions tools/fregrid/conserve_interp.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -84,10 +84,15 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles
//START NTILES_OUT
for(n=0; n<ntiles_out; n++) {

double *lon_out_min_list=NULL, *lon_out_max_list=NULL, *lon_out_avg=NULL;
double *lat_out_min_list=NULL, *lat_out_max_list=NULL;
double *lon_out_list=NULL, *lat_out_list=NULL;
int *n2_list=NULL;
Minmaxavg_lists out_minmaxavg_lists;
out_minmaxavg_lists.lon_list=NULL;
out_minmaxavg_lists.lat_list=NULL;
out_minmaxavg_lists.lon_min_list=NULL;
out_minmaxavg_lists.lat_min_list=NULL;
out_minmaxavg_lists.lon_max_list=NULL;
out_minmaxavg_lists.lat_max_list=NULL;
out_minmaxavg_lists.n_list=NULL;
out_minmaxavg_lists.lon_avg=NULL;

nx_out = grid_out[n].nxc;
ny_out = grid_out[n].nyc;
Expand All @@ -96,21 +101,24 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles
#pragma acc enter data copyin(grid_out[n].lonc[0:(nx_out+1)*(ny_out+1)], \
grid_out[n].latc[0:(nx_out+1)*(ny_out+1)])


//allocate memory for the lists
malloc_minmaxavg_lists(nx_out*ny_out, &lon_out_min_list, &lon_out_max_list,
&lat_out_min_list, &lat_out_max_list, &n2_list,
&lon_out_avg, &lon_out_list, &lat_out_list);
malloc_minmaxavg_lists(nx_out*ny_out, &out_minmaxavg_lists);

#define MAX_V 8
#pragma acc enter data create(lon_out_list[0:MAX_V*nx_out*ny_out], lat_out_list[0:MAX_V*nx_out*ny_out], \
lat_out_min_list[0:nx_out*ny_out], lat_out_max_list[0:nx_out*ny_out], \
lon_out_min_list[0:nx_out*ny_out], lon_out_max_list[0:nx_out*ny_out], \
lon_out_avg[0:nx_out*ny_out], n2_list[0:nx_out*ny_out])
#pragma acc enter data create(out_minmaxavg_lists)
#pragma acc enter data create(out_minmaxavg_lists.lon_list[0:MAX_V*nx_out*ny_out], \
out_minmaxavg_lists.lat_list[0:MAX_V*nx_out*ny_out], \
out_minmaxavg_lists.lon_min_list[0:nx_out*ny_out], \
out_minmaxavg_lists.lon_max_list[0:nx_out*ny_out], \
out_minmaxavg_lists.lat_min_list[0:nx_out*ny_out], \
out_minmaxavg_lists.lat_max_list[0:nx_out*ny_out], \
out_minmaxavg_lists.n_list[0:nx_out*ny_out], \
out_minmaxavg_lists.lon_avg[0:nx_out*ny_out] )


//compute the list values
get_minmaxavg_lists(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,
n2_list, lon_out_avg, lon_out_list, lat_out_list);
get_minmaxavg_lists(nx_out, ny_out, grid_out[n].lonc, grid_out[n].latc, &out_minmaxavg_lists);

//START NTILES_IN
for(m=0; m<ntiles_in; m++) {
Expand Down Expand Up @@ -161,6 +169,7 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles
mask, i_in, j_in, i_out, j_out, xgrid_area);
for(i=0; i<nxgrid; i++) j_in[i] += jstart;
free(mask);
#pragma acc exit data delete(mask)
} //opcode CONSERVE_ORDER1

else if(opcode & CONSERVE_ORDER2) {
Expand All @@ -177,10 +186,8 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles
#ifdef _OPENACC
nxgrid = create_xgrid_2dx2d_order2_acc(&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,
lon_out_min_list, lon_out_max_list, lon_out_avg,
lat_out_min_list, lat_out_max_list, n2_list,
lon_out_list, lat_out_list, mask,
i_in, j_in, i_out, j_out, xgrid_area, xgrid_clon, xgrid_clat);
&out_minmaxavg_lists, mask, i_in, j_in, i_out, j_out,
xgrid_area, xgrid_clon, xgrid_clat);
#else
nxgrid = 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,
Expand Down Expand Up @@ -281,11 +288,8 @@ void setup_conserve_interp(int ntiles_in, const Grid_config *grid_in, int ntiles
malloc_xgrid_arrays(zero, &i_in, &j_in, &i_out, &j_out, &xgrid_area, &xgrid_clon , &xgrid_clat);
#pragma acc exit data delete(grid_in[m].latc, grid_in[m].lonc)
} // ntiles_in
malloc_minmaxavg_lists(zero, &lon_out_min_list, &lon_out_max_list,
&lat_out_min_list, &lat_out_max_list, &n2_list,
&lon_out_avg, &lon_out_list, &lat_out_list);
#pragma acc exit data delete(lon_out_min_list, lon_out_max_list, lat_out_min_list, \
lat_out_max_list, n2_list, lon_out_avg, lon_out_list, lat_out_list)
malloc_minmaxavg_lists(zero, &out_minmaxavg_lists);
#pragma acc exit data delete(out_minmaxavg_lists)
#pragma acc exit data delete(grid_out[n].latc, grid_out[n].lonc)
} // ntimes_out

Expand Down
17 changes: 14 additions & 3 deletions tools/fregrid/globals.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -96,7 +96,7 @@ typedef struct {
char area_name[STRING];
int do_regrid;
int is_axis_data;
int dimsize[5];
int dimsize[5];
} Var_config;

typedef struct {
Expand All @@ -121,11 +121,11 @@ typedef struct {
int bndid;
int size;
nc_type type;
char cart;
char cart;
int bndtype;
int is_defined;
double *bnddata;
double *data;
double *data;
} Axis_config;

typedef struct {
Expand Down Expand Up @@ -236,4 +236,15 @@ typedef struct{
double *f_min;
} Monotone_config;

typedef struct{
double *lat_min_list;
double *lat_max_list;
double *lon_min_list;
double *lon_max_list;
double *lon_avg;
double *lon_list;
double *lat_list;
int *n_list;
} Minmaxavg_lists;

#endif
2 changes: 1 addition & 1 deletion tools/libfrencutils/Makefile.am
View file
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,7 @@ if WITH_OPENACC
libfrencutils_gpu_a_CFLAGS = $(CC_OPENACC_FLAGS) $(AM_CFLAGS)
endif

AM_CFLAGS = $(NETCDF_CFLAGS)
AM_CFLAGS = $(NETCDF_CFLAGS) -I$(top_srcdir)/tools/fregrid

libfrencutils_a_SOURCES = affinity.c \
constant.h \
Expand Down
Loading

0 comments on commit af6c817

Please sign in to comment.