1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
|
/* File produced by Kranc */
#define KRANC_C
#include <assert.h>
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "cctk.h"
#include "cctk_Arguments.h"
#include "cctk_Parameters.h"
#include "GenericFD.h"
#include "Differencing.h"
#include "cctk_Loop.h"
#include "loopcontrol.h"
#include "OpenCLRunTime.h"
#include "vectors.h"
/* Define macros used in calculations */
#define INITVALUE (42)
#define ScalarINV(x) ((CCTK_REAL)1.0 / (x))
#define ScalarSQR(x) ((x) * (x))
#define ScalarCUB(x) ((x) * ScalarSQR(x))
#define ScalarQAD(x) (ScalarSQR(ScalarSQR(x)))
#define INV(x) (kdiv(ToReal(1.0),x))
#define SQR(x) (kmul(x,x))
#define CUB(x) (kmul(x,SQR(x)))
#define QAD(x) (SQR(SQR(x)))
extern "C" void WT_CL_RHS_SelectBCs(CCTK_ARGUMENTS)
{
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
CCTK_INT ierr = 0;
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "ML_WaveToy_CL::WT_rhorhs","flat");
if (ierr < 0)
CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_rhorhs.");
ierr = Boundary_SelectGroupForBC(cctkGH, CCTK_ALL_FACES, GenericFD_GetBoundaryWidth(cctkGH), -1 /* no table */, "ML_WaveToy_CL::WT_urhs","flat");
if (ierr < 0)
CCTK_WARN(1, "Failed to register flat BC for ML_WaveToy_CL::WT_urhs.");
return;
}
static void WT_CL_RHS_Body(cGH const * restrict const cctkGH, int const dir, int const face, CCTK_REAL const normal[3], CCTK_REAL const tangentA[3], CCTK_REAL const tangentB[3], int const imin[3], int const imax[3], int const n_subblock_gfs, CCTK_REAL * restrict const subblock_gfs[])
{
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
char const *const source =
"\n"
"/* Include user-supplied include files */\n"
"\n"
"/* Initialise finite differencing variables */\n"
"ptrdiff_t const di = 1;\n"
"ptrdiff_t const dj = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
"ptrdiff_t const dk = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
"ptrdiff_t const cdi = sizeof(CCTK_REAL) * di;\n"
"ptrdiff_t const cdj = sizeof(CCTK_REAL) * dj;\n"
"ptrdiff_t const cdk = sizeof(CCTK_REAL) * dk;\n"
"CCTK_REAL_VEC const dx = ToReal(CCTK_DELTA_SPACE(0));\n"
"CCTK_REAL_VEC const dy = ToReal(CCTK_DELTA_SPACE(1));\n"
"CCTK_REAL_VEC const dz = ToReal(CCTK_DELTA_SPACE(2));\n"
"CCTK_REAL_VEC const dt = ToReal(CCTK_DELTA_TIME);\n"
"CCTK_REAL_VEC const t = ToReal(cctk_time);\n"
"CCTK_REAL_VEC const dxi = INV(dx);\n"
"CCTK_REAL_VEC const dyi = INV(dy);\n"
"CCTK_REAL_VEC const dzi = INV(dz);\n"
"CCTK_REAL_VEC const khalf = ToReal(0.5);\n"
"CCTK_REAL_VEC const kthird = ToReal(1.0/3.0);\n"
"CCTK_REAL_VEC const ktwothird = ToReal(2.0/3.0);\n"
"CCTK_REAL_VEC const kfourthird = ToReal(4.0/3.0);\n"
"CCTK_REAL_VEC const keightthird = ToReal(8.0/3.0);\n"
"CCTK_REAL_VEC const hdxi = kmul(ToReal(0.5), dxi);\n"
"CCTK_REAL_VEC const hdyi = kmul(ToReal(0.5), dyi);\n"
"CCTK_REAL_VEC const hdzi = kmul(ToReal(0.5), dzi);\n"
"\n"
"/* Initialize predefined quantities */\n"
"CCTK_REAL_VEC const p1o12dx = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n"
"CCTK_REAL_VEC const p1o12dy = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n"
"CCTK_REAL_VEC const p1o12dz = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n"
"CCTK_REAL_VEC const p1o144dxdy = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n"
"CCTK_REAL_VEC const p1o144dxdz = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n"
"CCTK_REAL_VEC const p1o144dydz = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n"
"CCTK_REAL_VEC const pm1o12dx2 = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n"
"CCTK_REAL_VEC const pm1o12dy2 = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n"
"CCTK_REAL_VEC const pm1o12dz2 = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dz,dz));\n"
"\n"
"/* Assign local copies of arrays functions */\n"
"\n"
"\n"
"\n"
"/* Calculate temporaries and arrays functions */\n"
"\n"
"/* Copy local copies back to grid functions */\n"
"\n"
"/* Loop over the grid points */\n"
"#pragma omp parallel\n"
"LC_LOOP3VEC(WT_CL_RHS,\n"
" i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n"
" cctk_ash[0],cctk_ash[1],cctk_ash[2],\n"
" CCTK_REAL_VEC_SIZE)\n"
"{\n"
" ptrdiff_t const index = di*i + dj*j + dk*k;\n"
" \n"
" /* Assign local copies of grid functions */\n"
" \n"
" CCTK_REAL_VEC rhoL = vec_load(rho[index]);\n"
" CCTK_REAL_VEC uL = vec_load(u[index]);\n"
" \n"
" \n"
" /* Include user supplied include files */\n"
" \n"
" /* Precompute derivatives */\n"
" CCTK_REAL_VEC const PDstandardNth11u = PDstandardNth11(&u[index]);\n"
" CCTK_REAL_VEC const PDstandardNth22u = PDstandardNth22(&u[index]);\n"
" CCTK_REAL_VEC const PDstandardNth33u = PDstandardNth33(&u[index]);\n"
" \n"
" /* Calculate temporaries and grid functions */\n"
" CCTK_REAL_VEC urhsL = rhoL;\n"
" \n"
" CCTK_REAL_VEC rhorhsL = \n"
" kadd(PDstandardNth11u,kadd(PDstandardNth22u,PDstandardNth33u));\n"
" \n"
" /* Copy local copies back to grid functions */\n"
" vec_store_partial_prepare(i,lc_imin,lc_imax);\n"
" vec_store_nta_partial(rhorhs[index],rhorhsL);\n"
" vec_store_nta_partial(urhs[index],urhsL);\n"
"}\n"
"LC_ENDLOOP3VEC(WT_CL_RHS);\n"
""
;
char const *const groups[] = {
"ML_WaveToy_CL::WT_rho",
"ML_WaveToy_CL::WT_rhorhs",
"ML_WaveToy_CL::WT_u",
"ML_WaveToy_CL::WT_urhs",
NULL};
static struct OpenCLKernel *kernel = NULL;
char const *const sources[] = {differencing, source, NULL};
OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_RHS",
sources, groups, NULL, NULL, NULL, -1,
imin, imax, &kernel);
}
extern "C" void WT_CL_RHS(CCTK_ARGUMENTS)
{
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
if (verbose > 1)
{
CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_RHS_Body");
}
if (cctk_iteration % WT_CL_RHS_calc_every != WT_CL_RHS_calc_offset)
{
return;
}
const char *const groups[] = {
"ML_WaveToy_CL::WT_rho",
"ML_WaveToy_CL::WT_rhorhs",
"ML_WaveToy_CL::WT_u",
"ML_WaveToy_CL::WT_urhs"};
GenericFD_AssertGroupStorage(cctkGH, "WT_CL_RHS", 4, groups);
GenericFD_EnsureStencilFits(cctkGH, "WT_CL_RHS", 2, 2, 2);
GenericFD_LoopOverInterior(cctkGH, WT_CL_RHS_Body);
if (verbose > 1)
{
CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_RHS_Body");
}
}
|
