Skip to content

Commit 2890a96

Browse files
davidrohrdavidrohr
committed
GPU: Support external instance for GPUTRDTracker class in GPUChainTracking
1 parent 4fc2ede commit 2890a96

5 files changed

Lines changed: 39 additions & 20 deletions

File tree

‎GPU/GPUTracking/Base/GPUReconstructionKernels.h‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -66,7 +66,8 @@ GPUCA_KRNL_LB((GPUTPCGMO2Output, prepare ), (simple), (), ())
6666
GPUCA_KRNL(( GPUTPCGMO2Output, sort ), (simple), (), ())
6767
GPUCA_KRNL_LB((GPUTPCGMO2Output, output ), (simple), (), ())
6868
GPUCA_KRNL(( GPUTPCGMO2Output, mc ), (simple), (), ())
69-
GPUCA_KRNL_LB((GPUTRDTrackerKernels ), (simple), (), ())
69+
GPUCA_KRNL_LB((GPUTRDTrackerKernels, gpuVersion ), (simple), (, GPUPtr1(GPUTRDTracker*, externalInstance)), (, GPUPtr2(GPUTRDTracker*, externalInstance)))
70+
GPUCA_KRNL_LB((GPUTRDTrackerKernels, o2Version ), (simple), (, GPUPtr1(GPUTRDTracker*, externalInstance)), (, GPUPtr2(GPUTRDTracker*, externalInstance)))
7071
GPUCA_KRNL_LB((GPUITSFitterKernel ), (simple), (), ())
7172
GPUCA_KRNL_LB((GPUTPCConvertKernel ), (simple), (), ())
7273
GPUCA_KRNL_LB((GPUTPCCompressionKernels, step0attached ), (simple), (), ())

‎GPU/GPUTracking/Definitions/GPUDefGPUParameters.h‎

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -237,8 +237,11 @@
237237
#ifndef GPUCA_LB_GPUTPCGlobalTracking
238238
#define GPUCA_LB_GPUTPCGlobalTracking 256
239239
#endif
240-
#ifndef GPUCA_LB_GPUTRDTrackerKernels
241-
#define GPUCA_LB_GPUTRDTrackerKernels 512
240+
#ifndef GPUCA_LB_GPUTRDTrackerKernels_gpuVersion
241+
#define GPUCA_LB_GPUTRDTrackerKernels_gpuVersion 512
242+
#endif
243+
#ifndef GPUCA_LB_GPUTRDTrackerKernels_o2Version
244+
#define GPUCA_LB_GPUTRDTrackerKernels_o2Version 512
242245
#endif
243246
#ifndef GPUCA_LB_GPUTPCConvertKernel
244247
#define GPUCA_LB_GPUTPCConvertKernel 256

‎GPU/GPUTracking/Global/GPUChainTrackingTRD.cxx‎

Lines changed: 15 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -69,18 +69,25 @@ int GPUChainTracking::DoTRDGPUTracking(GPUTRDTracker* externalInstance)
6969
{
7070
#ifdef GPUCA_HAVE_O2HEADERS
7171
bool doGPU = GetRecoStepsGPU() & RecoStep::TRDTracking;
72-
GPUTRDTrackerGPU& Tracker = processors()->trdTrackerGPU;
73-
GPUTRDTrackerGPU& TrackerShadow = doGPU ? processorsShadow()->trdTrackerGPU : Tracker;
72+
auto* Tracker = &processors()->getTRDTracker<I>();
73+
auto* TrackerShadow = doGPU ? &processorsShadow()->getTRDTracker<I>() : Tracker;
74+
if (externalInstance) {
75+
if constexpr (std::is_same_v<decltype(Tracker), decltype(externalInstance)>) {
76+
Tracker = externalInstance;
77+
} else {
78+
throw std::runtime_error("Must not provide external instance that does not match template type");
79+
}
80+
}
7481

7582
const auto& threadContext = GetThreadContext();
76-
SetupGPUProcessor(&Tracker, false);
77-
TrackerShadow.OverrideGPUGeometry(reinterpret_cast<GPUTRDGeometry*>(mFlatObjectsDevice.mCalibObjects.trdGeometry));
83+
SetupGPUProcessor(Tracker, false);
84+
TrackerShadow->OverrideGPUGeometry(reinterpret_cast<GPUTRDGeometry*>(mFlatObjectsDevice.mCalibObjects.trdGeometry));
7885

79-
WriteToConstantMemory(RecoStep::TRDTracking, (char*)&processors()->trdTrackerGPU - (char*)processors(), &TrackerShadow, sizeof(TrackerShadow), 0);
80-
TransferMemoryResourcesToGPU(RecoStep::TRDTracking, &Tracker, 0);
86+
WriteToConstantMemory(RecoStep::TRDTracking, (char*)&processors()->trdTrackerGPU - (char*)processors(), TrackerShadow, sizeof(*TrackerShadow), 0);
87+
TransferMemoryResourcesToGPU(RecoStep::TRDTracking, Tracker, 0);
8188

82-
runKernel<GPUTRDTrackerKernels>(GetGridAuto(0), krnlRunRangeNone);
83-
TransferMemoryResourcesToHost(RecoStep::TRDTracking, &Tracker, 0);
89+
runKernel<GPUTRDTrackerKernels>(GetGridAuto(0), krnlRunRangeNone, krnlEventNone, nullptr);
90+
TransferMemoryResourcesToHost(RecoStep::TRDTracking, Tracker, 0);
8491
SynchronizeStream(0);
8592

8693
if (GetProcessingSettings().debugLevel >= 2) {

‎GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx‎

Lines changed: 16 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -15,23 +15,31 @@
1515
#include "GPUTRDTrackerKernels.h"
1616
#include "GPUTRDGeometry.h"
1717
#include "GPUConstantMem.h"
18+
#include "GPUCommonTypeTraits.h"
1819
#if defined(WITH_OPENMP) && !defined(GPUCA_GPUCODE)
1920
#include "GPUReconstruction.h"
2021
#endif
2122

2223
using namespace GPUCA_NAMESPACE::gpu;
2324

2425
template <int I>
25-
GPUdii() void GPUTRDTrackerKernels::Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors)
26+
GPUdii() void GPUTRDTrackerKernels::Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTracker* externalInstance)
2627
{
27-
auto& trdTracker = processors.getTRDTracker<I>();
28-
GPUCA_OPENMP(parallel for if(!trdTracker.GetRec().GetProcessingSettings().ompKernels) num_threads(trdTracker.GetRec().GetProcessingSettings().ompThreads))
29-
for (int i = get_global_id(0); i < trdTracker.NTracks(); i += get_global_size(0)) {
30-
trdTracker.DoTrackingThread(i, get_global_id(0));
28+
auto* trdTracker = &processors.getTRDTracker<I>();
29+
#ifndef GPUCA_GPUCODE_DEVICE
30+
if constexpr (std::is_same_v<decltype(trdTracker), decltype(externalInstance)>) {
31+
if (externalInstance) {
32+
trdTracker = externalInstance;
33+
}
34+
}
35+
#endif
36+
GPUCA_OPENMP(parallel for if(!trdTracker->GetRec().GetProcessingSettings().ompKernels) num_threads(trdTracker->GetRec().GetProcessingSettings().ompThreads))
37+
for (int i = get_global_id(0); i < trdTracker->NTracks(); i += get_global_size(0)) {
38+
trdTracker->DoTrackingThread(i, get_global_id(0));
3139
}
3240
}
3341

34-
template GPUd() void GPUTRDTrackerKernels::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors);
42+
template GPUd() void GPUTRDTrackerKernels::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTracker* externalInstance);
3543
#ifdef GPUCA_HAVE_O2HEADERS
36-
template GPUd() void GPUTRDTrackerKernels::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors);
37-
#endif
44+
template GPUd() void GPUTRDTrackerKernels::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTracker* externalInstance);
45+
#endif // GPUCA_HAVE_O2HEADERS

‎GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ class GPUTRDTrackerKernels : public GPUKernelTemplate
3030
o2Version = 1 };
3131
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::TRDTracking; }
3232
template <int iKernel = defaultKernel>
33-
GPUd() static void Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors);
33+
GPUd() static void Thread(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTracker* externalInstance = nullptr);
3434
};
3535
} // namespace gpu
3636
} // namespace GPUCA_NAMESPACE

0 commit comments

Comments
 (0)