From d59fa46d24417b67554132bc666e45886160bd09 Mon Sep 17 00:00:00 2001 From: Weiqun Zhang Date: Wed, 24 Jul 2019 19:43:18 -0700 Subject: Reimplement AddPlasma. Commits related to AddPlasma in hackathon branch are squashed into one. --- Source/Initialization/InjectorMomentum.cpp | 50 ++++++++++++++++++++++++++++++ 1 file changed, 50 insertions(+) create mode 100644 Source/Initialization/InjectorMomentum.cpp (limited to 'Source/Initialization/InjectorMomentum.cpp') diff --git a/Source/Initialization/InjectorMomentum.cpp b/Source/Initialization/InjectorMomentum.cpp new file mode 100644 index 000000000..4e483fc33 --- /dev/null +++ b/Source/Initialization/InjectorMomentum.cpp @@ -0,0 +1,50 @@ +#include + +using namespace amrex; + +InjectorMomentum::~InjectorMomentum () +{ + switch (type) + { + case Type::parser: + { + object.parser.m_ux_parser.clear(); + object.parser.m_uy_parser.clear(); + object.parser.m_uz_parser.clear(); + break; + } + case Type::custom: + { + object.custom.clear(); + break; + } + } +} + +std::size_t +InjectorMomentum::sharedMemoryNeeded () const noexcept +{ + switch (type) + { + case Type::parser: + { + return amrex::Gpu::numThreadsPerBlockParallelFor() * sizeof(double); + } + default: + return 0; + } +} + +bool +InjectorMomentum::useRandom () const noexcept +{ + switch (type) + { + case Type::gaussian: + { + return true; + } + default: + return false; + } +} -- cgit v1.2.3 From a0bb6f3ee7863c97292c99715854821ede3bcdb5 Mon Sep 17 00:00:00 2001 From: Weiqun Zhang Date: Sun, 28 Jul 2019 18:21:34 -0700 Subject: No need to resize random number seeds anymore --- Source/Initialization/InjectorDensity.H | 1 - Source/Initialization/InjectorMomentum.H | 2 -- Source/Initialization/InjectorMomentum.cpp | 13 ------------- Source/Initialization/InjectorPosition.H | 2 -- Source/Initialization/InjectorPosition.cpp | 19 ------------------- Source/Initialization/Make.package | 1 - Source/Initialization/PlasmaInjector.H | 5 ----- Source/Particles/PhysicalParticleContainer.cpp | 9 --------- 8 files changed, 52 deletions(-) delete mode 100644 Source/Initialization/InjectorPosition.cpp (limited to 'Source/Initialization/InjectorMomentum.cpp') diff --git a/Source/Initialization/InjectorDensity.H b/Source/Initialization/InjectorDensity.H index 61471433f..a27affd95 100644 --- a/Source/Initialization/InjectorDensity.H +++ b/Source/Initialization/InjectorDensity.H @@ -119,7 +119,6 @@ struct InjectorDensity ~InjectorDensity (); std::size_t sharedMemoryNeeded () const noexcept; - bool useRandom () const noexcept { return false; } AMREX_GPU_HOST_DEVICE amrex::Real diff --git a/Source/Initialization/InjectorMomentum.H b/Source/Initialization/InjectorMomentum.H index 2434503cc..baf1e0e05 100644 --- a/Source/Initialization/InjectorMomentum.H +++ b/Source/Initialization/InjectorMomentum.H @@ -123,8 +123,6 @@ struct InjectorMomentum std::size_t sharedMemoryNeeded () const noexcept; - bool useRandom () const noexcept; - AMREX_GPU_HOST_DEVICE amrex::XDim3 getMomentum (amrex::Real x, amrex::Real y, amrex::Real z) const noexcept diff --git a/Source/Initialization/InjectorMomentum.cpp b/Source/Initialization/InjectorMomentum.cpp index 4e483fc33..f9070b29c 100644 --- a/Source/Initialization/InjectorMomentum.cpp +++ b/Source/Initialization/InjectorMomentum.cpp @@ -35,16 +35,3 @@ InjectorMomentum::sharedMemoryNeeded () const noexcept } } -bool -InjectorMomentum::useRandom () const noexcept -{ - switch (type) - { - case Type::gaussian: - { - return true; - } - default: - return false; - } -} diff --git a/Source/Initialization/InjectorPosition.H b/Source/Initialization/InjectorPosition.H index e74345ae5..83ea8aaf2 100644 --- a/Source/Initialization/InjectorPosition.H +++ b/Source/Initialization/InjectorPosition.H @@ -72,8 +72,6 @@ struct InjectorPosition std::size_t sharedMemoryNeeded () const noexcept { return 0; } - bool useRandom () const noexcept; - AMREX_GPU_HOST_DEVICE amrex::XDim3 getPositionUnitBox (int i_part, int ref_fac=1) const noexcept diff --git a/Source/Initialization/InjectorPosition.cpp b/Source/Initialization/InjectorPosition.cpp deleted file mode 100644 index b77f26920..000000000 --- a/Source/Initialization/InjectorPosition.cpp +++ /dev/null @@ -1,19 +0,0 @@ -#include - -using namespace amrex; - -bool -InjectorPosition::useRandom () const noexcept -{ - switch (type) - { - case Type::random: - { - return true; - } - default: - { - return false; - } - }; -} diff --git a/Source/Initialization/Make.package b/Source/Initialization/Make.package index 3a924f207..2c6458b6d 100644 --- a/Source/Initialization/Make.package +++ b/Source/Initialization/Make.package @@ -4,7 +4,6 @@ CEXE_sources += PlasmaInjector.cpp CEXE_headers += PlasmaInjector.H CEXE_headers += InjectorPosition.H -CEXE_sources += InjectorPosition.cpp CEXE_headers += InjectorDensity.H CEXE_sources += InjectorDensity.cpp diff --git a/Source/Initialization/PlasmaInjector.H b/Source/Initialization/PlasmaInjector.H index d5c064192..c0d74ab0c 100644 --- a/Source/Initialization/PlasmaInjector.H +++ b/Source/Initialization/PlasmaInjector.H @@ -71,11 +71,6 @@ public: InjectorDensity* getInjectorDensity (); InjectorMomentum* getInjectorMomentum (); - bool useRandom () const noexcept { - return inj_pos->useRandom() or inj_rho->useRandom() - or inj_mom->useRandom(); - } - std::size_t sharedMemoryNeeded () const noexcept { return amrex::max(inj_pos->sharedMemoryNeeded(), diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 31ffadddf..b205ba274 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -430,15 +430,6 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) overlap_realbox.lo(1), overlap_realbox.lo(2))}; -#ifdef WARPX_RZ - { -#else - if (plasma_injector->useRandom()) { -#endif - amrex::Gpu::streamSynchronize(); - amrex::CheckSeedArraySizeAndResize(max_new_particles); - } - std::size_t shared_mem_bytes = plasma_injector->sharedMemoryNeeded(); int lrrfac = rrfac; -- cgit v1.2.3 From 801665ed402def24b9953755b8b0befa196fa2bb Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Wed, 31 Jul 2019 10:51:41 -0700 Subject: fix memory needed in shared memory --- Source/Initialization/InjectorDensity.cpp | 2 +- Source/Initialization/InjectorMomentum.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) (limited to 'Source/Initialization/InjectorMomentum.cpp') diff --git a/Source/Initialization/InjectorDensity.cpp b/Source/Initialization/InjectorDensity.cpp index f796180fd..78846215b 100644 --- a/Source/Initialization/InjectorDensity.cpp +++ b/Source/Initialization/InjectorDensity.cpp @@ -31,7 +31,7 @@ InjectorDensity::sharedMemoryNeeded () const noexcept { case Type::parser: { - return amrex::Gpu::numThreadsPerBlockParallelFor() * sizeof(double); + return amrex::Gpu::numThreadsPerBlockParallelFor() * sizeof(double) * 3; } default: return 0; diff --git a/Source/Initialization/InjectorMomentum.cpp b/Source/Initialization/InjectorMomentum.cpp index f9070b29c..e43eece30 100644 --- a/Source/Initialization/InjectorMomentum.cpp +++ b/Source/Initialization/InjectorMomentum.cpp @@ -28,7 +28,7 @@ InjectorMomentum::sharedMemoryNeeded () const noexcept { case Type::parser: { - return amrex::Gpu::numThreadsPerBlockParallelFor() * sizeof(double); + return amrex::Gpu::numThreadsPerBlockParallelFor() * sizeof(double) * 3; } default: return 0; -- cgit v1.2.3 From 804620cdc9d8510c77341cc5741d96339c064f1e Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Wed, 31 Jul 2019 11:04:34 -0700 Subject: comments on all injectors --- Source/Initialization/InjectorDensity.H | 2 +- Source/Initialization/InjectorDensity.cpp | 6 ++++++ Source/Initialization/InjectorMomentum.H | 30 ++++++++++++++++++++++++++++++ Source/Initialization/InjectorMomentum.cpp | 3 +++ 4 files changed, 40 insertions(+), 1 deletion(-) (limited to 'Source/Initialization/InjectorMomentum.cpp') diff --git a/Source/Initialization/InjectorDensity.H b/Source/Initialization/InjectorDensity.H index 7e04515d2..b7f5c26eb 100644 --- a/Source/Initialization/InjectorDensity.H +++ b/Source/Initialization/InjectorDensity.H @@ -144,7 +144,7 @@ struct InjectorDensity std::size_t sharedMemoryNeeded () const noexcept; - // call getPositionUnitBox from the object stored in the union + // call getDensity from the object stored in the union // (the union is called Object, and the instance is called object). AMREX_GPU_HOST_DEVICE amrex::Real diff --git a/Source/Initialization/InjectorDensity.cpp b/Source/Initialization/InjectorDensity.cpp index 78846215b..7fed85b75 100644 --- a/Source/Initialization/InjectorDensity.cpp +++ b/Source/Initialization/InjectorDensity.cpp @@ -24,6 +24,7 @@ InjectorDensity::~InjectorDensity () } } +// Compute the amount of memory needed in GPU Shared Memory. std::size_t InjectorDensity::sharedMemoryNeeded () const noexcept { @@ -31,6 +32,8 @@ InjectorDensity::sharedMemoryNeeded () const noexcept { case Type::parser: { + // For parser injector, the 3D position of each particle + // is stored in shared memory. return amrex::Gpu::numThreadsPerBlockParallelFor() * sizeof(double) * 3; } default: @@ -45,6 +48,8 @@ InjectorDensityPredefined::InjectorDensityPredefined ( ParmParse pp(a_species_name); std::vector v; + // Read parameters for the predefined plasma profile, + // and store them in managed memory pp.getarr("predefined_profile_params", v); p = static_cast (amrex::The_Managed_Arena()->alloc(sizeof(amrex::Real)*v.size())); @@ -52,6 +57,7 @@ InjectorDensityPredefined::InjectorDensityPredefined ( p[i] = v[i]; } + // Parse predefined profile name, and update member variable profile. std::string which_profile_s; pp.query("predefined_profile_name", which_profile_s); std::transform(which_profile_s.begin(), which_profile_s.end(), diff --git a/Source/Initialization/InjectorMomentum.H b/Source/Initialization/InjectorMomentum.H index baf1e0e05..399ee7759 100644 --- a/Source/Initialization/InjectorMomentum.H +++ b/Source/Initialization/InjectorMomentum.H @@ -6,6 +6,7 @@ #include #include +// struct whose getMomentum returns constant momentum. struct InjectorMomentumConstant { InjectorMomentumConstant (amrex::Real a_ux, amrex::Real a_uy, amrex::Real a_uz) noexcept @@ -21,6 +22,8 @@ private: amrex::Real m_ux, m_uy, m_uz; }; +// struct whose getMomentum returns momentum for 1 particle, from random +// gaussian distribution. struct InjectorMomentumGaussian { InjectorMomentumGaussian (amrex::Real a_ux_m, amrex::Real a_uy_m, @@ -43,6 +46,8 @@ private: amrex::Real m_ux_th, m_uy_th, m_uz_th; }; +// struct whose getMomentum returns momentum for 1 particle, for +// radial expansion struct InjectorMomentumRadialExpansion { InjectorMomentumRadialExpansion (amrex::Real a_u_over_r) noexcept @@ -60,6 +65,7 @@ private: amrex::Real u_over_r; }; +// struct whose getMomentumm returns local momentum computed from parser. struct InjectorMomentumParser { InjectorMomentumParser (WarpXParser const& a_ux_parser, @@ -78,15 +84,29 @@ struct InjectorMomentumParser GpuParser m_ux_parser, m_uy_parser, m_uz_parser; }; +// Base struct for momentum injector. +// InjectorMomentum contains a union (called Object) that holds any one +// instance of: +// - InjectorMomentumConstant : to generate constant density; +// - InjectorMomentumGaussian : to generate gaussian distribution; +// - InjectorMomentumRadialExpansion: to generate radial expansion; +// - InjectorMomentumParser : to generate momentum from parser; +// The choice is made at runtime, depending in the constructor called. +// This mimics virtual functions, except the struct is stored in managed memory +// and member functions are made __host__ __device__ to run on CPU and GPU. +// This struct inherits from amrex::Gpu::Managed to provide new and delete +// operators in managed memory when running on GPU. Nothing special on CPU. struct InjectorMomentum : public amrex::Gpu::Managed { + // This constructor stores a InjectorMomentumConstant in union object. InjectorMomentum (InjectorMomentumConstant* t, amrex::Real a_ux, amrex::Real a_uy, amrex::Real a_uz) : type(Type::constant), object(t, a_ux, a_uy, a_uz) { } + // This constructor stores a InjectorMomentumParser in union object. InjectorMomentum (InjectorMomentumParser* t, WarpXParser const& a_ux_parser, WarpXParser const& a_uy_parser, @@ -95,6 +115,7 @@ struct InjectorMomentum object(t, a_ux_parser, a_uy_parser, a_uz_parser) { } + // This constructor stores a InjectorMomentumGaussian in union object. InjectorMomentum (InjectorMomentumGaussian* t, amrex::Real a_ux_m, amrex::Real a_uy_m, amrex::Real a_uz_m, amrex::Real a_ux_th, amrex::Real a_uy_th, amrex::Real a_uz_th) @@ -102,18 +123,22 @@ struct InjectorMomentum object(t,a_ux_m,a_uy_m,a_uz_m,a_ux_th,a_uy_th,a_uz_th) { } + // This constructor stores a InjectorMomentumCustom in union object. InjectorMomentum (InjectorMomentumCustom* t, std::string const& a_species_name) : type(Type::custom), object(t, a_species_name) { } + // This constructor stores a InjectorMomentumRadialExpansion in union object. InjectorMomentum (InjectorMomentumRadialExpansion* t, amrex::Real u_over_r) : type(Type::radial_expansion), object(t, u_over_r) { } + // Explicitly prevent the compiler from generating copy constructors + // and copy assignment operators. InjectorMomentum (InjectorMomentum const&) = delete; InjectorMomentum (InjectorMomentum&&) = delete; void operator= (InjectorMomentum const&) = delete; @@ -123,6 +148,8 @@ struct InjectorMomentum std::size_t sharedMemoryNeeded () const noexcept; + // call getMomentum from the object stored in the union + // (the union is called Object, and the instance is called object). AMREX_GPU_HOST_DEVICE amrex::XDim3 getMomentum (amrex::Real x, amrex::Real y, amrex::Real z) const noexcept @@ -161,6 +188,9 @@ private: enum struct Type { constant, custom, gaussian, radial_expansion, parser }; Type type; + // An instance of union Object constructs and stores any one of + // the objects declared (constant or custom or gaussian or + // radial_expansion or parser). union Object { Object (InjectorMomentumConstant*, amrex::Real a_ux, amrex::Real a_uy, amrex::Real a_uz) noexcept diff --git a/Source/Initialization/InjectorMomentum.cpp b/Source/Initialization/InjectorMomentum.cpp index e43eece30..a197b5bef 100644 --- a/Source/Initialization/InjectorMomentum.cpp +++ b/Source/Initialization/InjectorMomentum.cpp @@ -21,6 +21,7 @@ InjectorMomentum::~InjectorMomentum () } } +// Compute the amount of memory needed in GPU Shared Memory. std::size_t InjectorMomentum::sharedMemoryNeeded () const noexcept { @@ -28,6 +29,8 @@ InjectorMomentum::sharedMemoryNeeded () const noexcept { case Type::parser: { + // For parser injector, the 3D position of each particle + // is stored in shared memory. return amrex::Gpu::numThreadsPerBlockParallelFor() * sizeof(double) * 3; } default: -- cgit v1.2.3