diff --git a/.github/workflows/intel.yml b/.github/workflows/intel.yml index 3b1d6b546a4..8a7867fcbc8 100644 --- a/.github/workflows/intel.yml +++ b/.github/workflows/intel.yml @@ -115,7 +115,7 @@ jobs: cmake -S . -B build_sp \ -DCMAKE_CXX_FLAGS_RELEASE="-O1 -DNDEBUG" \ -DCMAKE_VERBOSE_MAKEFILE=ON \ - -DWarpX_EB=ON \ + -DWarpX_EB=OFF \ -DWarpX_PYTHON=ON \ -DWarpX_MPI=OFF \ -DWarpX_OPENPMD=ON \ diff --git a/.github/workflows/macos.yml b/.github/workflows/macos.yml index 00ac8f06b5d..dfd32f459f0 100644 --- a/.github/workflows/macos.yml +++ b/.github/workflows/macos.yml @@ -60,7 +60,7 @@ jobs: cmake -S . -B build_dp \ -DCMAKE_VERBOSE_MAKEFILE=ON \ - -DWarpX_EB=ON \ + -DWarpX_EB=OFF \ -DWarpX_OPENPMD=ON \ -DWarpX_openpmd_internal=OFF cmake --build build_dp -j 3 @@ -68,7 +68,7 @@ jobs: cmake -S . -B build_sp \ -DCMAKE_VERBOSE_MAKEFILE=ON \ -DPython_EXECUTABLE=$(which python3) \ - -DWarpX_EB=ON \ + -DWarpX_EB=OFF \ -DWarpX_PYTHON=ON \ -DWarpX_OPENPMD=ON \ -DWarpX_openpmd_internal=OFF \ diff --git a/.github/workflows/windows.yml b/.github/workflows/windows.yml index 2ef74cdb7f9..3c0faaf3636 100644 --- a/.github/workflows/windows.yml +++ b/.github/workflows/windows.yml @@ -91,7 +91,7 @@ jobs: -DCMAKE_BUILD_TYPE=Release ^ -DCMAKE_VERBOSE_MAKEFILE=ON ^ -DWarpX_COMPUTE=OMP ^ - -DWarpX_EB=ON ^ + -DWarpX_EB=OFF ^ -DWarpX_PYTHON=ON ^ -DWarpX_MPI=OFF ^ -DWarpX_OPENPMD=ON diff --git a/Docs/source/install/dependencies.rst b/Docs/source/install/dependencies.rst index 3bab32b7502..fadd38699e2 100644 --- a/Docs/source/install/dependencies.rst +++ b/Docs/source/install/dependencies.rst @@ -228,7 +228,7 @@ The `Advanced Package Tool (APT) ` .. code-block:: bash sudo apt update - sudo apt install build-essential ccache cmake g++ git libfftw3-mpi-dev libfftw3-dev libhdf5-openmpi-dev libopenmpi-dev pkg-config python3 python3-matplotlib python3-mpi4py python3-numpy python3-pandas python3-pip python3-scipy python3-venv + sudo apt install build-essential ccache cmake g++ git libfftw3-mpi-dev libfftw3-dev libhdf5-openmpi-dev libopenmpi-dev pkg-config python3 python3-dev python3-matplotlib python3-mpi4py python3-numpy python3-pandas python3-pip python3-scipy python3-venv # optional: # for CUDA, either install @@ -244,7 +244,7 @@ The `Advanced Package Tool (APT) ` .. code-block:: bash sudo apt update - sudo apt install build-essential ccache cmake g++ git libfftw3-dev libfftw3-dev libhdf5-dev pkg-config python3 python3-matplotlib python3-numpy python3-pandas python3-pip python3-scipy python3-venv + sudo apt install build-essential ccache cmake g++ git libfftw3-dev libfftw3-dev libhdf5-dev pkg-config python3 python3-dev python3-matplotlib python3-numpy python3-pandas python3-pip python3-scipy python3-venv # optional: # for CUDA, either install diff --git a/Examples/Tests/collision/analysis_collision_2d.py b/Examples/Tests/collision/analysis_collision_2d.py index 29482909b2e..c11b4cd9463 100755 --- a/Examples/Tests/collision/analysis_collision_2d.py +++ b/Examples/Tests/collision/analysis_collision_2d.py @@ -114,4 +114,4 @@ dim, species_name) test_name = os.path.split(os.getcwd())[1] -checksumAPI.evaluate_checksum(test_name, fn) +checksumAPI.evaluate_checksum(test_name, fn, rtol=1e-1) diff --git a/Examples/Tests/collision/analysis_collision_3d.py b/Examples/Tests/collision/analysis_collision_3d.py index 86a434caab2..98fd76cd704 100755 --- a/Examples/Tests/collision/analysis_collision_3d.py +++ b/Examples/Tests/collision/analysis_collision_3d.py @@ -111,4 +111,4 @@ dim, species_name) test_name = os.path.split(os.getcwd())[1] -checksumAPI.evaluate_checksum(test_name, fn) +checksumAPI.evaluate_checksum(test_name, fn, rtol=1e-1) diff --git a/Regression/WarpX-tests.ini b/Regression/WarpX-tests.ini index 6da53c6b760..f92ce4ac0e4 100644 --- a/Regression/WarpX-tests.ini +++ b/Regression/WarpX-tests.ini @@ -2436,21 +2436,6 @@ useOMP = 1 numthreads = 1 analysisRoutine = Examples/Tests/pml/analysis_pml_yee.py -[pml_x_yee_eb] -buildDir = . -inputFile = Examples/Tests/pml/inputs_2d -runtime_params = algo.maxwell_solver=yee chk.file_prefix=pml_x_yee_eb_chk chk.file_min_digits=5 -dim = 2 -addToCompileString = USE_EB=TRUE -cmakeSetupOpts = -DWarpX_DIMS=2 -DWarpX_EB=ON -restartTest = 1 -restartFileNum = 150 -useMPI = 1 -numprocs = 2 -useOMP = 1 -numthreads = 1 -analysisRoutine = Examples/Tests/pml/analysis_pml_yee.py - [projection_divb_cleaner_rz] buildDir = . inputFile = Examples/Tests/projection_divb_cleaner/inputs_rz diff --git a/Source/EmbeddedBoundary/CMakeLists.txt b/Source/EmbeddedBoundary/CMakeLists.txt index 3fa0ea0228f..a102087f8c9 100644 --- a/Source/EmbeddedBoundary/CMakeLists.txt +++ b/Source/EmbeddedBoundary/CMakeLists.txt @@ -5,5 +5,6 @@ foreach(D IN LISTS WarpX_DIMS) WarpXInitEB.cpp WarpXFaceExtensions.cpp WarpXFaceInfoBox.H + Enabled.cpp ) endforeach() diff --git a/Source/EmbeddedBoundary/Enabled.H b/Source/EmbeddedBoundary/Enabled.H new file mode 100644 index 00000000000..90ea5f35101 --- /dev/null +++ b/Source/EmbeddedBoundary/Enabled.H @@ -0,0 +1,18 @@ +/* Copyright 2024 Axel Huebl + * + * This file is part of WarpX. + * + * License: BSD-3-Clause-LBNL + */ +#ifndef WARPX_EB_ENABLED_H_ +#define WARPX_EB_ENABLED_H_ + +#include + +namespace EB +{ + /** Are embedded boundaries enabled? */ + bool enabled (); + +} // namespace EB +#endif // WARPX_EB_ENABLED_H_ diff --git a/Source/EmbeddedBoundary/Enabled.cpp b/Source/EmbeddedBoundary/Enabled.cpp new file mode 100644 index 00000000000..935c51c7c0f --- /dev/null +++ b/Source/EmbeddedBoundary/Enabled.cpp @@ -0,0 +1,39 @@ +/* Copyright 2024 Axel Huebl + * + * This file is part of WarpX. + * + * License: BSD-3-Clause-LBNL + */ +#include "Enabled.H" + +#ifdef AMREX_USE_EB +#include +#endif +#if defined(AMREX_USE_EB) && defined(WARPX_DIM_RZ) +#include +#endif + + +namespace EB +{ + bool enabled () + { +#ifndef AMREX_USE_EB + return false; +#else + amrex::ParmParse const pp_warpx("warpx"); + amrex::ParmParse const pp_eb2("eb2"); + + // test various runtime options to enable EBs + std::string eb_implicit_function; + bool eb_enabled = pp_warpx.query("eb_implicit_function", eb_implicit_function); + + // https://amrex-codes.github.io/amrex/docs_html/EB.html + std::string eb_stl; + eb_enabled |= pp_eb2.query("geom_type", eb_stl); + + return eb_enabled; +#endif + } + +} // namespace EB diff --git a/Source/EmbeddedBoundary/Make.package b/Source/EmbeddedBoundary/Make.package index 6ed944d19bd..76a20896f85 100644 --- a/Source/EmbeddedBoundary/Make.package +++ b/Source/EmbeddedBoundary/Make.package @@ -1,8 +1,10 @@ +CEXE_headers += Enabled.H CEXE_headers += ParticleScraper.H CEXE_headers += ParticleBoundaryProcess.H CEXE_headers += DistanceToEB.H CEXE_headers += WarpXFaceInfoBox.H +CEXE_sources += Enabled.cpp CEXE_sources += WarpXInitEB.cpp CEXE_sources += WarpXFaceExtensions.cpp diff --git a/Source/Make.WarpX b/Source/Make.WarpX index 502cb8a6714..57bac56e9a4 100644 --- a/Source/Make.WarpX +++ b/Source/Make.WarpX @@ -183,6 +183,8 @@ ifeq ($(USE_FFT),TRUE) INCLUDE_LOCATIONS += $(ROC_PATH)/rocfft/include LIBRARY_LOCATIONS += $(ROC_PATH)/rocfft/lib libraries += -lrocfft + else ifeq ($(USE_SYCL),TRUE) + # nothing else # Running on CPU # Use FFTW ifeq ($(PRECISION),FLOAT) diff --git a/Source/WarpX.cpp b/Source/WarpX.cpp index fe79ce9c771..7a93f014a55 100644 --- a/Source/WarpX.cpp +++ b/Source/WarpX.cpp @@ -14,6 +14,7 @@ #include "BoundaryConditions/PML.H" #include "Diagnostics/MultiDiagnostics.H" #include "Diagnostics/ReducedDiags/MultiReducedDiags.H" +#include "EmbeddedBoundary/Enabled.H" #include "EmbeddedBoundary/WarpXFaceInfoBox.H" #include "FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H" #include "FieldSolver/FiniteDifferenceSolver/MacroscopicProperties/MacroscopicProperties.H" @@ -789,6 +790,13 @@ WarpX::ReadParameters () potential_specified |= pp_boundary.query("potential_hi_z", m_poisson_boundary_handler.potential_zhi_str); #if defined(AMREX_USE_EB) potential_specified |= pp_warpx.query("eb_potential(x,y,z,t)", m_poisson_boundary_handler.potential_eb_str); + + if (!EB::enabled()) { + throw std::runtime_error( + "Currently, users MUST use EB if it was compiled in. " + "This will change with https://github.com/ECP-WarpX/WarpX/pull/4865 ." + ); + } #endif m_boundary_potential_specified = potential_specified; if (potential_specified & (WarpX::electromagnetic_solver_id == ElectromagneticSolverAlgo::HybridPIC)) { diff --git a/Source/ablastr/math/fft/AnyFFT.H b/Source/ablastr/math/fft/AnyFFT.H index 8a4c11c7654..237541f9379 100644 --- a/Source/ablastr/math/fft/AnyFFT.H +++ b/Source/ablastr/math/fft/AnyFFT.H @@ -10,6 +10,7 @@ #ifdef ABLASTR_USE_FFT # include +# include # include # if defined(AMREX_USE_CUDA) @@ -20,6 +21,8 @@ # else # include # endif +# elif defined(AMREX_USE_SYCL) +# include # else # include # endif @@ -62,6 +65,8 @@ namespace ablastr::math::anyfft # else using Complex = double2; # endif +# elif defined(AMREX_USE_SYCL) + using Complex = amrex::GpuComplex; # else # ifdef AMREX_USE_FLOAT using Complex = fftwf_complex; @@ -77,6 +82,15 @@ namespace ablastr::math::anyfft using VendorFFTPlan = cufftHandle; # elif defined(AMREX_USE_HIP) using VendorFFTPlan = rocfft_plan; +# elif defined(AMREX_USE_SYCL) + using VendorFFTPlan = oneapi::mkl::dft::descriptor< +# ifdef AMREX_USE_FLOAT + oneapi::mkl::dft::precision::SINGLE, +# else + oneapi::mkl::dft::precision::DOUBLE, +# endif + oneapi::mkl::dft::domain::REAL> *; + // dft::descriptor has no default ctor, so we have to use ptr. # else # ifdef AMREX_USE_FLOAT using VendorFFTPlan = fftwf_plan; @@ -99,6 +113,9 @@ namespace ablastr::math::anyfft VendorFFTPlan m_plan; /**< Vendor FFT plan */ direction m_dir; /**< direction (C2R or R2C) */ int m_dim; /**< Dimensionality of the FFT plan */ +#ifdef AMREX_USE_SYCL + amrex::gpuStream_t m_stream; +#endif }; /** Collection of FFT plans, one FFTplan per box */ diff --git a/Source/ablastr/math/fft/CMakeLists.txt b/Source/ablastr/math/fft/CMakeLists.txt index 913a912e1ee..b0ebc3c8050 100644 --- a/Source/ablastr/math/fft/CMakeLists.txt +++ b/Source/ablastr/math/fft/CMakeLists.txt @@ -5,6 +5,8 @@ foreach(D IN LISTS WarpX_DIMS) target_sources(ablastr_${SD} PRIVATE WrapCuFFT.cpp) elseif(WarpX_COMPUTE STREQUAL HIP) target_sources(ablastr_${SD} PRIVATE WrapRocFFT.cpp) + elseif(WarpX_COMPUTE STREQUAL SYCL) + target_sources(ablastr_${SD} PRIVATE WrapMklFFT.cpp) else() target_sources(ablastr_${SD} PRIVATE WrapFFTW.cpp) endif() diff --git a/Source/ablastr/math/fft/Make.package b/Source/ablastr/math/fft/Make.package index 63786cdc006..6e7dc2b3dd4 100644 --- a/Source/ablastr/math/fft/Make.package +++ b/Source/ablastr/math/fft/Make.package @@ -3,6 +3,8 @@ ifeq ($(USE_FFT),TRUE) CEXE_sources += WrapCuFFT.cpp else ifeq ($(USE_HIP),TRUE) CEXE_sources += WrapRocFFT.cpp + else ifeq ($(USE_SYCL),TRUE) + CEXE_sources += WrapMklFFT.cpp else CEXE_sources += WrapFFTW.cpp endif diff --git a/Source/ablastr/math/fft/WrapMklFFT.cpp b/Source/ablastr/math/fft/WrapMklFFT.cpp new file mode 100644 index 00000000000..ef2cb6c42fa --- /dev/null +++ b/Source/ablastr/math/fft/WrapMklFFT.cpp @@ -0,0 +1,96 @@ +/* Copyright 2019-2023 + * + * This file is part of ABLASTR. + * + * License: BSD-3-Clause-LBNL + */ + +#include "AnyFFT.H" + +#include "ablastr/utils/TextMsg.H" +#include "ablastr/profiler/ProfilerWrapper.H" + +#include + +namespace ablastr::math::anyfft +{ + + void setup () {/*nothing to do*/} + + void cleanup () {/*nothing to do*/} + + FFTplan CreatePlan (const amrex::IntVect& real_size, amrex::Real * const real_array, + Complex * const complex_array, const direction dir, const int dim) + { + FFTplan fft_plan; + ABLASTR_PROFILE("ablastr::math::anyfft::CreatePlan"); + + // Initialize fft_plan.m_plan with the vendor fft plan. + std::vector strides(dim+1); + if (dim == 3) { + fft_plan.m_plan = new std::remove_pointer_t( + {std::int64_t(real_size[2]), + std::int64_t(real_size[1]), + std::int64_t(real_size[0])}); + strides[0] = 0; + strides[1] = real_size[0] * real_size[1]; + strides[2] = real_size[0]; + strides[3] = 1; + } else if (dim == 2) { + fft_plan.m_plan = new std::remove_pointer_t( + {std::int64_t(real_size[1]), + std::int64_t(real_size[0])}); + strides[0] = 0; + strides[1] = real_size[0]; + strides[2] = 1; + } else if (dim == 1) { + strides[0] = 0; + strides[1] = 1; + fft_plan.m_plan = new std::remove_pointer_t( + std::int64_t(real_size[0])); + } else { + ABLASTR_ABORT_WITH_MESSAGE("only dim2 =1, dim=2 and dim=3 have been implemented"); + } + + fft_plan.m_plan->set_value(oneapi::mkl::dft::config_param::PLACEMENT, + DFTI_NOT_INPLACE); + fft_plan.m_plan->set_value(oneapi::mkl::dft::config_param::FWD_STRIDES, + strides.data()); + fft_plan.m_plan->commit(amrex::Gpu::Device::streamQueue()); + + // Store meta-data in fft_plan + fft_plan.m_real_array = real_array; + fft_plan.m_complex_array = complex_array; + fft_plan.m_dir = dir; + fft_plan.m_dim = dim; + fft_plan.m_stream = amrex::Gpu::gpuStream(); + + return fft_plan; + } + + void DestroyPlan (FFTplan& fft_plan) + { + delete fft_plan.m_plan; + } + + void Execute (FFTplan& fft_plan) + { + if (!(fft_plan.m_stream == amrex::Gpu::gpuStream())) { + amrex::Gpu::streamSynchronize(); + } + + sycl::event r; + if (fft_plan.m_dir == direction::R2C) { + r = oneapi::mkl::dft::compute_forward( + *fft_plan.m_plan, + fft_plan.m_real_array, + reinterpret_cast*>(fft_plan.m_complex_array)); + } else { + r = oneapi::mkl::dft::compute_backward( + *fft_plan.m_plan, + reinterpret_cast*>(fft_plan.m_complex_array), + fft_plan.m_real_array); + } + r.wait(); + } +} diff --git a/setup.py b/setup.py index 30f91d8fb4a..ecce858518d 100644 --- a/setup.py +++ b/setup.py @@ -17,7 +17,7 @@ def initialize_options(self): # clashes with directories many developers have in their source trees; # this can create confusing results with "pip install .", which clones # the whole source tree by default - self.build_base = '_tmppythonbuild' + self.build_base = os.path.join("_tmppythonbuild", "warpx") def run(self): # remove existing build directory