Regenerate examples

10 files changed, 76 insertions, 130 deletions

diff --git a/Examples/WaveCaKernel/cakernel.ccl b/Examples/WaveCaKernel/cakernel.ccl
index 3df4bbb..f756f91 100644
--- a/Examples/WaveCaKernel/cakernel.ccl
+++ b/Examples/WaveCaKernel/cakernel.ccl
@@ -1,90 +1,29 @@
 CCTK_CUDA_KERNEL calc_rhs_2 TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="1,1,1,1,1,1"
 {
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in
-  {
-    phi
-  }
-  "phi"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out
-  {
-    phirhs
-  }
-  "phirhs"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in
-  {
-    pi
-  }
-  "pi"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out
-  {
-    pirhs
-  }
-  "pirhs"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in {phi} "phi"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out {phirhs} "phirhs"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in {pi} "pi"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out {pirhs} "pirhs"
 }
 
 CCTK_CUDA_KERNEL calc_rhs_4 TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="2,2,2,2,2,2"
 {
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in
-  {
-    phi
-  }
-  "phi"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out
-  {
-    phirhs
-  }
-  "phirhs"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in
-  {
-    pi
-  }
-  "pi"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out
-  {
-    pirhs
-  }
-  "pirhs"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in {phi} "phi"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out {phirhs} "phirhs"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in {pi} "pi"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out {pirhs} "pirhs"
 }
 
 CCTK_CUDA_KERNEL calc_bound_rhs TYPE=gpu_cuda/boundary_s TILE="8,8,8" SHARECODE=yes
 {
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out
-  {
-    phirhs
-  }
-  "phirhs"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out
-  {
-    pirhs
-  }
-  "pirhs"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in
-  {
-    xCopy
-  }
-  "xCopy"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out {phirhs} "phirhs"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=out {pirhs} "pirhs"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=in {xCopy} "xCopy"
 }
 
 CCTK_CUDA_KERNEL copy_to_device TYPE=gpu_cuda/3dblock TILE="8,8,8" SHARECODE=yes STENCIL="0,0,0,0,0,0"
 {
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout
-  {
-    phi
-  }
-  "phi"
-  
-  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout
-  {
-    pi
-  }
-  "pi"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout {phi} "phi"
+  CCTK_CUDA_KERNEL_VARIABLE  cached=no intent=inout {pi} "pi"
 }
 
diff --git a/Examples/WaveCaKernel/configuration.ccl b/Examples/WaveCaKernel/configuration.ccl
index 46dbde5..6d94724 100644
--- a/Examples/WaveCaKernel/configuration.ccl
+++ b/Examples/WaveCaKernel/configuration.ccl
@@ -4,4 +4,4 @@ REQUIRES GenericFD
 OPTIONAL LoopControl
 {
 }
-REQUIRES CUDA
\ No newline at end of file
+REQUIRES CUDA MPI
diff --git a/Examples/WaveCaKernel/interface.ccl b/Examples/WaveCaKernel/interface.ccl
index 11c1416..1d883f5 100644
--- a/Examples/WaveCaKernel/interface.ccl
+++ b/Examples/WaveCaKernel/interface.ccl
@@ -2,13 +2,14 @@
 
 implements: WaveCaKernel
 
-inherits:   Grid GenericFD Boundary
+inherits:   Grid GenericFD Boundary Accelerator
 
 
 
 USES INCLUDE: GenericFD.h
 USES INCLUDE: Symmetry.h
 USES INCLUDE: sbp_calc_coeffs.h
+USES INCLUDE: CaCUDALib_driver_support.h
 USES INCLUDE: Boundary.h
 USES INCLUDE: loopcontrol.h
 
diff --git a/Examples/WaveCaKernel/schedule.ccl b/Examples/WaveCaKernel/schedule.ccl
index 74879b1..9dea53b 100644
--- a/Examples/WaveCaKernel/schedule.ccl
+++ b/Examples/WaveCaKernel/schedule.ccl
@@ -57,100 +57,100 @@ schedule WaveCaKernel_RegisterSymmetries in SymmetryRegister
 schedule initial_gaussian AT INITIAL
 {
   LANG: C
-  READS:    grid::coordinates
-  WRITES:   WaveCaKernel::phi_g
-  WRITES:   WaveCaKernel::pi_g
-  WRITES:   WaveCaKernel::xCopy_g
+  READS:    grid::x
+  WRITES:   WaveCaKernel::phi
+  WRITES:   WaveCaKernel::pi
+  WRITES:   WaveCaKernel::xCopy
 } "initial_gaussian"
 
 
 if (fdOrder == 2)
 {
-  schedule CAKERNEL_Launch_calc_rhs_2 in MoL_CalcRHS
+  schedule CAKERNEL_Launch_calc_rhs_2 as calc_rhs_2 in MoL_CalcRHS
   {
     LANG: C
     TAGS: Device=1
-    READS:    WaveCaKernel::phi_g
-    READS:    WaveCaKernel::pi_g
-    WRITES:   WaveCaKernel::phi_grhs
-    WRITES:   WaveCaKernel::pi_grhs
+    READS:    WaveCaKernel::phi
+    READS:    WaveCaKernel::pi
+    WRITES:   WaveCaKernel::phirhs
+    WRITES:   WaveCaKernel::pirhs
   } "calc_rhs_2"
 }
 
 
 if (fdOrder == 2)
 {
-  schedule CAKERNEL_Launch_calc_rhs_2 at ANALYSIS
+  schedule CAKERNEL_Launch_calc_rhs_2 as calc_rhs_2 at ANALYSIS
   {
     LANG: C
     SYNC: phi_grhs
     SYNC: pi_grhs
     TAGS: Device=1
-    READS:    WaveCaKernel::phi_g
-    READS:    WaveCaKernel::pi_g
-    WRITES:   WaveCaKernel::phi_grhs
-    WRITES:   WaveCaKernel::pi_grhs
+    READS:    WaveCaKernel::phi
+    READS:    WaveCaKernel::pi
+    WRITES:   WaveCaKernel::phirhs
+    WRITES:   WaveCaKernel::pirhs
   } "calc_rhs_2"
 }
 
 
 if (fdOrder == 4)
 {
-  schedule CAKERNEL_Launch_calc_rhs_4 in MoL_CalcRHS
+  schedule CAKERNEL_Launch_calc_rhs_4 as calc_rhs_4 in MoL_CalcRHS
   {
     LANG: C
     TAGS: Device=1
-    READS:    WaveCaKernel::phi_g
-    READS:    WaveCaKernel::pi_g
-    WRITES:   WaveCaKernel::phi_grhs
-    WRITES:   WaveCaKernel::pi_grhs
+    READS:    WaveCaKernel::phi
+    READS:    WaveCaKernel::pi
+    WRITES:   WaveCaKernel::phirhs
+    WRITES:   WaveCaKernel::pirhs
   } "calc_rhs_4"
 }
 
 
 if (fdOrder == 4)
 {
-  schedule CAKERNEL_Launch_calc_rhs_4 at ANALYSIS
+  schedule CAKERNEL_Launch_calc_rhs_4 as calc_rhs_4 at ANALYSIS
   {
     LANG: C
     SYNC: phi_grhs
     SYNC: pi_grhs
     TAGS: Device=1
-    READS:    WaveCaKernel::phi_g
-    READS:    WaveCaKernel::pi_g
-    WRITES:   WaveCaKernel::phi_grhs
-    WRITES:   WaveCaKernel::pi_grhs
+    READS:    WaveCaKernel::phi
+    READS:    WaveCaKernel::pi
+    WRITES:   WaveCaKernel::phirhs
+    WRITES:   WaveCaKernel::pirhs
   } "calc_rhs_4"
 }
 
-schedule CAKERNEL_Launch_calc_bound_rhs in MoL_RHSBoundaries
+schedule CAKERNEL_Launch_calc_bound_rhs as calc_bound_rhs in MoL_RHSBoundaries
 {
   LANG: C
   TAGS: Device=1
-  READS:    WaveCaKernel::xCopy_g
-  WRITES:   WaveCaKernel::phi_grhs
-  WRITES:   WaveCaKernel::pi_grhs
+  READS:    WaveCaKernel::xCopy
+  WRITES:   WaveCaKernel::phirhs
+  WRITES:   WaveCaKernel::pirhs
 } "calc_bound_rhs"
 
-schedule CAKERNEL_Launch_calc_bound_rhs at ANALYSIS
+schedule CAKERNEL_Launch_calc_bound_rhs as calc_bound_rhs at ANALYSIS
 {
   LANG: C
   SYNC: phi_grhs
   SYNC: pi_grhs
   TAGS: Device=1
-  READS:    WaveCaKernel::xCopy_g
-  WRITES:   WaveCaKernel::phi_grhs
-  WRITES:   WaveCaKernel::pi_grhs
+  READS:    WaveCaKernel::xCopy
+  WRITES:   WaveCaKernel::phirhs
+  WRITES:   WaveCaKernel::pirhs
 } "calc_bound_rhs"
 
-schedule CAKERNEL_Launch_copy_to_device at INITIAL after initial_gaussian
+schedule CAKERNEL_Launch_copy_to_device as copy_to_device at INITIAL after initial_gaussian
 {
   LANG: C
   TAGS: Device=1
-  READS:    WaveCaKernel::phi_g
-  READS:    WaveCaKernel::pi_g
-  WRITES:   WaveCaKernel::phi_g
-  WRITES:   WaveCaKernel::pi_g
+  READS:    WaveCaKernel::phi
+  READS:    WaveCaKernel::pi
+  WRITES:   WaveCaKernel::phi
+  WRITES:   WaveCaKernel::pi
 } "copy_to_device"
 
 schedule WaveCaKernel_SelectBoundConds in MoL_PostStep
@@ -173,6 +173,12 @@ schedule WaveCaKernel_RegisterVars in MoL_Register
   OPTIONS: meta
 } "Register Variables for MoL"
 
+schedule WaveCaKernel_Init in CCTK_BASEGRID after Accelerator_SetDevice
+{
+  LANG: C
+  OPTIONS: local
+} "Initialize CUDA Device"
+
 schedule group ApplyBCs as WaveCaKernel_ApplyBCs in MoL_PostStep after WaveCaKernel_SelectBoundConds
 {
   # no language specified
diff --git a/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code b/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code
index 0942be7..013ab6c 100644
--- a/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code
+++ b/Examples/WaveCaKernel/src/CaKernel__calc_bound_rhs.code
@@ -9,10 +9,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)))
 
 CAKERNEL_calc_bound_rhs_Begin
   
diff --git a/Examples/WaveCaKernel/src/CaKernel__calc_rhs_2.code b/Examples/WaveCaKernel/src/CaKernel__calc_rhs_2.code
index b85aad1..085e274 100644
--- a/Examples/WaveCaKernel/src/CaKernel__calc_rhs_2.code
+++ b/Examples/WaveCaKernel/src/CaKernel__calc_rhs_2.code
@@ -9,10 +9,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)))
 
 CAKERNEL_calc_rhs_2_Begin
   
diff --git a/Examples/WaveCaKernel/src/CaKernel__calc_rhs_4.code b/Examples/WaveCaKernel/src/CaKernel__calc_rhs_4.code
index a3f9bbc..c94216d 100644
--- a/Examples/WaveCaKernel/src/CaKernel__calc_rhs_4.code
+++ b/Examples/WaveCaKernel/src/CaKernel__calc_rhs_4.code
@@ -9,10 +9,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)))
 
 CAKERNEL_calc_rhs_4_Begin
   
diff --git a/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code b/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code
index 76cdc9a..40f0db4 100644
--- a/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code
+++ b/Examples/WaveCaKernel/src/CaKernel__copy_to_device.code
@@ -9,10 +9,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)))
 
 CAKERNEL_copy_to_device_Begin
   
diff --git a/Examples/WaveCaKernel/src/initial_gaussian.cc b/Examples/WaveCaKernel/src/initial_gaussian.cc
index ad73b93..28d41e0 100644
--- a/Examples/WaveCaKernel/src/initial_gaussian.cc
+++ b/Examples/WaveCaKernel/src/initial_gaussian.cc
@@ -17,10 +17,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)))
 
 static void initial_gaussian_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[])
 {
@@ -86,7 +86,7 @@ static void initial_gaussian_Body(cGH const * restrict const cctkGH, int const d
   #pragma omp parallel
   CCTK_LOOP3(initial_gaussian,
     i,j,k, imin[0],imin[1],imin[2], imax[0],imax[1],imax[2],
-    cctk_lsh[0],cctk_lsh[1],cctk_lsh[2])
+    cctk_ash[0],cctk_ash[1],cctk_ash[2])
   {
     ptrdiff_t const index = di*i + dj*j + dk*k;
     
diff --git a/Examples/WaveCaKernel/test b/Examples/WaveCaKernel/test
index 35022b1..7137e86 120000
--- a/Examples/WaveCaKernel/test
+++ b/Examples/WaveCaKernel/test
@@ -1 +1 @@
-../tests/WaveCaKernel/test
\ No newline at end of file
+../tests/WaveCaKernel
\ No newline at end of file