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

- Split the transport from the integration-related parts by creating a
  new library.
- 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

@@ -156,23 +156,21 @@ endif()
 #----------------------------------------------------------------------------#
 # Build Targets
 #----------------------------------------------------------------------------#
-set(ADEPT_G4_INTEGRATION_SRCS
-  src/AdePTTrackingManager.cc
-  src/AdePTTrackingManager.cu
-  src/AdePTPhysics.cc
-  src/HepEMPhysics.cc
-  src/AdePTGeant4Integration.cpp
-  src/AdePTConfigurationMessenger.cc
-)
-
 add_library(CopCore INTERFACE)
 target_include_directories(CopCore
   INTERFACE
     $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/include/AdePT/copcore/>
     $<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/>
 )
 
-add_library(AdePT_G4_integration SHARED ${ADEPT_G4_INTEGRATION_SRCS})
+add_library(AdePT_G4_integration SHARED
+  src/AdePTTrackingManager.cc
+  src/AdePTPhysics.cc
+  src/HepEMPhysics.cc
+  src/AdePTGeant4Integration.cpp
+  src/AdePTConfigurationMessenger.cc
+  src/AdePTConfiguration.cc
+)
 target_include_directories(AdePT_G4_integration
   PUBLIC
     $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>
@@ -182,17 +180,29 @@ target_link_libraries(AdePT_G4_integration
   PUBLIC
     CopCore
     VecGeom::vecgeom
-    VecGeom::vecgeomcuda_static
     VecGeom::vgdml
     ${Geant4_LIBRARIES}
     G4HepEm::g4HepEm
-    G4HepEm::g4HepEmData
-    G4HepEm::g4HepEmInit
-    G4HepEm::g4HepEmRun
-    CUDA::cudart
 )
 
-set_target_properties(AdePT_G4_integration
+add_library(AdePTTransport SHARED
+  src/AdePTTransport.cc
+  src/AdePTTransport.cu
+)
+target_include_directories(AdePTTransport
+  PUBLIC
+    $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>
+    $<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/>
+)
+target_link_libraries(AdePTTransport
+  PUBLIC
+    AdePT_G4_integration
+    VecGeom::vecgeomcuda_static
+    VecGeom::vecgeomcuda
+    G4HepEm::g4HepEm
+)
+
+set_target_properties(AdePTTransport
   PROPERTIES
     CUDA_SEPARABLE_COMPILATION ON
     CUDA_RESOLVE_DEVICE_SYMBOLS ON
@@ -223,7 +233,7 @@ configure_package_config_file(cmake/${PROJECT_NAME}Config.cmake.in
 )
 
 #Install the libraries
-install(TARGETS CopCore AdePT_G4_integration
+install(TARGETS CopCore AdePT_G4_integration AdePTTransport
   EXPORT ${PROJECT_NAME}Targets
   ARCHIVE DESTINATION "${CMAKE_INSTALL_LIBDIR}"
   LIBRARY DESTINATION "${CMAKE_INSTALL_LIBDIR}"
@@ -241,7 +251,7 @@ install(FILES "${PROJECT_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/AdePTConfig.cmake"
 )
 
 #Export the targets file
-export(TARGETS CopCore AdePT_G4_integration
+export(TARGETS CopCore AdePT_G4_integration AdePTTransport
   FILE "${PROJECT_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/${PROJECT_NAME}Targets.cmake"
 )
 

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}
+    AdePTTransport
 )
 
 # Install macros and geometry file

examples/IntegrationBenchmark/CMakeLists.txt

@@ -50,16 +50,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} 
+    AdePTTransport
+    ${HEPMC3_LIBRARIES}
     ${HEPMC3_FIO_LIBRARIES}
 )
 

include/AdePT/core/AdePTConfiguration.hh

@@ -4,14 +4,26 @@
 #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>
 
 class AdePTConfiguration {
 public:
-  AdePTConfiguration() { fAdePTConfigurationMessenger = new AdePTConfigurationMessenger(this); }
-  ~AdePTConfiguration() { delete fAdePTConfigurationMessenger; }
+  /// @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.
+  using AdeptFactoryFunction_t = std::shared_ptr<AdePTTransportInterface> (*)(
+      unsigned int nThread, unsigned int nTrackSlot, unsigned int nHitSlot, int verbosity,
+      std::vector<std::string> const *GPURegionNames, bool trackInAllRegions);
+
+  AdePTConfiguration();
+  ~AdePTConfiguration();
   void SetRandomSeed(int randomSeed) { fRandomSeed = randomSeed; }
   void SetTrackInAllRegions(bool trackInAllRegions) { fTrackInAllRegions = trackInAllRegions; }
   void AddGPURegionName(std::string name) { fGPURegionNames.push_back(name); }
@@ -22,6 +34,9 @@ public:
   void SetMillionsOfHitSlots(double millionSlots) { fMillionsOfHitSlots = millionSlots; }
   void SetHitBufferFlushThreshold(float threshold) { fHitBufferFlushThreshold = threshold; }
   void SetCUDAStackLimit(int limit) { fCUDAStackLimit = limit; }
+  /// Register a function to create AdePT instances. This function will be called on every thread that
+  /// needs an instance of AdePT.
+  static void SetAdePTFactoryFunction(AdeptFactoryFunction_t func) { sFactoryFunction = func; }
 
   // We temporarily load VecGeom geometry from GDML
   void SetVecGeomGDML(std::string filename) { fVecGeomGDML = filename; }
@@ -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,11 @@ 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;
+  inline static AdeptFactoryFunction_t sFactoryFunction = nullptr;
 };
 
 #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>
@@ -38,6 +41,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)
 {
@@ -65,7 +73,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;
@@ -491,3 +499,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