Skip to content

Commit 4491bd5

Browse files
mpucciompuccio
committed
ITS: simplifications
1 parent e9f275a commit 4491bd5

11 files changed

Lines changed: 99 additions & 121 deletions

File tree

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

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -91,7 +91,7 @@ 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 loadTrackExtensionStartStatesDevice();
94+
void loadTrackExtensionStartTracksDevice();
9595
void createTrackExtensionCandidatesDevice(const size_t);
9696
void createTrackExtensionScratchDevice(const int nThreads, const int beamWidth);
9797
void createTrackExtensionResultsDevice(const size_t);
@@ -134,7 +134,7 @@ class TimeFrameGPU : public TimeFrame<NLayers>
134134

135135
// Hybrid
136136
TrackITSExt* getDeviceTrackITSExt() { return mTrackITSExtDevice; }
137-
TrackExtensionStartState<NLayers>* getDeviceTrackExtensionStartStates() { return mTrackExtensionStartStatesDevice; }
137+
TrackITSExt* getDeviceTrackExtensionStartTracks() { return mTrackExtensionStartTracksDevice; }
138138
TrackExtensionCandidate<NLayers>* getDeviceTrackExtensionCandidates() { return mTrackExtensionCandidatesDevice; }
139139
int* getDeviceTrackExtensionCandidateOffsets() { return mTrackExtensionCandidateOffsetsDevice; }
140140
TrackExtensionHypothesis<NLayers>* getDeviceActiveTrackExtensionHypotheses() { return mActiveTrackExtensionHypothesesDevice; }
@@ -237,7 +237,7 @@ class TimeFrameGPU : public TimeFrame<NLayers>
237237
float** mCellSeedsChi2DeviceArray;
238238

239239
TrackITSExt* mTrackITSExtDevice;
240-
TrackExtensionStartState<NLayers>* mTrackExtensionStartStatesDevice{nullptr};
240+
TrackITSExt* mTrackExtensionStartTracksDevice{nullptr};
241241
TrackExtensionCandidate<NLayers>* mTrackExtensionCandidatesDevice{nullptr};
242242
int* mTrackExtensionCandidateOffsetsDevice{nullptr};
243243
TrackExtensionHypothesis<NLayers>* mActiveTrackExtensionHypothesesDevice{nullptr};
@@ -260,8 +260,7 @@ class TimeFrameGPU : public TimeFrame<NLayers>
260260

261261
// Temporary buffer for storing output tracks from GPU tracking
262262
bounded_vector<TrackITSExt> mTrackITSExt;
263-
// Temporary buffer for compact track states used by GPU track extension
264-
bounded_vector<TrackExtensionStartState<NLayers>> mTrackExtensionStartStates;
263+
bounded_vector<TrackITSExt> mTrackExtensionStartTracks;
265264
// Temporary buffer for compact track extension proposals from GPU tracking
266265
bounded_vector<TrackExtensionCandidate<NLayers>> mTrackExtensionCandidates;
267266
// Temporary buffer for fitted track extension proposals from GPU tracking

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -42,7 +42,7 @@ inline constexpr int kTrackExtensionLaunchThreadsPerBlock = 256;
4242
inline constexpr int kTrackExtensionLaunchThreads = kTrackExtensionLaunchBlocks * kTrackExtensionLaunchThreadsPerBlock;
4343

4444
template <int NLayers>
45-
void computeTrackExtensionCandidatesHandler(const TrackExtensionStartState<NLayers>* tracks,
45+
void computeTrackExtensionCandidatesHandler(const TrackITSExt* tracks,
4646
const IndexTableUtils<NLayers>* utils,
4747
const typename ROFMaskTable<NLayers>::View& rofMask,
4848
const typename ROFOverlapTable<NLayers>::View& rofOverlaps,
@@ -73,7 +73,7 @@ void computeTrackExtensionCandidatesHandler(const TrackExtensionStartState<NLaye
7373
gpu::Stream& stream);
7474

7575
template <int NLayers>
76-
void computeTrackExtensionResultsHandler(const TrackExtensionStartState<NLayers>* tracks,
76+
void computeTrackExtensionResultsHandler(const TrackITSExt* tracks,
7777
const TrackExtensionCandidate<NLayers>* candidates,
7878
const int* candidateOffsets,
7979
TrackExtensionResult<NLayers>* results,

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

Lines changed: 7 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -583,31 +583,17 @@ void TimeFrameGPU<NLayers>::createTrackITSExtDevice(const size_t nSeeds)
583583
}
584584

585585
template <int NLayers>
586-
void TimeFrameGPU<NLayers>::loadTrackExtensionStartStatesDevice()
586+
void TimeFrameGPU<NLayers>::loadTrackExtensionStartTracksDevice()
587587
{
588-
GPUTimer timer("loading track extension start states");
589-
GPULog("gpu-transfer: loading {} track extension start states, for {:.2f} MB.", this->mTracks.size(), this->mTracks.size() * sizeof(o2::its::TrackExtensionStartState<NLayers>) / constants::MB);
590-
mTrackExtensionStartStatesDevice = nullptr;
591-
mTrackExtensionStartStates = bounded_vector<TrackExtensionStartState<NLayers>>(this->mTracks.size(), {}, this->getMemoryPool().get());
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());
592592
if (this->mTracks.empty()) {
593593
return;
594594
}
595-
for (size_t iTrack{0}; iTrack < this->mTracks.size(); ++iTrack) {
596-
const auto& track = this->mTracks[iTrack];
597-
auto& state = mTrackExtensionStartStates[iTrack];
598-
state.paramIn = track.getParamIn();
599-
state.paramOut = track.getParamOut();
600-
state.time = track.getTimeStamp();
601-
state.chi2 = track.getChi2();
602-
state.nClusters = track.getNClusters();
603-
state.firstClusterLayer = static_cast<int>(track.getFirstClusterLayer());
604-
state.lastClusterLayer = static_cast<int>(track.getLastClusterLayer());
605-
for (int iLayer{0}; iLayer < NLayers; ++iLayer) {
606-
state.clusters[iLayer] = track.getClusterIndex(iLayer);
607-
}
608-
}
609-
allocMem(reinterpret_cast<void**>(&mTrackExtensionStartStatesDevice), mTrackExtensionStartStates.size() * sizeof(o2::its::TrackExtensionStartState<NLayers>), this->hasFrameworkAllocator(), (o2::gpu::GPUMemoryResource::MEMORY_GPU | o2::gpu::GPUMemoryResource::MEMORY_STACK));
610-
GPUChkErrS(cudaMemcpy(mTrackExtensionStartStatesDevice, mTrackExtensionStartStates.data(), mTrackExtensionStartStates.size() * sizeof(o2::its::TrackExtensionStartState<NLayers>), cudaMemcpyHostToDevice));
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));
611597
}
612598

613599
template <int NLayers>

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

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -422,7 +422,7 @@ void TrackerTraitsGPU<NLayers>::buildTrackExtensionCandidates(const int iteratio
422422
const auto nTracks = this->mTimeFrame->getTracks().size();
423423
const int beamWidth = std::max(1, this->mTrkParams[iteration].TrackFollowerBeamWidth);
424424
mTimeFrameGPU->syncStreams();
425-
mTimeFrameGPU->loadTrackExtensionStartStatesDevice();
425+
mTimeFrameGPU->loadTrackExtensionStartTracksDevice();
426426
mTimeFrameGPU->createTrackExtensionCandidatesDevice(nTracks);
427427
mTimeFrameGPU->createTrackExtensionScratchDevice(kTrackExtensionLaunchThreads, beamWidth);
428428
std::array<float, NLayers> layerRadii{};
@@ -431,7 +431,7 @@ void TrackerTraitsGPU<NLayers>::buildTrackExtensionCandidates(const int iteratio
431431
layerRadii[iLayer] = this->mTrkParams[iteration].LayerRadii[iLayer];
432432
layerxX0[iLayer] = this->mTrkParams[iteration].LayerxX0[iLayer];
433433
}
434-
computeTrackExtensionCandidatesHandler<NLayers>(mTimeFrameGPU->getDeviceTrackExtensionStartStates(),
434+
computeTrackExtensionCandidatesHandler<NLayers>(mTimeFrameGPU->getDeviceTrackExtensionStartTracks(),
435435
mTimeFrameGPU->getDeviceIndexTableUtils(),
436436
mTimeFrameGPU->getDeviceROFMaskTableView(),
437437
mTimeFrameGPU->getDeviceROFOverlapTableView(),
@@ -461,7 +461,7 @@ void TrackerTraitsGPU<NLayers>::buildTrackExtensionCandidates(const int iteratio
461461
this->mTrkParams[iteration].CorrType,
462462
mTimeFrameGPU->getStream(0));
463463
mTimeFrameGPU->createTrackExtensionResultsDevice(nTracks);
464-
computeTrackExtensionResultsHandler<NLayers>(mTimeFrameGPU->getDeviceTrackExtensionStartStates(),
464+
computeTrackExtensionResultsHandler<NLayers>(mTimeFrameGPU->getDeviceTrackExtensionStartTracks(),
465465
mTimeFrameGPU->getDeviceTrackExtensionCandidates(),
466466
mTimeFrameGPU->getDeviceTrackExtensionCandidateOffsets(),
467467
mTimeFrameGPU->getDeviceTrackExtensionResults(),

0 commit comments

Comments
 (0)