[mini] Store G4RegionId (#455)

This PR changes the fGPURegion to actually store the G4RegionId. Before,
it was 0 for non-GPU and 1 for GPU (should have been a bool), now it
actually stores the G4RegionId if it is used on the GPU, otherwise it is
set to -1. This is in preparation for Woodcock tracking on the GPU.

Author: SeverinDiederichs

include/AdePT/core/AdePTTransport.cuh

@@ -220,7 +220,7 @@ __global__ void InitTracks(adeptint::TrackData *trackinfo, int ntracks, int star
     // nextState is initialized as needed.
 
     int lvolID = track.navState.GetLogicalId();
-    assert(auxDataArray[lvolID].fGPUregion);
+    assert(auxDataArray[lvolID].fGPUregionId >= 0);
   }
 }
 

include/AdePT/core/CommonStruct.h

@@ -43,9 +43,9 @@ namespace adeptint {
 /// @brief Auxiliary logical volume data. This stores in the same structure the material-cuts couple index,
 /// the sensitive volume handler index and the flag if the region is active for AdePT.
 struct VolAuxData {
-  int fSensIndex{-1}; ///< index of handler for sensitive volumes (-1 means non-sensitive)
-  int fMCIndex{0};    ///< material-cut cuple index in G4HepEm
-  int fGPUregion{0};  ///< GPU region index (currently 1 or 0, meaning tracked on GPU or not)
+  int fSensIndex{-1};   ///< index of handler for sensitive volumes (-1 means non-sensitive)
+  int fMCIndex{0};      ///< material-cut cuple index in G4HepEm
+  int fGPUregionId{-1}; ///< GPU region index, corresponds to G4Region.instanceID if tracked on GPU, -1 otherwise
 };
 
 /// @brief Structure holding the arrays of auxiliary volume data on host and device

include/AdePT/kernels/electrons.cuh

@@ -502,7 +502,7 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
           const int nextlvolID          = nextState.GetLogicalId();
           VolAuxData const &nextauxData = gVolAuxData[nextlvolID];
           // track has left GPU region
-          if (nextauxData.fGPUregion <= 0) {
+          if (nextauxData.fGPUregionId < 0) {
             // To be safe, just push a bit the track exiting the GPU region to make sure
             // Geant4 does not relocate it again inside the same region
             pos += kPushDistance * dir;

include/AdePT/kernels/electrons_async.cuh

@@ -73,7 +73,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
     // TODO: Why do we have a specific AsyncAdePT VolAuxData?
     VolAuxData const &auxData = AsyncAdePT::gVolAuxData[lvolID];
 
-    assert(auxData.fGPUregion > 0); // make sure we don't get inconsistent region here
+    assert(auxData.fGPUregionId >= 0); // make sure we don't get inconsistent region here
     SlotManager &slotManager = IsElectron ? *secondaries.electrons.fSlotManager : *secondaries.positrons.fSlotManager;
 
     auto survive = [&](bool leak = false) {
@@ -87,7 +87,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
         activeQueue->push_back(slot);
     };
 
-    if (auxData.fGPUregion == 0) {
+    if (auxData.fGPUregionId < 0) {
       printf(__FILE__ ":%d Error: Should kick particle %d lvol %d (%15.12f %15.12f %15.12f) dir=(%f %f %f) "
                       "evt=%d thread=%d "
                       "e=%f safety=%f out of GPU\n",
@@ -337,7 +337,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
         const int nextlvolID = navState.GetLogicalId();
 #endif
         VolAuxData const &nextauxData = AsyncAdePT::gVolAuxData[nextlvolID];
-        if (nextauxData.fGPUregion > 0)
+        if (nextauxData.fGPUregionId >= 0)
           survive();
         else {
           // To be safe, just push a bit the track exiting the GPU region to make sure
@@ -386,7 +386,7 @@ static __device__ __forceinline__ void TransportElectrons(Track *electrons, cons
     //     const auto nextvolume         = navState.Top();
     //     const int nextlvolID          = nextvolume->GetLogicalVolume()->id();
     //     VolAuxData const &nextauxData = AsyncAdePT::gVolAuxData[nextlvolID];
-    //     if (nextauxData.fGPUregion > 0)
+    //     if (nextauxData.fGPUregionId >= 0)
     //       survive();
     //     else {
     //       // To be safe, just push a bit the track exiting the GPU region to make sure

include/AdePT/kernels/electrons_split.cuh

@@ -663,7 +663,7 @@ __global__ void ElectronRelocation(Track *electrons, Track *leaks, G4HepEmElectr
       // Check if the next volume belongs to the GPU region and push it to the appropriate queue
       const int nextlvolID          = currentTrack.navState.GetLogicalId();
       VolAuxData const &nextauxData = AsyncAdePT::gVolAuxData[nextlvolID];
-      if (nextauxData.fGPUregion > 0) {
+      if (nextauxData.fGPUregionId >= 0) {
         theTrack->SetMCIndex(nextauxData.fMCIndex);
         survive();
       } else {

include/AdePT/kernels/gammas.cuh

@@ -236,7 +236,7 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
         const int nextlvolID          = nextState.GetLogicalId();
         VolAuxData const &nextauxData = gVolAuxData[nextlvolID];
 
-        if (nextauxData.fGPUregion > 0) {
+        if (nextauxData.fGPUregionId >= 0) {
           surviveFlag = true;
         } else {
           // To be safe, just push a bit the track exiting the GPU region to make sure

include/AdePT/kernels/gammas_async.cuh

@@ -42,7 +42,7 @@ __global__ void __launch_bounds__(256, 1)
     auto navState              = currentTrack.navState;
     const auto preStepNavState = navState;
     VolAuxData const &auxData  = AsyncAdePT::gVolAuxData[currentTrack.navState.Top()->GetLogicalVolume()->id()];
-    assert(auxData.fGPUregion > 0); // make sure we don't get inconsistent region here
+    assert(auxData.fGPUregionId >= 0); // make sure we don't get inconsistent region here
     auto &slotManager = *secondaries.gammas.fSlotManager;
 
     // Write local variables back into track and enqueue
@@ -115,7 +115,7 @@ __global__ void __launch_bounds__(256, 1)
         const auto nextvolume         = navState.Top();
         const int nextlvolID          = nextvolume->GetLogicalVolume()->id();
         VolAuxData const &nextauxData = AsyncAdePT::gVolAuxData[nextlvolID];
-        if (nextauxData.fGPUregion > 0)
+        if (nextauxData.fGPUregionId >= 0)
           survive(activeQueue);
         else {
           // To be safe, just push a bit the track exiting the GPU region to make sure

include/AdePT/kernels/gammas_split.cuh

@@ -313,7 +313,7 @@ __global__ void GammaRelocation(Track *gammas, Track *leaks, G4HepEmGammaTrack *
       //  Check if the next volume belongs to the GPU region and push it to the appropriate queue
       const int nextlvolID          = currentTrack.navState.GetLogicalId();
       VolAuxData const &nextauxData = AsyncAdePT::gVolAuxData[nextlvolID];
-      if (nextauxData.fGPUregion > 0) {
+      if (nextauxData.fGPUregionId >= 0) {
         theTrack->SetMCIndex(nextauxData.fMCIndex);
         survive();
       } else {

src/AdePTGeant4Integration.cpp

@@ -324,11 +324,11 @@ void AdePTGeant4Integration::InitVolAuxData(adeptint::VolAuxData *volAuxData, G4
     if (!trackInAllRegions) {
       for (G4Region *gpuRegion : gpuRegions) {
         if (g4_lvol->GetRegion() == gpuRegion) {
-          volAuxData[vg_lvol->id()].fGPUregion = 1;
+          volAuxData[vg_lvol->id()].fGPUregionId = g4_lvol->GetRegion()->GetInstanceID();
         }
       }
     } else {
-      volAuxData[vg_lvol->id()].fGPUregion = 1;
+      volAuxData[vg_lvol->id()].fGPUregionId = g4_lvol->GetRegion()->GetInstanceID();
     }
 
     if (g4_lvol->GetSensitiveDetector() != nullptr) {