/* Copyright 2020 Revathi Jambunathan, Weiqun Zhang * * This file is part of WarpX. * * License: BSD-3-Clause-LBNL */ #ifndef WARPX_PARSER_WRAPPER_H_ #define WARPX_PARSER_WRAPPER_H_ #include "Parser/WarpXParser.H" #include "Parser/GpuParser.H" #include #include /** * \brief * The HostDeviceParser struct is non-owning and is suitable for being * value captured by device lamba. */ template struct HostDeviceParser { template AMREX_GPU_HOST_DEVICE std::enable_if_t::value, amrex::Real> operator() (Ts... var) const noexcept { #ifdef AMREX_USE_GPU #if AMREX_DEVICE_COMPILE // WarpX compiled for GPU, function compiled for __device__ amrex::GpuArray l_var{var...}; return wp_ast_eval<0>(m_gpu_parser_ast, l_var.data()); #else // WarpX compiled for GPU, function compiled for __host__ return (*m_gpu_parser)(var...); #endif #else // WarpX compiled for CPU return (*m_gpu_parser)(var...); #endif } #ifdef AMREX_USE_GPU struct wp_node * m_gpu_parser_ast = nullptr; #endif GpuParser const* m_gpu_parser = nullptr; }; /** * \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. */ template struct ParserWrapper : public GpuParser { using GpuParser::GpuParser; ParserWrapper (ParserWrapper const&) = delete; void operator= (ParserWrapper const&) = delete; ~ParserWrapper() { GpuParser::clear(); } void operator() () const = delete; // Hide GpuParser::operator() HostDeviceParser getParser () const { #ifdef AMREX_USE_GPU return HostDeviceParser{this->m_gpu_parser_ast, static_cast const*>(this)}; #else return HostDeviceParser{static_cast const*>(this)}; #endif } }; template HostDeviceParser getParser (ParserWrapper const* parser_wrapper) { return parser_wrapper ? parser_wrapper->getParser() : HostDeviceParser{}; } template HostDeviceParser getParser (std::unique_ptr > const& parser_wrapper) { return parser_wrapper ? parser_wrapper->getParser() : HostDeviceParser{}; } #endif