Skip to content

Commit f1d2bb7

Browse files
SeverinDiederichsSeverinDiederichs
authored
enable calling of UserSteppingAction and PostUserTrackingAction (apt-sim#356)
If the UserSteppingAction is required, we need to copy back every GPU step back to the G4 workers. This required to change the kernels, as we need to be able to record every step independent of edep or sensitive detector. Since copying back every step can lead to a very large amount of steps, this would quickly fill the buffer. Then, if the buffer gets too full before the Geant4 workers take care of their hits, the GPUStep Management Thread would start copying out the GPUSteps, as implemented in apt-sim#350. However, this copying is too slow, if every Step is recorded and leads to the GPU running out of HitSlots. Previously, the Geant4 workers would only take care of the GPU Steps after their transport has finished. However, this may be too late and the buffer may be too full, leading to copying. Therefore, the Geant4 workers must be able to process some of the steps already earlier. This is now done in the `AdePTTrackingManager`: Before a new track is processed, the `GPUStepProcessing` is called. This way, the GPU step buffer can be kept under control. To enable this, the processing of the GPUSteps is now encapsulated in a single function, that can be called from the `AdePTTrackingManager`. In the same manner, the PostUserTrackingAction is called. For this, the RecordHit also writes if it is the LastStep of a track. Note that it is straightforward to calso call the PreUserTrackingAction. This requires a StepCounter, which is availale in the B field update branch, so I will add it *after* the B field branch is merged. Both can be enabled via: ``` /adept/CallUserSteppingAction true /adept/CallPostUserTrackingAction true ``` Since this PR touches the kernels, below the physics validation at high statistics, which is as good as it should be: <img width="586" alt="Screenshot 2025-03-09 at 07 30 15" src="https://github.com/user-attachments/assets/8027a386-2680-4b22-9c4f-8da91a693ea3" />
1 parent d5e712f commit f1d2bb7

18 files changed

+227
-142
lines changed

‎include/AdePT/core/AdePTConfiguration.hh

+9
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,11 @@ public:
2929
~AdePTConfiguration() {}
3030
void SetNumThreads(int numThreads) { fNumThreads = numThreads; }
3131
void SetTrackInAllRegions(bool trackInAllRegions) { fTrackInAllRegions = trackInAllRegions; }
32+
void SetCallUserSteppingAction(bool callUserSteppingAction) { fCallUserSteppingAction = callUserSteppingAction; }
33+
void SetCallPostUserTrackingAction(bool callPostUserTrackingAction)
34+
{
35+
fCallPostUserTrackingAction = callPostUserTrackingAction;
36+
}
3237
void AddGPURegionName(std::string name) { fGPURegionNames.push_back(name); }
3338
void SetAdePTActivation(bool activateAdePT) { fAdePTActivated = activateAdePT; }
3439
void SetVerbosity(int verbosity) { fVerbosity = verbosity; };
@@ -43,6 +48,8 @@ public:
4348
void SetVecGeomGDML(std::string filename) { fVecGeomGDML = filename; }
4449

4550
bool GetTrackInAllRegions() { return fTrackInAllRegions; }
51+
bool GetCallUserSteppingAction() { return fCallUserSteppingAction; }
52+
bool GetCallPostUserTrackingAction() { return fCallPostUserTrackingAction; }
4653
bool IsAdePTActivated() { return fAdePTActivated; }
4754
int GetNumThreads() { return fNumThreads; };
4855
int GetVerbosity() { return fVerbosity; };
@@ -59,6 +66,8 @@ public:
5966

6067
private:
6168
bool fTrackInAllRegions{false};
69+
bool fCallUserSteppingAction{false};
70+
bool fCallPostUserTrackingAction{false};
6271
bool fAdePTActivated{true};
6372
int fNumThreads;
6473
int fVerbosity{0};

‎include/AdePT/core/AdePTScoringTemplate.cuh

+1-1
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);
26+
double aPostCharge, unsigned int eventId, short threadId, bool isLastStep);
2727

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

‎include/AdePT/core/AdePTTransport.h

+1-1
Original file line numberDiff line numberDiff line change
@@ -78,6 +78,7 @@ class AdePTTransport : public AdePTTransportInterface {
7878
void Cleanup();
7979
/// @brief Interface for transporting a buffer of tracks in AdePT.
8080
void Shower(int event, int threadId);
81+
void ProcessGPUSteps(int, int) {};
8182

8283
private:
8384
static inline G4HepEmState *fg4hepem_state{nullptr}; ///< The HepEm state singleton
@@ -106,7 +107,6 @@ class AdePTTransport : public AdePTTransportInterface {
106107
bool InitializeField(double bz);
107108
bool InitializeGeometry(const vecgeom::cxx::VPlacedVolume *world);
108109
bool InitializePhysics();
109-
void ProcessGPUHits();
110110
};
111111

112112
#include "AdePTTransport.icc"

‎include/AdePT/core/AdePTTransportInterface.hh

+3-2
Original file line numberDiff line numberDiff line change
@@ -47,8 +47,9 @@ public:
4747
/// @brief Initialize the ApplyCuts flag on device
4848
virtual bool InitializeApplyCuts(bool applycuts) = 0;
4949
/// @brief Interface for transporting a buffer of tracks in AdePT.
50-
virtual void Shower(int event, int threadId) = 0;
51-
virtual void Cleanup() = 0;
50+
virtual void Shower(int event, int threadId) = 0;
51+
virtual void Cleanup() = 0;
52+
virtual void ProcessGPUSteps(int threadId, int eventId) = 0;
5253
};
5354

5455
#endif

‎include/AdePT/core/AsyncAdePTTransport.cuh

+8-6
Original file line numberDiff line numberDiff line change
@@ -616,7 +616,8 @@ void HitProcessingLoop(HitProcessingContext *const context, GPUstate &gpuState,
616616

617617
void TransportLoop(int trackCapacity, int scoringCapacity, int numThreads, TrackBuffer &trackBuffer, GPUstate &gpuState,
618618
std::vector<std::atomic<EventState>> &eventStates, std::condition_variable &cvG4Workers,
619-
std::vector<AdePTScoring> &scoring, int adeptSeed, int debugLevel)
619+
std::vector<AdePTScoring> &scoring, int adeptSeed, int debugLevel, bool returnAllSteps,
620+
bool returnLastStep)
620621
{
621622
// NVTXTracer tracer{"TransportLoop"};
622623

@@ -791,7 +792,7 @@ void TransportLoop(int trackCapacity, int scoringCapacity, int numThreads, Track
791792
const auto [threads, blocks] = computeThreadsAndBlocks(particlesInFlight[ParticleType::Electron]);
792793
TransportElectrons<PerEventScoring><<<blocks, threads, 0, electrons.stream>>>(
793794
electrons.tracks, electrons.queues.currentlyActive, secondaries, electrons.queues.nextActive,
794-
electrons.queues.leakedTracksCurrent, gpuState.fScoring_dev);
795+
electrons.queues.leakedTracksCurrent, gpuState.fScoring_dev, returnAllSteps, returnLastStep);
795796

796797
COPCORE_CUDA_CHECK(cudaEventRecord(electrons.event, electrons.stream));
797798
COPCORE_CUDA_CHECK(cudaStreamWaitEvent(gpuState.stream, electrons.event, 0));
@@ -802,7 +803,7 @@ void TransportLoop(int trackCapacity, int scoringCapacity, int numThreads, Track
802803
const auto [threads, blocks] = computeThreadsAndBlocks(particlesInFlight[ParticleType::Positron]);
803804
TransportPositrons<PerEventScoring><<<blocks, threads, 0, positrons.stream>>>(
804805
positrons.tracks, positrons.queues.currentlyActive, secondaries, positrons.queues.nextActive,
805-
positrons.queues.leakedTracksCurrent, gpuState.fScoring_dev);
806+
positrons.queues.leakedTracksCurrent, gpuState.fScoring_dev, returnAllSteps, returnLastStep);
806807

807808
COPCORE_CUDA_CHECK(cudaEventRecord(positrons.event, positrons.stream));
808809
COPCORE_CUDA_CHECK(cudaStreamWaitEvent(gpuState.stream, positrons.event, 0));
@@ -813,7 +814,8 @@ void TransportLoop(int trackCapacity, int scoringCapacity, int numThreads, Track
813814
const auto [threads, blocks] = computeThreadsAndBlocks(particlesInFlight[ParticleType::Gamma]);
814815
TransportGammas<PerEventScoring><<<blocks, threads, 0, gammas.stream>>>(
815816
gammas.tracks, gammas.queues.currentlyActive, secondaries, gammas.queues.nextActive,
816-
gammas.queues.leakedTracksCurrent, gpuState.fScoring_dev); //, gpuState.gammaInteractions);
817+
gammas.queues.leakedTracksCurrent, gpuState.fScoring_dev, returnAllSteps,
818+
returnLastStep); //, gpuState.gammaInteractions);
817819

818820
// constexpr unsigned int intThreads = 128;
819821
// ApplyGammaInteractions<PerEventScoring><<<dim3(20, 3, 1), intThreads, 0, gammas.stream>>>(
@@ -1099,12 +1101,12 @@ void CloseGPUBuffer(unsigned int threadId, GPUstate &gpuState, GPUHit *begin, co
10991101
std::thread LaunchGPUWorker(int trackCapacity, int scoringCapacity, int numThreads, TrackBuffer &trackBuffer,
11001102
GPUstate &gpuState, std::vector<std::atomic<EventState>> &eventStates,
11011103
std::condition_variable &cvG4Workers, std::vector<AdePTScoring> &scoring, int adeptSeed,
1102-
int debugLevel)
1104+
int debugLevel, bool returnAllSteps, bool returnLastStep)
11031105
{
11041106
return std::thread{
11051107
&TransportLoop, trackCapacity, scoringCapacity, numThreads, std::ref(trackBuffer),
11061108
std::ref(gpuState), std::ref(eventStates), std::ref(cvG4Workers), std::ref(scoring), adeptSeed,
1107-
debugLevel};
1109+
debugLevel, returnAllSteps, returnLastStep};
11081110
}
11091111

11101112
void FreeGPU(std::unique_ptr<AsyncAdePT::GPUstate, AsyncAdePT::GPUstateDeleter> &gpuState, G4HepEmState &g4hepem_state,

‎include/AdePT/core/AsyncAdePTTransport.hh

+4
Original file line numberDiff line numberDiff line change
@@ -59,6 +59,9 @@ private:
5959
std::vector<double> fGPUNetEnergy;
6060
bool fTrackInAllRegions = false;
6161
std::vector<std::string> const *fGPURegionNames;
62+
// Flags for the kernels to return the last or all steps, needed for PostUserTrackingAction or UserSteppingAction
63+
bool fReturnAllSteps = false;
64+
bool fReturnLastStep = false;
6265

6366
void Initialize();
6467
void InitBVH();
@@ -105,6 +108,7 @@ public:
105108
void Shower(int event, int threadId) override { Flush(threadId, event); }
106109
/// Block until transport of the given event is done.
107110
void Flush(int threadId, int eventId);
111+
void ProcessGPUSteps(int threadId, int eventId) override;
108112
void Cleanup() override {}
109113
};
110114

‎include/AdePT/core/AsyncAdePTTransport.icc

+33-23
Original file line numberDiff line numberDiff line change
@@ -48,7 +48,7 @@ std::pair<GPUHit *, GPUHit *> GetGPUHitsFromBuffer(unsigned int, unsigned int, A
4848
void CloseGPUBuffer(unsigned int, AsyncAdePT::GPUstate &, GPUHit *, const bool);
4949
std::thread LaunchGPUWorker(int, int, int, AsyncAdePT::TrackBuffer &, AsyncAdePT::GPUstate &,
5050
std::vector<std::atomic<AsyncAdePT::EventState>> &, std::condition_variable &,
51-
std::vector<AdePTScoring> &, int, int);
51+
std::vector<AdePTScoring> &, int, int, bool, bool);
5252
std::unique_ptr<AsyncAdePT::GPUstate, AsyncAdePT::GPUstateDeleter> InitializeGPU(int trackCapacity, int scoringCapacity,
5353
int numThreads,
5454
AsyncAdePT::TrackBuffer &trackBuffer,
@@ -85,7 +85,8 @@ AsyncAdePTTransport<IntegrationLayer>::AsyncAdePTTransport(AdePTConfiguration &c
8585
fDebugLevel{configuration.GetVerbosity()}, fIntegrationLayerObjects(fNThread), fEventStates(fNThread),
8686
fGPUNetEnergy(fNThread, 0.0), fTrackInAllRegions{configuration.GetTrackInAllRegions()},
8787
fGPURegionNames{configuration.GetGPURegionNames()}, fCUDAStackLimit{configuration.GetCUDAStackLimit()},
88-
fCUDAHeapLimit{configuration.GetCUDAHeapLimit()}
88+
fCUDAHeapLimit{configuration.GetCUDAHeapLimit()}, fReturnAllSteps{configuration.GetCallUserSteppingAction()},
89+
fReturnLastStep{configuration.GetCallPostUserTrackingAction()}
8990
{
9091
if (fNThread > kMaxThreads)
9192
throw std::invalid_argument("AsyncAdePTTransport limited to " + std::to_string(kMaxThreads) + " threads");
@@ -244,7 +245,8 @@ void AsyncAdePTTransport<IntegrationLayer>::Initialize()
244245

245246
fGPUstate = async_adept_impl::InitializeGPU(fTrackCapacity, fScoringCapacity, fNThread, *fBuffer, fScoring);
246247
fGPUWorker = async_adept_impl::LaunchGPUWorker(fTrackCapacity, fScoringCapacity, fNThread, *fBuffer, *fGPUstate,
247-
fEventStates, fCV_G4Workers, fScoring, fAdePTSeed, fDebugLevel);
248+
fEventStates, fCV_G4Workers, fScoring, fAdePTSeed, fDebugLevel,
249+
fReturnAllSteps, fReturnLastStep);
248250
}
249251

250252
template <typename IntegrationLayer>
@@ -254,6 +256,33 @@ void AsyncAdePTTransport<IntegrationLayer>::InitBVH()
254256
vecgeom::cxx::BVHManager::DeviceInit();
255257
}
256258

259+
template <typename IntegrationLayer>
260+
void AsyncAdePTTransport<IntegrationLayer>::ProcessGPUSteps(int threadId, int eventId)
261+
{
262+
263+
AdePTGeant4Integration &integrationInstance = fIntegrationLayerObjects[threadId];
264+
std::pair<GPUHit *, GPUHit *> range;
265+
bool dataOnBuffer;
266+
267+
while ((range = async_adept_impl::GetGPUHitsFromBuffer(threadId, eventId, *fGPUstate, dataOnBuffer)).first !=
268+
nullptr) {
269+
for (auto it = range.first; it != range.second; ++it) {
270+
// important sanity check: thread should only process its own hits and only from the current event
271+
if (it->threadId != threadId)
272+
std::cerr << "\033[1;31mError, threadId doesn't match it->threadId " << it->threadId << " threadId " << threadId
273+
<< "\033[0m" << std::endl;
274+
if (it->fEventId != eventId) {
275+
std::cerr << "\033[1;31mError, eventId doesn't match it->fEventId " << it->fEventId << "eventId " << eventId
276+
<< " num hits to be processed " << (range.second - range.first) << " dataOnBuffer " << dataOnBuffer
277+
<< "state : " << static_cast<unsigned int>(fEventStates[threadId].load(std::memory_order_acquire))
278+
<< "\033[0m" << std::endl;
279+
}
280+
integrationInstance.ProcessGPUStep(*it, fReturnAllSteps, fReturnLastStep);
281+
}
282+
async_adept_impl::CloseGPUBuffer(threadId, *fGPUstate, range.first, dataOnBuffer);
283+
}
284+
}
285+
257286
template <typename IntegrationLayer>
258287
void AsyncAdePTTransport<IntegrationLayer>::Flush(G4int threadId, G4int eventId)
259288
{
@@ -268,31 +297,12 @@ void AsyncAdePTTransport<IntegrationLayer>::Flush(G4int threadId, G4int eventId)
268297

269298
while (fEventStates[threadId].load(std::memory_order_acquire) < EventState::DeviceFlushed) {
270299

271-
std::pair<GPUHit *, GPUHit *> range;
272-
bool dataOnBuffer;
273-
274300
{
275301
std::unique_lock lock{fMutex_G4Workers};
276302
fCV_G4Workers.wait(lock);
277303
}
278304

279-
while ((range = async_adept_impl::GetGPUHitsFromBuffer(threadId, eventId, *fGPUstate, dataOnBuffer)).first !=
280-
nullptr) {
281-
for (auto it = range.first; it != range.second; ++it) {
282-
// important sanity check: thread should only process its own hits and only from the current event
283-
if (it->threadId != threadId)
284-
std::cerr << "Error, threadId doesn't match it->threadId " << it->threadId << " threadId " << threadId
285-
<< std::endl;
286-
if (it->fEventId != eventId) {
287-
std::cerr << "Error, eventId doesn't match it->fEventId " << it->fEventId << "eventId " << eventId
288-
<< " num hits to be processed " << (range.second - range.first) << " dataOnBuffer " << dataOnBuffer
289-
<< "state : " << static_cast<unsigned int>(fEventStates[threadId].load(std::memory_order_acquire))
290-
<< std::endl;
291-
}
292-
integrationInstance.ProcessGPUHit(*it);
293-
}
294-
async_adept_impl::CloseGPUBuffer(threadId, *fGPUstate, range.first, dataOnBuffer);
295-
}
305+
ProcessGPUSteps(threadId, eventId);
296306
}
297307

298308
// Now device should be flushed, so retrieve the tracks:

‎include/AdePT/core/HostScoringImpl.cuh

+4-4
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)
158+
double aPostCharge, unsigned int, short, 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);
166+
aPostCharge, 0, 0, false);
167167
}
168168

169169
/// @brief Account for the number of produced secondaries
@@ -202,7 +202,7 @@ inline void EndOfIteration(HostScoring &hostScoring, HostScoring *hostScoring_de
202202
COPCORE_CUDA_CHECK(cudaStreamSynchronize(stream));
203203
// Process the hits on CPU
204204
for (const auto &hit : hostScoring) {
205-
integration.ProcessGPUHit(hit);
205+
integration.ProcessGPUStep(hit);
206206
}
207207
}
208208
}
@@ -220,7 +220,7 @@ inline void EndOfTransport(HostScoring &hostScoring, HostScoring *hostScoring_de
220220
COPCORE_CUDA_CHECK(cudaStreamSynchronize(stream));
221221
// Process the last hits on CPU
222222
for (const auto &hit : hostScoring) {
223-
integration.ProcessGPUHit(hit);
223+
integration.ProcessGPUStep(hit);
224224
}
225225
}
226226
} // namespace adept_scoring

‎include/AdePT/core/PerEventScoringImpl.cuh

+2-2
Original file line numberDiff line numberDiff line change
@@ -674,15 +674,15 @@ __device__ void RecordHit(AsyncAdePT::PerEventScoring * /*scoring*/, int aParent
674674
vecgeom::Vector3D<Precision> const &aPreMomentumDirection, double aPreEKin, double aPreCharge,
675675
vecgeom::NavigationState const &aPostState, vecgeom::Vector3D<Precision> const &aPostPosition,
676676
vecgeom::Vector3D<Precision> const &aPostMomentumDirection, double aPostEKin,
677-
double aPostCharge, unsigned int eventID, short threadID)
677+
double aPostCharge, unsigned int eventID, short threadID, bool isLastStep)
678678
{
679679
// Acquire a hit slot
680680
GPUHit &aGPUHit = AsyncAdePT::gHitScoringBuffer_dev.GetNextSlot(threadID);
681681

682682
// Fill the required data
683683
FillHit(aGPUHit, aParentID, aParticleType, aStepLength, aTotalEnergyDeposit, aPreState, aPrePosition,
684684
aPreMomentumDirection, aPreEKin, aPreCharge, aPostState, aPostPosition, aPostMomentumDirection, aPostEKin,
685-
aPostCharge, eventID, threadID);
685+
aPostCharge, eventID, threadID, isLastStep);
686686
}
687687

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

‎include/AdePT/core/ScoringCommons.hh

+8-7
Original file line numberDiff line numberDiff line change
@@ -21,18 +21,18 @@ struct GPUStepPoint {
2121
// Stores the necessary data to reconstruct GPU hits on the host , and
2222
// call the user-defined Geant4 sensitive detector code
2323
struct GPUHit {
24-
int fParentID{0}; // Track ID
24+
// Data needed to reconstruct pre-post step points
25+
GPUStepPoint fPreStepPoint;
26+
GPUStepPoint fPostStepPoint;
2527
// Data needed to reconstruct G4 Step
2628
double fStepLength{0};
2729
double fTotalEnergyDeposit{0};
2830
double fNonIonizingEnergyDeposit{0};
29-
// bool fFirstStepInVolume{false};
30-
// bool fLastStepInVolume{false};
31-
// Data needed to reconstruct pre-post step points
32-
GPUStepPoint fPreStepPoint;
33-
GPUStepPoint fPostStepPoint;
31+
int fParentID{0}; // Track ID
3432
unsigned int fEventId{0};
3533
short threadId{-1};
34+
// bool fFirstStepInVolume{false};
35+
bool fLastStepOfTrack{false};
3636
char fParticleType{0}; // Particle type ID
3737
};
3838

@@ -74,11 +74,12 @@ __device__ __forceinline__ void FillHit(GPUHit &aGPUHit, int aParentID, char aPa
7474
double aPreCharge, vecgeom::NavigationState const &aPostState,
7575
vecgeom::Vector3D<Precision> const &aPostPosition,
7676
vecgeom::Vector3D<Precision> const &aPostMomentumDirection, double aPostEKin,
77-
double aPostCharge, unsigned int eventID, short threadID)
77+
double aPostCharge, unsigned int eventID, short threadID, bool isLastStep)
7878
{
7979
aGPUHit.fEventId = eventID;
8080
aGPUHit.threadId = threadID;
8181

82+
aGPUHit.fLastStepOfTrack = isLastStep;
8283
// Fill the required data
8384
aGPUHit.fParentID = aParentID;
8485
aGPUHit.fParticleType = aParticleType;

‎include/AdePT/integration/AdePTConfigurationMessenger.hh

+2
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,8 @@ private:
3636
G4UIcmdWithAnInteger *fSetCUDAStackLimitCmd;
3737
G4UIcmdWithAnInteger *fSetCUDAHeapLimitCmd;
3838
G4UIcmdWithABool *fSetTrackInAllRegionsCmd;
39+
G4UIcmdWithABool *fSetCallUserSteppingActionCmd;
40+
G4UIcmdWithABool *fSetCallPostUserTrackingActionCmd;
3941
G4UIcmdWithAString *fAddRegionCmd;
4042
G4UIcmdWithABool *fActivateAdePTCmd;
4143
G4UIcmdWithAnInteger *fSetVerbosityCmd;

‎include/AdePT/integration/AdePTGeant4Integration.hh

+2-1
Original file line numberDiff line numberDiff line change
@@ -54,7 +54,8 @@ public:
5454
std::vector<G4LogicalVolume const *> &vecgeomLvToG4Map);
5555

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

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

‎include/AdePT/integration/AdePTTrackingManager.hh

+4
Original file line numberDiff line numberDiff line change
@@ -83,6 +83,10 @@ private:
8383
bool fAdePTInitialized{false};
8484
};
8585

86+
#ifdef ASYNC_MODE
87+
std::shared_ptr<AsyncAdePT::AsyncAdePTTransport<AdePTGeant4Integration>> GetAdePTInstance();
88+
#endif
89+
8690
//....oooOO0OOooo........oooOO0OOooo........oooOO0OOooo........oooOO0OOooo......
8791

8892
#endif

0 commit comments

Comments
 (0)