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
|
/* 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_Dirichlet_SelectBCs(CCTK_ARGUMENTS)
{
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
if (cctk_iteration % WT_CL_Dirichlet_calc_every != WT_CL_Dirichlet_calc_offset)
return;
CCTK_INT ierr CCTK_ATTRIBUTE_UNUSED = 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_Dirichlet_Body(const cGH* restrict const cctkGH, const int dir, const int face, const CCTK_REAL normal[3], const CCTK_REAL tangentA[3], const CCTK_REAL tangentB[3], const int imin[3], const int imax[3], const int n_subblock_gfs, CCTK_REAL* restrict const subblock_gfs[])
{
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
const char* const source =
"\n"
"/* Include user-supplied include files */\n"
"\n"
"/* Initialise finite differencing variables */\n"
"const ptrdiff_t di CCTK_ATTRIBUTE_UNUSED = 1;\n"
"const ptrdiff_t dj CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,1,0) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
"const ptrdiff_t dk CCTK_ATTRIBUTE_UNUSED = CCTK_GFINDEX3D(cctkGH,0,0,1) - CCTK_GFINDEX3D(cctkGH,0,0,0);\n"
"const ptrdiff_t cdi CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * di;\n"
"const ptrdiff_t cdj CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dj;\n"
"const ptrdiff_t cdk CCTK_ATTRIBUTE_UNUSED = sizeof(CCTK_REAL) * dk;\n"
"const CCTK_REAL_VEC dx CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(0));\n"
"const CCTK_REAL_VEC dy CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(1));\n"
"const CCTK_REAL_VEC dz CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_SPACE(2));\n"
"const CCTK_REAL_VEC dt CCTK_ATTRIBUTE_UNUSED = ToReal(CCTK_DELTA_TIME);\n"
"const CCTK_REAL_VEC t CCTK_ATTRIBUTE_UNUSED = ToReal(cctk_time);\n"
"const CCTK_REAL_VEC dxi CCTK_ATTRIBUTE_UNUSED = INV(dx);\n"
"const CCTK_REAL_VEC dyi CCTK_ATTRIBUTE_UNUSED = INV(dy);\n"
"const CCTK_REAL_VEC dzi CCTK_ATTRIBUTE_UNUSED = INV(dz);\n"
"const CCTK_REAL_VEC khalf CCTK_ATTRIBUTE_UNUSED = ToReal(0.5);\n"
"const CCTK_REAL_VEC kthird CCTK_ATTRIBUTE_UNUSED = ToReal(1.0/3.0);\n"
"const CCTK_REAL_VEC ktwothird CCTK_ATTRIBUTE_UNUSED = ToReal(2.0/3.0);\n"
"const CCTK_REAL_VEC kfourthird CCTK_ATTRIBUTE_UNUSED = ToReal(4.0/3.0);\n"
"const CCTK_REAL_VEC keightthird CCTK_ATTRIBUTE_UNUSED = ToReal(8.0/3.0);\n"
"const CCTK_REAL_VEC hdxi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dxi);\n"
"const CCTK_REAL_VEC hdyi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dyi);\n"
"const CCTK_REAL_VEC hdzi CCTK_ATTRIBUTE_UNUSED = kmul(ToReal(0.5), dzi);\n"
"\n"
"/* Initialize predefined quantities */\n"
"const CCTK_REAL_VEC p1o12dx CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dx);\n"
"const CCTK_REAL_VEC p1o12dy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dy);\n"
"const CCTK_REAL_VEC p1o12dz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.0833333333333333333333333333333),dz);\n"
"const CCTK_REAL_VEC p1o144dxdy CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dy,dx));\n"
"const CCTK_REAL_VEC p1o144dxdz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dx));\n"
"const CCTK_REAL_VEC p1o144dydz CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(0.00694444444444444444444444444444),kmul(dz,dy));\n"
"const CCTK_REAL_VEC pm1o12dx2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dx,dx));\n"
"const CCTK_REAL_VEC pm1o12dy2 CCTK_ATTRIBUTE_UNUSED = kdiv(ToReal(-0.0833333333333333333333333333333),kmul(dy,dy));\n"
"const CCTK_REAL_VEC pm1o12dz2 CCTK_ATTRIBUTE_UNUSED = 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"
"const int imin0=imin[0];\n"
"const int imin1=imin[1];\n"
"const int imin2=imin[2];\n"
"const int imax0=imax[0];\n"
"const int imax1=imax[1];\n"
"const int imax2=imax[2];\n"
"#pragma omp parallel // reduction(+: vec_iter_counter, vec_op_counter, vec_mem_counter)\n"
"CCTK_LOOP3STR(WT_CL_Dirichlet,\n"
" i,j,k, imin0,imin1,imin2, imax0,imax1,imax2,\n"
" cctk_ash[0],cctk_ash[1],cctk_ash[2],\n"
" vecimin,vecimax, CCTK_REAL_VEC_SIZE)\n"
"{\n"
" const ptrdiff_t index CCTK_ATTRIBUTE_UNUSED = di*i + dj*j + dk*k;\n"
" // vec_iter_counter+=CCTK_REAL_VEC_SIZE;\n"
" \n"
" /* Assign local copies of grid functions */\n"
" \n"
" \n"
" \n"
" /* Include user supplied include files */\n"
" \n"
" /* Precompute derivatives */\n"
" \n"
" /* Calculate temporaries and grid functions */\n"
" CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED urhsL = ToReal(0);\n"
" \n"
" CCTK_REAL_VEC CCTK_ATTRIBUTE_UNUSED rhorhsL = ToReal(0);\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"
"CCTK_ENDLOOP3STR(WT_CL_Dirichlet);\n"
""
;
const char* const groups[] = {
"ML_WaveToy_CL::WT_rhorhs",
"ML_WaveToy_CL::WT_urhs",
NULL};
static struct OpenCLKernel *kernel = NULL;
const char* const sources[] = {differencing, source, NULL};
OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "WT_CL_Dirichlet",
sources, groups, NULL, NULL, NULL, -1,
imin, imax, &kernel);
}
extern "C" void WT_CL_Dirichlet(CCTK_ARGUMENTS)
{
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
if (verbose > 1)
{
CCTK_VInfo(CCTK_THORNSTRING,"Entering WT_CL_Dirichlet_Body");
}
if (cctk_iteration % WT_CL_Dirichlet_calc_every != WT_CL_Dirichlet_calc_offset)
{
return;
}
const char* const groups[] = {
"ML_WaveToy_CL::WT_rhorhs",
"ML_WaveToy_CL::WT_urhs"};
GenericFD_AssertGroupStorage(cctkGH, "WT_CL_Dirichlet", 2, groups);
GenericFD_LoopOverBoundary(cctkGH, WT_CL_Dirichlet_Body);
if (verbose > 1)
{
CCTK_VInfo(CCTK_THORNSTRING,"Leaving WT_CL_Dirichlet_Body");
}
}
|
