diff options
author | Ian Hinder <ian.hinder@aei.mpg.de> | 2012-11-14 19:53:02 +0100 |
---|---|---|
committer | Ian Hinder <ian.hinder@aei.mpg.de> | 2012-11-14 19:53:02 +0100 |
commit | cda73baf0fbca06d75e51ad28ed708d7a087a849 (patch) | |
tree | 4fac3164d827c0713877954ac0730f166714a5ad /Examples/SimpleWaveOpenCL | |
parent | 72b0c57fbcbbc71f55464dfdf751dbdb0b45fcb8 (diff) |
Regenerate examples
Diffstat (limited to 'Examples/SimpleWaveOpenCL')
-rw-r--r-- | Examples/SimpleWaveOpenCL/cakernel.ccl | 1 | ||||
-rw-r--r-- | Examples/SimpleWaveOpenCL/schedule.ccl | 27 | ||||
-rw-r--r-- | Examples/SimpleWaveOpenCL/src/Differencing.h | 2 | ||||
-rw-r--r-- | Examples/SimpleWaveOpenCL/src/calc_rhs.cc | 36 | ||||
-rw-r--r-- | Examples/SimpleWaveOpenCL/src/initial_sine.cc | 40 | ||||
-rw-r--r-- | Examples/SimpleWaveOpenCL/src/make.code.defn | 2 |
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 |