Skip to content

Commit 871a671

Browse files
SeverinDiederichsSeverinDiederichs
authored
Enable PreUserTrackingAction and fixes related to leaked energy (#369)
1 parent 8ee04c2 commit 871a671

13 files changed

+66
-56
lines changed

‎include/AdePT/core/AdePTConfiguration.hh

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -30,10 +30,7 @@ public:
3030
void SetNumThreads(int numThreads) { fNumThreads = numThreads; }
3131
void SetTrackInAllRegions(bool trackInAllRegions) { fTrackInAllRegions = trackInAllRegions; }
3232
void SetCallUserSteppingAction(bool callUserSteppingAction) { fCallUserSteppingAction = callUserSteppingAction; }
33-
void SetCallPostUserTrackingAction(bool callPostUserTrackingAction)
34-
{
35-
fCallPostUserTrackingAction = callPostUserTrackingAction;
36-
}
33+
void SetCallUserTrackingAction(bool callUserTrackingAction) { fCallUserTrackingAction = callUserTrackingAction; }
3734
void AddGPURegionName(std::string name) { fGPURegionNames.push_back(name); }
3835
void SetAdePTActivation(bool activateAdePT) { fAdePTActivated = activateAdePT; }
3936
void SetVerbosity(int verbosity) { fVerbosity = verbosity; };
@@ -56,7 +53,7 @@ public:
5653

5754
bool GetTrackInAllRegions() { return fTrackInAllRegions; }
5855
bool GetCallUserSteppingAction() { return fCallUserSteppingAction; }
59-
bool GetCallPostUserTrackingAction() { return fCallPostUserTrackingAction; }
56+
bool GetCallUserTrackingAction() { return fCallUserTrackingAction; }
6057
bool GetSpeedOfLight() { return fSpeedOfLight; }
6158
bool IsAdePTActivated() { return fAdePTActivated; }
6259
int GetNumThreads() { return fNumThreads; };
@@ -78,7 +75,7 @@ public:
7875
private:
7976
bool fTrackInAllRegions{false};
8077
bool fCallUserSteppingAction{false};
81-
bool fCallPostUserTrackingAction{false};
78+
bool fCallUserTrackingAction{false};
8279
bool fSpeedOfLight{false};
8380
bool fAdePTActivated{true};
8481
int fNumThreads;

‎include/AdePT/core/AdePTScoringTemplate.cuh

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,7 +23,7 @@ __device__ void RecordHit(Scoring *scoring_dev, int aParentID, char aParticleTyp
2323
vecgeom::Vector3D<Precision> const &aPreMomentumDirection, double aPreEKin, double aPreCharge,
2424
vecgeom::NavigationState const &aPostState, vecgeom::Vector3D<Precision> const &aPostPosition,
2525
vecgeom::Vector3D<Precision> const &aPostMomentumDirection, double aPostEKin,
26-
double aPostCharge, unsigned int eventId, short threadId, bool isLastStep);
26+
double aPostCharge, unsigned int eventId, short threadId, bool isLastStep, bool isFirstStep);
2727

2828
template <typename Scoring>
2929
__device__ void AccountProduced(Scoring *scoring_dev, int num_ele, int num_pos, int num_gam);

‎include/AdePT/core/AsyncAdePTTransport.hh

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -63,8 +63,8 @@ private:
6363
bool fTrackInAllRegions = false;
6464
std::vector<std::string> const *fGPURegionNames;
6565
// Flags for the kernels to return the last or all steps, needed for PostUserTrackingAction or UserSteppingAction
66-
bool fReturnAllSteps = false;
67-
bool fReturnLastStep = false;
66+
bool fReturnAllSteps = false;
67+
bool fReturnFirstAndLastStep = false;
6868
std::string fBfieldFile{""}; ///< Path to magnetic field file (in the covfie format)
6969
GeneralMagneticField fMagneticField; ///< arbitrary magnetic field
7070
double fCPUCapacityFactor{

‎include/AdePT/core/AsyncAdePTTransport.icc

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -93,8 +93,8 @@ AsyncAdePTTransport<IntegrationLayer>::AsyncAdePTTransport(AdePTConfiguration &c
9393
fGPUNetEnergy(fNThread, 0.0), fTrackInAllRegions{configuration.GetTrackInAllRegions()},
9494
fGPURegionNames{configuration.GetGPURegionNames()}, fCUDAStackLimit{configuration.GetCUDAStackLimit()},
9595
fCUDAHeapLimit{configuration.GetCUDAHeapLimit()}, fReturnAllSteps{configuration.GetCallUserSteppingAction()},
96-
fReturnLastStep{configuration.GetCallPostUserTrackingAction()}, fBfieldFile{configuration.GetCovfieBfieldFile()},
97-
fLastNParticlesOnCPU{configuration.GetLastNParticlesOnCPU()},
96+
fReturnFirstAndLastStep{configuration.GetCallUserTrackingAction()},
97+
fBfieldFile{configuration.GetCovfieBfieldFile()}, fLastNParticlesOnCPU{configuration.GetLastNParticlesOnCPU()},
9898
fCPUCopyFraction{configuration.GetHitBufferFlushThreshold()},
9999
fCPUCapacityFactor{configuration.GetCPUCapacityFactor()}
100100
{
@@ -310,7 +310,7 @@ void AsyncAdePTTransport<IntegrationLayer>::Initialize()
310310
fCPUCapacityFactor, fCPUCopyFraction);
311311
fGPUWorker = async_adept_impl::LaunchGPUWorker(fTrackCapacity, fScoringCapacity, fNThread, *fBuffer, *fGPUstate,
312312
fEventStates, fCV_G4Workers, fScoring, fAdePTSeed, fDebugLevel,
313-
fReturnAllSteps, fReturnLastStep, fLastNParticlesOnCPU);
313+
fReturnAllSteps, fReturnFirstAndLastStep, fLastNParticlesOnCPU);
314314
}
315315

316316
template <typename IntegrationLayer>
@@ -341,7 +341,7 @@ void AsyncAdePTTransport<IntegrationLayer>::ProcessGPUSteps(int threadId, int ev
341341
<< " state : " << static_cast<unsigned int>(fEventStates[threadId].load(std::memory_order_acquire))
342342
<< "\033[0m" << std::endl;
343343
}
344-
integrationInstance.ProcessGPUStep(*it, fReturnAllSteps, fReturnLastStep);
344+
integrationInstance.ProcessGPUStep(*it, fReturnAllSteps, fReturnFirstAndLastStep);
345345
}
346346
async_adept_impl::CloseGPUBuffer(threadId, *fGPUstate, range.first, dataOnBuffer);
347347
}

‎include/AdePT/core/HostScoringImpl.cuh

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -155,15 +155,15 @@ __device__ void RecordHit(HostScoring *hostScoring_dev, int aParentID, char aPar
155155
vecgeom::Vector3D<Precision> const &aPreMomentumDirection, double aPreEKin, double aPreCharge,
156156
vecgeom::NavigationState const &aPostState, vecgeom::Vector3D<Precision> const &aPostPosition,
157157
vecgeom::Vector3D<Precision> const &aPostMomentumDirection, double aPostEKin,
158-
double aPostCharge, unsigned int, short, bool)
158+
double aPostCharge, unsigned int, short, bool, bool)
159159
{
160160
// Acquire a hit slot
161161
GPUHit &aGPUHit = *GetNextFreeHit(hostScoring_dev);
162162

163163
// Fill the required data
164164
FillHit(aGPUHit, aParentID, aParticleType, aStepLength, aTotalEnergyDeposit, aPreState, aPrePosition,
165165
aPreMomentumDirection, aPreEKin, aPreCharge, aPostState, aPostPosition, aPostMomentumDirection, aPostEKin,
166-
aPostCharge, 0, 0, false);
166+
aPostCharge, 0, 0, false, false);
167167
}
168168

169169
/// @brief Account for the number of produced secondaries

‎include/AdePT/core/PerEventScoringImpl.cuh

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -691,15 +691,15 @@ __device__ void RecordHit(AsyncAdePT::PerEventScoring * /*scoring*/, int aParent
691691
vecgeom::Vector3D<Precision> const &aPreMomentumDirection, double aPreEKin, double aPreCharge,
692692
vecgeom::NavigationState const &aPostState, vecgeom::Vector3D<Precision> const &aPostPosition,
693693
vecgeom::Vector3D<Precision> const &aPostMomentumDirection, double aPostEKin,
694-
double aPostCharge, unsigned int eventID, short threadID, bool isLastStep)
694+
double aPostCharge, unsigned int eventID, short threadID, bool isLastStep, bool isFirstStep)
695695
{
696696
// Acquire a hit slot
697697
GPUHit &aGPUHit = AsyncAdePT::gHitScoringBuffer_dev.GetNextSlot(threadID);
698698

699699
// Fill the required data
700700
FillHit(aGPUHit, aParentID, aParticleType, aStepLength, aTotalEnergyDeposit, aPreState, aPrePosition,
701701
aPreMomentumDirection, aPreEKin, aPreCharge, aPostState, aPostPosition, aPostMomentumDirection, aPostEKin,
702-
aPostCharge, eventID, threadID, isLastStep);
702+
aPostCharge, eventID, threadID, isLastStep, isLastStep);
703703
}
704704

705705
/// @brief Account for the number of produced secondaries

‎include/AdePT/core/ScoringCommons.hh

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -32,6 +32,7 @@ struct GPUHit {
3232
unsigned int fEventId{0};
3333
short threadId{-1};
3434
// bool fFirstStepInVolume{false};
35+
bool fFirstStepOfTrack{false};
3536
bool fLastStepOfTrack{false};
3637
char fParticleType{0}; // Particle type ID
3738
};
@@ -74,12 +75,14 @@ __device__ __forceinline__ void FillHit(GPUHit &aGPUHit, int aParentID, char aPa
7475
double aPreCharge, vecgeom::NavigationState const &aPostState,
7576
vecgeom::Vector3D<Precision> const &aPostPosition,
7677
vecgeom::Vector3D<Precision> const &aPostMomentumDirection, double aPostEKin,
77-
double aPostCharge, unsigned int eventID, short threadID, bool isLastStep)
78+
double aPostCharge, unsigned int eventID, short threadID, bool isLastStep,
79+
bool isFirstStep)
7880
{
7981
aGPUHit.fEventId = eventID;
8082
aGPUHit.threadId = threadID;
8183

82-
aGPUHit.fLastStepOfTrack = isLastStep;
84+
aGPUHit.fFirstStepOfTrack = isFirstStep;
85+
aGPUHit.fLastStepOfTrack = isLastStep;
8386
// Fill the required data
8487
aGPUHit.fParentID = aParentID;
8588
aGPUHit.fParticleType = aParticleType;

‎include/AdePT/integration/AdePTConfigurationMessenger.hh

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,7 @@ private:
3838
G4UIcmdWithAnInteger *fSetFinishOnCpuCmd;
3939
G4UIcmdWithABool *fSetTrackInAllRegionsCmd;
4040
G4UIcmdWithABool *fSetCallUserSteppingActionCmd;
41-
G4UIcmdWithABool *fSetCallPostUserTrackingActionCmd;
41+
G4UIcmdWithABool *fSetCallUserTrackingActionCmd;
4242
G4UIcmdWithABool *fSetSpeedOfLightCmd;
4343
G4UIcmdWithAString *fAddRegionCmd;
4444
G4UIcmdWithABool *fActivateAdePTCmd;

‎include/AdePT/integration/AdePTGeant4Integration.hh

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -55,7 +55,7 @@ public:
5555

5656
/// @brief Reconstructs GPU hits on host and calls the user-defined sensitive detector code
5757
void ProcessGPUStep(GPUHit const &hit, bool const callUserSteppingAction = false,
58-
bool const callPostUserTrackingaction = false);
58+
bool const callUserTrackingaction = false);
5959

6060
/// @brief Takes a range of tracks coming from the device and gives them back to Geant4
6161
template <typename Iterator>

‎include/AdePT/kernels/electrons.cuh

Lines changed: 17 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -654,6 +654,7 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
654654
}
655655

656656
if (stopped) {
657+
eKin = 0;
657658
if (!IsElectron) {
658659
// Annihilate the stopped positron into two gammas heading to opposite
659660
// directions (isotropic).
@@ -711,21 +712,22 @@ static __device__ __forceinline__ void TransportElectrons(adept::TrackManager<Tr
711712
// Record the step. Edep includes the continuous energy loss and edep from secondaries which were cut
712713
if ((energyDeposit > 0 && auxData.fSensIndex >= 0) || returnAllSteps || returnLastStep)
713714
adept_scoring::RecordHit(userScoring, currentTrack.parentId,
714-
static_cast<char>(IsElectron ? 0 : 1), // Particle type
715-
elTrack.GetPStepLength(), // Step length
716-
energyDeposit, // Total Edep
717-
navState, // Pre-step point navstate
718-
preStepPos, // Pre-step point position
719-
preStepDir, // Pre-step point momentum direction
720-
preStepEnergy, // Pre-step point kinetic energy
721-
IsElectron ? -1 : 1, // Pre-step point charge
722-
nextState, // Post-step point navstate
723-
pos, // Post-step point position
724-
dir, // Post-step point momentum direction
725-
eKin, // Post-step point kinetic energy
726-
IsElectron ? -1 : 1, // Post-step point charge
727-
currentTrack.eventId, currentTrack.threadId, // eventID and threadID
728-
returnLastStep); // whether this was the last step
715+
static_cast<char>(IsElectron ? 0 : 1), // Particle type
716+
elTrack.GetPStepLength(), // Step length
717+
energyDeposit, // Total Edep
718+
navState, // Pre-step point navstate
719+
preStepPos, // Pre-step point position
720+
preStepDir, // Pre-step point momentum direction
721+
preStepEnergy, // Pre-step point kinetic energy
722+
IsElectron ? -1 : 1, // Pre-step point charge
723+
nextState, // Post-step point navstate
724+
pos, // Post-step point position
725+
dir, // Post-step point momentum direction
726+
eKin, // Post-step point kinetic energy
727+
IsElectron ? -1 : 1, // Post-step point charge
728+
currentTrack.eventId, currentTrack.threadId, // eventID and threadID
729+
returnLastStep, // whether this was the last step
730+
currentTrack.stepCounter == 1 ? true : false); // whether this was the first step
729731
if (cross_boundary) {
730732
// Move to the next boundary now that the Step is recorded
731733
navState = nextState;

‎include/AdePT/kernels/gammas.cuh

Lines changed: 9 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -195,7 +195,6 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
195195

196196
#ifdef ADEPT_USE_SURF
197197
AdePTNavigator::RelocateToNextVolume(pos, dir, hitsurf_index, nextState);
198-
if (nextState.IsOutside()) continue;
199198
#else
200199
AdePTNavigator::RelocateToNextVolume(pos, dir, nextState);
201200
#endif
@@ -216,10 +215,12 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
216215
nextState, // Post-step point navstate
217216
pos, // Post-step point position
218217
dir, // Post-step point momentum direction
219-
0, // Post-step point kinetic energy
218+
eKin, // Post-step point kinetic energy
220219
0, // Post-step point charge
221220
currentTrack.eventId, currentTrack.threadId, // event and thread ID
222-
returnLastStep); // whether this is the last step of the track
221+
returnLastStep,
222+
currentTrack.stepCounter == 1 ? true
223+
: false); // whether this is the last step of the track
223224

224225
// Move to the next boundary.
225226
navState = nextState;
@@ -266,10 +267,11 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
266267
nextState, // Post-step point navstate
267268
pos, // Post-step point position
268269
dir, // Post-step point momentum direction
269-
0, // Post-step point kinetic energy
270+
eKin, // Post-step point kinetic energy
270271
0, // Post-step point charge
271272
currentTrack.eventId, currentTrack.threadId, // event and thread ID
272-
returnLastStep); // whether this is the last step of the track
273+
returnLastStep, // whether this is the last step of the track
274+
currentTrack.stepCounter == 1 ? true : false); // whether this is the first step
273275
}
274276
continue;
275277
} else {
@@ -485,7 +487,8 @@ __global__ void TransportGammas(adept::TrackManager<Track> *gammas, Secondaries
485487
newEnergyGamma, // Post-step point kinetic energy
486488
0, // Post-step point charge
487489
currentTrack.eventId, currentTrack.threadId, // event and thread ID
488-
returnLastStep); // whether this is the last step of the track
490+
returnLastStep, // whether this is the last step of the track
491+
currentTrack.stepCounter == 1 ? true : false); // whether this is the first step
489492
}
490493
}
491494
}

‎src/AdePTConfigurationMessenger.cc

Lines changed: 5 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -33,13 +33,8 @@ AdePTConfigurationMessenger::AdePTConfigurationMessenger(AdePTConfiguration *ade
3333
"secondary before the primary has finished its track."
3434
" NOTE: This means that every single step is recorded on GPU and send back to CPU, which can impact performance");
3535

36-
fSetCallPostUserTrackingActionCmd = new G4UIcmdWithABool("/adept/CallPostUserTrackingAction", this);
37-
fSetCallPostUserTrackingActionCmd->SetGuidance(
38-
"If true, the PostUserTrackingAction is called for on every track. NOTE: This "
39-
"means that the last step of every track is recorded on GPU and send back to CPU");
40-
41-
fSetCallPostUserTrackingActionCmd = new G4UIcmdWithABool("/adept/CallPostUserTrackingAction", this);
42-
fSetCallPostUserTrackingActionCmd->SetGuidance(
36+
fSetCallUserTrackingActionCmd = new G4UIcmdWithABool("/adept/CallUserTrackingAction", this);
37+
fSetCallUserTrackingActionCmd->SetGuidance(
4338
"If true, the PostUserTrackingAction is called for on every track. NOTE: This "
4439
"means that the last step of every track is recorded on GPU and send back to CPU");
4540

@@ -109,7 +104,7 @@ AdePTConfigurationMessenger::~AdePTConfigurationMessenger()
109104
delete fSetCUDAHeapLimitCmd;
110105
delete fSetTrackInAllRegionsCmd;
111106
delete fSetCallUserSteppingActionCmd;
112-
delete fSetCallPostUserTrackingActionCmd;
107+
delete fSetCallUserTrackingActionCmd;
113108
delete fAddRegionCmd;
114109
delete fActivateAdePTCmd;
115110
delete fSetVerbosityCmd;
@@ -133,8 +128,8 @@ void AdePTConfigurationMessenger::SetNewValue(G4UIcommand *command, G4String new
133128
fAdePTConfiguration->SetTrackInAllRegions(fSetTrackInAllRegionsCmd->GetNewBoolValue(newValue));
134129
} else if (command == fSetCallUserSteppingActionCmd) {
135130
fAdePTConfiguration->SetCallUserSteppingAction(newValue);
136-
} else if (command == fSetCallPostUserTrackingActionCmd) {
137-
fAdePTConfiguration->SetCallPostUserTrackingAction(newValue);
131+
} else if (command == fSetCallUserTrackingActionCmd) {
132+
fAdePTConfiguration->SetCallUserTrackingAction(newValue);
138133
} else if (command == fSetSpeedOfLightCmd) {
139134
fAdePTConfiguration->SetSpeedOfLightCmd(newValue);
140135
} else if (command == fAddRegionCmd) {

0 commit comments

Comments
 (0)