ApplyCuts (rebase) (#334)

This PR adds the `ApplyCuts` used by G4 in AdePT. It is based and
supersedes PR #317. The `ApplyCuts` value is read from Geant4 and can be
toggled via `/process/em/applyCuts true`

Note that currently the positron cut is not yet available in G4HepEm
(but will be shortly), and we use just use the electron cut. This will
be fixed soon.

The restructuring of the scoring is now fully functional.

The run time shows that the `ApplyCuts` works as intended and
accelerates the simulation. Since the scoring was also effected, the
table also shows the results with `ApplyCuts = false`, which shows just
the new scoring implementation
| Setting | 4 ttbar 1 threads CMS (no Bfield) | 1e4 10 GeV electrons Em3
|
|-------------|-------------|---------------|
| master |  90.4422 s | 76.5938 s|
| new scoring | 90.3233 s| 75.4082 s |
| new scoring + ApplyCuts | 87.3564 s | 60.8005 s|

Additionally, the update of the time for gammas was incorrect (as it
would have been skipped if a boundary was hit), this is also fixed.

The physics validation looks really good using the `ApplyCuts` both in
G4 and AdePT (for 10^7 primaries

Co-authored-by: JuanGonzalezCaminero <[email protected]>

include/AdePT/core/AdePTTransport.cuh

@@ -56,6 +56,7 @@ inline __constant__ __device__ struct G4HepEmData g4HepEmData;
 
 inline __constant__ __device__ adeptint::VolAuxData *gVolAuxData = nullptr;
 inline __constant__ __device__ double BzFieldValue               = 0;
+inline __constant__ __device__ bool ApplyCuts                    = false;
 
 bool InitializeVolAuxArray(adeptint::VolAuxArray &array)
 {
@@ -108,6 +109,13 @@ G4HepEmState *InitG4HepEm()
   return state;
 }
 
+bool InitializeApplyCuts(bool applycuts)
+{
+  // Initialize ApplyCut
+  COPCORE_CUDA_CHECK(cudaMemcpyToSymbol(ApplyCuts, &applycuts, sizeof(bool)));
+  return true;
+}
+
 // Kernel function to initialize tracks comming from a Geant4 buffer
 __global__ void InitTracks(adeptint::TrackData *trackinfo, int ntracks, int startTrack, int event,
                            Secondaries secondaries, const vecgeom::VPlacedVolume *world, AdeptScoring *userScoring,

include/AdePT/core/AdePTTransport.h

@@ -70,6 +70,8 @@ class AdePTTransport : public AdePTTransportInterface {
   /// @brief Create material-cut couple index array
   /// @brief Initialize service and copy geometry & physics data on device
   void Initialize(bool common_data = false);
+  /// @brief Initializes the ApplyCut flag. Can only be called after G4 Physics is build
+  bool InitializeApplyCuts(bool applycuts);
   /// @brief Final cleanup
   void Cleanup();
   /// @brief Interface for transporting a buffer of tracks in AdePT.

include/AdePT/core/AdePTTransport.icc

@@ -25,6 +25,7 @@ namespace adept_impl {
 /// Forward declarations for methods implemented in AdePTTransport.cu
 using TrackBuffer = adeptint::TrackBuffer;
 bool InitializeField(double);
+bool InitializeApplyCuts(bool);
 bool InitializeVolAuxArray(adeptint::VolAuxArray &);
 void FreeVolAuxArray(adeptint::VolAuxArray &);
 G4HepEmState *InitG4HepEm();
@@ -60,6 +61,12 @@ bool AdePTTransport<IntegrationLayer>::InitializeField(double bz)
   return adept_impl::InitializeField(bz);
 }
 
+template <typename IntegrationLayer>
+bool AdePTTransport<IntegrationLayer>::InitializeApplyCuts(bool applycuts)
+{
+  return adept_impl::InitializeApplyCuts(applycuts);
+}
+
 template <typename IntegrationLayer>
 void AdePTTransport<IntegrationLayer>::AddTrack(int pdg, int parent_id, double energy, double x, double y, double z,
                                                 double dirx, double diry, double dirz, double globalTime,

include/AdePT/core/AdePTTransportInterface.hh

@@ -41,6 +41,8 @@ public:
   virtual void SetCUDAStackLimit(int limit)                                   = 0;
   /// @brief Initialize service and copy geometry & physics data on device
   virtual void Initialize(bool common_data = false) = 0;
+  /// @brief Initialize the ApplyCuts flag on device
+  virtual bool InitializeApplyCuts(bool applycuts) = 0;
   /// @brief Interface for transporting a buffer of tracks in AdePT.
   virtual void Shower(int event, int threadId) = 0;
   virtual void Cleanup()                       = 0;

include/AdePT/core/AdePTTransportStruct.cuh

@@ -109,5 +109,6 @@ extern __constant__ struct G4HepEmData g4HepEmData;
 
 extern __constant__ __device__ adeptint::VolAuxData *gVolAuxData;
 extern __constant__ __device__ double BzFieldValue;
+extern __constant__ __device__ bool ApplyCuts;
 } // namespace adept_impl
 #endif