Skip to content

Commit a0e63ef

Browse files
davidrohrdavidrohr
committed
GPU TPC: Drop unused TrackletConstructor_allSectors kernel
1 parent 6d54cfc commit a0e63ef

File tree

5 files changed

+4
-89
lines changed

5 files changed

+4
-89
lines changed

‎GPU/GPUTracking/Definitions/GPUDefParametersDefault.h‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -515,8 +515,6 @@
515515
#define GPUCA_LB_GPUTPCCFStreamCompaction_scanTop GPUCA_THREAD_COUNT_SCAN
516516
#define GPUCA_LB_GPUTPCCFStreamCompaction_scanDown GPUCA_THREAD_COUNT_SCAN
517517
#define GPUCA_LB_GPUTPCCFStreamCompaction_compactDigits GPUCA_THREAD_COUNT_SCAN
518-
#define GPUCA_LB_GPUTPCTrackletConstructor_singleSector GPUCA_LB_GPUTPCTrackletConstructor
519-
#define GPUCA_LB_GPUTPCTrackletConstructor_allSectors GPUCA_LB_GPUTPCTrackletConstructor
520518
#define GPUCA_LB_GPUTPCCompressionGatherKernels_unbuffered GPUCA_LB_COMPRESSION_GATHER
521519
#define GPUCA_LB_GPUTPCCompressionGatherKernels_buffered32 GPUCA_LB_COMPRESSION_GATHER
522520
#define GPUCA_LB_GPUTPCCompressionGatherKernels_buffered64 GPUCA_LB_COMPRESSION_GATHER

‎GPU/GPUTracking/Global/GPUChainTrackingSectorTracker.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ int32_t GPUChainTracking::RunTPCTrackingSectors_internal()
105105
for (uint32_t iSector = 0; iSector < NSECTORS; iSector++) {
106106
processorsShadow()->tpcTrackers[iSector].GPUParametersConst()->gpumem = (char*)mRec->DeviceMemoryBase();
107107
// Initialize Startup Constants
108-
processors()->tpcTrackers[iSector].GPUParameters()->nextStartHit = (((getKernelProperties<GPUTPCTrackletConstructor, GPUTPCTrackletConstructor::allSectors>().minBlocks * BlockCount()) + NSECTORS - 1 - iSector) / NSECTORS) * getKernelProperties<GPUTPCTrackletConstructor, GPUTPCTrackletConstructor::allSectors>().nThreads;
108+
processors()->tpcTrackers[iSector].GPUParameters()->nextStartHit = (((getKernelProperties<GPUTPCTrackletConstructor>().minBlocks * BlockCount()) + NSECTORS - 1 - iSector) / NSECTORS) * getKernelProperties<GPUTPCTrackletConstructor>().nThreads;
109109
processorsShadow()->tpcTrackers[iSector].SetGPUTextureBase(mRec->DeviceMemoryBase());
110110
}
111111

‎GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.cxx‎

Lines changed: 1 addition & 74 deletions
Original file line numberDiff line numberDiff line change
@@ -476,7 +476,7 @@ GPUdic(2, 1) void GPUTPCTrackletConstructor::DoTracklet(GPUconstantref() GPUTPCT
476476
}
477477

478478
template <>
479-
GPUdii() void GPUTPCTrackletConstructor::Thread<GPUTPCTrackletConstructor::singleSector>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& sMem, processorType& GPUrestrict() tracker)
479+
GPUdii() void GPUTPCTrackletConstructor::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& sMem, processorType& GPUrestrict() tracker)
480480
{
481481
if (get_local_id(0) == 0) {
482482
sMem.mNStartHits = *tracker.NStartHits();
@@ -491,79 +491,6 @@ GPUdii() void GPUTPCTrackletConstructor::Thread<GPUTPCTrackletConstructor::singl
491491
}
492492
}
493493

494-
template <>
495-
GPUdii() void GPUTPCTrackletConstructor::Thread<GPUTPCTrackletConstructor::allSectors>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& sMem, processorType& GPUrestrict() tracker0)
496-
{
497-
GPUconstantref() GPUTPCTracker* GPUrestrict() pTracker = &tracker0;
498-
#ifdef GPUCA_GPUCODE
499-
int32_t mySector = get_group_id(0) % GPUCA_NSECTORS;
500-
int32_t currentSector = -1;
501-
502-
if (get_local_id(0) == 0) {
503-
sMem.mNextStartHitFirstRun = 1;
504-
}
505-
GPUCA_UNROLL(, U())
506-
for (uint32_t iSector = 0; iSector < GPUCA_NSECTORS; iSector++) {
507-
GPUconstantref() GPUTPCTracker& GPUrestrict() tracker = pTracker[mySector];
508-
509-
GPUTPCThreadMemory rMem;
510-
511-
while ((rMem.mISH = FetchTracklet(tracker, sMem)) != -2) {
512-
if (rMem.mISH >= 0 && get_local_id(0) < GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCTrackletConstructor)) {
513-
rMem.mISH += get_local_id(0);
514-
} else {
515-
rMem.mISH = -1;
516-
}
517-
518-
if (mySector != currentSector) {
519-
if (get_local_id(0) == 0) {
520-
sMem.mNStartHits = *tracker.NStartHits();
521-
}
522-
CA_SHARED_CACHE(&sMem.mRows[0], tracker.TrackingDataRows(), GPUCA_ROW_COUNT * sizeof(GPUTPCRow));
523-
GPUbarrier();
524-
currentSector = mySector;
525-
}
526-
527-
if (rMem.mISH >= 0 && rMem.mISH < sMem.mNStartHits) {
528-
rMem.mGo = true;
529-
DoTracklet(tracker, sMem, rMem);
530-
}
531-
}
532-
if (++mySector >= GPUCA_NSECTORS) {
533-
mySector = 0;
534-
}
535-
}
536-
#else
537-
for (int32_t iSector = 0; iSector < GPUCA_NSECTORS; iSector++) {
538-
Thread<singleSector>(nBlocks, nThreads, iBlock, iThread, sMem, pTracker[iSector]);
539-
}
540-
#endif
541-
}
542-
543-
#ifdef GPUCA_GPUCODE
544-
545-
GPUd() int32_t GPUTPCTrackletConstructor::FetchTracklet(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUSharedMemory& sMem)
546-
{
547-
const uint32_t nStartHit = *tracker.NStartHits();
548-
GPUbarrier();
549-
if (get_local_id(0) == 0) {
550-
int32_t firstStartHit = -2;
551-
if (sMem.mNextStartHitFirstRun == 1) {
552-
firstStartHit = (get_group_id(0) - tracker.ISector()) / GPUCA_NSECTORS * GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCTrackletConstructor);
553-
sMem.mNextStartHitFirstRun = 0;
554-
} else {
555-
if (tracker.GPUParameters()->nextStartHit < nStartHit) {
556-
firstStartHit = CAMath::AtomicAdd<uint32_t>(&tracker.GPUParameters()->nextStartHit, GPUCA_GET_THREAD_COUNT(GPUCA_LB_GPUTPCTrackletConstructor));
557-
}
558-
}
559-
sMem.mNextStartHitFirst = firstStartHit < (int32_t)nStartHit ? firstStartHit : -2;
560-
}
561-
GPUbarrier();
562-
return (sMem.mNextStartHitFirst);
563-
}
564-
565-
#endif // GPUCA_GPUCODE
566-
567494
template <> // FIXME: GPUgeneric() needed to make the clang spirv output link correctly
568495
GPUd() int32_t GPUTPCTrackletConstructor::GPUTPCTrackletConstructorExtrapolationTracking<GPUgeneric() GPUTPCExtrapolationTracking::GPUSharedMemory>(GPUconstantref() GPUTPCTracker& GPUrestrict() tracker, GPUsharedref() GPUTPCExtrapolationTracking::GPUSharedMemory& sMem, GPUTPCTrackParam& GPUrestrict() tParam, int32_t row, int32_t increment, int32_t iTracklet, calink* rowHits)
569496
{

‎GPU/GPUTracking/SectorTracker/GPUTPCTrackletConstructor.h‎

Lines changed: 1 addition & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -28,14 +28,9 @@ namespace o2::gpu
2828
*/
2929
class GPUTPCTracker;
3030

31-
class GPUTPCTrackletConstructor
31+
class GPUTPCTrackletConstructor : public GPUKernelTemplate
3232
{
3333
public:
34-
enum K {
35-
singleSector = 0,
36-
allSectors = 1
37-
};
38-
3934
class GPUTPCThreadMemory
4035
{
4136
friend class GPUTPCTrackletConstructor; //! friend class
@@ -89,10 +84,6 @@ class GPUTPCTrackletConstructor
8984

9085
GPUd() static void DoTracklet(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() GPUTPCTrackletConstructor::GPUSharedMemory& sMem, GPUTPCThreadMemory& rMem);
9186

92-
#ifdef GPUCA_GPUCODE
93-
GPUd() static int32_t FetchTracklet(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() GPUSharedMemory& sMem);
94-
#endif // GPUCA_GPUCODE
95-
9687
template <class T>
9788
GPUd() static int32_t GPUTPCTrackletConstructorExtrapolationTracking(GPUconstantref() GPUTPCTracker& tracker, GPUsharedref() T& sMem, GPUTPCTrackParam& tParam, int32_t startrow, int32_t increment, int32_t iTracklet, calink* rowHits);
9889

‎GPU/GPUTracking/kernels.cmake‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -34,8 +34,7 @@ o2_gpu_add_kernel("GPUTPCNeighboursFinder" "= TPCTRAC
3434
o2_gpu_add_kernel("GPUTPCNeighboursCleaner" "= TPCTRACKER" LB)
3535
o2_gpu_add_kernel("GPUTPCStartHitsFinder" "= TPCTRACKER" LB)
3636
o2_gpu_add_kernel("GPUTPCStartHitsSorter" "= TPCTRACKER" LB)
37-
o2_gpu_add_kernel("GPUTPCTrackletConstructor, singleSector" "= TPCTRACKER" LB)
38-
o2_gpu_add_kernel("GPUTPCTrackletConstructor, allSectors" "= TPCTRACKER" LB)
37+
o2_gpu_add_kernel("GPUTPCTrackletConstructor" "= TPCTRACKER" LB)
3938
o2_gpu_add_kernel("GPUTPCTrackletSelector" "= TPCTRACKER" LB)
4039
o2_gpu_add_kernel("GPUMemClean16" "GPUGeneralKernels" NO void* ptr "uint64_t" size)
4140
o2_gpu_add_kernel("GPUitoa" "GPUGeneralKernels" NO int32_t* ptr "uint64_t" size)

0 commit comments

Comments
 (0)