XcalableACC: An Integration of XcalableMP and OpenACC
137
Fig. 14 Lattice QCD pseudo code
1 typedef struct Quark {
2 double v[4][3][2];
3 } Quark_t;
4 typedef struct Gluon {
5 double v[3][3][2];
6 } Gluon_t;
7 Quark_t v[NT][NZ][NY][NX], tmp_v[NT][NZ][NY][NX];
8 Gluon_t u[4][NT][NZ][NY][NX];
9
10 #pragma xmp template t[NT][NZ]
11 #pragma xmp nodes n[NODES_T][NODES_Z]
12 #pragma xmp distribute t[block][block] onto n
13 #pragma xmp align v[i][j][*][*] with t[i][j]
14 #pragma xmp align tmp_v[i][j][*][*] with t[i][j]
15 #pragma xmp align u[*][i][j][*][*] with t[i][j]
16 #pragma xmp shadow v[1:1][1:1][0][0]
17 #pragma xmp shadow tmp_v[1:1][1:1][0][0]
18 #pragma xmp shadow u[0][1:1][1:1][0][0]
19 ...
20 int main(){
21 ...
22 #pragma acc enter data copyin(v, tmp_v, u)
Fig. 15 Declaration of distributed arrays for Lattice QCD
each quark and gluon element is affected by its neighboring orthogonal elements.
Note that “0” in the shadow directive means that no halo region exists. In line 22, the
enter data directive transfers the distributed arrays from host memory to accelerator
memory.
Figure 16 shows how to call WD(). The reflect directives are inserted before
WD() in order to update own halo region. In line 2, “1:0” in width clause means
only the lower halo region is updated because only it is needed in WD(). The u is
not updated before the second WD() function because values of u are not updated