diff options
Diffstat (limited to 'Source/Parser/GpuParser.H')
| -rw-r--r-- | Source/Parser/GpuParser.H | 14 |
1 files changed, 10 insertions, 4 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 }; |