Bridge++  Ver.2.1.3
index_eo_alt_openacc-inc.h
Go to the documentation of this file.
1 
10 #include "lib_alt_Accel/inline/define_params.h"
11 #include "lib_alt_Accel/inline/define_index.h"
12 
13 //====================================================================
14 void split(real_t* RESTRICT ve, real_t* RESTRICT vo,
15  real_t* RESTRICT w,
16  int ieo_origin, int nin, int* Nsize)
17 {
18  int Nx = Nsize[0];
19  int Ny = Nsize[1];
20  int Nz = Nsize[2];
21  int Nt = Nsize[3];
22 
23  int Nvol = Nx * Ny * Nz * Nt;
24  int Nvol2 = Nvol/2;
25 
26  int Nvol_pad = CEIL_NWP(Nvol);
27  int Nvol2_pad = CEIL_NWP(Nvol2);
28 
29  int nv = nin * Nvol_pad;
30  int nv2 = nin * Nvol2_pad;
31 
32 #pragma acc data present(ve[0:nv2], vo[0:nv2], w[0:nv]), \
33  copyin(ieo_origin, nin, Nx, Ny, Nz, Nt, Nvol, Nvol2_pad)
34  {
35 
36 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
37  {
38 #pragma acc loop gang worker vector
39  for(int ist2 = Nvol2_pad - NWP; ist2 < Nvol2_pad; ++ist2){
40  for(int in = 0; in < nin; ++in){
41  ve[IDX2(nin, in, ist2)] = 0.0;
42  vo[IDX2(nin, in, ist2)] = 0.0;
43  }
44  }
45  }
46 
47 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
48  {
49 #pragma acc loop gang worker vector
50  for(int ist = 0; ist < Nvol; ++ist){
51  int ix = ist % Nx;
52  int iy = (ist/Nx) % Ny;
53  int iz = (ist/(Nx * Ny)) % Nz;
54  int it = ist/(Nx * Ny * Nz);
55  int ieo = (ix + iy + iz + it + ieo_origin) % 2;
56  int ist2 = ist/2;
57  if(ieo == 0){
58  for(int in = 0; in < nin; ++in){
59  ve[IDX2(nin, in, ist2)] = w[IDX2(nin, in, ist)];
60  }
61  }else{
62  for(int in = 0; in < nin; ++in){
63  vo[IDX2(nin, in, ist2)] = w[IDX2(nin, in, ist)];
64  }
65  }
66  }
67  }
68 
69  } // acc data
70 
71 }
72 
73 //====================================================================
74 void merge(real_t* RESTRICT v,
75  real_t* RESTRICT we, real_t* RESTRICT wo,
76  int ieo_origin, int nin, int* Nsize)
77 {
78  int Nx = Nsize[0];
79  int Ny = Nsize[1];
80  int Nz = Nsize[2];
81  int Nt = Nsize[3];
82 
83  int Nvol = Nx * Ny * Nz * Nt;
84  int Nvol2 = Nvol/2;
85 
86  int Nvol_pad = CEIL_NWP(Nvol);
87  int Nvol2_pad = CEIL_NWP(Nvol2);
88 
89  int nv = nin * Nvol_pad;
90  int nv2 = nin * Nvol2_pad;
91 
92 #pragma acc data present(we[0:nv2], wo[0:nv2], v[0:nv]), \
93  copyin(ieo_origin, nin, Nx, Ny, Nz, Nt, Nvol)
94 #pragma acc parallel num_workers(NUM_WORKERS) vector_length(VECTOR_LENGTH)
95  {
96 
97 #pragma acc loop gang worker vector
98  for(int ist = 0; ist < Nvol_pad; ++ist){
99  int ix = ist % Nx;
100  int iy = (ist/Nx) % Ny;
101  int iz = (ist/(Nx * Ny)) % Nz;
102  int it = ist/(Nx * Ny * Nz);
103  int ieo = (ix + iy + iz + it + ieo_origin) % 2;
104  int ist2 = ist/2;
105  if(ist < Nvol){
106  if(ieo == 0){
107  for(int in = 0; in < nin; ++in){
108  v[IDX2(nin, in, ist)] = we[IDX2(nin, in, ist2)];
109  }
110  }else{
111  for(int in = 0; in < nin; ++in){
112  v[IDX2(nin, in, ist)] = wo[IDX2(nin, in, ist2)];
113  }
114  }
115  }else{
116  for(int in = 0; in < nin; ++in){
117  v[IDX2(nin, in, ist)] = 0.0;
118  }
119  }
120  }
121  }
122 
123 }
124 
125 //============================================================END=====
merge
void merge(real_t *RESTRICT v, real_t *RESTRICT we, real_t *RESTRICT wo, int ieo_origin, int nin, int *Nsize)
Definition: index_eo_alt_openacc-inc.h:74
iy
int iy
Definition: mult_Wilson_xyz_openacc-inc.h:239
CEIL_NWP
#define CEIL_NWP(nst)
Definition: define_params.h:47
ix
int ix
Definition: mult_Wilson_xyz_openacc-inc.h:16
split
void split(real_t *RESTRICT ve, real_t *RESTRICT vo, real_t *RESTRICT w, int ieo_origin, int nin, int *Nsize)
Definition: index_eo_alt_openacc-inc.h:14
it
int it
Definition: mult_Wilson_xyz_openacc-inc.h:461
real_t
double real_t
Definition: bridgeACC_AField_double.cpp:14
IDX2
#define IDX2(nin, in, ist)
Definition: define_index.h:28
define_index.h
NWP
#define NWP
Definition: define_params.h:32
iz
int iz
Definition: mult_Wilson_xyz_openacc-inc.h:462
define_params.h