AdePT compiles and runs, but there are issues with the cleanup

CMakeLists.txt

@@ -124,7 +124,7 @@ add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CONFIG:RelWithDebInfo>>:
 # - For Debug, generate full debug information - this completely disables optimizations!
 add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<CONFIG:Debug>>:--device-debug>")
 # - For both, interleave the source in PTX to enhance the debugging experience.
-add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<OR:$<CONFIG:RelWithDebInfo>,$<CONFIG:Debug>>>:--source-in-ptx>")
+# add_compile_options("$<$<AND:$<COMPILE_LANGUAGE:CUDA>,$<OR:$<CONFIG:RelWithDebInfo>,$<CONFIG:Debug>>>:--source-in-ptx>")
 
 # Disable warnings from the CUDA frontend about unknown GCC pragmas - let the compiler decide what it likes.
 add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcudafe;--diag_suppress=unrecognized_gcc_pragma>")

include/AdePT/core/AdePTScoring.cuh → include/AdePT/core/AdePTScoringTemplate.cuh

@@ -7,7 +7,7 @@
 namespace adept_scoring
 {
   template <typename Scoring>
-  Scoring* InitializeOnGPU(Scoring &scoring){ return nullptr; }
+  Scoring* InitializeOnGPU(Scoring *scoring){}
 
   template <typename Scoring>
   void FreeGPU(Scoring *scoring_dev){}

include/AdePT/core/AdePTTransport.cuh

@@ -1,9 +1,9 @@
 // SPDX-FileCopyrightText: 2022 CERN
 // SPDX-License-Identifier: Apache-2.0
 
-#include <AdePT/core/AdePTScoring.cuh>
-#include <AdePT/core/HostScoring.cuh>
-#include <AdePT/core/HostScoringImpl.cuh>
+#include <AdePT/core/AdePTScoringTemplate.cuh>
+#include <AdePT/core/HostScoringStruct.cuh>
+#include <AdePT/core/HostScoringCudaImpl.cuh>
 
 #include <AdePT/core/AdePTTransportStruct.cuh>
 #include <AdePT/base/Atomic.h>
@@ -212,15 +212,12 @@ void PrepareLeakedBuffers(int numLeaked, adeptint::TrackBuffer &buffer, GPUstate
   }
 }
 
-GPUstate *InitializeGPU(adeptint::TrackBuffer &buffer, int capacity, int maxbatch, AdeptScoring *scoring, AdeptScoring *scoring_dev)
+GPUstate *InitializeGPU(adeptint::TrackBuffer &buffer, int capacity, int maxbatch)
 {
   using TrackData    = adeptint::TrackData;
   auto gpuState_ptr  = new GPUstate;
   auto &gpuState      = *gpuState_ptr;
 
-  // Initialize Scoring
-  scoring_dev = adept_scoring::InitializeOnGPU(*scoring);
-
   // Allocate track managers, streams and synchronization events.
   const size_t kQueueSize = MParrayTracks::SizeOfInstance(capacity);
   // Create a stream to synchronize kernels of all particle types.
@@ -249,6 +246,12 @@ GPUstate *InitializeGPU(adeptint::TrackBuffer &buffer, int capacity, int maxbatc
   return gpuState_ptr;
 }
 
+AdeptScoring *InitializeScoringGPU(AdeptScoring *scoring)
+{
+  // Initialize Scoring
+  return adept_scoring::InitializeOnGPU(scoring);
+}
+
 void FreeGPU(GPUstate &gpuState, G4HepEmState *g4hepem_state)
 {
   // Free resources.

include/AdePT/core/AdePTTransport.h

@@ -17,8 +17,8 @@
 #include <G4HepEmState.hh>
 
 #include "CommonStruct.h"
-#include <AdePT/core/AdePTScoring.cuh>
-#include <AdePT/core/HostScoring.cuh>
+#include <AdePT/core/AdePTScoringTemplate.cuh>
+#include <AdePT/core/HostScoringStruct.cuh>
 
 class G4Region;
 struct GPUstate;

include/AdePT/core/AdePTTransport.icc

@@ -24,7 +24,8 @@ bool InitializeField(double);
 bool InitializeVolAuxArray(adeptint::VolAuxArray &);
 void FreeVolAuxArray(adeptint::VolAuxArray &);
 G4HepEmState *InitG4HepEm();
-GPUstate *InitializeGPU(TrackBuffer &, int, int, AdeptScoring *, AdeptScoring *);
+GPUstate *InitializeGPU(TrackBuffer &, int, int);
+AdeptScoring *InitializeScoringGPU(AdeptScoring *scoring);
 void FreeGPU(GPUstate &, G4HepEmState *);
 template <typename IntegrationLayer>
 void ShowerGPU(IntegrationLayer &integration, int event, TrackBuffer &buffer, GPUstate &gpuState, AdeptScoring *scoring,
@@ -137,12 +138,12 @@ void AdePTTransport<IntegrationLayer>::Initialize(bool common_data)
   std::cout << "=== AdePTTransport: initializing transport engine for thread: " << fIntegrationLayer.GetThreadID()
             << std::endl;
 
-  // Initialize user scoring data
+  // Initialize user scoring data on Host
   fScoring     = new AdeptScoring(fHitBufferCapacity);
-  //fScoring_dev = adept_scoring::InitializeOnGPU(fScoring);
 
   // Initialize the transport engine for the current thread
-  fGPUstate = adept_impl::InitializeGPU(fBuffer, fCapacity, fMaxBatch, fScoring, fScoring_dev);
+  fGPUstate = adept_impl::InitializeGPU(fBuffer, fCapacity, fMaxBatch);
+  fScoring_dev = adept_impl::InitializeScoringGPU(fScoring);
 
   fInit = true;
 }

include/AdePT/core/AdePTTransportStruct.cuh

@@ -5,7 +5,7 @@
 #define ADEPT_TRANSPORT_STRUCT_CUH
 
 #include <AdePT/core/CommonStruct.h>
-#include <AdePT/core/HostScoring.cuh>
+#include <AdePT/core/HostScoringStruct.cuh>
 
 #include "Track.cuh"
 #include <AdePT/base/TrackManager.cuh>

include/AdePT/core/HostScoringImpl.cuh → include/AdePT/core/HostScoringCudaImpl.cuh

@@ -1,4 +1,4 @@
-#include <AdePT/core/HostScoring.cuh>
+#include <AdePT/core/HostScoringStruct.cuh>
 
 // CUDA Methods specific to HostScoring
 
@@ -96,35 +96,35 @@ void CopyGlobalCountersToHost(HostScoring &hostScoring, cudaStream_t stream)
                                      cudaMemcpyDeviceToHost, stream));
 }
 
-#include <AdePT/core/AdePTScoring.cuh>
 
-// Specialization of CUDA Methods for HostScoring
 
+// Specialization of CUDA Methods for HostScoring
+#include <AdePT/core/AdePTScoringTemplate.cuh>
 namespace adept_scoring
 {
   /// @brief Allocate and initialize data structures on device
   template <>
-  HostScoring* InitializeOnGPU(HostScoring &hostScoring)
+  HostScoring* InitializeOnGPU(HostScoring *hostScoring)
   {
     // Allocate space for the hits buffer
-    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring.fGPUHitsBuffer_device, sizeof(GPUHit) * hostScoring.fBufferCapacity));
+    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring->fGPUHitsBuffer_device, sizeof(GPUHit) * hostScoring->fBufferCapacity));
 
     // Allocate space for the global counters
-    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring.fGlobalCounters_dev, sizeof(GlobalCounters)));
+    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring->fGlobalCounters_dev, sizeof(GlobalCounters)));
 
     // Allocate space for the atomic variables on device
-    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring.fUsedSlots_d, sizeof(adept::Atomic_t<unsigned int>)));
-    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring.fNextFreeHit_d, sizeof(adept::Atomic_t<unsigned int>)));
+    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring->fUsedSlots_d, sizeof(adept::Atomic_t<unsigned int>)));
+    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring->fNextFreeHit_d, sizeof(adept::Atomic_t<unsigned int>)));
 
     // Allocate space for the stats on device
     // Allocate space for the global counters
-    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring.fStats_device, sizeof(HostScoring::Stats)));
+    COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring->fStats_device, sizeof(HostScoring::Stats)));
 
     // Allocate space for the instance on GPU and copy the data members from the host
     // Now allocate space for the BasicScoring placeholder on device and copy the device pointers of components
     HostScoring *hostScoring_dev = nullptr;
     COPCORE_CUDA_CHECK(cudaMalloc(&hostScoring_dev, sizeof(HostScoring)));
-    COPCORE_CUDA_CHECK(cudaMemcpy(hostScoring_dev, &hostScoring, sizeof(HostScoring), cudaMemcpyHostToDevice));
+    COPCORE_CUDA_CHECK(cudaMemcpy(hostScoring_dev, hostScoring, sizeof(HostScoring), cudaMemcpyHostToDevice));
 
     return hostScoring_dev;
   }

include/AdePT/integration/AdePTGeant4Integration.hh

@@ -14,7 +14,7 @@
 #include <G4HepEmState.hh>
 
 #include <AdePT/core/CommonStruct.h>
-#include <AdePT/core/HostScoring.cuh>
+#include <AdePT/core/HostScoringStruct.cuh>
 
 #include <G4VPhysicalVolume.hh>
 #include <G4LogicalVolume.hh>

include/AdePT/integration/AdePTTrackingManager.hh

@@ -33,11 +33,11 @@ public:
   void SetVerbosity(int verbosity) { fVerbosity = verbosity; }
 
   // Set the AdePTTransport instance, also initializes the GPU region list
-  void SetAdePTTransport(AdePTTransport<AdePTGeant4Integration> *adept)
+  void SetAdePTTransport(AdePTTransport<AdePTGeant4Integration> *adeptTransport)
   {
-    fAdept = adept;
-    if (!adept->GetTrackInAllRegions()) {
-      for (std::string regionName : *(adept->GetGPURegionNames())) {
+    fAdeptTransport = adeptTransport;
+    if (!adeptTransport->GetTrackInAllRegions()) {
+      for (std::string regionName : *(adeptTransport->GetGPURegionNames())) {
         G4cout << "AdePTTrackingManager: Marking " << regionName << " as a GPU Region" << G4endl;
         G4Region *region = G4RegionStore::GetInstance()->GetRegion(regionName);
         if (region != nullptr)
@@ -61,7 +61,7 @@ private:
   void StepInHostRegion(G4Track *aTrack);
 
   std::vector<G4Region *> fGPURegions{};
-  AdePTTransport<AdePTGeant4Integration> *fAdept;
+  AdePTTransport<AdePTGeant4Integration> *fAdeptTransport;
   int fVerbosity{0};
   G4double ProductionCut = 0.7 * copcore::units::mm;
   int MCIndex[100];

src/AdePTTrackingManager.cc

@@ -21,7 +21,7 @@ AdePTTrackingManager::AdePTTrackingManager() {}
 
 AdePTTrackingManager::~AdePTTrackingManager()
 {
-  fAdept->Cleanup();
+  fAdeptTransport->Cleanup();
 }
 
 //....oooOO0OOooo........oooOO0OOooo........oooOO0OOooo........oooOO0OOooo......
@@ -78,9 +78,9 @@ void AdePTTrackingManager::FlushEvent()
 
   if (fVerbosity > 0)
     G4cout << "No more particles on the stack, triggering shower to flush the AdePT buffer with "
-           << fAdept->GetNtoDevice() << " particles left." << G4endl;
+           << fAdeptTransport->GetNtoDevice() << " particles left." << G4endl;
 
-  fAdept->Shower(G4EventManager::GetEventManager()->GetConstCurrentEvent()->GetEventID());
+  fAdeptTransport->Shower(G4EventManager::GetEventManager()->GetConstCurrentEvent()->GetEventID());
 }
 
 void AdePTTrackingManager::ProcessTrack(G4Track *aTrack)
@@ -119,7 +119,7 @@ void AdePTTrackingManager::ProcessTrack(G4Track *aTrack)
     G4Region *region = aTrack->GetVolume()->GetLogicalVolume()->GetRegion();
     // Check if the particle is in a GPU region
     bool isGPURegion = false;
-    if (fAdept->GetTrackInAllRegions()) {
+    if (fAdeptTransport->GetTrackInAllRegions()) {
       isGPURegion = true;
     } else {
       for (G4Region *gpuRegion : fGPURegions) {
@@ -138,7 +138,7 @@ void AdePTTrackingManager::ProcessTrack(G4Track *aTrack)
       G4double properTime    = aTrack->GetProperTime();
       auto pdg               = aTrack->GetParticleDefinition()->GetPDGEncoding();
 
-      fAdept->AddTrack(pdg, energy, particlePosition[0], particlePosition[1], particlePosition[2], particleDirection[0],
+      fAdeptTransport->AddTrack(pdg, energy, particlePosition[0], particlePosition[1], particlePosition[2], particleDirection[0],
                        particleDirection[1], particleDirection[2], globalTime, localTime, properTime);
 
       // The track dies from the point of view of Geant4
@@ -180,7 +180,7 @@ void AdePTTrackingManager::StepInHostRegion(G4Track *aTrack)
       aTrack->SetTouchableHandle(aTrack->GetNextTouchableHandle());
       G4Region *region = aTrack->GetVolume()->GetLogicalVolume()->GetRegion();
       // This should never be true if this flag is set, as all particles would be sent to AdePT
-      assert(fAdept->GetTrackInAllRegions() == false);
+      assert(fAdeptTransport->GetTrackInAllRegions() == false);
       // Check whether the particle has entered a GPU region
       for (G4Region *gpuRegion : fGPURegions) {
         if (region == gpuRegion) {