Changing kernel names

apt-sim · Sep 24, 2024 · a819871 · a819871
1 parent c4f2c4f
commit a819871
Show file tree

Hide file tree

Showing 5 changed files with 16 additions and 28 deletions.
diff --git a/examples/Example1/CMakeLists.txt b/examples/Example1/CMakeLists.txt
@@ -74,8 +74,6 @@ configure_file("macros/example1_ttbar_noadept.mac.in" "${PROJECT_BINARY_DIR}/exa
 
 # Tests
 
-string(APPEND CMAKE_CUDA_FLAGS " -Xptxas=-v") 
-
 add_test(NAME example1
   COMMAND $<TARGET_FILE:example1> -m ${PROJECT_BINARY_DIR}/example1_large_stack.mac
 )
diff --git a/examples/Example1/macros/example1_ttbar_LHCb.mac.in b/examples/Example1/macros/example1_ttbar_LHCb.mac.in
@@ -26,7 +26,7 @@
 /adept/setMillionsOfTrackSlots 4
 /adept/setMillionsOfHitSlots 1
 ## Device stack limit
-/adept/setCUDAStackLimit 4096
+# /adept/setCUDAStackLimit 4096
 
 # If true, particles are transported on the GPU across the whole geometry, GPU regions are ignored
 /adept/setTrackInAllRegions true

diff --git a/include/AdePT/core/AdePTTransport.cuh b/include/AdePT/core/AdePTTransport.cuh
@@ -407,12 +407,12 @@ void ShowerGPU(IntegrationLayer &integration, int event, adeptint::TrackBuffer &
           electrons.trackmgr, secondaries, electrons.leakedTracks, scoring_dev,
           VolAuxArray::GetInstance().fAuxData_dev);
 #else
-      ElectronPhysics1<true><<<transportBlocks, TransportThreads, 0, electrons.stream>>>(
+      ElectronHowFar<true><<<transportBlocks, TransportThreads, 0, electrons.stream>>>(
         electrons.trackmgr, gpuState.hepEMBuffers_d.electronsHepEm, VolAuxArray::GetInstance().fAuxData_dev
       );
-      ElectronTransport1<true><<<transportBlocks, TransportThreads, 0, electrons.stream>>>(
+      ElectronPropagation<true><<<transportBlocks, TransportThreads, 0, electrons.stream>>>(
         electrons.trackmgr, gpuState.hepEMBuffers_d.electronsHepEm);
-      MSC1<true><<<transportBlocks, TransportThreads, 0, electrons.stream>>>(
+      ElectronMSC<true><<<transportBlocks, TransportThreads, 0, electrons.stream>>>(
         electrons.trackmgr, gpuState.hepEMBuffers_d.electronsHepEm);
       ElectronRelocation<true><<<transportBlocks, TransportThreads, 0, electrons.stream>>>(
         electrons.trackmgr);
@@ -436,12 +436,12 @@ void ShowerGPU(IntegrationLayer &integration, int event, adeptint::TrackBuffer &
           positrons.trackmgr, secondaries, positrons.leakedTracks, scoring_dev,
           VolAuxArray::GetInstance().fAuxData_dev);
 #else
-      ElectronPhysics1<false><<<transportBlocks, TransportThreads, 0, positrons.stream>>>(
+      ElectronHowFar<false><<<transportBlocks, TransportThreads, 0, positrons.stream>>>(
         positrons.trackmgr, gpuState.hepEMBuffers_d.positronsHepEm, VolAuxArray::GetInstance().fAuxData_dev
       );
-      ElectronTransport1<false><<<transportBlocks, TransportThreads, 0, positrons.stream>>>(
+      ElectronPropagation<false><<<transportBlocks, TransportThreads, 0, positrons.stream>>>(
         positrons.trackmgr, gpuState.hepEMBuffers_d.positronsHepEm);
-      MSC1<false><<<transportBlocks, TransportThreads, 0, positrons.stream>>>(
+      ElectronMSC<false><<<transportBlocks, TransportThreads, 0, positrons.stream>>>(
         positrons.trackmgr, gpuState.hepEMBuffers_d.positronsHepEm);
       ElectronRelocation<false><<<transportBlocks, TransportThreads, 0, positrons.stream>>>(
         positrons.trackmgr);
@@ -464,16 +464,16 @@ void ShowerGPU(IntegrationLayer &integration, int event, adeptint::TrackBuffer &
       TransportGammas<AdeptScoring><<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
           gammas.trackmgr, secondaries, gammas.leakedTracks, scoring_dev, VolAuxArray::GetInstance().fAuxData_dev);
 #else
-      GammaPhysics1<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+      GammaHowFar<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
         gammas.trackmgr, gpuState.hepEMBuffers_d.gammasHepEm, VolAuxArray::GetInstance().fAuxData_dev
       );
-      GammaTransport1<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+      GammaPropagation<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
         gammas.trackmgr, gpuState.hepEMBuffers_d.gammasHepEm, VolAuxArray::GetInstance().fAuxData_dev
       );
       GammaRelocation<<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
         gammas.trackmgr, gammas.leakedTracks, VolAuxArray::GetInstance().fAuxData_dev
       );
-      GammaPhysics2<AdeptScoring><<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
+      GammaInteractions<AdeptScoring><<<transportBlocks, TransportThreads, 0, gammas.stream>>>(
         gammas.trackmgr,  gpuState.hepEMBuffers_d.gammasHepEm, secondaries, scoring_dev, VolAuxArray::GetInstance().fAuxData_dev
       );
 #endif

diff --git a/include/AdePT/kernels/electrons_split.cuh b/include/AdePT/kernels/electrons_split.cuh
@@ -35,7 +35,7 @@ __device__ double GetVelocity(double eKin)
 }
 
 template <bool IsElectron>
-__global__ void ElectronPhysics1(adept::TrackManager<Track> *electrons, G4HepEmElectronTrack *hepEMTracks, VolAuxData const *auxDataArray)
+__global__ void ElectronHowFar(adept::TrackManager<Track> *electrons, G4HepEmElectronTrack *hepEMTracks, VolAuxData const *auxDataArray)
 {
   constexpr int Charge               = IsElectron ? -1 : 1;
   constexpr double restMass          = copcore::units::kElectronMassC2;
@@ -128,7 +128,7 @@ __global__ void ElectronPhysics1(adept::TrackManager<Track> *electrons, G4HepEmE
 }
 
 template <bool IsElectron>
-static __global__ void ElectronTransport1(adept::TrackManager<Track> *electrons, G4HepEmElectronTrack *hepEMTracks)
+static __global__ void ElectronPropagation(adept::TrackManager<Track> *electrons, G4HepEmElectronTrack *hepEMTracks)
 {
 #ifdef VECGEOM_FLOAT_PRECISION
   const Precision kPush = 10 * vecgeom::kTolerance;
@@ -184,7 +184,7 @@ static __global__ void ElectronTransport1(adept::TrackManager<Track> *electrons,
 
 // May work well to separate MSC1 (Continuous effects) and MSC2 (Checks + safety)
 template <bool IsElectron>
-static __global__ void MSC1(adept::TrackManager<Track> *electrons, G4HepEmElectronTrack *hepEMTracks)
+static __global__ void ElectronMSC(adept::TrackManager<Track> *electrons, G4HepEmElectronTrack *hepEMTracks)
 {
   constexpr double restMass          = copcore::units::kElectronMassC2;
   int activeSize = electrons->fActiveTracks->size();

diff --git a/include/AdePT/kernels/gammas_split.cuh b/include/AdePT/kernels/gammas_split.cuh
@@ -20,7 +20,7 @@
 
 using VolAuxData = adeptint::VolAuxData;
 
-__global__ void GammaPhysics1(adept::TrackManager<Track> *gammas, G4HepEmGammaTrack *hepEMTracks, VolAuxData const *auxDataArray)
+__global__ void GammaHowFar(adept::TrackManager<Track> *gammas, G4HepEmGammaTrack *hepEMTracks, VolAuxData const *auxDataArray)
 {
   int activeSize = gammas->fActiveTracks->size();
   for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < activeSize; i += blockDim.x * gridDim.x) {
@@ -58,20 +58,10 @@ __global__ void GammaPhysics1(adept::TrackManager<Track> *gammas, G4HepEmGammaTr
 
     // Call G4HepEm to compute the physics step limit.
     G4HepEmGammaManager::HowFar(&g4HepEmData, &g4HepEmPars, &gammaTrack);
-
-    // hepEMTracks[slot] = gammaTrack;
-
-    // Save the info we need from the G4HepEM track
-    // currentTrack.geometricalStepLengthFromPhysics = theTrack->GetGStepLength();
-    // currentTrack.winnerProcessIndex               = theTrack->GetWinnerProcessIndex();
-    // currentTrack.PEmxSec                          = gammaTrack.GetPEmxSec();
-    // currentTrack.preStepMFPs[0]                   = theTrack->GetMFP(0);
-    // currentTrack.preStepMFPs[1]                   = theTrack->GetMFP(1);
-    // currentTrack.preStepMFPs[2]                   = theTrack->GetMFP(2);
   }
 }
 
-__global__ void GammaTransport1(adept::TrackManager<Track> *gammas, G4HepEmGammaTrack *hepEMTracks, VolAuxData const *auxDataArray)
+__global__ void GammaPropagation(adept::TrackManager<Track> *gammas, G4HepEmGammaTrack *hepEMTracks, VolAuxData const *auxDataArray)
 {
 #ifdef VECGEOM_FLOAT_PRECISION
   const Precision kPush = 10 * vecgeom::kTolerance;
@@ -199,7 +189,7 @@ __global__ void GammaRelocation(adept::TrackManager<Track> *gammas, MParrayTrack
 }
 
 template <typename Scoring>
-__global__ void GammaPhysics2(adept::TrackManager<Track> *gammas, G4HepEmGammaTrack *hepEMTracks, Secondaries secondaries, Scoring *userScoring,
+__global__ void GammaInteractions(adept::TrackManager<Track> *gammas, G4HepEmGammaTrack *hepEMTracks, Secondaries secondaries, Scoring *userScoring,
                          VolAuxData const *auxDataArray)
 {
   int activeSize = gammas->fActiveTracks->size();