[core,integration] Introduce an interface for using different "AdePT"s.

- Split the transport from the integration-related parts introducing a
  new source file. The integration-related parts can be reused in a
  different transport implementation.
- Create a transport abstraction, so AdePTTrackingManager is independent
  of the transport implementation.
- Add thread and event IDs to the transport interface. These are
  necessary for the async transport implementation.
- Start to enumerate tracks in the tracking manager. This can be used to
  reproducibly seed the AdePT random sequences.
- Add some const declarations for the default AdePT implementation.
- Use a factory function to instantiate AdePT. Like this, different
  AdePT implementations can be used without changing code in the tracking
  manager or in AdePTPhysics.
- Replace a few includes with forward declarations.
- Fix device link errors that can show when using a symbol in multiple
  cuda translation units.

CMakeLists.txt

@@ -164,11 +164,11 @@ endif()
 #----------------------------------------------------------------------------#
 set(ADEPT_G4_INTEGRATION_SRCS
   src/AdePTTrackingManager.cc
-  src/AdePTTrackingManager.cu
   src/AdePTPhysics.cc
   src/HepEMPhysics.cc
   src/AdePTGeant4Integration.cpp
   src/AdePTConfigurationMessenger.cc
+  src/AdePTConfiguration.cc
 )
 
 add_library(CopCore INTERFACE)
@@ -178,7 +178,11 @@ target_include_directories(CopCore
     $<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/>
 )
 
-add_library(AdePT_G4_integration SHARED ${ADEPT_G4_INTEGRATION_SRCS})
+add_library(AdePT_G4_integration SHARED
+  ${ADEPT_G4_INTEGRATION_SRCS}
+  src/AdePTTransport.cc
+  src/AdePTTransport.cu
+)
 target_include_directories(AdePT_G4_integration
   PUBLIC
     $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>

examples/Example1/CMakeLists.txt

@@ -51,17 +51,17 @@ endif()
 # example1
 add_executable(example1 example1.cpp ${sources_g4} ${sources_hepmc3})
 target_include_directories(example1
-  PRIVATE 
-    ${PROJECT_SOURCE_DIR}/examples/Example1/include 
+  PRIVATE
+    ${PROJECT_SOURCE_DIR}/examples/Example1/include
     ${PROJECT_SOURCE_DIR}/examples/Example1
     ${PROJECT_SOURCE_DIR}/examples/common/include
     ${HEPMC3_INCLUDE_DIR}
 )
 target_link_libraries(example1
   PRIVATE
-    AdePT_G4_integration
-    ${HEPMC3_LIBRARIES} 
+    ${HEPMC3_LIBRARIES}
     ${HEPMC3_FIO_LIBRARIES}
+    AdePT_G4_integration
 )
 
 # Install macros and geometry file

examples/IntegrationBenchmark/CMakeLists.txt

@@ -52,16 +52,16 @@ endif()
 # integrationBenchmark
 add_executable(integrationBenchmark integrationBenchmark.cpp ${sources_g4} ${sources_hepmc3})
 target_include_directories(integrationBenchmark
-  PRIVATE 
-    ${PROJECT_SOURCE_DIR}/examples/IntegrationBenchmark/include 
+  PRIVATE
+    ${PROJECT_SOURCE_DIR}/examples/IntegrationBenchmark/include
     ${PROJECT_SOURCE_DIR}/examples/IntegrationBenchmark
     ${PROJECT_SOURCE_DIR}/examples/common/include
     ${HEPMC3_INCLUDE_DIR}
 )
 target_link_libraries(integrationBenchmark
   PRIVATE
-    AdePT_G4_integration  
-    ${HEPMC3_LIBRARIES} 
+    AdePT_G4_integration
+    ${HEPMC3_LIBRARIES}
     ${HEPMC3_FIO_LIBRARIES}
 )
 

include/AdePT/core/AdePTConfiguration.hh

@@ -4,14 +4,29 @@
 #ifndef ADEPT_CONFIGURATION_HH
 #define ADEPT_CONFIGURATION_HH
 
+#include <AdePT/integration/AdePTConfigurationMessenger.hh>
+#include <AdePT/core/AdePTTransportInterface.hh>
+
+#include <memory>
 #include <string>
 #include <vector>
-#include <AdePT/integration/AdePTConfigurationMessenger.hh>
 
+/// @brief Factory function to create AdePT instances.
+/// Every AdePT transport implementation needs to provide this function to create
+/// instances of the transport implementation. These might either be one instance
+/// per thread, or share one instance across many threads. This is up to the
+/// transport implementation.
+std::shared_ptr<AdePTTransportInterface> AdePTTransportFactory(unsigned int nThread, unsigned int nTrackSlot,
+                                                               unsigned int nHitSlot, int verbosity,
+                                                               std::vector<std::string> const *GPURegionNames,
+                                                               bool trackInAllRegions);
+
+/// @brief Create and configure instances of an AdePT transport implementation.
+///
 class AdePTConfiguration {
 public:
-  AdePTConfiguration() { fAdePTConfigurationMessenger = new AdePTConfigurationMessenger(this); }
-  ~AdePTConfiguration() { delete fAdePTConfigurationMessenger; }
+  AdePTConfiguration();
+  ~AdePTConfiguration();
   void SetRandomSeed(int randomSeed) { fRandomSeed = randomSeed; }
   void SetTrackInAllRegions(bool trackInAllRegions) { fTrackInAllRegions = trackInAllRegions; }
   void AddGPURegionName(std::string name) { fGPURegionNames.push_back(name); }
@@ -36,6 +51,8 @@ public:
   double GetMillionsOfHitSlots() { return fMillionsOfHitSlots; }
   std::vector<std::string> *GetGPURegionNames() { return &fGPURegionNames; }
 
+  std::shared_ptr<AdePTTransportInterface> CreateAdePTInstance(unsigned int nThread);
+
   // Temporary
   std::string GetVecGeomGDML() { return fVecGeomGDML; }
 
@@ -50,10 +67,10 @@ private:
   double fMillionsOfTrackSlots{1};
   double fMillionsOfHitSlots{1};
   std::vector<std::string> fGPURegionNames{};
+  int fNThread = -1;
 
   std::string fVecGeomGDML{""};
-
-  AdePTConfigurationMessenger *fAdePTConfigurationMessenger;
+  std::unique_ptr<AdePTConfigurationMessenger> fAdePTConfigurationMessenger;
 };
 
 #endif

include/AdePT/core/AdePTTransport.cuh

@@ -1,6 +1,9 @@
 // SPDX-FileCopyrightText: 2022 CERN
 // SPDX-License-Identifier: Apache-2.0
 
+#ifndef ADEPT_TRANSPORT_CUH
+#define ADEPT_TRANSPORT_CUH
+
 #include <AdePT/core/AdePTScoringTemplate.cuh>
 #include <AdePT/core/HostScoringStruct.cuh>
 #include <AdePT/core/HostScoringImpl.cuh>
@@ -48,6 +51,11 @@
 #include <algorithm>
 
 namespace adept_impl {
+inline __constant__ __device__ struct G4HepEmParameters g4HepEmPars;
+inline __constant__ __device__ struct G4HepEmData g4HepEmData;
+
+inline __constant__ __device__ adeptint::VolAuxData *gVolAuxData = nullptr;
+inline __constant__ __device__ double BzFieldValue               = 0;
 
 bool InitializeVolAuxArray(adeptint::VolAuxArray &array)
 {
@@ -75,7 +83,7 @@ G4HepEmState *InitG4HepEm()
 
   // Copy to GPU.
   CopyG4HepEmDataToGPU(state->fData);
-  COPCORE_CUDA_CHECK(cudaMemcpyToSymbol(g4HepEmPars, state->fParameters, sizeof(G4HepEmParameters)));
+  COPCORE_CUDA_CHECK(cudaMemcpyToSymbol(adept_impl::g4HepEmPars, state->fParameters, sizeof(G4HepEmParameters)));
 
   // Create G4HepEmData with the device pointers.
   G4HepEmData dataOnDevice;
@@ -552,3 +560,5 @@ void ShowerGPU(IntegrationLayer &integration, int event, adeptint::TrackBuffer &
   adept_scoring::EndOfTransport<IntegrationLayer>(*scoring, scoring_dev, gpuState.stream, integration);
 }
 } // namespace adept_impl
+
+#endif

include/AdePT/core/AdePTTransport.h

@@ -8,27 +8,25 @@
 #ifndef ADEPT_INTEGRATION_H
 #define ADEPT_INTEGRATION_H
 
+#include "AdePTTransportInterface.hh"
+
 #include <unordered_map>
 #include <VecGeom/base/Config.h>
 #ifdef VECGEOM_ENABLE_CUDA
 #include <VecGeom/management/CudaManager.h> // forward declares vecgeom::cxx::VPlacedVolume
 #endif
 
-#include <G4HepEmState.hh>
-
 #include "CommonStruct.h"
 #include <AdePT/core/AdePTScoringTemplate.cuh>
 #include <AdePT/core/HostScoringStruct.cuh>
 
 class G4Region;
 struct GPUstate;
 class G4VPhysicalVolume;
-
-template <class TTag>
-class TestManager;
+struct G4HepEmState;
 
 template <typename IntegrationLayer>
-class AdePTTransport {
+class AdePTTransport : public AdePTTransportInterface {
 public:
   static constexpr int kMaxThreads = 256;
   using TrackBuffer                = adeptint::TrackBuffer;
@@ -43,8 +41,9 @@ class AdePTTransport {
   int GetNfromDevice() const { return fBuffer.fromDevice.size(); }
 
   /// @brief Adds a track to the buffer
-  void AddTrack(int pdg, int parentID, double energy, double x, double y, double z, double dirx, double diry, double dirz,
-                double globalTime, double localTime, double properTime);
+  void AddTrack(int pdg, int parentID, double energy, double x, double y, double z, double dirx, double diry,
+                double dirz, double globalTime, double localTime, double properTime, int threadId, unsigned int eventId,
+                unsigned int trackIndex);
 
   void SetTrackCapacity(size_t capacity) { fCapacity = capacity; }
   /// @brief Get the track capacity on GPU
@@ -61,19 +60,19 @@ class AdePTTransport {
   void SetDebugLevel(int level) { fDebugLevel = level; }
   /// @brief Set whether AdePT should transport particles across the whole geometry
   void SetTrackInAllRegions(bool trackInAllRegions) { fTrackInAllRegions = trackInAllRegions; }
-  bool GetTrackInAllRegions() { return fTrackInAllRegions; }
+  bool GetTrackInAllRegions() const { return fTrackInAllRegions; }
   /// @brief Set Geant4 region to which it applies
-  void SetGPURegionNames(std::vector<std::string> *regionNames) { fGPURegionNames = regionNames; }
+  void SetGPURegionNames(std::vector<std::string> const *regionNames) { fGPURegionNames = regionNames; }
   /// @brief Set CUDA device stack limit
   void SetCUDAStackLimit(int limit) { fCUDAStackLimit = limit; }
-  std::vector<std::string> *GetGPURegionNames() { return fGPURegionNames; }
+  std::vector<std::string> const *GetGPURegionNames() { return fGPURegionNames; }
   /// @brief Create material-cut couple index array
   /// @brief Initialize service and copy geometry & physics data on device
   void Initialize(bool common_data = false);
   /// @brief Final cleanup
   void Cleanup();
   /// @brief Interface for transporting a buffer of tracks in AdePT.
-  void Shower(int event);
+  void Shower(int event, int threadId);
 
 private:
   static inline G4HepEmState *fg4hepem_state{nullptr}; ///< The HepEm state singleton
@@ -90,7 +89,7 @@ class AdePTTransport {
   AdeptScoring *fScoring{nullptr};                     ///< User scoring object
   AdeptScoring *fScoring_dev{nullptr};                 ///< Device ptr for scoring data
   TrackBuffer fBuffer;                                 ///< Vector of buffers of tracks to/from device (per thread)
-  std::vector<std::string> *fGPURegionNames{};         ///< Region to which applies
+  std::vector<std::string> const *fGPURegionNames{};   ///< Region to which applies
   IntegrationLayer fIntegrationLayer; ///< Provides functionality needed for integration with the simulation toolkit
   bool fInit{false};                  ///< Service initialized flag
   bool fTrackInAllRegions;            ///< Whether the whole geometry is a GPU region

include/AdePT/core/AdePTTransport.icc

@@ -47,7 +47,8 @@ bool AdePTTransport<IntegrationLayer>::InitializeField(double bz)
 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,
-                                                double localTime, double properTime)
+                                                double localTime, double properTime, int /*threadId*/, unsigned int eventId,
+                                                unsigned int /*trackIndex*/)
 {
   fBuffer.toDevice.emplace_back(pdg, parent_id, energy, x, y, z, dirx, diry, dirz, globalTime, localTime, properTime);
   if (pdg == 11)
@@ -60,7 +61,7 @@ void AdePTTransport<IntegrationLayer>::AddTrack(int pdg, int parent_id, double e
   if (fBuffer.toDevice.size() >= fBufferThreshold) {
     if (fDebugLevel > 0)
       std::cout << "Reached the threshold of " << fBufferThreshold << " triggering the shower" << std::endl;
-    this->Shower(fIntegrationLayer.GetEventID());
+    this->Shower(eventId, 0);
   }
 }
 
@@ -199,7 +200,7 @@ void AdePTTransport<IntegrationLayer>::Cleanup()
 }
 
 template <typename IntegrationLayer>
-void AdePTTransport<IntegrationLayer>::Shower(int event)
+void AdePTTransport<IntegrationLayer>::Shower(int event, int /*threadId*/)
 {
   int tid = fIntegrationLayer.GetThreadID();
   if (fDebugLevel > 0 && fBuffer.toDevice.size() == 0) {

include/AdePT/core/AdePTTransportInterface.hh

@@ -0,0 +1,45 @@
+// SPDX-FileCopyrightText: 2022 CERN
+// SPDX-License-Identifier: Apache-2.0
+
+#ifndef ADEPT_TRANSPORT_INTERFACE_H
+#define ADEPT_TRANSPORT_INTERFACE_H
+
+#include <memory>
+#include <string>
+#include <vector>
+
+class AdePTTransportInterface {
+public:
+  virtual ~AdePTTransportInterface() {}
+
+  /// @brief Adds a track to the buffer
+  virtual void AddTrack(int pdg, int id, double energy, double x, double y, double z, double dirx, double diry,
+                        double dirz, double globalTime, double localTime, double properTime, int threadId,
+                        unsigned int eventId, unsigned int trackIndex) = 0;
+
+  /// @brief Set capacity of on-GPU track buffer.
+  virtual void SetTrackCapacity(size_t capacity) = 0;
+  /// @brief Set Hit buffer capacity on GPU and Host
+  virtual void SetHitBufferCapacity(size_t capacity) = 0;
+  /// @brief Set maximum batch size
+  virtual void SetMaxBatch(int npart) = 0;
+  /// @brief Set buffer threshold
+  virtual void SetBufferThreshold(int limit) = 0;
+  /// @brief Set debug level for transport
+  virtual void SetDebugLevel(int level) = 0;
+  /// @brief Set whether AdePT should transport particles across the whole geometry
+  virtual void SetTrackInAllRegions(bool trackInAllRegions) = 0;
+  /// @brief Check whether AdePT should transport particles across the whole geometry
+  virtual bool GetTrackInAllRegions() const = 0;
+  /// @brief Set Geant4 region to which it applies
+  virtual void SetGPURegionNames(std::vector<std::string> const *regionNames) = 0;
+  virtual std::vector<std::string> const *GetGPURegionNames()                 = 0;
+  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 Interface for transporting a buffer of tracks in AdePT.
+  virtual void Shower(int event, int threadId) = 0;
+  virtual void Cleanup()                       = 0;
+};
+
+#endif

include/AdePT/core/AdePTTransportStruct.cuh

@@ -102,21 +102,12 @@ struct GPUstate {
   Stats *stats{nullptr};              ///< statistics object pointer on host
 };
 
-// Constant data structures from G4HepEm accessed by the kernels.
-// (defined in TestEm3.cu)
-extern __constant__ __device__ struct G4HepEmParameters g4HepEmPars;
-extern __constant__ __device__ struct G4HepEmData g4HepEmData;
+namespace adept_impl {
+constexpr double kPush = 1.e-8 * copcore::units::cm;
+extern __constant__ struct G4HepEmParameters g4HepEmPars;
+extern __constant__ struct G4HepEmData g4HepEmData;
 
-// Pointer for array of volume auxiliary data on device
 extern __constant__ __device__ adeptint::VolAuxData *gVolAuxData;
-
-// constexpr float BzFieldValue = 0.1 * copcore::units::tesla;
 extern __constant__ __device__ double BzFieldValue;
-constexpr double kPush = 1.e-8 * copcore::units::cm;
-__constant__ __device__ struct G4HepEmParameters g4HepEmPars;
-__constant__ __device__ struct G4HepEmData g4HepEmData;
-
-__constant__ __device__ adeptint::VolAuxData *gVolAuxData = nullptr;
-__constant__ __device__ double BzFieldValue               = 0;
-
+} // namespace adept_impl
 #endif

include/AdePT/integration/AdePTGeant4Integration.hh

@@ -12,13 +12,13 @@
 #include <AdePT/core/CommonStruct.h>
 #include <AdePT/core/HostScoringStruct.cuh>
 
-#include <G4HepEmState.hh>
-
 #include <G4EventManager.hh>
 #include <G4Event.hh>
 
 #include <unordered_map>
 
+struct G4HepEmState;
+
 namespace AdePTGeant4Integration_detail {
 struct ScoringObjects;
 struct Deleter {

include/AdePT/integration/AdePTPhysics.hh

@@ -6,8 +6,9 @@
 
 #include "G4VPhysicsConstructor.hh"
 #include "globals.hh"
-#include <AdePT/integration/AdePTTrackingManager.hh>
-#include <AdePT/core/AdePTConfiguration.hh>
+
+class AdePTTrackingManager;
+class AdePTConfiguration;
 
 class AdePTPhysics : public G4VPhysicsConstructor {
 public: