aboutsummaryrefslogtreecommitdiff
path: root/Examples/SimpleWaveOpenCL
diff options
context:
space:
mode:
authorIan Hinder <ian.hinder@aei.mpg.de>2012-11-14 19:53:02 +0100
committerIan Hinder <ian.hinder@aei.mpg.de>2012-11-14 19:53:02 +0100
commitcda73baf0fbca06d75e51ad28ed708d7a087a849 (patch)
tree4fac3164d827c0713877954ac0730f166714a5ad /Examples/SimpleWaveOpenCL
parent72b0c57fbcbbc71f55464dfdf751dbdb0b45fcb8 (diff)
Regenerate examples
Diffstat (limited to 'Examples/SimpleWaveOpenCL')
-rw-r--r--Examples/SimpleWaveOpenCL/cakernel.ccl1
-rw-r--r--Examples/SimpleWaveOpenCL/schedule.ccl27
-rw-r--r--Examples/SimpleWaveOpenCL/src/Differencing.h2
-rw-r--r--Examples/SimpleWaveOpenCL/src/calc_rhs.cc36
-rw-r--r--Examples/SimpleWaveOpenCL/src/initial_sine.cc40
-rw-r--r--Examples/SimpleWaveOpenCL/src/make.code.defn2
6 files changed, 61 insertions, 47 deletions
diff --git a/Examples/SimpleWaveOpenCL/cakernel.ccl b/Examples/SimpleWaveOpenCL/cakernel.ccl
deleted file mode 100644
index 32dc04a..0000000
--- a/Examples/SimpleWaveOpenCL/cakernel.ccl
+++ /dev/null
@@ -1 +0,0 @@
-KrancThorn`Private`cakernel$438 \ No newline at end of file
diff --git a/Examples/SimpleWaveOpenCL/schedule.ccl b/Examples/SimpleWaveOpenCL/schedule.ccl
index 13a096e..ec20325 100644
--- a/Examples/SimpleWaveOpenCL/schedule.ccl
+++ b/Examples/SimpleWaveOpenCL/schedule.ccl
@@ -33,12 +33,6 @@ schedule SimpleWaveOpenCL_Startup at STARTUP
OPTIONS: meta
} "create banner"
-schedule SimpleWaveOpenCL_RegisterVars in MoL_Register
-{
- LANG: C
- OPTIONS: meta
-} "Register Variables for MoL"
-
schedule SimpleWaveOpenCL_RegisterSymmetries in SymmetryRegister
{
LANG: C
@@ -48,17 +42,20 @@ schedule SimpleWaveOpenCL_RegisterSymmetries in SymmetryRegister
schedule initial_sine AT INITIAL
{
LANG: C
- # TAGS: OpenCL=1
- READS: grid::coordinates
- WRITES: SimpleWaveOpenCL::evolved_group
+ TAGS: Device=1
+ READS: grid::x
+ WRITES: SimpleWaveOpenCL::phi
+ WRITES: SimpleWaveOpenCL::pi
} "initial_sine"
schedule calc_rhs in MoL_CalcRHS
{
LANG: C
- # TAGS: OpenCL=1
- READS: SimpleWaveOpenCL::evolved_group
- WRITES: SimpleWaveOpenCL::evolved_grouprhs
+ TAGS: Device=1
+ READS: SimpleWaveOpenCL::phi
+ READS: SimpleWaveOpenCL::pi
+ WRITES: SimpleWaveOpenCL::phirhs
+ WRITES: SimpleWaveOpenCL::pirhs
} "calc_rhs"
schedule SimpleWaveOpenCL_SelectBoundConds in MoL_PostStep
@@ -74,6 +71,12 @@ schedule SimpleWaveOpenCL_CheckBoundaries at BASEGRID
OPTIONS: meta
} "check boundaries treatment"
+schedule SimpleWaveOpenCL_RegisterVars in MoL_Register
+{
+ LANG: C
+ OPTIONS: meta
+} "Register Variables for MoL"
+
schedule group ApplyBCs as SimpleWaveOpenCL_ApplyBCs in MoL_PostStep after SimpleWaveOpenCL_SelectBoundConds
{
# no language specified
diff --git a/Examples/SimpleWaveOpenCL/src/Differencing.h b/Examples/SimpleWaveOpenCL/src/Differencing.h
index 476bbe1..7d48129 100644
--- a/Examples/SimpleWaveOpenCL/src/Differencing.h
+++ b/Examples/SimpleWaveOpenCL/src/Differencing.h
@@ -1,4 +1,4 @@
-static char const * const differencing =
+static char const *const differencing =
"#ifndef KRANC_DIFF_FUNCTIONS\n"
"# define PDstandard2nd1(u) ((-KRANC_GFOFFSET3D(u,-1,0,0) + KRANC_GFOFFSET3D(u,1,0,0))*p1o2dx)\n"
"#else\n"
diff --git a/Examples/SimpleWaveOpenCL/src/calc_rhs.cc b/Examples/SimpleWaveOpenCL/src/calc_rhs.cc
index f463547..0344c0d 100644
--- a/Examples/SimpleWaveOpenCL/src/calc_rhs.cc
+++ b/Examples/SimpleWaveOpenCL/src/calc_rhs.cc
@@ -18,10 +18,10 @@
/* Define macros used in calculations */
#define INITVALUE (42)
-#define QAD(x) (SQR(SQR(x)))
-#define INV(x) ((1.0) / (x))
+#define INV(x) ((CCTK_REAL)1.0 / (x))
#define SQR(x) ((x) * (x))
-#define CUB(x) ((x) * (x) * (x))
+#define CUB(x) ((x) * SQR(x))
+#define QAD(x) (SQR(SQR(x)))
extern "C" void calc_rhs_SelectBCs(CCTK_ARGUMENTS)
{
@@ -40,7 +40,7 @@ static void calc_rhs_Body(cGH const * restrict const cctkGH, int const dir, int
DECLARE_CCTK_ARGUMENTS;
DECLARE_CCTK_PARAMETERS;
- char const * const source =
+ char const *const source =
"\n"
"/* Include user-supplied include files */\n"
"\n"
@@ -86,9 +86,9 @@ static void calc_rhs_Body(cGH const * restrict const cctkGH, int const dir, int
"\n"
"/* Loop over the grid points */\n"
"#pragma omp parallel\n"
- "CCTK_LOOP3 (calc_rhs,\n"
+ "CCTK_LOOP3(calc_rhs,\n"
" i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n"
- " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2])\n"
+ " cctk_ash[0],cctk_ash[1],cctk_ash[2])\n"
"{\n"
" ptrdiff_t const index = di*i + dj*j + dk*k;\n"
" \n"
@@ -112,20 +112,24 @@ static void calc_rhs_Body(cGH const * restrict const cctkGH, int const dir, int
" PDstandard2nd33phi;\n"
" \n"
" /* Copy local copies back to grid functions */\n"
+ " vec_store_partial_prepare(i,lc_imin,lc_imax);\n"
" vec_store_nta_partial(phirhs[index],phirhsL);\n"
" vec_store_nta_partial(pirhs[index],pirhsL);\n"
"}\n"
- "CCTK_ENDLOOP3 (calc_rhs);\n"
+ "CCTK_ENDLOOP3(calc_rhs);\n"
""
;
- char const * const groups[] = {"SimpleWaveOpenCL::evolved_group","SimpleWaveOpenCL::evolved_grouprhs",NULL};
+ char const *const groups[] = {
+ "SimpleWaveOpenCL::evolved_group",
+ "SimpleWaveOpenCL::evolved_grouprhs",
+ NULL};
- static struct OpenCLKernel * kernel = NULL;
- char const * const sources[] = {differencing, source, NULL};
- OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "calc_rhs",
- sources, groups, NULL, NULL, NULL, -1,
- imin, imax, &kernel);
+ static struct OpenCLKernel *kernel = NULL;
+ char const *const sources[] = {differencing, source, NULL};
+ OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "calc_rhs",
+ sources, groups, NULL, NULL, NULL, -1,
+ imin, imax, &kernel);
}
@@ -145,12 +149,14 @@ extern "C" void calc_rhs(CCTK_ARGUMENTS)
return;
}
- const char *groups[] = {"SimpleWaveOpenCL::evolved_group","SimpleWaveOpenCL::evolved_grouprhs"};
+ const char *const groups[] = {
+ "SimpleWaveOpenCL::evolved_group",
+ "SimpleWaveOpenCL::evolved_grouprhs"};
GenericFD_AssertGroupStorage(cctkGH, "calc_rhs", 2, groups);
GenericFD_EnsureStencilFits(cctkGH, "calc_rhs", 1, 1, 1);
- GenericFD_LoopOverInterior(cctkGH, &calc_rhs_Body);
+ GenericFD_LoopOverInterior(cctkGH, calc_rhs_Body);
if (verbose > 1)
{
diff --git a/Examples/SimpleWaveOpenCL/src/initial_sine.cc b/Examples/SimpleWaveOpenCL/src/initial_sine.cc
index fa553dd..b29c7b5 100644
--- a/Examples/SimpleWaveOpenCL/src/initial_sine.cc
+++ b/Examples/SimpleWaveOpenCL/src/initial_sine.cc
@@ -18,17 +18,17 @@
/* Define macros used in calculations */
#define INITVALUE (42)
-#define QAD(x) (SQR(SQR(x)))
-#define INV(x) ((1.0) / (x))
+#define INV(x) ((CCTK_REAL)1.0 / (x))
#define SQR(x) ((x) * (x))
-#define CUB(x) ((x) * (x) * (x))
+#define CUB(x) ((x) * SQR(x))
+#define QAD(x) (SQR(SQR(x)))
static void initial_sine_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 =
+ char const *const source =
"\n"
"/* Include user-supplied include files */\n"
"\n"
@@ -74,9 +74,9 @@ static void initial_sine_Body(cGH const * restrict const cctkGH, int const dir,
"\n"
"/* Loop over the grid points */\n"
"#pragma omp parallel\n"
- "CCTK_LOOP3 (initial_sine,\n"
+ "CCTK_LOOP3(initial_sine,\n"
" i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],\n"
- " cctk_lsh[0],cctk_lsh[1],cctk_lsh[2])\n"
+ " cctk_ash[0],cctk_ash[1],cctk_ash[2])\n"
"{\n"
" ptrdiff_t const index = di*i + dj*j + dk*k;\n"
" \n"
@@ -90,25 +90,29 @@ static void initial_sine_Body(cGH const * restrict const cctkGH, int const dir,
" /* Precompute derivatives */\n"
" \n"
" /* Calculate temporaries and grid functions */\n"
- " CCTK_REAL phiL = Sin(2*Pi*(xL - t));\n"
+ " CCTK_REAL phiL = sin(2*Pi*(xL - t));\n"
" \n"
- " CCTK_REAL piL = -2*Pi*Cos(2*Pi*(xL - t));\n"
+ " CCTK_REAL piL = -2*Pi*cos(2*Pi*(xL - t));\n"
" \n"
" /* Copy local copies back to grid functions */\n"
+ " vec_store_partial_prepare(i,lc_imin,lc_imax);\n"
" vec_store_nta_partial(phi[index],phiL);\n"
" vec_store_nta_partial(pi[index],piL);\n"
"}\n"
- "CCTK_ENDLOOP3 (initial_sine);\n"
+ "CCTK_ENDLOOP3(initial_sine);\n"
""
;
- char const * const groups[] = {"SimpleWaveOpenCL::evolved_group","grid::coordinates",NULL};
+ char const *const groups[] = {
+ "SimpleWaveOpenCL::evolved_group",
+ "grid::coordinates",
+ NULL};
- static struct OpenCLKernel * kernel = NULL;
- char const * const sources[] = {differencing, source, NULL};
- OpenCLRunTime_CallKernel (cctkGH, CCTK_THORNSTRING, "initial_sine",
- sources, groups, NULL, NULL, NULL, -1,
- imin, imax, &kernel);
+ static struct OpenCLKernel *kernel = NULL;
+ char const *const sources[] = {differencing, source, NULL};
+ OpenCLRunTime_CallKernel(cctkGH, CCTK_THORNSTRING, "initial_sine",
+ sources, groups, NULL, NULL, NULL, -1,
+ imin, imax, &kernel);
}
@@ -128,11 +132,13 @@ extern "C" void initial_sine(CCTK_ARGUMENTS)
return;
}
- const char *groups[] = {"SimpleWaveOpenCL::evolved_group","grid::coordinates"};
+ const char *const groups[] = {
+ "SimpleWaveOpenCL::evolved_group",
+ "grid::coordinates"};
GenericFD_AssertGroupStorage(cctkGH, "initial_sine", 2, groups);
- GenericFD_LoopOverEverything(cctkGH, &initial_sine_Body);
+ GenericFD_LoopOverEverything(cctkGH, initial_sine_Body);
if (verbose > 1)
{
diff --git a/Examples/SimpleWaveOpenCL/src/make.code.defn b/Examples/SimpleWaveOpenCL/src/make.code.defn
index 535f5d0..3332924 100644
--- a/Examples/SimpleWaveOpenCL/src/make.code.defn
+++ b/Examples/SimpleWaveOpenCL/src/make.code.defn
@@ -1,3 +1,3 @@
# File produced by Kranc
-SRCS = Startup.cc RegisterMoL.cc RegisterSymmetries.cc initial_sine.cc calc_rhs.cc Boundaries.cc
+SRCS = Startup.cc RegisterSymmetries.cc RegisterMoL.cc initial_sine.cc calc_rhs.cc Boundaries.cc