diff --git a/src/AbsBeamline/Component.h b/src/AbsBeamline/Component.h index 6914f43d7e753b59fc5ac35caf6884e8b2c5aa37..fadc16e665edea917a7f4a82352d15e3b33c6e25 100644 --- a/src/AbsBeamline/Component.h +++ b/src/AbsBeamline/Component.h @@ -22,6 +22,30 @@ // $Author: fci $ // // ------------------------------------------------------------------------ + +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "AbsBeamline/ElementBase.h" #include "Fields/EMField.h" #include "OPALTypes.h" diff --git a/src/AbsBeamline/Drift.h b/src/AbsBeamline/Drift.h index 257478635f00e6a95c17361d57f08e864d57d3bf..9e53d6d789e0c1d135248e3aecd1882346d6ddce 100644 --- a/src/AbsBeamline/Drift.h +++ b/src/AbsBeamline/Drift.h @@ -23,6 +23,29 @@ #include "AbsBeamline/Component.h" +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + // Class Drift // ------------------------------------------------------------------------ /// Interface for drift space. @@ -70,4 +93,4 @@ inline int Drift::getRequiredNumberOfTimeSteps() const { return 1; } -#endif // CLASSIC_Drift_HH \ No newline at end of file +#endif // CLASSIC_Drift_HH diff --git a/src/AbsBeamline/ElementBase.cpp b/src/AbsBeamline/ElementBase.cpp index 956f17a01bbee32fc7bff1a18159a5d72b4ae542..3b1cbd6d58dad3f73ad2ca4382c2f28e0bafda50 100644 --- a/src/AbsBeamline/ElementBase.cpp +++ b/src/AbsBeamline/ElementBase.cpp @@ -64,8 +64,54 @@ #include "Channels/Channel.h" +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include <boost/filesystem.hpp> +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + const std::map<ElementType, std::string> ElementBase::elementTypeToString_s = { {ElementType::ANY, "Any"}, {ElementType::BEAMLINE, "Beamline"}, diff --git a/src/AbsBeamline/ElementBase.h b/src/AbsBeamline/ElementBase.h index 5820a22be762c06d9c882662477869d5536aeaa8..bfc92afe02b8ba05f43a6c724e89254e54297677 100644 --- a/src/AbsBeamline/ElementBase.h +++ b/src/AbsBeamline/ElementBase.h @@ -63,6 +63,29 @@ #ifndef CLASSIC_ElementBase_HH #define CLASSIC_ElementBase_HH +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "AbsBeamline/AttributeSet.h" #include "Algorithms/CoordinateSystemTrafo.h" #include "Algorithms/Quaternion.hpp" diff --git a/src/AbsBeamline/Multipole.h b/src/AbsBeamline/Multipole.h index 9a94e95fbef5e2ca924bb2c521e657e20fcead0a..450829ff429a181b5f3d8622abea7cabbf338932 100644 --- a/src/AbsBeamline/Multipole.h +++ b/src/AbsBeamline/Multipole.h @@ -23,6 +23,29 @@ #include "BeamlineGeometry/StraightGeometry.h" #include "Fields/BMultipoleField.h" +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + class Fieldmap; // Class Multipole @@ -160,4 +183,4 @@ inline size_t Multipole::getMaxSkewComponentIndex() const { return SkewComponents.size(); } -#endif // CLASSIC_Multipole_HH \ No newline at end of file +#endif // CLASSIC_Multipole_HH diff --git a/src/AbsBeamline/MultipoleTBase.cpp b/src/AbsBeamline/MultipoleTBase.cpp index 2ceda734d86337ed8f8811880f9a358e5bad46f4..9d106f062e7e5db3f8e83c571a94088169e6f6aa 100644 --- a/src/AbsBeamline/MultipoleTBase.cpp +++ b/src/AbsBeamline/MultipoleTBase.cpp @@ -250,4 +250,4 @@ double MultipoleTBase::getFnDerivS(const std::size_t &n, deriv += -1. * getFn(n, x, s + 2. * stepSize); deriv /= 12 * stepSize; return deriv; -} \ No newline at end of file +} diff --git a/src/AbsBeamline/MultipoleTBase.h b/src/AbsBeamline/MultipoleTBase.h index 845fa6f65529afed239595579d1f1d36b5d47471..e0c24d47c92268d9cd4689f4582ca2bac050017b 100644 --- a/src/AbsBeamline/MultipoleTBase.h +++ b/src/AbsBeamline/MultipoleTBase.h @@ -79,6 +79,30 @@ #include "AbsBeamline/Component.h" #include "AbsBeamline/EndFieldModel/Tanh.h" +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + + #include "Fields/BMultipoleField.h" class MultipoleTBase : public Component { diff --git a/src/AbsBeamline/MultipoleTCurvedConstRadius.h b/src/AbsBeamline/MultipoleTCurvedConstRadius.h index 951b719c0c1695b90cb5d473bc91adaf024eee50..cdd66af3686b7a217f98d461e13be1dcb2e31b15 100644 --- a/src/AbsBeamline/MultipoleTCurvedConstRadius.h +++ b/src/AbsBeamline/MultipoleTCurvedConstRadius.h @@ -75,6 +75,29 @@ * --------------------------------------------------------------------- */ +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "BeamlineGeometry/PlanarArcGeometry.h" #include "AbsBeamline/MultipoleTBase.h" #include "AbsBeamline/MultipoleTFunctions/RecursionRelation.h" diff --git a/src/AbsBeamline/MultipoleTCurvedVarRadius.cpp b/src/AbsBeamline/MultipoleTCurvedVarRadius.cpp index 8903bb59a0c6edee250c84ddded72cf760b15c29..de585699159532688d743c02af07113baf1fa883 100644 --- a/src/AbsBeamline/MultipoleTCurvedVarRadius.cpp +++ b/src/AbsBeamline/MultipoleTCurvedVarRadius.cpp @@ -26,7 +26,7 @@ * POSSIBILITY OF SUCH DAMAGE. */ - +#include <experimental/source_location> #include <vector> #include "gsl/gsl_sf_pow_int.h" #include "MultipoleTCurvedVarRadius.h" diff --git a/src/AbsBeamline/MultipoleTCurvedVarRadius.h b/src/AbsBeamline/MultipoleTCurvedVarRadius.h index 4f5f32f4db4be10f6183028285ee8c8485e82743..589d34af507a7a63179b8121a0b3ebb58eaef71b 100644 --- a/src/AbsBeamline/MultipoleTCurvedVarRadius.h +++ b/src/AbsBeamline/MultipoleTCurvedVarRadius.h @@ -76,6 +76,31 @@ * --------------------------------------------------------------------- */ + + +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "BeamlineGeometry/VarRadiusGeometry.h" #include "AbsBeamline/MultipoleTBase.h" diff --git a/src/AbsBeamline/MultipoleTStraight.cpp b/src/AbsBeamline/MultipoleTStraight.cpp index 152b6c2235e3507f0ca64c6b6c51a473e15c16a9..b5b8dd99897308795f2c351bf5e35067887d410d 100644 --- a/src/AbsBeamline/MultipoleTStraight.cpp +++ b/src/AbsBeamline/MultipoleTStraight.cpp @@ -26,7 +26,6 @@ * POSSIBILITY OF SUCH DAMAGE. */ - #include "gsl/gsl_sf_gamma.h" #include "gsl/gsl_sf_pow_int.h" #include "MultipoleTStraight.h" diff --git a/src/AbsBeamline/MultipoleTStraight.h b/src/AbsBeamline/MultipoleTStraight.h index 562f42a4ac0b38b2db10a75436cb726d865cc80d..e4663886fc6585e33198f78a115d5d8a69969b3e 100644 --- a/src/AbsBeamline/MultipoleTStraight.h +++ b/src/AbsBeamline/MultipoleTStraight.h @@ -79,6 +79,29 @@ #include "AbsBeamline/MultipoleTBase.h" #include <vector> +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + class MultipoleTStraight: public MultipoleTBase { public: /** Constructor diff --git a/src/AbsBeamline/ScalingFFAMagnet.cpp b/src/AbsBeamline/ScalingFFAMagnet.cpp index 74a11e1f9255ad9dc6bc49b0880968ece7ca9a27..e71af8cc98f586a9e59a8c9bf19683e519c9b5e5 100644 --- a/src/AbsBeamline/ScalingFFAMagnet.cpp +++ b/src/AbsBeamline/ScalingFFAMagnet.cpp @@ -26,7 +26,6 @@ */ #include <cmath> - #include "AbsBeamline/BeamlineVisitor.h" #include "AbsBeamline/ScalingFFAMagnet.h" #include "PartBunch/PartBunch.hpp" diff --git a/src/AbsBeamline/ScalingFFAMagnet.h b/src/AbsBeamline/ScalingFFAMagnet.h index 850573aead0e01701e51438d74409859187a1442..0b8a547f5e7e6c6c6873937984f087c34968ee32 100644 --- a/src/AbsBeamline/ScalingFFAMagnet.h +++ b/src/AbsBeamline/ScalingFFAMagnet.h @@ -30,6 +30,29 @@ #include "AbsBeamline/EndFieldModel/EndFieldModel.h" #include "AbsBeamline/Component.h" +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #ifndef ABSBEAMLINE_ScalingFFAMagnet_H #define ABSBEAMLINE_ScalingFFAMagnet_H diff --git a/src/AbsBeamline/Solenoid.h b/src/AbsBeamline/Solenoid.h index 130b406b8d95d6ed1d661c9374f621c3d7caadf0..ac0d1e55a6c3189338fd0f9426dce6f8d29da555 100644 --- a/src/AbsBeamline/Solenoid.h +++ b/src/AbsBeamline/Solenoid.h @@ -21,6 +21,29 @@ // // ------------------------------------------------------------------------ +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "AbsBeamline/Component.h" #include "Algorithms/CoordinateSystemTrafo.h" diff --git a/src/AbstractObjects/OpalData.h b/src/AbstractObjects/OpalData.h index 3f7b2b4c72153b3740566de376e1a5bbafa6b0ec..1260d0d8f7b32d647fabbac2d8eaf8d5ee748bb7 100644 --- a/src/AbstractObjects/OpalData.h +++ b/src/AbstractObjects/OpalData.h @@ -22,6 +22,29 @@ #ifndef OPAL_OpalData_HH #define OPAL_OpalData_HH +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include <iosfwd> #include <map> #include <stack> diff --git a/src/Algorithms/DistributionMoments.cpp b/src/Algorithms/DistributionMoments.cpp index 3b92719b84be569bca127823169809dcb24a4fd4..9b1c6ff0eca708336849f69f4d766562a0eef097 100644 --- a/src/Algorithms/DistributionMoments.cpp +++ b/src/Algorithms/DistributionMoments.cpp @@ -279,7 +279,7 @@ void DistributionMoments::computeMoments(ippl::ParticleAttrib<Vector_t<double,3> } double betaGamma = std::sqrt(std::pow(meanGamma_m, 2) - 1.0); - geometricEps_m = normalizedEps_m / Vector_t(betaGamma); + geometricEps_m = normalizedEps_m / Vector_t<double,3>(betaGamma); } diff --git a/src/Algorithms/DistributionMoments.h b/src/Algorithms/DistributionMoments.h index bf893cb2579fd0395f2caf3e5647ecaf3968e3a6..d5f1eee09ea60bafa363c1767abab2956b3bc532 100644 --- a/src/Algorithms/DistributionMoments.h +++ b/src/Algorithms/DistributionMoments.h @@ -18,6 +18,29 @@ #ifndef DISTRIBUTIONMOMENTS_H #define DISTRIBUTIONMOMENTS_H +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "Ippl.h" #include <Kokkos_Core.hpp> #include "Algorithms/BoostMatrix.h" @@ -95,12 +118,12 @@ public: double getTotalMass() const; double getTotalNumParticles() const; -private: - bool isParticleExcluded(const OpalParticle&) const; void computeMeans(ippl::ParticleAttrib<Vector_t<double,3>>::view_type& Rview, ippl::ParticleAttrib<Vector_t<double,3>>::view_type& Pview, ippl::ParticleAttrib<double>::view_type& Mview, size_t Np); +private: + bool isParticleExcluded(const OpalParticle&) const; //template <class InputIt> //void computeMeans(const InputIt&, const InputIt&); diff --git a/src/Algorithms/ParallelTracker.h b/src/Algorithms/ParallelTracker.h index fc5ab41e83540f71258d4328dacda21584860b8c..c76ba87f91429d517b62075fcc0798e519358ffe 100644 --- a/src/Algorithms/ParallelTracker.h +++ b/src/Algorithms/ParallelTracker.h @@ -180,6 +180,18 @@ public: /// Apply the algorithm to a vertical FFA magnet. virtual void visitVerticalFFAMagnet(const VerticalFFAMagnet& bend); + // made following public: __host__ __device__ lambda cannot have private or protected access within its class + void kickParticles(const BorisPusher& pusher); + + void pushParticles(const BorisPusher& pusher); + + void timeIntegration2(BorisPusher& pusher); + + void changeDT(bool backTrack = false); + + void computeSpaceChargeFields(unsigned long long step); + + void setTime(); private: // Not implemented. ParallelTracker(); @@ -221,8 +233,6 @@ private: /********************** END VARIABLES ***********************************/ - void kickParticles(const BorisPusher& pusher); - void pushParticles(const BorisPusher& pusher); void updateReferenceParticle(const BorisPusher& pusher); void writePhaseSpace(const long long step, bool psDump, bool statDump); @@ -237,22 +247,18 @@ private: void prepareSections(); void timeIntegration1(BorisPusher& pusher); - void timeIntegration2(BorisPusher& pusher); void selectDT(bool backTrack = false); - void changeDT(bool backTrack = false); void emitParticles(long long step); void computeExternalFields(OrbitThreader& oth); void computeWakefield(IndexMap::value_t& elements); void computeParticleMatterInteraction(IndexMap::value_t elements, OrbitThreader& oth); - void computeSpaceChargeFields(unsigned long long step); // void prepareOpalBeamlineSections(); void dumpStats(long long step, bool psDump, bool statDump); void setOptionalVariables(); bool hasEndOfLineReached(const BoundingBox& globalBoundingBox); void handleRestartRun(); - void setTime(); void doBinaryRepartition(); void transformBunch(const CoordinateSystemTrafo& trafo); @@ -341,4 +347,4 @@ inline void ParallelTracker::pushParticles(const BorisPusher& pusher) { ippl::Comm->barrier(); } -#endif // OPAL_ParallelTracker_HH \ No newline at end of file +#endif // OPAL_ParallelTracker_HH diff --git a/src/Algorithms/PartData.h b/src/Algorithms/PartData.h index 73636d2a3bc2b7ca35380b052a5a33f985e10e4c..49619e68b6b8cfa3797cca0679ae6861bc99aa10 100644 --- a/src/Algorithms/PartData.h +++ b/src/Algorithms/PartData.h @@ -32,6 +32,8 @@ // by the compiler perform the correct operation. For speed reasons // they are not implemented. +#include <Kokkos_Core.hpp> + class PartData { public: @@ -48,10 +50,10 @@ public: PartData(); /// The constant charge per particle. - double getQ() const; + KOKKOS_INLINE_FUNCTION double getQ() const; /// The constant mass per particle. - double getM() const; + KOKKOS_INLINE_FUNCTION double getM() const; /// The constant reference momentum per particle. double getP() const; @@ -101,12 +103,12 @@ protected: // Inline functions. // ------------------------------------------------------------------------ -inline double PartData::getQ() const { +KOKKOS_INLINE_FUNCTION double PartData::getQ() const { return charge; } -inline double PartData::getM() const { +KOKKOS_INLINE_FUNCTION double PartData::getM() const { return mass; } diff --git a/src/BasicActions/Option.h b/src/BasicActions/Option.h index 6db1c8bdc3029da3110d869f19ba548b541f2632..8562165c3cbdc06df1b29104b518bb9ba71082ab 100644 --- a/src/BasicActions/Option.h +++ b/src/BasicActions/Option.h @@ -20,6 +20,29 @@ #ifndef OPAL_Option_HH #define OPAL_Option_HH +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "AbstractObjects/Action.h" #include "Utilities/Options.h" diff --git a/src/Distribution/Distribution.cpp b/src/Distribution/Distribution.cpp index d8c1b23e8ad0efb661192e95abbe5f456777f81d..bceabbde8f2285289cc1253f9a4287f1b247cd3c 100644 --- a/src/Distribution/Distribution.cpp +++ b/src/Distribution/Distribution.cpp @@ -59,8 +59,6 @@ using Base = ippl::ParticleBase<ippl::ParticleSpatialLayout<T, Dim>>; using view_type = typename ippl::detail::ViewType<ippl::Vector<double, Dim>, 1>::view_type; -constexpr double SMALLESTCUTOFF = 1e-12; - namespace DISTRIBUTION { enum { TYPE, FNAME, SIGMAX, SIGMAY, SIGMAZ, SIGMAPX, SIGMAPY, SIGMAPZ, SIZE, CUTOFFPX, CUTOFFPY, CUTOFFPZ, CUTOFFX, CUTOFFY, CUTOFFLONG }; } diff --git a/src/Main.cpp b/src/Main.cpp index 78fdd2d1592a9b2b12a2d4ba282e4552f51861e6..9784cc7795447ed7d0f63bd4f972fb333905db70 100644 --- a/src/Main.cpp +++ b/src/Main.cpp @@ -15,6 +15,29 @@ // along with OPAL. If not, see <https://www.gnu.org/licenses/>. // +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "H5hut.h" #include "AbstractObjects/OpalData.h" diff --git a/src/OPALTypes.h b/src/OPALTypes.h index b1f8f980da77ba13236c161759ae3e4a6880a5be..6c66ac99e12cd8fdfcbdc505d623754dac05ef12 100644 --- a/src/OPALTypes.h +++ b/src/OPALTypes.h @@ -69,7 +69,7 @@ template <typename T = double, unsigned Dim = 3, class... ViewArgs> using VField_t = Field<Vector_t<T, Dim>, Dim, ViewArgs...>; */ -template <typename T, unsigned Dim = 3> +template <typename T, unsigned Dim> using Vector_t = ippl::Vector<T, Dim>; typedef typename std::pair<Vector_t<double, 3>, Vector_t<double, 3>> VectorPair_t; @@ -81,13 +81,13 @@ enum UnitState_t { units = 0, unitless = 1 }; // euclidean norm template <class T, unsigned D> -inline double euclidean_norm(const Vector_t<T, D>& v) { - return std::sqrt(dot(v, v).apply()); +KOKKOS_INLINE_FUNCTION double euclidean_norm(const Vector_t<T, D>& v) { + return Kokkos::sqrt(dot(v, v).apply()); } // dot products template <class T, unsigned D> -inline double dot(const Vector_t<T, D>& v, const Vector_t<T, D>& w) { +KOKKOS_INLINE_FUNCTION double dot(const Vector_t<T, D>& v, const Vector_t<T, D>& w) { double res = 0.0; for (unsigned i = 0; i < D; i++) res += v(i) * w(i); @@ -95,11 +95,11 @@ inline double dot(const Vector_t<T, D>& v, const Vector_t<T, D>& w) { } template <class T, unsigned D> -inline double dot(const Vector_t<T, D>& v) { +KOKKOS_INLINE_FUNCTION double dot(const Vector_t<T, D>& v) { double res = 0.0; for (unsigned i = 0; i < D; i++) res += v(i) * v(i); return res; } -#endif \ No newline at end of file +#endif diff --git a/src/PartBunch/#[-Wunused-variable]# b/src/PartBunch/#[-Wunused-variable]# deleted file mode 100644 index 091276e8d88504da43085e1b5da7541ad81073bd..0000000000000000000000000000000000000000 --- a/src/PartBunch/#[-Wunused-variable]# +++ /dev/null @@ -1,2 +0,0 @@ - 499 | int tag = 101; - \ No newline at end of file diff --git a/src/PartBunch/PartBunch.hpp b/src/PartBunch/PartBunch.hpp index 5a9069d3be284858866c3eca92ac3ab099daaef7..41de9eb2b33e320c2f40329d230adc755021c857 100644 --- a/src/PartBunch/PartBunch.hpp +++ b/src/PartBunch/PartBunch.hpp @@ -43,7 +43,6 @@ diagnostics(): calculate statistics and maybe write tp h5 and stat files #include "Algorithms/BoostMatrix.h" #include "Algorithms/CoordinateSystemTrafo.h" -//#include "Algorithms/DistributionMoments.h" #include "Attributes/Attributes.h" #include "Distribution/Distribution.h" #include "Manager/BaseManager.h" @@ -86,18 +85,33 @@ public: double time_m; - Vector_t<int, Dim> nr_m; - size_type totalP_m; int nt_m; + double lbt_m; + double dt_m; + int it_m; std::string integration_method_m; + std::string solver_m; + bool isFirstRepartition_m; + +private: + double qi_m; + + double mi_m; + + double rmsDensity_m; + + +public: + Vector_t<int, Dim> nr_m; + Vector_t<double, Dim> origin_m; Vector_t<double, Dim> rmin_m; Vector_t<double, Dim> rmax_m; @@ -112,7 +126,6 @@ public: ippl::NDIndex<Dim> domain_m; std::array<bool, Dim> decomp_m; - bool isFirstRepartition_m; /* Up to here it is like the opaltest @@ -178,15 +191,17 @@ private: PartData* reference_m; double couplingConstant_m; - double qi_m; - double mi_m; - double rmsDensity_m; /// step in a TRACK command long long localTrackStep_m; /// if multiple TRACK commands long long globalTrackStep_m; + + + std::shared_ptr<Distribution> OPALdist_m; + + std::shared_ptr<FieldSolverCmd> OPALFieldSolver_m; // unit state of PartBunch // UnitState_t unit_state_m; @@ -204,13 +219,11 @@ private: bool dcBeam_m; double periodLength_m; - Distribution* OPALdist_m; - FieldSolverCmd* OPALFieldSolver_m; - public: PartBunch( double qi, double mi, size_t totalP, int nt, double lbt, std::string integration_method, - Distribution* OPALdistribution, FieldSolverCmd* OPALFieldSolver) + std::shared_ptr<Distribution> &OPALdistribution, + std::shared_ptr<FieldSolverCmd> &OPALFieldSolver) : ippl::PicManager< T, Dim, ParticleContainer<T, Dim>, FieldContainer<T, Dim>, LoadBalancer<T, Dim>>(), time_m(0.0), @@ -221,7 +234,7 @@ public: it_m(0), integration_method_m(integration_method), solver_m(""), - isFirstRepartition_m(true), + isFirstRepartition_m(true), qi_m(qi), mi_m(mi), rmsDensity_m(0.0), @@ -229,9 +242,8 @@ public: globalTrackStep_m(0), OPALdist_m(OPALdistribution), OPALFieldSolver_m(OPALFieldSolver){ - - Inform m("PartBunch() "); + Inform m("PartBunch() "); static IpplTimings::TimerRef gatherInfoPartBunch = IpplTimings::getTimer("gatherInfoPartBunch"); IpplTimings::startTimer(gatherInfoPartBunch); @@ -263,12 +275,10 @@ public: this->origin_m = -3.0; this->dt_m = 0.5 / this->nr_m[2]; - using ParticleContainer_t = ParticleContainer<T, Dim>; - using FieldContainer_t = FieldContainer<T, Dim>; + rmin_m = origin_m; + rmax_m = origin_m + length; - this->setFieldContainer(std::make_shared<FieldContainer_t>( - this->hr_m, this->rmin_m, this->rmax_m, this->decomp_m, this->domain_m, this->origin_m, - isAllPeriodic)); + this->setFieldContainer( std::make_shared<FieldContainer_t>(hr_m, rmin_m, rmax_m, decomp_m, domain_m, origin_m, isAllPeriodic) ); this->setParticleContainer(std::make_shared<ParticleContainer_t>( this->fcontainer_m->getMesh(), this->fcontainer_m->getFL())); @@ -347,7 +357,13 @@ public: } void gatherCIC() { - gather(this->pcontainer_m->E, this->fcontainer_m->getE(), this->pcontainer_m->R); +/* + using Base = ippl::ParticleBase<ippl::ParticleSpatialLayout<T, Dim>>; + typename Base::particle_position_type* Ep = &this->pcontainer_m->E; + typename Base::particle_position_type* R = &this->pcontainer_m->R; + VField_t<T, Dim>* Ef = &this->fcontainer_m->getE(); + gather(*Ep, *Ef, *R); +*/ } void scatterCIC(); @@ -411,8 +427,10 @@ public: } double getGamma(int i) const { + return 0.0; } double getBeta(int i) const { + return 0.0; } void actT() { @@ -449,7 +467,7 @@ public: bool isGridFixed() { - return false; + return false; } void boundp() { @@ -490,6 +508,7 @@ public: } size_t calcNumPartsOutside(Vector_t<double, Dim> x) { + return 0; } void calcLineDensity( @@ -500,6 +519,7 @@ public: } Vector_t<double, Dim> getEExtrema() { + return Vector_t<double, Dim>(0); } void computeSelfFields(); @@ -512,15 +532,20 @@ public: } bool getFieldSolverType() { + return false; } bool getIfBeamEmitting() { + return false; } int getLastEmittedEnergyBin() { + return 0; } size_t getNumberOfEmissionSteps() { + return 0; } int getNumberOfEnergyBins() { + return 0; } void Rebin() { @@ -548,12 +573,14 @@ public: void rebin() { } int getLastemittedBin() { + return 0; } void setLocalBinCount(size_t num, int bin) { } void calcGammas() { } double getBinGamma(int bin) { + return 0.0; } bool hasBinning() { return false; @@ -563,31 +590,43 @@ public: void setBinCharge(int bin) { } double calcMeanPhi() { + return 0.0; } bool resetPartBinID2(const double eta) { return false; } bool resetPartBinBunch() { + return false; } double getPx(int i) { + return 0.0; } double getPy(int i) { + return 0.0; } double getPz(int i) { + return 0.0; } double getPx0(int i) { + return 0.0; } double getPy0(int i) { + return 0.0; } double getX(int i) { + return 0.0; } double getY(int i) { + return 0.0; } double getZ(int i) { + return 0.0; } double getX0(int i) { + return 0.0; } double getY0(int i) { + return 0.0; } void setZ(int i, double zcoo) { diff --git a/src/Physics/ParticleProperties.h b/src/Physics/ParticleProperties.h index 8143e4317314ce16abd7c7f8aec121e5f4b6ab7f..b3cbd6384552137759d0ea775f25b9ff86dd3f31 100644 --- a/src/Physics/ParticleProperties.h +++ b/src/Physics/ParticleProperties.h @@ -18,6 +18,29 @@ #ifndef PARTICLEPROPERTIES_H #define PARTICLEPROPERTIES_H +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include <boost/bimap.hpp> #include <map> diff --git a/src/Steppers/BorisPusher.h b/src/Steppers/BorisPusher.h index 1f547d7c9d0cc9280030226bd20529d95b979995..cc24294da6e48fe0ee18eeafc811475523f0a829 100644 --- a/src/Steppers/BorisPusher.h +++ b/src/Steppers/BorisPusher.h @@ -112,7 +112,7 @@ KOKKOS_INLINE_FUNCTION void BorisPusher::kick( P += cross(P, t); */ - double gamma = sqrt(1.0 + dot(P, P)); + double gamma = Kokkos::sqrt(1.0 + dot(P, P)); Vector_t<double, 3> const t = 0.5 * dt * charge * Physics::c * Physics::c / (gamma * mass) * Bf; Vector_t<double, 3> const w = P + cross(P, t); Vector_t<double, 3> const s = 2.0 / (1.0 + dot(t, t)) * t; @@ -135,7 +135,7 @@ KOKKOS_INLINE_FUNCTION void BorisPusher::push( * R[i] += 0.5 * P[i] * recpgamma; * \endcode */ - R += 0.5 * P / std::sqrt(1.0 + dot(P)); + R += 0.5 * P / Kokkos::sqrt(1.0 + dot(P)); } #endif diff --git a/src/Structure/SDDSColumn.h b/src/Structure/SDDSColumn.h index a23cec4760706aeb598721491763333e7b9dc2ae..1c95c31dc6b946cda3d3a81ee21f25f4f3da52a5 100644 --- a/src/Structure/SDDSColumn.h +++ b/src/Structure/SDDSColumn.h @@ -18,6 +18,29 @@ #ifndef SDDSWRITERCOLUMN_H #define SDDSWRITERCOLUMN_H +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include <boost/variant.hpp> #include <ostream> @@ -79,4 +102,4 @@ void SDDSColumn::addValue(const T& val) { std::ostream& operator<<(std::ostream& os, const SDDSColumn& col); -#endif \ No newline at end of file +#endif diff --git a/src/Structure/SDDSColumnSet.h b/src/Structure/SDDSColumnSet.h index 406ab8b6e2a699c48a477ac1cfb9d14295dd12b3..225ccfb4817cbe8c383e44bf364185706d6f534a 100644 --- a/src/Structure/SDDSColumnSet.h +++ b/src/Structure/SDDSColumnSet.h @@ -18,6 +18,29 @@ #ifndef SDDSWRITERCOLUMNSET_H #define SDDSWRITERCOLUMNSET_H +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "Structure/SDDSColumn.h" #include "Utilities/OpalException.h" @@ -80,4 +103,4 @@ bool SDDSColumnSet::hasColumns() const { } -#endif \ No newline at end of file +#endif diff --git a/src/Track/TrackRun.cpp b/src/Track/TrackRun.cpp index fb05d592b71e5b8303ed8e66d7f641dc71ca5a9e..9fe64b704eb7474ff28fdccd84721269de11d4ce 100644 --- a/src/Track/TrackRun.cpp +++ b/src/Track/TrackRun.cpp @@ -228,12 +228,13 @@ void TrackRun::execute() { // \todo here we can loop over several distributions - dist_m = Distribution::find(distributionArray[0]); + dist_m = std::shared_ptr<Distribution>(Distribution::find(distributionArray[0])); *gmsg << *dist_m << endl; - fs_m = FieldSolverCmd::find(Attributes::getString(itsAttr[TRACKRUN::FIELDSOLVER])); + fs_m = std::shared_ptr<FieldSolverCmd>(FieldSolverCmd::find(Attributes::getString(itsAttr[TRACKRUN::FIELDSOLVER]))); *gmsg << *fs_m << endl; + Beam* beam = Beam::find(Attributes::getString(itsAttr[TRACKRUN::BEAM])); *gmsg << *beam << endl; @@ -249,7 +250,7 @@ void TrackRun::execute() { // There's a change of units for particle mass that seems strange -> gives consistent Kinetic Energy - bunch_m = std::make_unique<bunch_type>(macrocharge_m, + bunch_m = std::make_shared<bunch_type>(macrocharge_m, beam->getMass()*1e9*Units::eV2MeV, beam->getNumberOfParticles(), 10, 1.0, "LF2", dist_m, fs_m); bunch_m->setT(0.0); diff --git a/src/Track/TrackRun.h b/src/Track/TrackRun.h index 3ba1f1e5aaf41824c98ba0f889d42ebdb383dee7..42ee950ab40ebd6d2d562ded32836f1336afb20d 100644 --- a/src/Track/TrackRun.h +++ b/src/Track/TrackRun.h @@ -75,13 +75,13 @@ private: Tracker* itsTracker_m; - Distribution* dist_m; + std::shared_ptr<Distribution> dist_m; std::vector<Distribution*> distrs_m; std::shared_ptr<SamplingBase> sampler_m; - FieldSolverCmd* fs_m; + std::shared_ptr<FieldSolverCmd> fs_m; DataSink* ds_m; @@ -95,7 +95,7 @@ private: */ using bunch_type = PartBunch_t; - std::unique_ptr<bunch_type> bunch_m; + std::shared_ptr<bunch_type> bunch_m; bool isFollowupTrack_m; diff --git a/src/Utilities/MSLang/matheval.hpp b/src/Utilities/MSLang/matheval.hpp index 0644d15766eaf069b0c304520dee579be3a550a9..e7915beec12ffe6dd20f9a66d0e844ff207576cc 100644 --- a/src/Utilities/MSLang/matheval.hpp +++ b/src/Utilities/MSLang/matheval.hpp @@ -14,6 +14,29 @@ // http://www.boost.org/LICENSE_1_0.txt) #pragma once +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include <algorithm> #include <cmath> #include <functional> diff --git a/src/Utilities/SDDSParser.h b/src/Utilities/SDDSParser.h index e4925675500382fc5a97e7a41fa68f7b61b03ddc..703d4c7bb7b5881f99574f36ebcc50f51830ec0a 100644 --- a/src/Utilities/SDDSParser.h +++ b/src/Utilities/SDDSParser.h @@ -15,6 +15,30 @@ // You should have received a copy of the GNU General Public License // along with OPAL. If not, see <https://www.gnu.org/licenses/>. // + +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "SDDSParser/ast.hpp" #include "SDDSParser/file.hpp" #include "SDDSParser/skipper.hpp" @@ -239,4 +263,4 @@ namespace SDDS { } } -#endif \ No newline at end of file +#endif diff --git a/src/Utilities/SDDSParser/array.hpp b/src/Utilities/SDDSParser/array.hpp index baaae4bb377ea1dfb74a6a96c5b4fb7ea59d3035..d336f92c52ced3d6627da3b90c25896b1fc44b7c 100644 --- a/src/Utilities/SDDSParser/array.hpp +++ b/src/Utilities/SDDSParser/array.hpp @@ -17,6 +17,29 @@ #ifndef ARRAY_HPP_ #define ARRAY_HPP_ +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "ast.hpp" #include "skipper.hpp" #include "error_handler.hpp" diff --git a/src/Utilities/SDDSParser/associate.hpp b/src/Utilities/SDDSParser/associate.hpp index e4eabb22b98348d1cf6a3061db27c10531f2bd4e..6c7410e8efd7e061c78bb9c0a5a37f94b9eb572c 100644 --- a/src/Utilities/SDDSParser/associate.hpp +++ b/src/Utilities/SDDSParser/associate.hpp @@ -17,6 +17,29 @@ #ifndef ASSOCIATE_HPP_ #define ASSOCIATE_HPP_ +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include "ast.hpp" #include "skipper.hpp" #include "error_handler.hpp" diff --git a/src/Utilities/SDDSParser/ast.hpp b/src/Utilities/SDDSParser/ast.hpp index 4a75f840d37f19db571918b376b9d7dab96300d5..c27eeb83db124a54772f7398d0c1fc185f191c34 100644 --- a/src/Utilities/SDDSParser/ast.hpp +++ b/src/Utilities/SDDSParser/ast.hpp @@ -17,6 +17,29 @@ #ifndef AST_HPP_ #define AST_HPP_ +#ifdef __CUDACC__ +#pragma push_macro("__cpp_consteval") +#pragma push_macro("_NODISCARD") +#pragma push_macro("__builtin_LINE") + +#define __cpp_consteval 201811L + +#ifdef _NODISCARD + #undef _NODISCARD + #define _NODISCARD +#endif + +#define consteval constexpr + +#include <source_location> + +#undef consteval +#pragma pop_macro("__cpp_consteval") +#pragma pop_macro("_NODISCARD") +#else +#include <source_location> +#endif + #include <boost/variant.hpp> #include <boost/spirit/include/qi.hpp> @@ -92,4 +115,4 @@ namespace SDDS { } } -#endif // AST_HPP_ \ No newline at end of file +#endif // AST_HPP_