diff options
Diffstat (limited to 'Source/Parser/wp_parser_c.h')
| -rw-r--r-- | Source/Parser/wp_parser_c.h | 44 |
1 files changed, 19 insertions, 25 deletions
diff --git a/Source/Parser/wp_parser_c.h b/Source/Parser/wp_parser_c.h index 2cf0e2c00..c9c0d82ac 100644 --- a/Source/Parser/wp_parser_c.h +++ b/Source/Parser/wp_parser_c.h @@ -23,16 +23,10 @@ extern "C" { AMREX_GPU_HOST_DEVICE inline amrex_real -wp_ast_eval (struct wp_node* node) +wp_ast_eval (struct wp_node* node, amrex_real const* x) { amrex_real result; -#ifdef AMREX_DEVICE_COMPILE - extern __shared__ amrex_real extern_xyz[]; - int tid = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*(blockDim.x*blockDim.y); - amrex_real* x = extern_xyz + tid*4; // parser assumes 4 independent variables (x,y,z,t) -#endif - switch (node->type) { case WP_NUMBER: @@ -42,7 +36,7 @@ wp_ast_eval (struct wp_node* node) } case WP_SYMBOL: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i =((struct wp_symbol*)node)->ip.i; result = x[i]; #else @@ -52,45 +46,45 @@ wp_ast_eval (struct wp_node* node) } case WP_ADD: { - result = wp_ast_eval(node->l) + wp_ast_eval(node->r); + result = wp_ast_eval(node->l,x) + wp_ast_eval(node->r,x); break; } case WP_SUB: { - result = wp_ast_eval(node->l) - wp_ast_eval(node->r); + result = wp_ast_eval(node->l,x) - wp_ast_eval(node->r,x); break; } case WP_MUL: { - result = wp_ast_eval(node->l) * wp_ast_eval(node->r); + result = wp_ast_eval(node->l,x) * wp_ast_eval(node->r,x); break; } case WP_DIV: { - result = wp_ast_eval(node->l) / wp_ast_eval(node->r); + result = wp_ast_eval(node->l,x) / wp_ast_eval(node->r,x); break; } case WP_NEG: { - result = -wp_ast_eval(node->l); + result = -wp_ast_eval(node->l,x); break; } case WP_F1: { result = wp_call_f1(((struct wp_f1*)node)->ftype, - wp_ast_eval(((struct wp_f1*)node)->l)); + wp_ast_eval(((struct wp_f1*)node)->l,x)); 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)); + wp_ast_eval(((struct wp_f2*)node)->l,x), + wp_ast_eval(((struct wp_f2*)node)->r,x)); break; } case WP_ADD_VP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->rip.i; result = node->lvp.v + x[i]; #else @@ -100,7 +94,7 @@ wp_ast_eval (struct wp_node* node) } case WP_ADD_PP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->lvp.ip.i; int j = node->rip.i; result = x[i] + x[j]; @@ -111,7 +105,7 @@ wp_ast_eval (struct wp_node* node) } case WP_SUB_VP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->rip.i; result = node->lvp.v - x[i]; #else @@ -121,7 +115,7 @@ wp_ast_eval (struct wp_node* node) } case WP_SUB_PP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->lvp.ip.i; int j = node->rip.i; result = x[i] - x[j]; @@ -132,7 +126,7 @@ wp_ast_eval (struct wp_node* node) } case WP_MUL_VP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->rip.i; result = node->lvp.v * x[i]; #else @@ -142,7 +136,7 @@ wp_ast_eval (struct wp_node* node) } case WP_MUL_PP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->lvp.ip.i; int j = node->rip.i; result = x[i] * x[j]; @@ -153,7 +147,7 @@ wp_ast_eval (struct wp_node* node) } case WP_DIV_VP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->rip.i; result = node->lvp.v / x[i]; #else @@ -163,7 +157,7 @@ wp_ast_eval (struct wp_node* node) } case WP_DIV_PP: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->lvp.ip.i; int j = node->rip.i; result = x[i] / x[j]; @@ -174,7 +168,7 @@ wp_ast_eval (struct wp_node* node) } case WP_NEG_P: { -#ifdef AMREX_DEVICE_COMPILE +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) int i = node->rip.i; result = -x[i]; #else |