Skip to content

Commit 2a2b4b9

Browse files
mpucciompuccio
committed
ITS: implement code review
1 parent d4e7f5a commit 2a2b4b9

20 files changed

Lines changed: 550 additions & 1087 deletions

‎DataFormats/Detectors/ITSMFT/ITS/include/DataFormatsITS/TrackITS.h‎

Lines changed: 39 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,11 @@ namespace its
3535

3636
class TrackITS : public o2::track::TrackParCov
3737
{
38+
public:
39+
static constexpr unsigned int ExtendedPatternShift = 24;
40+
static constexpr int MaxLayersInTrackPattern = 8;
41+
42+
private:
3843
enum UserBits {
3944
kSharedClusters = 1 << 28
4045
};
@@ -106,16 +111,47 @@ class TrackITS : public o2::track::TrackParCov
106111
GPUhdi() uint32_t getPattern() const { return mPattern; }
107112
bool hasHitOnLayer(uint32_t i) const { return mPattern & (0x1 << i); }
108113
bool isFakeOnLayer(uint32_t i) const { return !(mPattern & (0x1 << (16 + i))); }
109-
bool isExtendedOnLayer(uint32_t i) const { return (mPattern & (0x1 << (24 + i))); } // only correct if getNClusters <= 8 on layers <= 8
110-
uint32_t getLastClusterLayer() const
114+
bool isExtendedOnLayer(uint32_t i) const { return (mPattern & (0x1 << (ExtendedPatternShift + i))); } // only correct if getNClusters <= 8 on layers <= 8
115+
template <int NLayers>
116+
GPUhdi() static constexpr uint32_t getLayerPatternMask()
117+
{
118+
return (NLayers >= 32) ? 0xffffffffu : ((1u << NLayers) - 1u);
119+
}
120+
template <int NLayers>
121+
GPUhdi() void setExtendedLayerPattern(uint32_t pattern)
122+
{
123+
pattern &= getLayerPatternMask<NLayers>();
124+
setUserField(static_cast<uint16_t>(pattern));
125+
if constexpr (NLayers <= MaxLayersInTrackPattern) {
126+
setPattern(getPattern() | (pattern << ExtendedPatternShift));
127+
}
128+
}
129+
template <int NLayers>
130+
GPUhdi() uint32_t getExtendedLayerPattern() const
131+
{
132+
const auto mask = getLayerPatternMask<NLayers>();
133+
if constexpr (NLayers <= MaxLayersInTrackPattern) {
134+
const auto pattern = (getPattern() >> ExtendedPatternShift) & mask;
135+
if (pattern) {
136+
return pattern;
137+
}
138+
}
139+
return getUserField() & mask;
140+
}
141+
GPUhdi() void clearExtendedLayerPattern()
142+
{
143+
setUserField(0);
144+
getParamOut().setUserField(0);
145+
}
146+
GPUhdi() uint32_t getLastClusterLayer() const
111147
{
112148
uint32_t r{0}, v{mPattern & ((1 << 16) - 1)};
113149
while (v >>= 1) {
114150
r++;
115151
}
116152
return r;
117153
}
118-
uint32_t getFirstClusterLayer() const
154+
GPUhdi() uint32_t getFirstClusterLayer() const
119155
{
120156
int s{0};
121157
while (!(mPattern & (1 << s))) {

‎Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h‎

Lines changed: 1 addition & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919
#include "ITStracking/BoundedAllocator.h"
2020
#include "ITStracking/TimeFrame.h"
2121
#include "ITStracking/Configuration.h"
22-
#include "ITStracking/TrackExtensionCandidate.h"
22+
#include "ITStracking/TrackExtensionHypothesis.h"
2323
#include "ITStrackingGPU/Utils.h"
2424

2525
namespace o2::its::gpu
@@ -91,13 +91,9 @@ class TimeFrameGPU : public TimeFrame<NLayers>
9191
void createNeighboursDevice(const unsigned int layer);
9292
void createNeighboursLUTDevice(const int, const unsigned int);
9393
void createTrackITSExtDevice(const size_t);
94-
void loadTrackExtensionStartTracksDevice();
95-
void createTrackExtensionCandidatesDevice(const size_t);
9694
void createTrackExtensionScratchDevice(const int nThreads, const int beamWidth);
97-
void createTrackExtensionResultsDevice(const size_t);
9895
void downloadTrackITSExtDevice();
9996
void downloadCellsNeighboursDevice(std::vector<bounded_vector<CellNeighbour>>&, const int);
100-
void downloadTrackExtensionResultsDevice();
10197
void downloadNeighboursLUTDevice(bounded_vector<int>&, const int);
10298
void downloadCellsDevice();
10399
void downloadCellsLUTDevice();
@@ -124,20 +120,15 @@ class TimeFrameGPU : public TimeFrame<NLayers>
124120
const auto getDeviceTrackingTopologyView() const { return mDeviceTrackingTopologyView; }
125121
int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; }
126122
auto& getTrackITSExt() { return mTrackITSExt; }
127-
auto& getTrackExtensionResults() { return mTrackExtensionResults; }
128123
Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; }
129124
int* getDeviceROFramesPV() { return mROFramesPVDevice; }
130125
unsigned char* getDeviceUsedClusters(const int);
131126
const o2::base::Propagator* getChainPropagator();
132127

133128
// Hybrid
134129
TrackITSExt* getDeviceTrackITSExt() { return mTrackITSExtDevice; }
135-
TrackITSExt* getDeviceTrackExtensionStartTracks() { return mTrackExtensionStartTracksDevice; }
136-
TrackExtensionCandidate<NLayers>* getDeviceTrackExtensionCandidates() { return mTrackExtensionCandidatesDevice; }
137-
int* getDeviceTrackExtensionCandidateOffsets() { return mTrackExtensionCandidateOffsetsDevice; }
138130
TrackExtensionHypothesis<NLayers>* getDeviceActiveTrackExtensionHypotheses() { return mActiveTrackExtensionHypothesesDevice; }
139131
TrackExtensionHypothesis<NLayers>* getDeviceNextTrackExtensionHypotheses() { return mNextTrackExtensionHypothesesDevice; }
140-
TrackExtensionResult<NLayers>* getDeviceTrackExtensionResults() { return mTrackExtensionResultsDevice; }
141132
int* getDeviceNeighboursLUT(const int layer) { return mNeighboursLUTDevice[layer]; }
142133
gsl::span<int*> getDeviceNeighboursLUTs() { return mNeighboursLUTDevice; }
143134
CellNeighbour** getDeviceArrayNeighbours() { return mNeighboursDeviceArray; }
@@ -235,13 +226,8 @@ class TimeFrameGPU : public TimeFrame<NLayers>
235226
float** mCellSeedsChi2DeviceArray;
236227

237228
TrackITSExt* mTrackITSExtDevice;
238-
TrackITSExt* mTrackExtensionStartTracksDevice{nullptr};
239-
TrackExtensionCandidate<NLayers>* mTrackExtensionCandidatesDevice{nullptr};
240-
int* mTrackExtensionCandidateOffsetsDevice{nullptr};
241229
TrackExtensionHypothesis<NLayers>* mActiveTrackExtensionHypothesesDevice{nullptr};
242230
TrackExtensionHypothesis<NLayers>* mNextTrackExtensionHypothesesDevice{nullptr};
243-
TrackExtensionResult<NLayers>* mTrackExtensionResultsDevice{nullptr};
244-
unsigned int mNTrackExtensionResults{0};
245231
std::array<CellNeighbour*, MaxCells> mNeighboursDevice{};
246232
CellNeighbour** mNeighboursDeviceArray{nullptr};
247233
std::array<TrackingFrameInfo*, NLayers> mTrackingFrameInfoDevice;
@@ -258,9 +244,6 @@ class TimeFrameGPU : public TimeFrame<NLayers>
258244

259245
// Temporary buffer for storing output tracks from GPU tracking
260246
bounded_vector<TrackITSExt> mTrackITSExt;
261-
bounded_vector<TrackITSExt> mTrackExtensionStartTracks;
262-
// Temporary buffer for fitted track extension proposals from GPU tracking
263-
bounded_vector<TrackExtensionResult<NLayers>> mTrackExtensionResults;
264247
};
265248

266249
template <int NLayers>

‎Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h‎

Lines changed: 1 addition & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919
#include "ITStracking/BoundedAllocator.h"
2020
#include "ITStracking/ROFLookupTables.h"
2121
#include "ITStracking/TrackingTopology.h"
22-
#include "ITStracking/TrackExtensionCandidate.h"
22+
#include "ITStracking/TrackExtensionHypothesis.h"
2323
#include "ITStrackingGPU/Utils.h"
2424
#include "DetectorsBase/Propagator.h"
2525

@@ -37,58 +37,6 @@ class Cluster;
3737
class TrackITSExt;
3838
class ExternalAllocator;
3939

40-
inline constexpr int kTrackExtensionLaunchBlocks = 60;
41-
inline constexpr int kTrackExtensionLaunchThreadsPerBlock = 256;
42-
inline constexpr int kTrackExtensionLaunchThreads = kTrackExtensionLaunchBlocks * kTrackExtensionLaunchThreadsPerBlock;
43-
44-
template <int NLayers>
45-
void computeTrackExtensionCandidatesHandler(const TrackITSExt* tracks,
46-
const IndexTableUtils<NLayers>* utils,
47-
const typename ROFMaskTable<NLayers>::View& rofMask,
48-
const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
49-
const Cluster** clusters,
50-
const unsigned char** usedClusters,
51-
const int** clustersIndexTables,
52-
const int** ROFClusters,
53-
const TrackingFrameInfo** trackingFrameInfo,
54-
TrackExtensionCandidate<NLayers>* candidates,
55-
int* candidateOffsets,
56-
TrackExtensionHypothesis<NLayers>* activeHypotheses,
57-
TrackExtensionHypothesis<NLayers>* nextHypotheses,
58-
const std::array<float, NLayers> layerRadii,
59-
const std::array<float, NLayers> layerxX0,
60-
const int nTracks,
61-
const int nLayers,
62-
const int phiBins,
63-
const int beamWidth,
64-
const bool extendTop,
65-
const bool extendBot,
66-
const float bz,
67-
const float maxChi2ClusterAttachment,
68-
const float maxChi2NDF,
69-
const float nSigmaCutPhi,
70-
const float nSigmaCutZ,
71-
const o2::base::Propagator* propagator,
72-
const o2::base::PropagatorF::MatCorrType matCorrType,
73-
gpu::Stream& stream);
74-
75-
template <int NLayers>
76-
void computeTrackExtensionResultsHandler(const TrackITSExt* tracks,
77-
const TrackExtensionCandidate<NLayers>* candidates,
78-
const int* candidateOffsets,
79-
TrackExtensionResult<NLayers>* results,
80-
const TrackingFrameInfo** trackingFrameInfo,
81-
const std::array<float, NLayers> layerxX0,
82-
const int nTracks,
83-
const int nLayers,
84-
const float bz,
85-
const float maxChi2ClusterAttachment,
86-
const float maxChi2NDF,
87-
const o2::base::Propagator* propagator,
88-
const o2::base::PropagatorF::MatCorrType matCorrType,
89-
const bool shiftRefToCluster,
90-
gpu::Stream& stream);
91-
9240
template <int NLayers>
9341
void countTrackletsInROFsHandler(const IndexTableUtils<NLayers>* utils,
9442
const typename ROFMaskTable<NLayers>::View& rofMask,

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu‎

Lines changed: 0 additions & 62 deletions
Original file line numberDiff line numberDiff line change
@@ -582,35 +582,6 @@ void TimeFrameGPU<NLayers>::createTrackITSExtDevice(const size_t nSeeds)
582582
GPUChkErrS(cudaMemset(mTrackITSExtDevice, 0, mNTracks * sizeof(o2::its::TrackITSExt)));
583583
}
584584

585-
template <int NLayers>
586-
void TimeFrameGPU<NLayers>::loadTrackExtensionStartTracksDevice()
587-
{
588-
GPUTimer timer("loading track extension start tracks");
589-
GPULog("gpu-transfer: loading {} track extension start tracks, for {:.2f} MB.", this->mTracks.size(), this->mTracks.size() * sizeof(o2::its::TrackITSExt) / constants::MB);
590-
mTrackExtensionStartTracksDevice = nullptr;
591-
mTrackExtensionStartTracks = bounded_vector<TrackITSExt>(this->mTracks.begin(), this->mTracks.end(), this->getMemoryPool().get());
592-
if (this->mTracks.empty()) {
593-
return;
594-
}
595-
allocMem(reinterpret_cast<void**>(&mTrackExtensionStartTracksDevice), mTrackExtensionStartTracks.size() * sizeof(o2::its::TrackITSExt), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK));
596-
GPUChkErrS(cudaMemcpy(mTrackExtensionStartTracksDevice, mTrackExtensionStartTracks.data(), mTrackExtensionStartTracks.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyHostToDevice));
597-
}
598-
599-
template <int NLayers>
600-
void TimeFrameGPU<NLayers>::createTrackExtensionCandidatesDevice(const size_t nTracks)
601-
{
602-
GPUTimer timer("reserving track extension candidates");
603-
const size_t nCandidates = nTracks * MaxTrackExtensionCandidatesPerTrack;
604-
GPULog("gpu-allocation: reserving {} track extension candidates, for {:.2f} MB.", nCandidates, nCandidates * sizeof(o2::its::TrackExtensionCandidate<NLayers>) / constants::MB);
605-
mTrackExtensionCandidatesDevice = nullptr;
606-
mTrackExtensionCandidateOffsetsDevice = nullptr;
607-
if (nCandidates == 0) {
608-
return;
609-
}
610-
allocMem(reinterpret_cast<void**>(&mTrackExtensionCandidatesDevice), nCandidates * sizeof(o2::its::TrackExtensionCandidate<NLayers>), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK));
611-
allocMem(reinterpret_cast<void**>(&mTrackExtensionCandidateOffsetsDevice), (nTracks + 1) * sizeof(int), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK));
612-
}
613-
614585
template <int NLayers>
615586
void TimeFrameGPU<NLayers>::createTrackExtensionScratchDevice(const int nThreads, const int beamWidth)
616587
{
@@ -626,28 +597,6 @@ void TimeFrameGPU<NLayers>::createTrackExtensionScratchDevice(const int nThreads
626597
allocMem(reinterpret_cast<void**>(&mNextTrackExtensionHypothesesDevice), nHypotheses * sizeof(o2::its::TrackExtensionHypothesis<NLayers>), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK));
627598
}
628599

629-
template <int NLayers>
630-
void TimeFrameGPU<NLayers>::createTrackExtensionResultsDevice(const size_t nTracks)
631-
{
632-
GPUTimer timer("reserving fitted track extension results");
633-
mNTrackExtensionResults = 0;
634-
if (nTracks == 0 || mTrackExtensionCandidateOffsetsDevice == nullptr) {
635-
mTrackExtensionResults = bounded_vector<TrackExtensionResult<NLayers>>(0, {}, this->getMemoryPool().get());
636-
mTrackExtensionResultsDevice = nullptr;
637-
return;
638-
}
639-
int nResults{0};
640-
GPUChkErrS(cudaMemcpy(&nResults, mTrackExtensionCandidateOffsetsDevice + nTracks, sizeof(int), cudaMemcpyDeviceToHost));
641-
mNTrackExtensionResults = nResults;
642-
GPULog("gpu-allocation: reserving {} fitted track extension results, for {:.2f} MB.", mNTrackExtensionResults, mNTrackExtensionResults * sizeof(o2::its::TrackExtensionResult<NLayers>) / constants::MB);
643-
mTrackExtensionResults = bounded_vector<TrackExtensionResult<NLayers>>(mNTrackExtensionResults, {}, this->getMemoryPool().get());
644-
mTrackExtensionResultsDevice = nullptr;
645-
if (mTrackExtensionResults.empty()) {
646-
return;
647-
}
648-
allocMem(reinterpret_cast<void**>(&mTrackExtensionResultsDevice), mNTrackExtensionResults * sizeof(o2::its::TrackExtensionResult<NLayers>), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK));
649-
}
650-
651600
template <int NLayers>
652601
void TimeFrameGPU<NLayers>::downloadCellsDevice()
653602
{
@@ -694,17 +643,6 @@ void TimeFrameGPU<NLayers>::downloadTrackITSExtDevice()
694643
GPUChkErrS(cudaMemcpy(mTrackITSExt.data(), mTrackITSExtDevice, mTrackITSExt.size() * sizeof(o2::its::TrackITSExt), cudaMemcpyDeviceToHost));
695644
}
696645

697-
template <int NLayers>
698-
void TimeFrameGPU<NLayers>::downloadTrackExtensionResultsDevice()
699-
{
700-
GPUTimer timer("downloading fitted track extension results");
701-
GPULog("gpu-transfer: downloading {} fitted track extension results, for {:.2f} MB.", mTrackExtensionResults.size(), mTrackExtensionResults.size() * sizeof(o2::its::TrackExtensionResult<NLayers>) / constants::MB);
702-
if (mTrackExtensionResults.empty()) {
703-
return;
704-
}
705-
GPUChkErrS(cudaMemcpy(mTrackExtensionResults.data(), mTrackExtensionResultsDevice, mTrackExtensionResults.size() * sizeof(o2::its::TrackExtensionResult<NLayers>), cudaMemcpyDeviceToHost));
706-
}
707-
708646
template <int NLayers>
709647
void TimeFrameGPU<NLayers>::unregisterHostMemory(const int maxLayers)
710648
{

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx‎

Lines changed: 6 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,10 @@
2121

2222
namespace o2::its
2323
{
24+
namespace
25+
{
26+
constexpr int trackExtensionLaunchThreads = 60 * 256;
27+
}
2428

2529
template <int NLayers>
2630
void TrackerTraitsGPU<NLayers>::initialiseTimeFrame(const int iteration)
@@ -309,8 +313,6 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
309313
const bool extendTop = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerTop];
310314
const bool extendBot = this->mTrkParams[iteration].PassFlags[IterationStep::TrackFollowerBot];
311315
const bool extendTracks = extendTop || extendBot;
312-
size_t nExtendedTracks{0};
313-
size_t nExtendedClusters{0};
314316
for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) {
315317
bounded_vector<TrackSeed<NLayers>> trackSeeds(this->getMemoryPool().get());
316318
for (int startCellTopologyId{0}; startCellTopologyId < hostTopology.nCells; ++startCellTopologyId) {
@@ -369,7 +371,7 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
369371
mTimeFrameGPU->getFrameworkAllocator());
370372
mTimeFrameGPU->createTrackITSExtDevice(trackSeeds.size());
371373
if (extendTracks) {
372-
mTimeFrameGPU->createTrackExtensionScratchDevice(kTrackExtensionLaunchThreads, this->mTrkParams[iteration].TrackFollowerBeamWidth);
374+
mTimeFrameGPU->createTrackExtensionScratchDevice(trackExtensionLaunchThreads, this->mTrkParams[iteration].TrackFollowerBeamWidth);
373375
}
374376
computeTrackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(),
375377
mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
@@ -409,12 +411,9 @@ void TrackerTraitsGPU<NLayers>::findRoads(const int iteration)
409411
mTimeFrameGPU->downloadTrackITSExtDevice();
410412

411413
auto& tracks = mTimeFrameGPU->getTrackITSExt();
412-
this->acceptTracks(iteration, tracks, firstClusters, nExtendedTracks, nExtendedClusters);
414+
this->acceptTracks(iteration, tracks, firstClusters);
413415
mTimeFrameGPU->loadUsedClustersDevice();
414416
}
415-
if (extendTracks) {
416-
LOGP(info, "Integrated track extension accepted {} tracks using {} clusters in iteration {}", nExtendedTracks, nExtendedClusters, iteration);
417-
}
418417
this->markTracks(iteration);
419418
// wipe the artefact memory
420419
mTimeFrameGPU->popMemoryStack(iteration);

0 commit comments

Comments
 (0)