diff options
Diffstat (limited to 'Source/Parser')
| -rw-r--r-- | Source/Parser/GpuParser.H | 72 | ||||
| -rw-r--r-- | Source/Parser/GpuParser.cpp | 73 | ||||
| -rw-r--r-- | Source/Parser/Make.package | 2 | ||||
| -rw-r--r-- | Source/Parser/WarpXParser.H | 4 | ||||
| -rw-r--r-- | Source/Parser/wp_parser_c.h | 122 | ||||
| -rw-r--r-- | Source/Parser/wp_parser_y.c | 129 | ||||
| -rw-r--r-- | Source/Parser/wp_parser_y.h | 22 |
7 files changed, 377 insertions, 47 deletions
diff --git a/Source/Parser/GpuParser.H b/Source/Parser/GpuParser.H new file mode 100644 index 000000000..1533ee6b9 --- /dev/null +++ b/Source/Parser/GpuParser.H @@ -0,0 +1,72 @@ +#ifndef WARPX_GPU_PARSER_H_ +#define WARPX_GPU_PARSER_H_ + +#include <WarpXParser.H> +#include <AMReX_Gpu.H> + +// When compiled for CPU, wrap WarpXParser and enable threading. +// When compiled for GPU, store one copy of the parser in +// CUDA managed memory for __device__ code, and one copy of the parser +// in CUDA managed memory for __host__ code. This way, the parser can be +// efficiently called from both host and device. +class GpuParser +{ +public: + GpuParser (WarpXParser const& wp); + void clear (); + + AMREX_GPU_HOST_DEVICE + double + operator() (double x, double y, double z) const noexcept + { +#ifdef AMREX_USE_GPU + +#ifdef AMREX_DEVICE_COMPILE +// WarpX compiled for GPU, function compiled for __device__ + // the 3D position of each particle is stored in shared memory. + amrex::Gpu::SharedMemory<double> gsm; + double* p = gsm.dataPtr(); + int tid = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*(blockDim.x*blockDim.y); + p[tid*3] = x; + p[tid*3+1] = y; + p[tid*3+2] = z; + return wp_ast_eval(m_gpu_parser.ast); +#else +// WarpX compiled for GPU, function compiled for __host__ + m_var.x = x; + m_var.y = y; + m_var.z = z; + return wp_ast_eval(m_cpu_parser.ast); +#endif + +#else +// WarpX compiled for CPU +#ifdef _OPENMP + int tid = omp_get_thread_num(); +#else + int tid = 0; +#endif + m_var[tid].x = x; + m_var[tid].y = y; + m_var[tid].z = z; + return wp_ast_eval(m_parser[tid]->ast); +#endif + } + +private: + +#ifdef AMREX_USE_GPU + // Copy of the parser running on __device__ + struct wp_parser m_gpu_parser; + // Copy of the parser running on __host__ + struct wp_parser m_cpu_parser; + mutable amrex::XDim3 m_var; +#else + // Only one parser + struct wp_parser** m_parser; + mutable amrex::XDim3* m_var; + int nthreads; +#endif +}; + +#endif diff --git a/Source/Parser/GpuParser.cpp b/Source/Parser/GpuParser.cpp new file mode 100644 index 000000000..97b96d465 --- /dev/null +++ b/Source/Parser/GpuParser.cpp @@ -0,0 +1,73 @@ +#include <GpuParser.H> + +GpuParser::GpuParser (WarpXParser const& wp) +{ +#ifdef AMREX_USE_GPU + + struct wp_parser* a_wp = wp.m_parser; + // Initialize GPU parser: allocate memory in CUDA managed memory, + // copy all data needed on GPU to m_gpu_parser + m_gpu_parser.sz_mempool = wp_ast_size((struct wp_node*)a_wp); + m_gpu_parser.p_root = (struct wp_node*) + amrex::The_Managed_Arena()->alloc(m_gpu_parser.sz_mempool); + m_gpu_parser.p_free = m_gpu_parser.p_root; + // 0: don't free the source + m_gpu_parser.ast = wp_parser_ast_dup(&m_gpu_parser, a_wp->ast, 0); + wp_parser_regvar_gpu(&m_gpu_parser, "x", 0); + wp_parser_regvar_gpu(&m_gpu_parser, "y", 1); + wp_parser_regvar_gpu(&m_gpu_parser, "z", 2); + + // Initialize CPU parser: allocate memory in CUDA managed memory, + // copy all data needed on CPU to m_cpu_parser + m_cpu_parser.sz_mempool = wp_ast_size((struct wp_node*)a_wp); + m_cpu_parser.p_root = (struct wp_node*) + amrex::The_Managed_Arena()->alloc(m_cpu_parser.sz_mempool); + m_cpu_parser.p_free = m_cpu_parser.p_root; + // 0: don't free the source + m_cpu_parser.ast = wp_parser_ast_dup(&m_cpu_parser, a_wp->ast, 0); + wp_parser_regvar(&m_cpu_parser, "x", &(m_var.x)); + wp_parser_regvar(&m_cpu_parser, "y", &(m_var.y)); + wp_parser_regvar(&m_cpu_parser, "z", &(m_var.z)); + +#else // not defined AMREX_USE_GPU + +#ifdef _OPENMP + nthreads = omp_get_max_threads(); +#else // _OPENMP + nthreads = 1; +#endif // _OPENMP + + m_parser = ::new struct wp_parser*[nthreads]; + m_var = ::new amrex::XDim3[nthreads]; + + for (int tid = 0; tid < nthreads; ++tid) + { +#ifdef _OPENMP + m_parser[tid] = wp_parser_dup(wp.m_parser[tid]); +#else // _OPENMP + m_parser[tid] = wp_parser_dup(wp.m_parser); +#endif // _OPENMP + wp_parser_regvar(m_parser[tid], "x", &(m_var[tid].x)); + wp_parser_regvar(m_parser[tid], "y", &(m_var[tid].y)); + wp_parser_regvar(m_parser[tid], "z", &(m_var[tid].z)); + } + +#endif // AMREX_USE_GPU +} + +void +GpuParser::clear () +{ +#ifdef AMREX_USE_GPU + amrex::The_Managed_Arena()->free(m_gpu_parser.ast); + amrex::The_Managed_Arena()->free(m_cpu_parser.ast); +#else + for (int tid = 0; tid < nthreads; ++tid) + { + wp_parser_delete(m_parser[tid]); + } + ::delete[] m_parser; + ::delete[] m_var; +#endif +} + diff --git a/Source/Parser/Make.package b/Source/Parser/Make.package index 26ef4fb43..5ce02cbda 100644 --- a/Source/Parser/Make.package +++ b/Source/Parser/Make.package @@ -3,6 +3,8 @@ cEXE_sources += wp_parser_y.c wp_parser.tab.c wp_parser.lex.c wp_parser_c.c cEXE_headers += wp_parser_y.h wp_parser.tab.h wp_parser.lex.h wp_parser_c.h CEXE_sources += WarpXParser.cpp CEXE_headers += WarpXParser.H +CEXE_headers += GpuParser.H +CEXE_sources += GpuParser.cpp INCLUDE_LOCATIONS += $(WARPX_HOME)/Source/Parser VPATH_LOCATIONS += $(WARPX_HOME)/Source/Parser diff --git a/Source/Parser/WarpXParser.H b/Source/Parser/WarpXParser.H index 046491e29..ffa61e457 100644 --- a/Source/Parser/WarpXParser.H +++ b/Source/Parser/WarpXParser.H @@ -13,6 +13,8 @@ #include <omp.h> #endif +class GpuParser; + class WarpXParser { public: @@ -46,6 +48,8 @@ public: std::set<std::string> symbols () const; + friend class GpuParser; + private: void clear (); 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); } diff --git a/Source/Parser/wp_parser_y.c b/Source/Parser/wp_parser_y.c index 46cb199db..259f9368b 100644 --- a/Source/Parser/wp_parser_y.c +++ b/Source/Parser/wp_parser_y.c @@ -6,6 +6,8 @@ #include "wp_parser_y.h" #include "wp_parser.tab.h" +#include <AMReX_GpuQualifiers.H> + static struct wp_node* wp_root = NULL; /* This is called by a bison rule to store the original AST in a @@ -33,7 +35,7 @@ wp_makesymbol (char* name) struct wp_symbol* symbol = (struct wp_symbol*) malloc(sizeof(struct wp_symbol)); symbol->type = WP_SYMBOL; symbol->name = strdup(name); - symbol->pointer = NULL; + symbol->ip.p = NULL; return symbol; } @@ -74,13 +76,19 @@ wp_newf2 (enum wp_f2_t ftype, struct wp_node* l, struct wp_node* r) return (struct wp_node*) tmp; } +AMREX_GPU_HOST_DEVICE void yyerror (char const *s, ...) { va_list vl; va_start(vl, s); +#ifdef AMREX_DEVICE_COMPILE + printf(s,"\n"); + assert(0); +#else vfprintf(stderr, s, vl); fprintf(stderr, "\n"); +#endif va_end(vl); } @@ -97,7 +105,7 @@ wp_parser_new (void) my_parser->ast = wp_parser_ast_dup(my_parser, wp_root,1); /* 1: free the source wp_root */ - if (my_parser->p_root + my_parser->sz_mempool != my_parser->p_free) { + if ((char*)my_parser->p_root + my_parser->sz_mempool != (char*)my_parser->p_free) { yyerror("wp_parser_new: error in memory size"); exit(1); } @@ -145,6 +153,7 @@ wp_parser_dup (struct wp_parser* source) return dest; } +AMREX_GPU_HOST_DEVICE double wp_call_f1 (enum wp_f1_t type, double a) { @@ -175,6 +184,7 @@ wp_call_f1 (enum wp_f1_t type, double a) } } +AMREX_GPU_HOST_DEVICE double wp_call_f2 (enum wp_f2_t type, double a, double b) { @@ -346,23 +356,23 @@ wp_parser_ast_dup (struct wp_parser* my_parser, struct wp_node* node, int move) #define WP_MOVEUP_R(node, v) \ struct wp_node* n = node->r->r; \ - double* p = node->r->rp; \ + double* p = node->r->rip.p; \ node->r = n; \ node->lvp.v = v; \ - node->rp = p; + node->rip.p = p; #define WP_MOVEUP_L(node, v) \ struct wp_node* n = node->l->r; \ - double* p = node->l->rp; \ + double* p = node->l->rip.p; \ node->r = n; \ node->lvp.v = v; \ - node->rp = p; + node->rip.p = p; #define WP_EVAL_R(node) node->r->lvp.v #define WP_EVAL_L(node) node->l->lvp.v #define WP_NEG_MOVEUP(node) \ node->r = node->l->r; \ node->lvp.v = -node->l->lvp.v; \ - node->rp = node->l->rp; + node->rip.p = node->l->rip.p; void wp_ast_optimize (struct wp_node* node) @@ -391,22 +401,22 @@ wp_ast_optimize (struct wp_node* node) node->r->type == WP_SYMBOL) { node->lvp.v = ((struct wp_number*)(node->l))->value; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_ADD_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_NUMBER) { node->lvp.v = ((struct wp_number*)(node->r))->value; - node->rp = ((struct wp_symbol*)(node->l))->pointer; + node->rip.p = ((struct wp_symbol*)(node->l))->ip.p; node->r = node->l; node->type = WP_ADD_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_SYMBOL) { - node->lvp.p = ((struct wp_symbol*)(node->l))->pointer; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->lvp.ip.p = ((struct wp_symbol*)(node->l))->ip.p; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_ADD_PP; } else if (node->l->type == WP_NUMBER && @@ -454,22 +464,22 @@ wp_ast_optimize (struct wp_node* node) node->r->type == WP_SYMBOL) { node->lvp.v = ((struct wp_number*)(node->l))->value; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_SUB_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_NUMBER) { node->lvp.v = -((struct wp_number*)(node->r))->value; - node->rp = ((struct wp_symbol*)(node->l))->pointer; + node->rip.p = ((struct wp_symbol*)(node->l))->ip.p; node->r = node->l; node->type = WP_ADD_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_SYMBOL) { - node->lvp.p = ((struct wp_symbol*)(node->l))->pointer; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->lvp.ip.p = ((struct wp_symbol*)(node->l))->ip.p; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_SUB_PP; } else if (node->l->type == WP_NUMBER && @@ -517,22 +527,22 @@ wp_ast_optimize (struct wp_node* node) node->r->type == WP_SYMBOL) { node->lvp.v = ((struct wp_number*)(node->l))->value; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_MUL_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_NUMBER) { node->lvp.v = ((struct wp_number*)(node->r))->value; - node->rp = ((struct wp_symbol*)(node->l))->pointer; + node->rip.p = ((struct wp_symbol*)(node->l))->ip.p; node->r = node->l; node->type = WP_MUL_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_SYMBOL) { - node->lvp.p = ((struct wp_symbol*)(node->l))->pointer; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->lvp.ip.p = ((struct wp_symbol*)(node->l))->ip.p; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_MUL_PP; } else if (node->l->type == WP_NUMBER && @@ -580,22 +590,22 @@ wp_ast_optimize (struct wp_node* node) node->r->type == WP_SYMBOL) { node->lvp.v = ((struct wp_number*)(node->l))->value; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_DIV_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_NUMBER) { node->lvp.v = 1./((struct wp_number*)(node->r))->value; - node->rp = ((struct wp_symbol*)(node->l))->pointer; + node->rip.p = ((struct wp_symbol*)(node->l))->ip.p; node->r = node->l; node->type = WP_MUL_VP; } else if (node->l->type == WP_SYMBOL && node->r->type == WP_SYMBOL) { - node->lvp.p = ((struct wp_symbol*)(node->l))->pointer; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->lvp.ip.p = ((struct wp_symbol*)(node->l))->ip.p; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; node->type = WP_DIV_PP; } else if (node->l->type == WP_NUMBER && @@ -637,7 +647,7 @@ wp_ast_optimize (struct wp_node* node) } else if (node->l->type == WP_SYMBOL) { - node->lvp.p = ((struct wp_symbol*)(node->l))->pointer; + node->lvp.ip.p = ((struct wp_symbol*)(node->l))->ip.p; node->type = WP_NEG_P; } else if (node->l->type == WP_ADD_VP) @@ -936,7 +946,7 @@ wp_ast_regvar (struct wp_node* node, char const* name, double* p) break; case WP_SYMBOL: if (strcmp(name, ((struct wp_symbol*)node)->name) == 0) { - ((struct wp_symbol*)node)->pointer = p; + ((struct wp_symbol*)node)->ip.p = p; } break; case WP_ADD: @@ -961,11 +971,11 @@ wp_ast_regvar (struct wp_node* node, char const* name, double* p) case WP_MUL_VP: case WP_DIV_VP: wp_ast_regvar(node->r, name, p); - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; break; case WP_NEG_P: wp_ast_regvar(node->l, name, p); - node->lvp.p = ((struct wp_symbol*)(node->l))->pointer; + node->lvp.ip.p = ((struct wp_symbol*)(node->l))->ip.p; break; case WP_ADD_PP: case WP_SUB_PP: @@ -973,8 +983,8 @@ wp_ast_regvar (struct wp_node* node, char const* name, double* p) case WP_DIV_PP: wp_ast_regvar(node->l, name, p); wp_ast_regvar(node->r, name, p); - node->lvp.p = ((struct wp_symbol*)(node->l))->pointer; - node->rp = ((struct wp_symbol*)(node->r))->pointer; + node->lvp.ip.p = ((struct wp_symbol*)(node->l))->ip.p; + node->rip.p = ((struct wp_symbol*)(node->r))->ip.p; break; default: yyerror("wp_ast_regvar: unknown node type %d\n", node->type); @@ -982,6 +992,61 @@ wp_ast_regvar (struct wp_node* node, char const* name, double* p) } } +void +wp_ast_regvar_gpu (struct wp_node* node, char const* name, int i) +{ + switch (node->type) + { + case WP_NUMBER: + break; + case WP_SYMBOL: + if (strcmp(name, ((struct wp_symbol*)node)->name) == 0) { + ((struct wp_symbol*)node)->ip.i = i; + } + break; + case WP_ADD: + case WP_SUB: + case WP_MUL: + case WP_DIV: + wp_ast_regvar_gpu(node->l, name, i); + wp_ast_regvar_gpu(node->r, name, i); + break; + case WP_NEG: + wp_ast_regvar_gpu(node->l, name, i); + break; + case WP_F1: + wp_ast_regvar_gpu(node->l, name, i); + break; + case WP_F2: + wp_ast_regvar_gpu(node->l, name, i); + wp_ast_regvar_gpu(node->r, name, i); + break; + case WP_ADD_VP: + case WP_SUB_VP: + case WP_MUL_VP: + case WP_DIV_VP: + wp_ast_regvar_gpu(node->r, name, i); + node->rip.i = ((struct wp_symbol*)(node->r))->ip.i; + break; + case WP_NEG_P: + wp_ast_regvar_gpu(node->l, name, i); + node->lvp.ip.i = ((struct wp_symbol*)(node->l))->ip.i; + break; + case WP_ADD_PP: + case WP_SUB_PP: + case WP_MUL_PP: + case WP_DIV_PP: + wp_ast_regvar_gpu(node->l, name, i); + wp_ast_regvar_gpu(node->r, name, i); + node->lvp.ip.i = ((struct wp_symbol*)(node->l))->ip.i; + node->rip.i = ((struct wp_symbol*)(node->r))->ip.i; + break; + default: + yyerror("wp_ast_regvar_gpu: unknown node type %d\n", node->type); + exit(1); + } +} + void wp_ast_setconst (struct wp_node* node, char const* name, double c) { switch (node->type) @@ -1040,6 +1105,12 @@ wp_parser_regvar (struct wp_parser* parser, char const* name, double* p) } void +wp_parser_regvar_gpu (struct wp_parser* parser, char const* name, int i) +{ + wp_ast_regvar_gpu(parser->ast, name, i); +} + +void wp_parser_setconst (struct wp_parser* parser, char const* name, double c) { wp_ast_setconst(parser->ast, name, c); diff --git a/Source/Parser/wp_parser_y.h b/Source/Parser/wp_parser_y.h index 4a3aeda40..8c9f8e4e4 100644 --- a/Source/Parser/wp_parser_y.h +++ b/Source/Parser/wp_parser_y.h @@ -1,6 +1,8 @@ #ifndef WP_PARSER_Y_H_ #define WP_PARSER_Y_H_ +#include <AMReX_GpuQualifiers.H> + #ifdef __cplusplus #include <cstdlib> extern "C" { @@ -73,17 +75,22 @@ enum wp_node_t { * wp_node_t type can be safely checked to determine their real type. */ -union wp_vp { - double v; +union wp_ip { + int i; double* p; }; +union wp_vp { + double v; + union wp_ip ip; +}; + struct wp_node { enum wp_node_t type; struct wp_node* l; struct wp_node* r; union wp_vp lvp; // After optimization, this may store left value/pointer. - double* rp; // this may store right pointer. + union wp_ip rip; // this may store right pointer. }; struct wp_number { @@ -94,7 +101,7 @@ struct wp_number { struct wp_symbol { enum wp_node_t type; char* name; - double* pointer; + union wp_ip ip; }; struct wp_f1 { /* Builtin functions with one argument */ @@ -124,6 +131,7 @@ struct wp_node* wp_newf1 (enum wp_f1_t ftype, struct wp_node* l); struct wp_node* wp_newf2 (enum wp_f2_t ftype, struct wp_node* l, struct wp_node* r); +AMREX_GPU_HOST_DEVICE void yyerror (char const *s, ...); /*******************************************************************/ @@ -146,6 +154,7 @@ struct wp_parser* wp_parser_dup (struct wp_parser* source); struct wp_node* wp_parser_ast_dup (struct wp_parser* parser, struct wp_node* src, int move); void wp_parser_regvar (struct wp_parser* parser, char const* name, double* p); +void wp_parser_regvar_gpu (struct wp_parser* parser, char const* name, int i); void wp_parser_setconst (struct wp_parser* parser, char const* name, double c); /* We need to walk the tree in these functions */ @@ -153,10 +162,11 @@ void wp_ast_optimize (struct wp_node* node); size_t wp_ast_size (struct wp_node* node); void wp_ast_print (struct wp_node* node); void wp_ast_regvar (struct wp_node* node, char const* name, double* p); +void wp_ast_regvar_gpu (struct wp_node* node, char const* name, int i); void wp_ast_setconst (struct wp_node* node, char const* name, double c); -double wp_call_f1 (enum wp_f1_t type, double a); -double wp_call_f2 (enum wp_f2_t type, double a, double b); +AMREX_GPU_HOST_DEVICE double wp_call_f1 (enum wp_f1_t type, double a); +AMREX_GPU_HOST_DEVICE double wp_call_f2 (enum wp_f2_t type, double a, double b); #ifdef __cplusplus } |