diff --git a/examples/Example1/CMakeLists.txt b/examples/Example1/CMakeLists.txt index adc85a54..595d455f 100644 --- 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 $ -m ${PROJECT_BINARY_DIR}/example1_large_stack.mac ) \ No newline at end of file diff --git a/examples/Example1/macros/example1_ttbar_LHCb.mac.in b/examples/Example1/macros/example1_ttbar_LHCb.mac.in index 7af59167..051736cc 100644 --- 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 index 4b0a4c99..225e157c 100644 --- 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<<>>( + ElectronHowFar<<>>( electrons.trackmgr, gpuState.hepEMBuffers_d.electronsHepEm, VolAuxArray::GetInstance().fAuxData_dev ); - ElectronTransport1<<>>( + ElectronPropagation<<>>( electrons.trackmgr, gpuState.hepEMBuffers_d.electronsHepEm); - MSC1<<>>( + ElectronMSC<<>>( electrons.trackmgr, gpuState.hepEMBuffers_d.electronsHepEm); ElectronRelocation<<>>( 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<<>>( + ElectronHowFar<<>>( positrons.trackmgr, gpuState.hepEMBuffers_d.positronsHepEm, VolAuxArray::GetInstance().fAuxData_dev ); - ElectronTransport1<<>>( + ElectronPropagation<<>>( positrons.trackmgr, gpuState.hepEMBuffers_d.positronsHepEm); - MSC1<<>>( + ElectronMSC<<>>( positrons.trackmgr, gpuState.hepEMBuffers_d.positronsHepEm); ElectronRelocation<<>>( positrons.trackmgr); @@ -464,16 +464,16 @@ void ShowerGPU(IntegrationLayer &integration, int event, adeptint::TrackBuffer & TransportGammas<<>>( gammas.trackmgr, secondaries, gammas.leakedTracks, scoring_dev, VolAuxArray::GetInstance().fAuxData_dev); #else - GammaPhysics1<<>>( + GammaHowFar<<>>( gammas.trackmgr, gpuState.hepEMBuffers_d.gammasHepEm, VolAuxArray::GetInstance().fAuxData_dev ); - GammaTransport1<<>>( + GammaPropagation<<>>( gammas.trackmgr, gpuState.hepEMBuffers_d.gammasHepEm, VolAuxArray::GetInstance().fAuxData_dev ); GammaRelocation<<>>( gammas.trackmgr, gammas.leakedTracks, VolAuxArray::GetInstance().fAuxData_dev ); - GammaPhysics2<<>>( + GammaInteractions<<>>( 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 index 2a210bac..816ad237 100644 --- a/include/AdePT/kernels/electrons_split.cuh +++ b/include/AdePT/kernels/electrons_split.cuh @@ -35,7 +35,7 @@ __device__ double GetVelocity(double eKin) } template -__global__ void ElectronPhysics1(adept::TrackManager *electrons, G4HepEmElectronTrack *hepEMTracks, VolAuxData const *auxDataArray) +__global__ void ElectronHowFar(adept::TrackManager *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 *electrons, G4HepEmE } template -static __global__ void ElectronTransport1(adept::TrackManager *electrons, G4HepEmElectronTrack *hepEMTracks) +static __global__ void ElectronPropagation(adept::TrackManager *electrons, G4HepEmElectronTrack *hepEMTracks) { #ifdef VECGEOM_FLOAT_PRECISION const Precision kPush = 10 * vecgeom::kTolerance; @@ -184,7 +184,7 @@ static __global__ void ElectronTransport1(adept::TrackManager *electrons, // May work well to separate MSC1 (Continuous effects) and MSC2 (Checks + safety) template -static __global__ void MSC1(adept::TrackManager *electrons, G4HepEmElectronTrack *hepEMTracks) +static __global__ void ElectronMSC(adept::TrackManager *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 index 9b85e095..f70d0bf2 100644 --- 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 *gammas, G4HepEmGammaTrack *hepEMTracks, VolAuxData const *auxDataArray) +__global__ void GammaHowFar(adept::TrackManager *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 *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 *gammas, G4HepEmGammaTrack *hepEMTracks, VolAuxData const *auxDataArray) +__global__ void GammaPropagation(adept::TrackManager *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 *gammas, MParrayTrack } template -__global__ void GammaPhysics2(adept::TrackManager *gammas, G4HepEmGammaTrack *hepEMTracks, Secondaries secondaries, Scoring *userScoring, +__global__ void GammaInteractions(adept::TrackManager *gammas, G4HepEmGammaTrack *hepEMTracks, Secondaries secondaries, Scoring *userScoring, VolAuxData const *auxDataArray) { int activeSize = gammas->fActiveTracks->size();