aboutsummaryrefslogtreecommitdiff
path: root/Source/Parser
diff options
context:
space:
mode:
Diffstat (limited to 'Source/Parser')
-rw-r--r--Source/Parser/GpuParser.H21
-rw-r--r--Source/Parser/GpuParser.cpp11
-rw-r--r--Source/Parser/Make.package1
-rw-r--r--Source/Parser/WarpXParser.H6
-rw-r--r--Source/Parser/WarpXParser.cpp6
-rw-r--r--Source/Parser/WarpXParserWrapper.H41
-rw-r--r--Source/Parser/wp_parser_c.h2
7 files changed, 83 insertions, 5 deletions
diff --git a/Source/Parser/GpuParser.H b/Source/Parser/GpuParser.H
index c158ee314..c6d870800 100644
--- a/Source/Parser/GpuParser.H
+++ b/Source/Parser/GpuParser.H
@@ -1,3 +1,10 @@
+/* Copyright 2019-2020 Maxence Thevenet, Revathi Jambunathan, Weiqun Zhang
+ *
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
#ifndef WARPX_GPU_PARSER_H_
#define WARPX_GPU_PARSER_H_
@@ -17,7 +24,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 +34,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 +58,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 +72,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..22fab6313 100644
--- a/Source/Parser/GpuParser.cpp
+++ b/Source/Parser/GpuParser.cpp
@@ -1,3 +1,10 @@
+/* Copyright 2019-2020 Maxence Thevenet, Revathi Jambunathan, Weiqun Zhang
+ *
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
#include <GpuParser.H>
GpuParser::GpuParser (WarpXParser const& wp)
@@ -16,6 +23,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 +36,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 +48,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 +60,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/Make.package b/Source/Parser/Make.package
index 5ce02cbda..15115c138 100644
--- a/Source/Parser/Make.package
+++ b/Source/Parser/Make.package
@@ -5,6 +5,7 @@ CEXE_sources += WarpXParser.cpp
CEXE_headers += WarpXParser.H
CEXE_headers += GpuParser.H
CEXE_sources += GpuParser.cpp
+CEXE_headers += WarpXParserWrapper.H
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 8c1d854d8..863b35fb8 100644
--- a/Source/Parser/WarpXParser.H
+++ b/Source/Parser/WarpXParser.H
@@ -1,3 +1,9 @@
+/* Copyright 2019 Weiqun Zhang
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
#ifndef WARPX_PARSER_H_
#define WARPX_PARSER_H_
diff --git a/Source/Parser/WarpXParser.cpp b/Source/Parser/WarpXParser.cpp
index ced536327..8c8be7ecb 100644
--- a/Source/Parser/WarpXParser.cpp
+++ b/Source/Parser/WarpXParser.cpp
@@ -1,3 +1,9 @@
+/* Copyright 2019 Weiqun Zhang
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
#include <algorithm>
#include "WarpXParser.H"
diff --git a/Source/Parser/WarpXParserWrapper.H b/Source/Parser/WarpXParserWrapper.H
new file mode 100644
index 000000000..2c76d97a3
--- /dev/null
+++ b/Source/Parser/WarpXParserWrapper.H
@@ -0,0 +1,41 @@
+/* Copyright 2020 Revathi Jambunathan
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
+#ifndef WARPX_PARSER_WRAPPER_H_
+#define WARPX_PARSER_WRAPPER_H_
+
+#include <WarpXParser.H>
+#include <AMReX_Gpu.H>
+#include <GpuParser.H>
+/**
+ * \brief
+ * The ParserWrapper struct is constructed to safely use the GpuParser class
+ * that can essentially be though of as a raw pointer. The GpuParser does
+ * not have an explicit destructor and the AddPlasma subroutines handle the GpuParser
+ * in a safe way. The ParserWrapper struct is used to avoid memory leak
+ * in the EB parser functions.
+ */
+struct ParserWrapper
+ : public amrex::Gpu::Managed
+{
+ ParserWrapper (WarpXParser const& a_parser) noexcept
+ : m_parser(a_parser) {}
+
+ ~ParserWrapper() {
+ m_parser.clear();
+ }
+
+ AMREX_GPU_HOST_DEVICE
+ amrex::Real
+ getField (amrex::Real x, amrex::Real y, amrex::Real z, amrex::Real t=0.0) const noexcept
+ {
+ return m_parser(x,y,z,t);
+ }
+
+ GpuParser m_parser;
+};
+
+#endif
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)
generated by cgit v1.2.3 (git 2.46.0) at 2025-08-31 19:55:02 +0000