Merge branch 'gpu-compile' into 'master'

GPU compilable

See merge request OPAL/opal-x/src!7

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"

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
+#endif  // CLASSIC_Drift_HH
\ No newline at end of file

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"},

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"

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
+#endif  // CLASSIC_Multipole_HH
\ No newline at end of file

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

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 {

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"

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"

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"

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"

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

src/AbsBeamline/ScalingFFAMagnet.cpp

@@ -26,7 +26,6 @@
  */
 
 #include <cmath>
-
 #include "AbsBeamline/BeamlineVisitor.h"
 #include "AbsBeamline/ScalingFFAMagnet.h"
 #include "PartBunch/PartBunch.hpp"

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
 

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"
 

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>

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);
 
 }
 

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&);

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
+#endif  // OPAL_ParallelTracker_HH
\ No newline at end of file

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;
 }