diff options
author | 2020-01-10 20:57:16 -0800 | |
---|---|---|
committer | 2020-01-10 21:36:51 -0800 | |
commit | 6166de191d0a2b468ea56bacec481e1ff27fcae1 (patch) | |
tree | 63994693f1df766d98ff4b4f045740f0efb3fa04 /Source/FieldSolver/FiniteDifferenceSolver | |
parent | c049e742014fca3bba78534b7afcc6bcd477aa0e (diff) | |
download | WarpX-6166de191d0a2b468ea56bacec481e1ff27fcae1.tar.gz WarpX-6166de191d0a2b468ea56bacec481e1ff27fcae1.tar.zst WarpX-6166de191d0a2b468ea56bacec481e1ff27fcae1.zip |
Start fixing compilation errors
Diffstat (limited to 'Source/FieldSolver/FiniteDifferenceSolver')
4 files changed, 63 insertions, 55 deletions
diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp index 33253c986..a70548817 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp @@ -1,28 +1,32 @@ +#include <WarpXAlgorithmSelection.H> #include<FiniteDifferenceAlgorithms/YeeAlgorithm.H> #include<FiniteDifferenceSolver.H> +#include<AMReX_Gpu.H> -FiniteDifferenceSolver::EvolveB ( VectorField Bfield, - ConstVectorField Efield, - amrex::Real dt ) { +using namespace amrex; + +void FiniteDifferenceSolver::EvolveB ( VectorField& Bfield, + VectorField& Efield, + amrex::Real dt ) { // Select algorithm (The choice of algorithm is a runtime option, // but we compile code for each algorithm, using templates) - if (fdtd_algo == MaxwellSolverAlgo::Yee){ + if (m_fdtd_algo == MaxwellSolverAlgo::Yee){ EvolveBwithAlgo <YeeAlgorithm> ( Bfield, Efield, dt ); // } else if (fdtd_algo == MaxwellSolverAlgo::CKC) { // EvolveBwithAlgo <CKCAlgorithm> ( Bfield, Efield, dt ); } else { amrex::Abort("Unknown algorithm"); } -) +} template<typename algo> -FiniteDifferenceSolver::EvolveBwithAlgo ( VectorField Bfield, - ConstVectorField Efield, - amrex::Real dt ) { +void FiniteDifferenceSolver::EvolveBwithAlgo ( VectorField& Bfield, + VectorField& Efield, + amrex::Real dt ) { // Loop through the grids, and over the tiles within each grid #ifdef _OPENMP -#pragma omp parallel if (Gpu::notInLaunchRegion()) +#pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*Bfield[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) { @@ -66,4 +70,4 @@ FiniteDifferenceSolver::EvolveBwithAlgo ( VectorField Bfield, } -}; +} diff --git a/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/YeeAlgorithm.H b/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/YeeAlgorithm.H index 6ea2f440d..7e32aebb4 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/YeeAlgorithm.H +++ b/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/YeeAlgorithm.H @@ -1,9 +1,14 @@ #ifndef WARPX_FINITE_DIFFERENCE_ALGORITHM_YEE_H_ #define WARPX_FINITE_DIFFERENCE_ALGORITHM_YEE_H_ +#include <AMReX_REAL.H> +#include <AMReX_Array4.H> +#include <AMReX_Gpu.H> + struct YeeAlgorithm { - void InitializeStencilCoefficients( std::array<Real,3> cell_size, + static void InitializeStencilCoefficients( + std::array<amrex::Real,3> cell_size, amrex::Gpu::ManagedVector<amrex::Real> stencil_coefs_x, amrex::Gpu::ManagedVector<amrex::Real> stencil_coefs_y, amrex::Gpu::ManagedVector<amrex::Real> stencil_coefs_z ) { @@ -18,30 +23,33 @@ struct YeeAlgorithm { } AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE - amrex::Real UpwardDx( - Array4 const F, int const i, int const j, int const k, - Real const* coefs_x, int const n_coefs_x ) { + static amrex::Real UpwardDx( + amrex::Array4<amrex::Real> const F, + int const i, int const j, int const k, + amrex::Real const* coefs_x, int const n_coefs_x ) { - amrex::Real const inv_dx = coefs_x[0]; + amrex::Real inv_dx = coefs_x[0]; return inv_dx*( F(i+1,j,k) - F(i,j,k) ); }; AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE - amrex::Real DownwardDx( - Array4 const F, int const i, int const j, int const k, - Real const* coefs_x, int const n_coefs_x ) { + static amrex::Real DownwardDx( + amrex::Array4<amrex::Real> const F, + int const i, int const j, int const k, + amrex::Real const* coefs_x, int const n_coefs_x ) { - amrex::Real const inv_dx = coefs_x[0]; + amrex::Real inv_dx = coefs_x[0]; return inv_dx*( F(i,j,k) - F(i-1,j,k) ); }; AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE - amrex::Real UpwardDy( - Array4 const F, int const i, int const j, int const k, - Real const* coefs_y, int const n_coefs_y ) { + static amrex::Real UpwardDy( + amrex::Array4<amrex::Real> const F, + int const i, int const j, int const k, + amrex::Real const* coefs_y, int const n_coefs_y ) { #if defined WARPX_DIM_3D - amrex::Real const inv_dy = coefs_y[0]; + amrex::Real inv_dy = coefs_y[0]; return inv_dy*( F(i,j+1,k) - F(i,j,k) ); #elif (defined WARPX_DIM_XZ) return 0; // 2D Cartesian: derivative along y is 0 @@ -49,12 +57,13 @@ struct YeeAlgorithm { }; AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE - amrex::Real DownwardDy( - Array4 const F, int const i, int const j, int const k, - Real const* coefs_y, int const n_coefs_y ) { + static amrex::Real DownwardDy( + amrex::Array4<amrex::Real> const F, + int const i, int const j, int const k, + amrex::Real const* coefs_y, int const n_coefs_y ) { #if defined WARPX_DIM_3D - amrex::Real const inv_dy = coefs_y[0]; + amrex::Real inv_dy = coefs_y[0]; return inv_dy*( F(i,j,k) - F(i,j-1,k) ); #elif (defined WARPX_DIM_XZ) return 0; // 2D Cartesian: derivative along y is 0 @@ -62,11 +71,12 @@ struct YeeAlgorithm { }; AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE - amrex::Real UpwardDz( - Array4 const F, int const i, int const j, int const k, - Real const* coefs_z, int const n_coefs_z ) { + static amrex::Real UpwardDz( + amrex::Array4<amrex::Real> const F, + int const i, int const j, int const k, + amrex::Real const* coefs_z, int const n_coefs_z ) { - amrex::Real const inv_dz = coefs_z[0]; + amrex::Real inv_dz = coefs_z[0]; #if defined WARPX_DIM_3D return inv_dz*( F(i,j,k+1) - F(i,j,k) ); #elif (defined WARPX_DIM_XZ) @@ -74,19 +84,6 @@ struct YeeAlgorithm { #endif }; - AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE - amrex::Real DownwardDz( - Array4 const F, int const i, int const j, int const k, - Real const* coefs_z, int const n_coefs_z ) { - - amrex::Real const inv_dz = coefs_z[0]; - #if defined WARPX_DIM_3D - return inv_dz*( F(i,j,k) - F(i,j,k-1) ); - #elif (defined WARPX_DIM_XZ) - return inv_dz*( F(i,j,k) - F(i,j-1,k) ); - #endif - }; - -} +}; #endif // WARPX_FINITE_DIFFERENCE_ALGORITHM_YEE_H_ diff --git a/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H b/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H index 1aa9c7e38..5496adff9 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H +++ b/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H @@ -1,6 +1,8 @@ #ifndef WARPX_FINITE_DIFFERENCE_SOLVER_H_ #define WARPX_FINITE_DIFFERENCE_SOLVER_H_ +#include <AMReX_MultiFab.H> +#include <WarpXAlgorithmSelection.H> #include<FiniteDifferenceAlgorithms/YeeAlgorithm.H> /** @@ -13,13 +15,13 @@ class FiniteDifferenceSolver public: using VectorField = std::array< std::unique_ptr<amrex::MultiFab>, 3 >; - using ConstVectorField = std::array< std::unique_ptr<amrex::MultiFab const>, 3 >; // Constructor - void FiniteDifferenceSolver::FiniteDifferenceSolver ( std::array<Real,3> cell_size ) { + FiniteDifferenceSolver ( int fdtd_algo, + std::array<amrex::Real,3> cell_size ) { // Register the type of finite-difference algorithm - fdtd_algo = WarpX::maxwell_fdtd_solver_id; + m_fdtd_algo = fdtd_algo; // Calculate coefficients of finite-difference stencil if (fdtd_algo == MaxwellSolverAlgo::Yee){ @@ -33,21 +35,21 @@ class FiniteDifferenceSolver } }; - void EvolveB ( VectorField Bfield, - ConstVectorField Efield, - amrex::Real dt ) const; + void EvolveB ( VectorField& Bfield, + VectorField& Efield, + amrex::Real dt ); private: - int fdtd_algo; + int m_fdtd_algo; amrex::Gpu::ManagedVector<amrex::Real> stencil_coefs_x; amrex::Gpu::ManagedVector<amrex::Real> stencil_coefs_y; amrex::Gpu::ManagedVector<amrex::Real> stencil_coefs_z; template< typename fdtd_algo > - void EvolveBwithAlgo ( VectorField Bfield, - ConstVectorField Efield, - amrex::Real dt ) const; + void EvolveBwithAlgo ( VectorField& Bfield, + VectorField& Efield, + amrex::Real dt ); }; diff --git a/Source/FieldSolver/FiniteDifferenceSolver/Make.package b/Source/FieldSolver/FiniteDifferenceSolver/Make.package new file mode 100644 index 000000000..27ac23846 --- /dev/null +++ b/Source/FieldSolver/FiniteDifferenceSolver/Make.package @@ -0,0 +1,5 @@ +CEXE_headers += FiniteDifferenceSolver.H +CEXE_sources += EvolveB.cpp + +INCLUDE_LOCATIONS += $(WARPX_HOME)/Source/FieldSolver/FiniteDifferenceSolver +VPATH_LOCATIONS += $(WARPX_HOME)/Source/FieldSolver/FiniteDifferenceSolver |