diff options
Diffstat (limited to 'Source/Parser')
| -rw-r--r-- | Source/Parser/GpuParser.H | 14 | ||||
| -rw-r--r-- | Source/Parser/GpuParser.cpp | 4 | ||||
| -rw-r--r-- | Source/Parser/WarpXParserWrapper.H | 2 | ||||
| -rw-r--r-- | Source/Parser/wp_parser_c.h | 2 |
4 files changed, 16 insertions, 6 deletions
diff --git a/Source/Parser/GpuParser.H b/Source/Parser/GpuParser.H index c158ee314..ff855d275 100644 --- a/Source/Parser/GpuParser.H +++ b/Source/Parser/GpuParser.H @@ -17,7 +17,7 @@ public: AMREX_GPU_HOST_DEVICE amrex::Real - operator() (amrex::Real x, amrex::Real y, amrex::Real z) const noexcept + operator() (amrex::Real x, amrex::Real y, amrex::Real z, amrex::Real t=0.0) const noexcept { #ifdef AMREX_USE_GPU @@ -27,15 +27,17 @@ public: amrex::Gpu::SharedMemory<amrex::Real> gsm; amrex::Real* 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; + p[tid*4] = x; + p[tid*4+1] = y; + p[tid*4+2] = z; + p[tid*4+3] = t; 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; + m_t = t; return wp_ast_eval(m_cpu_parser.ast); #endif @@ -49,10 +51,12 @@ public: m_var[tid].x = x; m_var[tid].y = y; m_var[tid].z = z; + m_t[tid] = t; return wp_ast_eval(m_parser[tid]->ast); #endif } + private: #ifdef AMREX_USE_GPU @@ -61,10 +65,12 @@ private: // Copy of the parser running on __host__ struct wp_parser m_cpu_parser; mutable amrex::XDim3 m_var; + mutable amrex::Real m_t; #else // Only one parser struct wp_parser** m_parser; mutable amrex::XDim3* m_var; + mutable amrex::Real* m_t; int nthreads; #endif }; diff --git a/Source/Parser/GpuParser.cpp b/Source/Parser/GpuParser.cpp index 5078b498b..ba904666b 100644 --- a/Source/Parser/GpuParser.cpp +++ b/Source/Parser/GpuParser.cpp @@ -16,6 +16,7 @@ GpuParser::GpuParser (WarpXParser const& wp) 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); + wp_parser_regvar_gpu(&m_gpu_parser, "t", 3); // Initialize CPU parser: allocate memory in CUDA managed memory, // copy all data needed on CPU to m_cpu_parser @@ -28,6 +29,7 @@ GpuParser::GpuParser (WarpXParser const& wp) 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)); + wp_parser_regvar(&m_cpu_parser, "t", &(m_t)); #else // not defined AMREX_USE_GPU @@ -39,6 +41,7 @@ GpuParser::GpuParser (WarpXParser const& wp) m_parser = ::new struct wp_parser*[nthreads]; m_var = ::new amrex::XDim3[nthreads]; + m_t = ::new amrex::Real[nthreads]; for (int tid = 0; tid < nthreads; ++tid) { @@ -50,6 +53,7 @@ GpuParser::GpuParser (WarpXParser const& wp) 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)); + wp_parser_regvar(m_parser[tid], "t", &(m_t[tid])); } #endif // AMREX_USE_GPU diff --git a/Source/Parser/WarpXParserWrapper.H b/Source/Parser/WarpXParserWrapper.H index 2dd7f72c7..2a4ff6fd2 100644 --- a/Source/Parser/WarpXParserWrapper.H +++ b/Source/Parser/WarpXParserWrapper.H @@ -24,7 +24,7 @@ struct ParserWrapper AMREX_GPU_HOST_DEVICE amrex::Real - getField (amrex::Real x, amrex::Real y, amrex::Real z) const noexcept + getField (amrex::Real x, amrex::Real y, amrex::Real z, amrex::Real t=0.0) const noexcept { return m_parser(x,y,z); } diff --git a/Source/Parser/wp_parser_c.h b/Source/Parser/wp_parser_c.h index 970d6b355..2cf0e2c00 100644 --- a/Source/Parser/wp_parser_c.h +++ b/Source/Parser/wp_parser_c.h @@ -30,7 +30,7 @@ wp_ast_eval (struct wp_node* node) #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*3; + amrex_real* x = extern_xyz + tid*4; // parser assumes 4 independent variables (x,y,z,t) #endif switch (node->type) |