Creating the exchange grid gpu

Creating the exchange grid

The get_upbound_nxcells_2dx2d_gpu function needs to be called before create_xgrid_2dx2d_conserve1_gpu or create_xgrid_2dx2d_conserve2_gpu. Get_upbound_nxcells_2dx2d_gpu provides the size of the arrays that need to be malloc-ed for create_xgrid_2dx2d_gpu; the starting and ending indices for the inner loop execution in create_xgrid_2dx2d_gpu; and array indices for storing exchange grid information. The following outlines the algorithm for get_upbound_nxcells_2dx2d_gpu and create_xgrid_2dx2d_order1. The logic of create_xgrid_2dx2d_order2 is very similar to that in create_xgrid_2dx2d_order1 and contains additional computation for centroid points. Create_xgrid_2dx2d_oder2 is currently being optimized for better performance and is not outlined here.

$${\color{#fc6c85}\large{<><><>}}$$

Get_upbound_nxcells_2dx2d_gpu

$$\color{#fc6c85}\large{\text{\#pragma acc data present/copy/copyin/copyout ...}}$$
$$\color{#fc6c85}\large{\text{\#pragma acc parallel loop reduction(sum:upbound\_nxcells) ...}}$$
for each input cell ij1

• get input cell vertices, area computed from the cell vertices, cell longitudinal “center”, and leftmost (lon_min), rightmost (lon_max), bottomost (lat_min), and topmost (lat_max) coordinate values of the cell.
  
  $$\color{#fc6c85}\large{\text{\#pragma acc parallel loop reduction(+:upbound\_nxcells) reduction(min:ij2\_min)}}$$ $$\color{#fc6c85}\large{\text{\qquad\qquad\qquad\qquad\qquad\qquad\qquad reduction(max:ij2\_max) reduction(+:i\_approx\_xcells\_per\_ij1...)}}$$
  for each output cell ij2

  • get ij2 cell information (same as above) stored in the output_grid_cells struct
  • check if cells overlap (see diagram)
  • if cells overlap
    • upbound_nxcells++
    • i_approx_xcells_per_ij1++ : increment the count of exchange grid cells created from input cell ij1
    • ij2_min = min(ij2_min, ij2: keep track of the first output cell ij2 that overlaps with input cell ij1
    • ij2_max = max(ij2_max, ij: keep track of the last output cell ij2 that overlaps with input cell ij1

• Store information: $$\space$$ approx_nxcells_per_ij1[ij1] = i_approx_xcells_per_ij1, $$\space$$ ij2_start[ij1] = ij2_min, $$\space$$ ij2_end[ij1] = ij2_max
  

return upbound_nxcells;

$${\color{#fc6c85}\large{<><><>}}$$

Create_xgrid_2dx2d_gpu

$$\color{#fc6c85}\large{\text{\#pragma acc data present/copy/copyin/copyout ...}}$$
$$\color{#fc6c85}\large{\text{\#pragma acc parallel loop...}}$$
for each input cell ij1

• int ixcell=0: initialize ixcell that will keep track of the number of exchange grid cells formed from cell ij1
• get input cell information (as listed in get_upbound_nxcells)
• compute approx_nxcells_b4_ij1. This value is approximate to the array index of the first exchange grid cell formed from ij1 and is the cumulative sum of approx_nxgrid_per_ij1 values over all cells before ij1.
  
  $$\color{#fc6c85}\large{\text{\#pragma acc parallel seq...}}$$
  for each output cell ij2 where ij2_start[ij1] $$\le$$ ij2 $$\le$$ ij2_end[ij1]
  • get ij2 cell information from output_grid_cell struct
  • check if cells overlap (see diagram)
  • if cells overlap
    • clip_2dx2d_gpu: get exchange grid cell vertices
    • poly_area_gpu: get exchange grid cell area
    • if the area of the exchange grid cell is larger than AREA_RATIO_THRESH
      • store exchange grid cell area: $$\space$$ store_xcell_area[approx_nxcells_b4_ij1+ixcell] = xcell_area
      • store input parent cell index: $$\space$$ parent_input_index[approx_nxcells_b4_ij1+ixcell] = ij1
      • store output parent cell idnex: $$\space$$ parent_output_index[approx_nxcells_b4_ij1+ixcell] = ij2
      • ixcell++

return nxgrid;

$${\color{#fc6c85}\large{<><><>}}$$

Additional information

The approximated number of exchange grid cells from get_upbound_nxcells_2dx2d_gpu will be an upper bound to the actual nxgrid. Get_upbound_nxcells_2dx2d_gpu only checks to see if the input and output grid cells overlap and does not check to see if the overlapping area meets the minimal overlap threshold. Thus not all exchange grid cells counted in upbound_nxgrid_2dx2d_gpu will be considered an exchange in create_xgrid_2dx2d_gpu and an additional routine called copy_data_to_interp_on_device needs to be called to correctly set the data in Interp.

$${\color{#fc6c85}\large{<><><>}}$$

Diagram illustrating the overlap check as done here

Diagram illustrating how the rotate variable is used here

$${\color{#fc6c85}\large{<><><>}}$$