1 files changed, 110 insertions, 12 deletions

diff --git a/Source/Parser/wp_parser_c.h b/Source/Parser/wp_parser_c.h
index d810bd685..3aafdec65 100644
--- a/Source/Parser/wp_parser_c.h
+++ b/Source/Parser/wp_parser_c.h
@@ -2,6 +2,8 @@
 #define WP_PARSER_C_H_
 
 #include "wp_parser_y.h"
+#include <AMReX_GpuQualifiers.H>
+#include <AMReX_Extension.H>
 
 #ifdef __cplusplus
 extern "C" {
@@ -18,71 +20,167 @@ extern "C" {
 #include <set>
 #include <string>
 
-inline
-double
+AMREX_GPU_HOST_DEVICE
+inline double
 wp_ast_eval (struct wp_node* node)
 {
     double result;
 
+#ifdef AMREX_DEVICE_COMPILE
+    extern __shared__ double extern_xyz[];
+    int tid = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*(blockDim.x*blockDim.y);
+    double* x = extern_xyz + tid*3;
+#endif
+
     switch (node->type)
     {
     case WP_NUMBER:
+    {
         result = ((struct wp_number*)node)->value;
         break;
+    }
     case WP_SYMBOL:
-        result = *(((struct wp_symbol*)node)->pointer);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i =((struct wp_symbol*)node)->ip.i;
+        result = x[i];
+#else
+        result = *(((struct wp_symbol*)node)->ip.p);
+#endif
         break;
+    }
     case WP_ADD:
+    {
         result = wp_ast_eval(node->l) + wp_ast_eval(node->r);
         break;
+    }
     case WP_SUB:
+    {
         result = wp_ast_eval(node->l) - wp_ast_eval(node->r);
         break;
+    }
     case WP_MUL:
+    {
         result = wp_ast_eval(node->l) * wp_ast_eval(node->r);
         break;
+    }
     case WP_DIV:
+    {
         result = wp_ast_eval(node->l) / wp_ast_eval(node->r);
         break;
+    }
     case WP_NEG:
+    {
         result = -wp_ast_eval(node->l);
         break;
+    }
     case WP_F1:
+    {
         result = wp_call_f1(((struct wp_f1*)node)->ftype,
                 wp_ast_eval(((struct wp_f1*)node)->l));
         break;
+    }
     case WP_F2:
+    {
         result = wp_call_f2(((struct wp_f2*)node)->ftype,
                 wp_ast_eval(((struct wp_f2*)node)->l),
                 wp_ast_eval(((struct wp_f2*)node)->r));
         break;
+    }
     case WP_ADD_VP:
-        result = node->lvp.v + *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->rip.i;
+        result = node->lvp.v + x[i];
+#else
+        result = node->lvp.v + *(node->rip.p);
+#endif
         break;
+    }
     case WP_ADD_PP:
-        result = *(node->lvp.p) + *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->lvp.ip.i;
+        int j = node->rip.i;
+        result = x[i] + x[j];
+#else
+        result = *(node->lvp.ip.p) + *(node->rip.p);
+#endif
         break;
+    }
     case WP_SUB_VP:
-        result = node->lvp.v - *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->rip.i;
+        result = node->lvp.v - x[i];
+#else
+        result = node->lvp.v - *(node->rip.p);
+#endif
         break;
+    }
     case WP_SUB_PP:
-        result = *(node->lvp.p) - *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->lvp.ip.i;
+        int j = node->rip.i;
+        result = x[i] - x[j];
+#else
+        result = *(node->lvp.ip.p) - *(node->rip.p);
+#endif
         break;
+    }
     case WP_MUL_VP:
-        result = node->lvp.v * *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->rip.i;
+        result = node->lvp.v * x[i];
+#else
+        result = node->lvp.v * *(node->rip.p);
+#endif
         break;
+    }
     case WP_MUL_PP:
-        result = *(node->lvp.p) * *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->lvp.ip.i;
+        int j = node->rip.i;
+        result = x[i] * x[j];
+#else
+        result = *(node->lvp.ip.p) * *(node->rip.p);
+#endif
         break;
+    }
     case WP_DIV_VP:
-        result = node->lvp.v / *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->rip.i;
+        result = node->lvp.v / x[i];
+#else
+        result = node->lvp.v / *(node->rip.p);
+#endif
         break;
+    }
     case WP_DIV_PP:
-        result = *(node->lvp.p) / *(node->rp);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->lvp.ip.i;
+        int j = node->rip.i;
+        result = x[i] / x[j];
+#else
+        result = *(node->lvp.ip.p) / *(node->rip.p);
+#endif
         break;
+    }
     case WP_NEG_P:
-        result = -*(node->lvp.p);
+    {
+#ifdef AMREX_DEVICE_COMPILE
+        int i = node->rip.i;
+        result = -x[i];
+#else
+        result = -*(node->lvp.ip.p);
+#endif
         break;
+    }
     default:
         yyerror("wp_ast_eval: unknown node type %d\n", node->type);
     }