diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index 88666cdfdb7fb..14aa0ea5d67fb 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -30,7 +30,7 @@ class DefaultGPUAllocator : public ExternalAllocator }; template -class TimeFrameGPU : public TimeFrame +class TimeFrameGPU : public TimeFrame { public: TimeFrameGPU(); @@ -205,14 +205,14 @@ class TimeFrameGPU : public TimeFrame template inline int TimeFrameGPU::getNClustersInRofSpan(const int rofIdstart, const int rofSpanSize, const int layerId) const { - return static_cast(mROFramesClusters[layerId][(rofIdstart + rofSpanSize) < mROFramesClusters.size() ? rofIdstart + rofSpanSize : mROFramesClusters.size() - 1] - mROFramesClusters[layerId][rofIdstart]); + return static_cast(this->mROFramesClusters[layerId][(rofIdstart + rofSpanSize) < this->mROFramesClusters.size() ? rofIdstart + rofSpanSize : this->mROFramesClusters.size() - 1] - this->mROFramesClusters[layerId][rofIdstart]); } template inline std::vector TimeFrameGPU::getClusterSizes() { - std::vector sizes(mUnsortedClusters.size()); - std::transform(mUnsortedClusters.begin(), mUnsortedClusters.end(), sizes.begin(), + std::vector sizes(this->mUnsortedClusters.size()); + std::transform(this->mUnsortedClusters.begin(), this->mUnsortedClusters.end(), sizes.begin(), [](const auto& v) { return static_cast(v.size()); }); return sizes; } diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h index c765307473749..1654f8cc8cf94 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h @@ -16,19 +16,17 @@ #include "ITStracking/TrackerTraits.h" #include "ITStrackingGPU/TimeFrameGPU.h" -namespace o2 -{ -namespace its +namespace o2::its { template -class TrackerTraitsGPU final : public TrackerTraits +class TrackerTraitsGPU final : public TrackerTraits { public: TrackerTraitsGPU() = default; ~TrackerTraitsGPU() override = default; - void adoptTimeFrame(TimeFrame* tf) final; + void adoptTimeFrame(TimeFrame* tf) final; void initialiseTimeFrame(const int iteration) final; void computeLayerTracklets(const int iteration, int, int) final; @@ -54,13 +52,6 @@ class TrackerTraitsGPU final : public TrackerTraits gpu::TimeFrameGPU<7>* mTimeFrameGPU; }; -template -inline void TrackerTraitsGPU::adoptTimeFrame(TimeFrame* tf) -{ - mTimeFrameGPU = static_cast*>(tf); - mTimeFrame = static_cast(tf); -} -} // namespace its -} // namespace o2 +} // namespace o2::its -#endif \ No newline at end of file +#endif diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 21b14fd9292d2..aabd3a940a532 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -73,8 +73,8 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, const float NSigmaCut, std::vector& phiCuts, const float resolutionPV, - std::vector& minR, - std::vector& maxR, + std::array& minR, + std::array& maxR, std::vector& resolutions, std::vector& radii, std::vector& mulScatAng, @@ -106,8 +106,8 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const float NSigmaCut, std::vector& phiCuts, const float resolutionPV, - std::vector& minR, - std::vector& maxR, + std::array& minR, + std::array& maxR, std::vector& resolutions, std::vector& radii, std::vector& mulScatAng, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/VertexerTraitsGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/VertexerTraitsGPU.h index eb4dc2179cdb4..a5c3709081a82 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/VertexerTraitsGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/VertexerTraitsGPU.h @@ -37,37 +37,30 @@ class ROframe; using constants::its2::InversePhiBinSize; -class VertexerTraitsGPU : public VertexerTraits +class VertexerTraitsGPU final : public VertexerTraits { public: - VertexerTraitsGPU(); - ~VertexerTraitsGPU() = default; - void initialise(const TrackingParameters&, const int iteration = 0) override; - void adoptTimeFrame(TimeFrame*) override; - void computeTracklets(const int iteration = 0) override; - void computeTrackletMatching(const int iteration = 0) override; - void computeVertices(const int iteration = 0) override; - void updateVertexingParameters(const std::vector&, const TimeFrameGPUParameters&) override; - - // Hybrid - void initialiseHybrid(const TrackingParameters& pars, const int iteration = 0) override { VertexerTraits::initialise(pars, iteration); } - void computeTrackletsHybrid(const int iteration = 0) override { VertexerTraits::computeTracklets(iteration); } - void computeTrackletMatchingHybrid(const int iteration = 0) override { VertexerTraits::computeTrackletMatching(iteration); } - void computeVerticesHybrid(const int iteration = 0) override { VertexerTraits::computeVertices(iteration); } - void adoptTimeFrameHybrid(TimeFrame* tf) override { VertexerTraits::adoptTimeFrame(tf); } - + void initialise(const TrackingParameters&, const int iteration = 0) final; + void adoptTimeFrame(TimeFrame<7>*) noexcept final; + void computeTracklets(const int iteration = 0) final; + void computeTrackletMatching(const int iteration = 0) final; + void computeVertices(const int iteration = 0) final; + void updateVertexingParameters(const std::vector&, const TimeFrameGPUParameters&) final; void computeVerticesHist(); + bool isGPU() const noexcept final { return true; } + const char* getName() const noexcept final { return "GPU"; } + protected: IndexTableUtils* mDeviceIndexTableUtils; gpu::TimeFrameGPU<7>* mTimeFrameGPU; TimeFrameGPUParameters mTfGPUParams; }; -inline void VertexerTraitsGPU::adoptTimeFrame(TimeFrame* tf) +inline void VertexerTraitsGPU::adoptTimeFrame(TimeFrame<7>* tf) noexcept { mTimeFrameGPU = static_cast*>(tf); - mTimeFrame = static_cast(tf); + mTimeFrame = static_cast*>(tf); } } // namespace its diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index ee99955ba6669..5a1e9148c0548 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -94,7 +94,7 @@ void* DefaultGPUAllocator::allocate(size_t size) template TimeFrameGPU::TimeFrameGPU() { - mIsGPU = true; + this->mIsGPU = true; } template @@ -104,7 +104,7 @@ template void TimeFrameGPU::allocMemAsync(void** ptr, size_t size, Stream* strPtr, bool extAllocator) { if (extAllocator) { - *ptr = mAllocator->allocate(size); + *ptr = this->mAllocator->allocate(size); } else { LOGP(debug, "Calling default CUDA allocator"); GPUChkErrS(cudaMallocAsync(reinterpret_cast(ptr), size, strPtr->get())); @@ -114,7 +114,7 @@ void TimeFrameGPU::allocMemAsync(void** ptr, size_t size, Stream* strPt template void TimeFrameGPU::setDevicePropagator(const o2::base::PropagatorImpl* propagator) { - mPropagatorDevice = propagator; + this->mPropagatorDevice = propagator; } template @@ -123,10 +123,10 @@ void TimeFrameGPU::loadIndexTableUtils(const int iteration) START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading indextable utils"); if (!iteration) { LOGP(debug, "gpu-allocation: allocating IndexTableUtils buffer, for {} MB.", sizeof(IndexTableUtils) / MB); - allocMemAsync(reinterpret_cast(&mIndexTableUtilsDevice), sizeof(IndexTableUtils), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mIndexTableUtilsDevice), sizeof(IndexTableUtils), nullptr, this->getExtAllocator()); } LOGP(debug, "gpu-transfer: loading IndexTableUtils object, for {} MB.", sizeof(IndexTableUtils) / MB); - GPUChkErrS(cudaMemcpyAsync(mIndexTableUtilsDevice, &mIndexTableUtils, sizeof(IndexTableUtils), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + GPUChkErrS(cudaMemcpyAsync(mIndexTableUtilsDevice, &(this->mIndexTableUtils), sizeof(IndexTableUtils), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -136,12 +136,12 @@ void TimeFrameGPU::loadUnsortedClustersDevice(const int iteration) if (!iteration) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading unsorted clusters"); for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - LOGP(debug, "gpu-transfer: loading {} unsorted clusters on layer {}, for {} MB.", mUnsortedClusters[iLayer].size(), iLayer, mUnsortedClusters[iLayer].size() * sizeof(Cluster) / MB); - allocMemAsync(reinterpret_cast(&mUnsortedClustersDevice[iLayer]), mUnsortedClusters[iLayer].size() * sizeof(Cluster), nullptr, getExtAllocator()); - GPUChkErrS(cudaHostRegister(mUnsortedClusters[iLayer].data(), mUnsortedClusters[iLayer].size() * sizeof(Cluster), cudaHostRegisterPortable)); - GPUChkErrS(cudaMemcpyAsync(mUnsortedClustersDevice[iLayer], mUnsortedClusters[iLayer].data(), mUnsortedClusters[iLayer].size() * sizeof(Cluster), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} unsorted clusters on layer {}, for {} MB.", this->mUnsortedClusters[iLayer].size(), iLayer, this->mUnsortedClusters[iLayer].size() * sizeof(Cluster) / MB); + allocMemAsync(reinterpret_cast(&mUnsortedClustersDevice[iLayer]), this->mUnsortedClusters[iLayer].size() * sizeof(Cluster), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaHostRegister(this->mUnsortedClusters[iLayer].data(), this->mUnsortedClusters[iLayer].size() * sizeof(Cluster), cudaHostRegisterPortable)); + GPUChkErrS(cudaMemcpyAsync(mUnsortedClustersDevice[iLayer], this->mUnsortedClusters[iLayer].data(), this->mUnsortedClusters[iLayer].size() * sizeof(Cluster), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mUnsortedClustersDeviceArray), nLayers * sizeof(Cluster*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mUnsortedClustersDeviceArray), nLayers * sizeof(Cluster*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaHostRegister(mUnsortedClustersDevice.data(), nLayers * sizeof(Cluster*), cudaHostRegisterPortable)); GPUChkErrS(cudaMemcpyAsync(mUnsortedClustersDeviceArray, mUnsortedClustersDevice.data(), nLayers * sizeof(Cluster*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); @@ -154,12 +154,12 @@ void TimeFrameGPU::loadClustersDevice(const int iteration) if (!iteration) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading sorted clusters"); for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - LOGP(debug, "gpu-transfer: loading {} clusters on layer {}, for {} MB.", mClusters[iLayer].size(), iLayer, mClusters[iLayer].size() * sizeof(Cluster) / MB); - allocMemAsync(reinterpret_cast(&mClustersDevice[iLayer]), mClusters[iLayer].size() * sizeof(Cluster), nullptr, getExtAllocator()); - GPUChkErrS(cudaHostRegister(mClusters[iLayer].data(), mClusters[iLayer].size() * sizeof(Cluster), cudaHostRegisterPortable)); - GPUChkErrS(cudaMemcpyAsync(mClustersDevice[iLayer], mClusters[iLayer].data(), mClusters[iLayer].size() * sizeof(Cluster), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} clusters on layer {}, for {} MB.", this->mClusters[iLayer].size(), iLayer, this->mClusters[iLayer].size() * sizeof(Cluster) / MB); + allocMemAsync(reinterpret_cast(&mClustersDevice[iLayer]), this->mClusters[iLayer].size() * sizeof(Cluster), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaHostRegister(this->mClusters[iLayer].data(), this->mClusters[iLayer].size() * sizeof(Cluster), cudaHostRegisterPortable)); + GPUChkErrS(cudaMemcpyAsync(mClustersDevice[iLayer], this->mClusters[iLayer].data(), this->mClusters[iLayer].size() * sizeof(Cluster), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mClustersDeviceArray), nLayers * sizeof(Cluster*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mClustersDeviceArray), nLayers * sizeof(Cluster*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaHostRegister(mClustersDevice.data(), nLayers * sizeof(Cluster*), cudaHostRegisterPortable)); GPUChkErrS(cudaMemcpyAsync(mClustersDeviceArray, mClustersDevice.data(), nLayers * sizeof(Cluster*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); @@ -172,11 +172,11 @@ void TimeFrameGPU::loadClustersIndexTables(const int iteration) if (!iteration) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading sorted clusters"); for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - LOGP(debug, "gpu-transfer: loading clusters indextable for layer {} with {} elements, for {} MB.", iLayer, mIndexTables[iLayer].size(), mIndexTables[iLayer].size() * sizeof(int) / MB); - allocMemAsync(reinterpret_cast(&mClustersIndexTablesDevice[iLayer]), mIndexTables[iLayer].size() * sizeof(int), nullptr, getExtAllocator()); - GPUChkErrS(cudaMemcpyAsync(mClustersIndexTablesDevice[iLayer], mIndexTables[iLayer].data(), mIndexTables[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading clusters indextable for layer {} with {} elements, for {} MB.", iLayer, this->mIndexTables[iLayer].size(), this->mIndexTables[iLayer].size() * sizeof(int) / MB); + allocMemAsync(reinterpret_cast(&mClustersIndexTablesDevice[iLayer]), this->mIndexTables[iLayer].size() * sizeof(int), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaMemcpyAsync(mClustersIndexTablesDevice[iLayer], this->mIndexTables[iLayer].data(), this->mIndexTables[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mClustersIndexTablesDeviceArray), nLayers * sizeof(int), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mClustersIndexTablesDeviceArray), nLayers * sizeof(int), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemcpyAsync(mClustersIndexTablesDeviceArray, mClustersIndexTablesDevice.data(), nLayers * sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -188,11 +188,11 @@ void TimeFrameGPU::createUsedClustersDevice(const int iteration) if (!iteration) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "creating used clusters flags"); for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - LOGP(debug, "gpu-transfer: creating {} used clusters flags on layer {}, for {} MB.", mUsedClusters[iLayer].size(), iLayer, mUsedClusters[iLayer].size() * sizeof(unsigned char) / MB); - allocMemAsync(reinterpret_cast(&mUsedClustersDevice[iLayer]), mUsedClusters[iLayer].size() * sizeof(unsigned char), nullptr, getExtAllocator()); - GPUChkErrS(cudaMemsetAsync(mUsedClustersDevice[iLayer], 0, mUsedClusters[iLayer].size() * sizeof(unsigned char), mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: creating {} used clusters flags on layer {}, for {} MB.", this->mUsedClusters[iLayer].size(), iLayer, this->mUsedClusters[iLayer].size() * sizeof(unsigned char) / MB); + allocMemAsync(reinterpret_cast(&mUsedClustersDevice[iLayer]), this->mUsedClusters[iLayer].size() * sizeof(unsigned char), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaMemsetAsync(mUsedClustersDevice[iLayer], 0, this->mUsedClusters[iLayer].size() * sizeof(unsigned char), mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mUsedClustersDeviceArray), nLayers * sizeof(unsigned char*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mUsedClustersDeviceArray), nLayers * sizeof(unsigned char*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemcpyAsync(mUsedClustersDeviceArray, mUsedClustersDevice.data(), nLayers * sizeof(unsigned char*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -203,8 +203,8 @@ void TimeFrameGPU::loadUsedClustersDevice() { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading used clusters flags"); for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - LOGP(debug, "gpu-transfer: loading {} used clusters flags on layer {}, for {} MB.", mUsedClusters[iLayer].size(), iLayer, mClusters[iLayer].size() * sizeof(unsigned char) / MB); - GPUChkErrS(cudaMemcpyAsync(mUsedClustersDevice[iLayer], mUsedClusters[iLayer].data(), mUsedClusters[iLayer].size() * sizeof(unsigned char), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} used clusters flags on layer {}, for {} MB.", this->mUsedClusters[iLayer].size(), iLayer, this->mClusters[iLayer].size() * sizeof(unsigned char) / MB); + GPUChkErrS(cudaMemcpyAsync(mUsedClustersDevice[iLayer], this->mUsedClusters[iLayer].data(), this->mUsedClusters[iLayer].size() * sizeof(unsigned char), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -215,11 +215,11 @@ void TimeFrameGPU::loadROframeClustersDevice(const int iteration) if (!iteration) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading ROframe clusters"); for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - LOGP(debug, "gpu-transfer: loading {} ROframe clusters info on layer {}, for {} MB.", mROFramesClusters[iLayer].size(), iLayer, mROFramesClusters[iLayer].size() * sizeof(int) / MB); - allocMemAsync(reinterpret_cast(&mROFramesClustersDevice[iLayer]), mROFramesClusters[iLayer].size() * sizeof(int), nullptr, getExtAllocator()); - GPUChkErrS(cudaMemcpyAsync(mROFramesClustersDevice[iLayer], mROFramesClusters[iLayer].data(), mROFramesClusters[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} ROframe clusters info on layer {}, for {} MB.", this->mROFramesClusters[iLayer].size(), iLayer, this->mROFramesClusters[iLayer].size() * sizeof(int) / MB); + allocMemAsync(reinterpret_cast(&mROFramesClustersDevice[iLayer]), this->mROFramesClusters[iLayer].size() * sizeof(int), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaMemcpyAsync(mROFramesClustersDevice[iLayer], this->mROFramesClusters[iLayer].data(), this->mROFramesClusters[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mROFrameClustersDeviceArray), nLayers * sizeof(int*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mROFrameClustersDeviceArray), nLayers * sizeof(int*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemcpyAsync(mROFrameClustersDeviceArray, mROFramesClustersDevice.data(), nLayers * sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -231,12 +231,12 @@ void TimeFrameGPU::loadTrackingFrameInfoDevice(const int iteration) START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading trackingframeinfo"); if (!iteration) { for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - LOGP(debug, "gpu-transfer: loading {} tfinfo on layer {}, for {} MB.", mTrackingFrameInfo[iLayer].size(), iLayer, mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo) / MB); - allocMemAsync(reinterpret_cast(&mTrackingFrameInfoDevice[iLayer]), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), nullptr, getExtAllocator()); - GPUChkErrS(cudaHostRegister(mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable)); - GPUChkErrS(cudaMemcpyAsync(mTrackingFrameInfoDevice[iLayer], mTrackingFrameInfo[iLayer].data(), mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} tfinfo on layer {}, for {} MB.", this->mTrackingFrameInfo[iLayer].size(), iLayer, this->mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo) / MB); + allocMemAsync(reinterpret_cast(&mTrackingFrameInfoDevice[iLayer]), this->mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaHostRegister(this->mTrackingFrameInfo[iLayer].data(), this->mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaHostRegisterPortable)); + GPUChkErrS(cudaMemcpyAsync(mTrackingFrameInfoDevice[iLayer], this->mTrackingFrameInfo[iLayer].data(), this->mTrackingFrameInfo[iLayer].size() * sizeof(TrackingFrameInfo), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mTrackingFrameInfoDeviceArray), nLayers * sizeof(TrackingFrameInfo*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mTrackingFrameInfoDeviceArray), nLayers * sizeof(TrackingFrameInfo*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaHostRegister(mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaHostRegisterPortable)); GPUChkErrS(cudaMemcpyAsync(mTrackingFrameInfoDeviceArray, mTrackingFrameInfoDevice.data(), nLayers * sizeof(TrackingFrameInfo*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } @@ -248,9 +248,9 @@ void TimeFrameGPU::loadMultiplicityCutMask(const int iteration) { if (!iteration) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading multiplicity cut mask"); - LOGP(debug, "gpu-transfer: loading multiplicity cut mask with {} elements, for {} MB.", mMultiplicityCutMask.size(), mMultiplicityCutMask.size() * sizeof(bool) / MB); - allocMemAsync(reinterpret_cast(&mMultMaskDevice), mMultiplicityCutMask.size() * sizeof(uint8_t), nullptr, getExtAllocator()); - GPUChkErrS(cudaMemcpyAsync(mMultMaskDevice, mMultiplicityCutMask.data(), mMultiplicityCutMask.size() * sizeof(uint8_t), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading multiplicity cut mask with {} elements, for {} MB.", this->mMultiplicityCutMask.size(), this->mMultiplicityCutMask.size() * sizeof(bool) / MB); + allocMemAsync(reinterpret_cast(&mMultMaskDevice), this->mMultiplicityCutMask.size() * sizeof(uint8_t), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaMemcpyAsync(mMultMaskDevice, this->mMultiplicityCutMask.data(), this->mMultiplicityCutMask.size() * sizeof(uint8_t), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } } @@ -260,12 +260,12 @@ void TimeFrameGPU::loadVertices(const int iteration) { if (!iteration) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading seeding vertices"); - LOGP(debug, "gpu-transfer: loading {} ROframes vertices, for {} MB.", mROFramesPV.size(), mROFramesPV.size() * sizeof(int) / MB); - allocMemAsync(reinterpret_cast(&mROFramesPVDevice), mROFramesPV.size() * sizeof(int), nullptr, getExtAllocator()); - GPUChkErrS(cudaMemcpyAsync(mROFramesPVDevice, mROFramesPV.data(), mROFramesPV.size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); - LOGP(debug, "gpu-transfer: loading {} seeding vertices, for {} MB.", mPrimaryVertices.size(), mPrimaryVertices.size() * sizeof(Vertex) / MB); - allocMemAsync(reinterpret_cast(&mPrimaryVerticesDevice), mPrimaryVertices.size() * sizeof(Vertex), nullptr, getExtAllocator()); - GPUChkErrS(cudaMemcpyAsync(mPrimaryVerticesDevice, mPrimaryVertices.data(), mPrimaryVertices.size() * sizeof(Vertex), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} ROframes vertices, for {} MB.", this->mROFramesPV.size(), this->mROFramesPV.size() * sizeof(int) / MB); + allocMemAsync(reinterpret_cast(&mROFramesPVDevice), this->mROFramesPV.size() * sizeof(int), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaMemcpyAsync(mROFramesPVDevice, this->mROFramesPV.data(), this->mROFramesPV.size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} seeding vertices, for {} MB.", this->mPrimaryVertices.size(), this->mPrimaryVertices.size() * sizeof(Vertex) / MB); + allocMemAsync(reinterpret_cast(&mPrimaryVerticesDevice), this->mPrimaryVertices.size() * sizeof(Vertex), nullptr, this->getExtAllocator()); + GPUChkErrS(cudaMemcpyAsync(mPrimaryVerticesDevice, this->mPrimaryVertices.data(), this->mPrimaryVertices.size() * sizeof(Vertex), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } } @@ -276,13 +276,13 @@ void TimeFrameGPU::createTrackletsLUTDevice(const int iteration) START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "creating tracklets LUTs"); for (auto iLayer{0}; iLayer < nLayers - 1; ++iLayer) { if (!iteration) { - LOGP(debug, "gpu-transfer: creating tracklets LUT for {} elements on layer {}, for {} MB.", mClusters[iLayer].size() + 1, iLayer, (mClusters[iLayer].size() + 1) * sizeof(int) / MB); - allocMemAsync(reinterpret_cast(&mTrackletsLUTDevice[iLayer]), (mClusters[iLayer].size() + 1) * sizeof(int), nullptr, getExtAllocator()); + LOGP(debug, "gpu-transfer: creating tracklets LUT for {} elements on layer {}, for {} MB.", this->mClusters[iLayer].size() + 1, iLayer, (this->mClusters[iLayer].size() + 1) * sizeof(int) / MB); + allocMemAsync(reinterpret_cast(&mTrackletsLUTDevice[iLayer]), (this->mClusters[iLayer].size() + 1) * sizeof(int), nullptr, this->getExtAllocator()); } - GPUChkErrS(cudaMemsetAsync(mTrackletsLUTDevice[iLayer], 0, (mClusters[iLayer].size() + 1) * sizeof(int), mGpuStreams[0]->get())); + GPUChkErrS(cudaMemsetAsync(mTrackletsLUTDevice[iLayer], 0, (this->mClusters[iLayer].size() + 1) * sizeof(int), mGpuStreams[0]->get())); } if (!iteration) { - allocMemAsync(reinterpret_cast(&mTrackletsLUTDeviceArray), (nLayers - 1) * sizeof(int*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mTrackletsLUTDeviceArray), (nLayers - 1) * sizeof(int*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemcpyAsync(mTrackletsLUTDeviceArray, mTrackletsLUTDevice.data(), mTrackletsLUTDevice.size() * sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); @@ -294,11 +294,11 @@ void TimeFrameGPU::createTrackletsBuffers() START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "creating cells buffers"); for (auto iLayer{0}; iLayer < nLayers - 1; ++iLayer) { mNTracklets[iLayer] = 0; - GPUChkErrS(cudaMemcpyAsync(&mNTracklets[iLayer], mTrackletsLUTDevice[iLayer] + mClusters[iLayer].size(), sizeof(int), cudaMemcpyDeviceToHost)); + GPUChkErrS(cudaMemcpyAsync(&mNTracklets[iLayer], mTrackletsLUTDevice[iLayer] + this->mClusters[iLayer].size(), sizeof(int), cudaMemcpyDeviceToHost)); LOGP(debug, "gpu-transfer: creating tracklets buffer for {} elements on layer {}, for {} MB.", mNTracklets[iLayer], iLayer, mNTracklets[iLayer] * sizeof(Tracklet) / MB); - allocMemAsync(reinterpret_cast(&mTrackletsDevice[iLayer]), mNTracklets[iLayer] * sizeof(Tracklet), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mTrackletsDevice[iLayer]), mNTracklets[iLayer] * sizeof(Tracklet), nullptr, this->getExtAllocator()); } - allocMemAsync(reinterpret_cast(&mTrackletsDeviceArray), (nLayers - 1) * sizeof(Tracklet*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mTrackletsDeviceArray), (nLayers - 1) * sizeof(Tracklet*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaHostRegister(mTrackletsDevice.data(), (nLayers - 1) * sizeof(Tracklet*), cudaHostRegisterPortable)); GPUChkErrS(cudaMemcpyAsync(mTrackletsDeviceArray, mTrackletsDevice.data(), (nLayers - 1) * sizeof(Tracklet*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); @@ -309,9 +309,9 @@ void TimeFrameGPU::loadTrackletsDevice() { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading tracklets"); for (auto iLayer{0}; iLayer < nLayers - 1; ++iLayer) { - LOGP(debug, "gpu-transfer: loading {} tracklets on layer {}, for {} MB.", mTracklets[iLayer].size(), iLayer, mTracklets[iLayer].size() * sizeof(Tracklet) / MB); - GPUChkErrS(cudaHostRegister(mTracklets[iLayer].data(), mTracklets[iLayer].size() * sizeof(Tracklet), cudaHostRegisterPortable)); - GPUChkErrS(cudaMemcpyAsync(mTrackletsDevice[iLayer], mTracklets[iLayer].data(), mTracklets[iLayer].size() * sizeof(Tracklet), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} tracklets on layer {}, for {} MB.", this->mTracklets[iLayer].size(), iLayer, this->mTracklets[iLayer].size() * sizeof(Tracklet) / MB); + GPUChkErrS(cudaHostRegister(this->mTracklets[iLayer].data(), this->mTracklets[iLayer].size() * sizeof(Tracklet), cudaHostRegisterPortable)); + GPUChkErrS(cudaMemcpyAsync(mTrackletsDevice[iLayer], this->mTracklets[iLayer].data(), this->mTracklets[iLayer].size() * sizeof(Tracklet), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -321,9 +321,9 @@ void TimeFrameGPU::loadTrackletsLUTDevice() { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading tracklets"); for (auto iLayer{0}; iLayer < nLayers - 2; ++iLayer) { - LOGP(debug, "gpu-transfer: loading tracklets LUT for {} elements on layer {}, for {} MB", mTrackletsLookupTable[iLayer].size(), iLayer + 1, mTrackletsLookupTable[iLayer].size() * sizeof(int) / MB); - GPUChkErrS(cudaHostRegister(mTrackletsLookupTable[iLayer].data(), mTrackletsLookupTable[iLayer].size() * sizeof(int), cudaHostRegisterPortable)); - GPUChkErrS(cudaMemcpyAsync(mTrackletsLUTDevice[iLayer + 1], mTrackletsLookupTable[iLayer].data(), mTrackletsLookupTable[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice)); + LOGP(debug, "gpu-transfer: loading tracklets LUT for {} elements on layer {}, for {} MB", this->mTrackletsLookupTable[iLayer].size(), iLayer + 1, this->mTrackletsLookupTable[iLayer].size() * sizeof(int) / MB); + GPUChkErrS(cudaHostRegister(this->mTrackletsLookupTable[iLayer].data(), this->mTrackletsLookupTable[iLayer].size() * sizeof(int), cudaHostRegisterPortable)); + GPUChkErrS(cudaMemcpyAsync(mTrackletsLUTDevice[iLayer + 1], this->mTrackletsLookupTable[iLayer].data(), this->mTrackletsLookupTable[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice)); } GPUChkErrS(cudaHostRegister(mTrackletsLUTDevice.data(), (nLayers - 1) * sizeof(int*), cudaHostRegisterPortable)); GPUChkErrS(cudaMemcpyAsync(mTrackletsLUTDeviceArray, mTrackletsLUTDevice.data(), (nLayers - 1) * sizeof(int*), cudaMemcpyHostToDevice)); @@ -335,12 +335,12 @@ void TimeFrameGPU::createNeighboursIndexTablesDevice() { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "creating cells neighbours"); // Here we do also the creation of the CellsDeviceArray, as the cells buffers are populated separately in the previous steps. - allocMemAsync(reinterpret_cast(&mCellsDeviceArray), (nLayers - 2) * sizeof(CellSeed*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mCellsDeviceArray), (nLayers - 2) * sizeof(CellSeed*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaHostRegister(mCellsDevice.data(), (nLayers - 2) * sizeof(CellSeed*), cudaHostRegisterPortable)); GPUChkErrS(cudaMemcpyAsync(mCellsDeviceArray, mCellsDevice.data(), (nLayers - 2) * sizeof(CellSeed*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); for (auto iLayer{0}; iLayer < nLayers - 2; ++iLayer) { LOGP(debug, "gpu-transfer: loading neighbours LUT for {} elements on layer {}, for {} MB.", mNCells[iLayer], iLayer, mNCells[iLayer] * sizeof(CellSeed) / MB); - allocMemAsync(reinterpret_cast(&mNeighboursIndexTablesDevice[iLayer]), (mNCells[iLayer] + 1) * sizeof(int), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mNeighboursIndexTablesDevice[iLayer]), (mNCells[iLayer] + 1) * sizeof(int), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemsetAsync(mNeighboursIndexTablesDevice[iLayer], 0, (mNCells[iLayer] + 1) * sizeof(int), mGpuStreams[0]->get())); } STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); @@ -351,7 +351,7 @@ void TimeFrameGPU::createNeighboursLUTDevice(const int layer, const uns { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "reserving neighboursLUT"); LOGP(debug, "gpu-allocation: reserving neighbours LUT for {} elements on layer {} , for {} MB.", nCells + 1, layer, (nCells + 1) * sizeof(int) / MB); - allocMemAsync(reinterpret_cast(&mNeighboursLUTDevice[layer]), (nCells + 1) * sizeof(int), nullptr, getExtAllocator()); // We need one element more to move exc -> inc + allocMemAsync(reinterpret_cast(&mNeighboursLUTDevice[layer]), (nCells + 1) * sizeof(int), nullptr, this->getExtAllocator()); // We need one element more to move exc -> inc GPUChkErrS(cudaMemsetAsync(mNeighboursLUTDevice[layer], 0, (nCells + 1) * sizeof(int), mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -361,13 +361,13 @@ void TimeFrameGPU::loadCellsDevice() { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading cell seeds"); for (auto iLayer{0}; iLayer < nLayers - 2; ++iLayer) { - LOGP(debug, "gpu-transfer: loading {} cell seeds on layer {}, for {} MB.", mCells[iLayer].size(), iLayer, mCells[iLayer].size() * sizeof(CellSeed) / MB); - allocMemAsync(reinterpret_cast(&mCellsDevice[iLayer]), mCells[iLayer].size() * sizeof(CellSeed), nullptr, getExtAllocator()); - allocMemAsync(reinterpret_cast(&mNeighboursIndexTablesDevice[iLayer]), (mCells[iLayer].size() + 1) * sizeof(int), nullptr, getExtAllocator()); // accessory for the neigh. finding. - GPUChkErrS(cudaMemsetAsync(mNeighboursIndexTablesDevice[iLayer], 0, (mCells[iLayer].size() + 1) * sizeof(int), mGpuStreams[0]->get())); - GPUChkErrS(cudaMemcpyAsync(mCellsDevice[iLayer], mCells[iLayer].data(), mCells[iLayer].size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} cell seeds on layer {}, for {} MB.", this->mCells[iLayer].size(), iLayer, this->mCells[iLayer].size() * sizeof(CellSeed) / MB); + allocMemAsync(reinterpret_cast(&mCellsDevice[iLayer]), this->mCells[iLayer].size() * sizeof(CellSeed), nullptr, this->getExtAllocator()); + allocMemAsync(reinterpret_cast(&mNeighboursIndexTablesDevice[iLayer]), (this->mCells[iLayer].size() + 1) * sizeof(int), nullptr, this->getExtAllocator()); // accessory for the neigh. finding. + GPUChkErrS(cudaMemsetAsync(mNeighboursIndexTablesDevice[iLayer], 0, (this->mCells[iLayer].size() + 1) * sizeof(int), mGpuStreams[0]->get())); + GPUChkErrS(cudaMemcpyAsync(mCellsDevice[iLayer], this->mCells[iLayer].data(), this->mCells[iLayer].size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mCellsDeviceArray), (nLayers - 2) * sizeof(CellSeed*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mCellsDeviceArray), (nLayers - 2) * sizeof(CellSeed*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemcpyAsync(mCellsDeviceArray, mCellsDevice.data(), (nLayers - 2) * sizeof(CellSeed*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -378,10 +378,10 @@ void TimeFrameGPU::createCellsLUTDevice() START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "creating cells LUTs"); for (auto iLayer{0}; iLayer < nLayers - 2; ++iLayer) { LOGP(debug, "gpu-transfer: creating cell LUT for {} elements on layer {}, for {} MB.", mNTracklets[iLayer] + 1, iLayer, (mNTracklets[iLayer] + 1) * sizeof(int) / MB); - allocMemAsync(reinterpret_cast(&mCellsLUTDevice[iLayer]), (mNTracklets[iLayer] + 1) * sizeof(int), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mCellsLUTDevice[iLayer]), (mNTracklets[iLayer] + 1) * sizeof(int), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemsetAsync(mCellsLUTDevice[iLayer], 0, (mNTracklets[iLayer] + 1) * sizeof(int), mGpuStreams[0]->get())); } - allocMemAsync(reinterpret_cast(&mCellsLUTDeviceArray), (nLayers - 2) * sizeof(int*), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mCellsLUTDeviceArray), (nLayers - 2) * sizeof(int*), nullptr, this->getExtAllocator()); GPUChkErrS(cudaMemcpyAsync(mCellsLUTDeviceArray, mCellsLUTDevice.data(), mCellsLUTDevice.size() * sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -393,7 +393,7 @@ void TimeFrameGPU::createCellsBuffers(const int layer) mNCells[layer] = 0; GPUChkErrS(cudaMemcpyAsync(&mNCells[layer], mCellsLUTDevice[layer] + mNTracklets[layer], sizeof(int), cudaMemcpyDeviceToHost)); LOGP(debug, "gpu-transfer: creating cell buffer for {} elements on layer {}, for {} MB.", mNCells[layer], layer, mNCells[layer] * sizeof(CellSeed) / MB); - allocMemAsync(reinterpret_cast(&mCellsDevice[layer]), mNCells[layer] * sizeof(CellSeed), nullptr, getExtAllocator()); + allocMemAsync(reinterpret_cast(&mCellsDevice[layer]), mNCells[layer] * sizeof(CellSeed), nullptr, this->getExtAllocator()); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -403,9 +403,9 @@ void TimeFrameGPU::loadCellsLUTDevice() { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading cells LUTs"); for (auto iLayer{0}; iLayer < nLayers - 3; ++iLayer) { - LOGP(debug, "gpu-transfer: loading cell LUT for {} elements on layer {}, for {} MB.", mCellsLookupTable[iLayer].size(), iLayer, mCellsLookupTable[iLayer].size() * sizeof(int) / MB); - GPUChkErrS(cudaHostRegister(mCellsLookupTable[iLayer].data(), mCellsLookupTable[iLayer].size() * sizeof(int), cudaHostRegisterPortable)); - GPUChkErrS(cudaMemcpyAsync(mCellsLUTDevice[iLayer + 1], mCellsLookupTable[iLayer].data(), mCellsLookupTable[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading cell LUT for {} elements on layer {}, for {} MB.", this->mCellsLookupTable[iLayer].size(), iLayer, this->mCellsLookupTable[iLayer].size() * sizeof(int) / MB); + GPUChkErrS(cudaHostRegister(this->mCellsLookupTable[iLayer].data(), this->mCellsLookupTable[iLayer].size() * sizeof(int), cudaHostRegisterPortable)); + GPUChkErrS(cudaMemcpyAsync(mCellsLUTDevice[iLayer + 1], this->mCellsLookupTable[iLayer].data(), this->mCellsLookupTable[iLayer].size() * sizeof(int), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -413,10 +413,10 @@ void TimeFrameGPU::loadCellsLUTDevice() template void TimeFrameGPU::loadRoadsDevice() { - LOGP(debug, "gpu-transfer: loading {} roads, for {} MB.", mRoads.size(), mRoads.size() * sizeof(Road) / MB); - allocMemAsync(reinterpret_cast(&mRoadsDevice), mRoads.size() * sizeof(Road), mGpuStreams[0], getExtAllocator()); - GPUChkErrS(cudaHostRegister(mRoads.data(), mRoads.size() * sizeof(Road), cudaHostRegisterPortable)); - GPUChkErrS(cudaMemcpyAsync(mRoadsDevice, mRoads.data(), mRoads.size() * sizeof(Road), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); + LOGP(debug, "gpu-transfer: loading {} roads, for {} MB.", this->mRoads.size(), this->mRoads.size() * sizeof(Road) / MB); + allocMemAsync(reinterpret_cast(&mRoadsDevice), this->mRoads.size() * sizeof(Road), mGpuStreams[0], this->getExtAllocator()); + GPUChkErrS(cudaHostRegister(this->mRoads.data(), this->mRoads.size() * sizeof(Road), cudaHostRegisterPortable)); + GPUChkErrS(cudaMemcpyAsync(mRoadsDevice, this->mRoads.data(), this->mRoads.size() * sizeof(Road), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); } template @@ -424,7 +424,7 @@ void TimeFrameGPU::loadTrackSeedsDevice(std::vector& seeds) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "loading track seeds"); LOGP(debug, "gpu-transfer: loading {} track seeds, for {} MB.", seeds.size(), seeds.size() * sizeof(CellSeed) / MB); - allocMemAsync(reinterpret_cast(&mTrackSeedsDevice), seeds.size() * sizeof(CellSeed), mGpuStreams[0], getExtAllocator()); + allocMemAsync(reinterpret_cast(&mTrackSeedsDevice), seeds.size() * sizeof(CellSeed), mGpuStreams[0], this->getExtAllocator()); GPUChkErrS(cudaHostRegister(seeds.data(), seeds.size() * sizeof(CellSeed), cudaHostRegisterPortable)); GPUChkErrS(cudaMemcpyAsync(mTrackSeedsDevice, seeds.data(), seeds.size() * sizeof(CellSeed), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); @@ -435,10 +435,10 @@ void TimeFrameGPU::createNeighboursDevice(const unsigned int layer, con { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "reserving neighbours"); LOGP(debug, "gpu-allocation: reserving {} neighbours (pairs), for {} MB.", nNeighbours, nNeighbours * sizeof(gpuPair) / MB); - allocMemAsync(reinterpret_cast(&mNeighbourPairsDevice[layer]), nNeighbours * sizeof(gpuPair), mGpuStreams[0], getExtAllocator()); + allocMemAsync(reinterpret_cast(&mNeighbourPairsDevice[layer]), nNeighbours * sizeof(gpuPair), mGpuStreams[0], this->getExtAllocator()); GPUChkErrS(cudaMemsetAsync(mNeighbourPairsDevice[layer], -1, nNeighbours * sizeof(gpuPair), mGpuStreams[0]->get())); LOGP(debug, "gpu-allocation: reserving {} neighbours, for {} MB.", nNeighbours, nNeighbours * sizeof(gpuPair) / MB); - allocMemAsync(reinterpret_cast(&mNeighboursDevice[layer]), nNeighbours * sizeof(int), mGpuStreams[0], getExtAllocator()); + allocMemAsync(reinterpret_cast(&mNeighboursDevice[layer]), nNeighbours * sizeof(int), mGpuStreams[0], this->getExtAllocator()); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -446,13 +446,13 @@ template void TimeFrameGPU::createNeighboursDevice(const unsigned int layer, std::vector>& neighbours) { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "reserving neighbours"); - mCellsNeighbours[layer].clear(); - mCellsNeighbours[layer].resize(neighbours.size()); + this->mCellsNeighbours[layer].clear(); + this->mCellsNeighbours[layer].resize(neighbours.size()); LOGP(debug, "gpu-allocation: reserving {} neighbours (pairs), for {} MB.", neighbours.size(), neighbours.size() * sizeof(gpuPair) / MB); - allocMemAsync(reinterpret_cast(&mNeighbourPairsDevice[layer]), neighbours.size() * sizeof(gpuPair), mGpuStreams[0], getExtAllocator()); + allocMemAsync(reinterpret_cast(&mNeighbourPairsDevice[layer]), neighbours.size() * sizeof(gpuPair), mGpuStreams[0], this->getExtAllocator()); GPUChkErrS(cudaMemsetAsync(mNeighbourPairsDevice[layer], -1, neighbours.size() * sizeof(gpuPair), mGpuStreams[0]->get())); LOGP(debug, "gpu-allocation: reserving {} neighbours, for {} MB.", neighbours.size(), neighbours.size() * sizeof(gpuPair) / MB); - allocMemAsync(reinterpret_cast(&mNeighboursDevice[layer]), neighbours.size() * sizeof(int), mGpuStreams[0], getExtAllocator()); + allocMemAsync(reinterpret_cast(&mNeighboursDevice[layer]), neighbours.size() * sizeof(int), mGpuStreams[0], this->getExtAllocator()); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -460,7 +460,7 @@ template void TimeFrameGPU::createNeighboursDeviceArray() { START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "reserving neighbours"); - allocMemAsync(reinterpret_cast(&mNeighboursDeviceArray), (nLayers - 2) * sizeof(int*), mGpuStreams[0], getExtAllocator()); + allocMemAsync(reinterpret_cast(&mNeighboursDeviceArray), (nLayers - 2) * sizeof(int*), mGpuStreams[0], this->getExtAllocator()); GPUChkErrS(cudaMemcpyAsync(mNeighboursDeviceArray, mNeighboursDevice.data(), (nLayers - 2) * sizeof(int*), cudaMemcpyHostToDevice, mGpuStreams[0]->get())); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -472,7 +472,7 @@ void TimeFrameGPU::createTrackITSExtDevice(std::vector& seeds mTrackITSExt.clear(); mTrackITSExt.resize(seeds.size()); LOGP(debug, "gpu-allocation: reserving {} tracks, for {} MB.", seeds.size(), seeds.size() * sizeof(o2::its::TrackITSExt) / MB); - allocMemAsync(reinterpret_cast(&mTrackITSExtDevice), seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0], getExtAllocator()); + allocMemAsync(reinterpret_cast(&mTrackITSExtDevice), seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0], this->getExtAllocator()); GPUChkErrS(cudaMemsetAsync(mTrackITSExtDevice, 0, seeds.size() * sizeof(o2::its::TrackITSExt), mGpuStreams[0]->get())); GPUChkErrS(cudaHostRegister(mTrackITSExt.data(), seeds.size() * sizeof(o2::its::TrackITSExt), cudaHostRegisterPortable)); STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); @@ -484,8 +484,8 @@ void TimeFrameGPU::downloadCellsDevice() START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "downloading cells"); for (int iLayer{0}; iLayer < nLayers - 2; ++iLayer) { LOGP(debug, "gpu-transfer: downloading {} cells on layer: {}, for {} MB.", mNCells[iLayer], iLayer, mNCells[iLayer] * sizeof(CellSeed) / MB); - mCells[iLayer].resize(mNCells[iLayer]); - GPUChkErrS(cudaMemcpyAsync(mCells[iLayer].data(), mCellsDevice[iLayer], mNCells[iLayer] * sizeof(CellSeed), cudaMemcpyDeviceToHost, mGpuStreams[0]->get())); + this->mCells[iLayer].resize(mNCells[iLayer]); + GPUChkErrS(cudaMemcpyAsync(this->mCells[iLayer].data(), this->mCellsDevice[iLayer], mNCells[iLayer] * sizeof(CellSeed), cudaMemcpyDeviceToHost, mGpuStreams[0]->get())); } STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -496,8 +496,8 @@ void TimeFrameGPU::downloadCellsLUTDevice() START_GPU_STREAM_TIMER(mGpuStreams[0]->get(), "downloading cell luts"); for (auto iLayer{0}; iLayer < nLayers - 3; ++iLayer) { LOGP(debug, "gpu-transfer: downloading cells lut on layer {} for {} elements", iLayer, (mNTracklets[iLayer + 1] + 1)); - mCellsLookupTable[iLayer].resize(mNTracklets[iLayer + 1] + 1); - GPUChkErrS(cudaMemcpyAsync(mCellsLookupTable[iLayer].data(), mCellsLUTDevice[iLayer + 1], (mNTracklets[iLayer + 1] + 1) * sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[0]->get())); + this->mCellsLookupTable[iLayer].resize(mNTracklets[iLayer + 1] + 1); + GPUChkErrS(cudaMemcpyAsync(this->mCellsLookupTable[iLayer].data(), mCellsLUTDevice[iLayer + 1], (mNTracklets[iLayer + 1] + 1) * sizeof(int), cudaMemcpyDeviceToHost, mGpuStreams[0]->get())); } STOP_GPU_STREAM_TIMER(mGpuStreams[0]->get()); } @@ -545,9 +545,9 @@ template void TimeFrameGPU::unregisterHostMemory(const int maxLayers) { for (auto iLayer{0}; iLayer < nLayers; ++iLayer) { - GPUChkErrS(cudaHostUnregister(mUnsortedClusters[iLayer].data())); - GPUChkErrS(cudaHostUnregister(mClusters[iLayer].data())); - GPUChkErrS(cudaHostUnregister(mTrackingFrameInfo[iLayer].data())); + GPUChkErrS(cudaHostUnregister(this->mUnsortedClusters[iLayer].data())); + GPUChkErrS(cudaHostUnregister(this->mClusters[iLayer].data())); + GPUChkErrS(cudaHostUnregister(this->mTrackingFrameInfo[iLayer].data())); } GPUChkErrS(cudaHostUnregister(mTrackingFrameInfoDevice.data())); GPUChkErrS(cudaHostUnregister(mUnsortedClustersDevice.data())); @@ -565,7 +565,7 @@ void TimeFrameGPU::initialise(const int iteration, for (auto& str : mGpuStreams) { str = new Stream(); } - o2::its::TimeFrame::initialise(iteration, trkParam, maxLayers); + o2::its::TimeFrame::initialise(iteration, trkParam, maxLayers); } template class TimeFrameGPU<7>; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx index f3b62ec8a6108..8741b42a697a8 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx @@ -25,7 +25,7 @@ constexpr int UnusedIndex{-1}; template void TrackerTraitsGPU::initialiseTimeFrame(const int iteration) { - mTimeFrameGPU->initialise(iteration, mTrkParams[iteration], nLayers); + mTimeFrameGPU->initialise(iteration, this->mTrkParams[iteration], nLayers); mTimeFrameGPU->loadClustersDevice(iteration); mTimeFrameGPU->loadUnsortedClustersDevice(iteration); mTimeFrameGPU->loadClustersIndexTables(iteration); @@ -37,23 +37,30 @@ void TrackerTraitsGPU::initialiseTimeFrame(const int iteration) mTimeFrameGPU->loadIndexTableUtils(iteration); } +template +void TrackerTraitsGPU::adoptTimeFrame(TimeFrame* tf) +{ + mTimeFrameGPU = static_cast*>(tf); + this->mTimeFrame = static_cast*>(tf); +} + template void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int iROFslice, int iVertex) { auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); mTimeFrameGPU->createTrackletsLUTDevice(iteration); - const Vertex diamondVert({mTrkParams[iteration].Diamond[0], mTrkParams[iteration].Diamond[1], mTrkParams[iteration].Diamond[2]}, {25.e-6f, 0.f, 0.f, 25.e-6f, 0.f, 36.f}, 1, 1.f); + const Vertex diamondVert({this->mTrkParams[iteration].Diamond[0], this->mTrkParams[iteration].Diamond[1], this->mTrkParams[iteration].Diamond[2]}, {25.e-6f, 0.f, 0.f, 25.e-6f, 0.f, 36.f}, 1, 1.f); gsl::span diamondSpan(&diamondVert, 1); - int startROF{mTrkParams[iteration].nROFsPerIterations > 0 ? iROFslice * mTrkParams[iteration].nROFsPerIterations : 0}; - int endROF{o2::gpu::CAMath::Min(mTrkParams[iteration].nROFsPerIterations > 0 ? (iROFslice + 1) * mTrkParams[iteration].nROFsPerIterations + mTrkParams[iteration].DeltaROF : mTimeFrameGPU->getNrof(), mTimeFrameGPU->getNrof())}; + int startROF{this->mTrkParams[iteration].nROFsPerIterations > 0 ? iROFslice * this->mTrkParams[iteration].nROFsPerIterations : 0}; + int endROF{o2::gpu::CAMath::Min(this->mTrkParams[iteration].nROFsPerIterations > 0 ? (iROFslice + 1) * this->mTrkParams[iteration].nROFsPerIterations + this->mTrkParams[iteration].DeltaROF : mTimeFrameGPU->getNrof(), mTimeFrameGPU->getNrof())}; countTrackletsInROFsHandler(mTimeFrameGPU->getDeviceIndexTableUtils(), mTimeFrameGPU->getDeviceMultCutMask(), startROF, endROF, mTimeFrameGPU->getNrof(), - mTrkParams[iteration].DeltaROF, + this->mTrkParams[iteration].DeltaROF, iVertex, mTimeFrameGPU->getDeviceVertices(), mTimeFrameGPU->getDeviceROFramesPV(), @@ -66,13 +73,13 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i mTimeFrameGPU->getDeviceArrayTrackletsLUT(), mTimeFrameGPU->getDeviceTrackletsLUTs(), // Required for the exclusive sums iteration, - mTrkParams[iteration].NSigmaCut, + this->mTrkParams[iteration].NSigmaCut, mTimeFrameGPU->getPhiCuts(), - mTrkParams[iteration].PVres, + this->mTrkParams[iteration].PVres, mTimeFrameGPU->getMinRs(), mTimeFrameGPU->getMaxRs(), mTimeFrameGPU->getPositionResolutions(), - mTrkParams[iteration].LayerRadii, + this->mTrkParams[iteration].LayerRadii, mTimeFrameGPU->getMSangles(), conf.nBlocks, conf.nThreads); @@ -82,7 +89,7 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i startROF, endROF, mTimeFrameGPU->getNrof(), - mTrkParams[iteration].DeltaROF, + this->mTrkParams[iteration].DeltaROF, iVertex, mTimeFrameGPU->getDeviceVertices(), mTimeFrameGPU->getDeviceROFramesPV(), @@ -98,13 +105,13 @@ void TrackerTraitsGPU::computeLayerTracklets(const int iteration, int i mTimeFrameGPU->getDeviceArrayTrackletsLUT(), mTimeFrameGPU->getDeviceTrackletsLUTs(), iteration, - mTrkParams[iteration].NSigmaCut, + this->mTrkParams[iteration].NSigmaCut, mTimeFrameGPU->getPhiCuts(), - mTrkParams[iteration].PVres, + this->mTrkParams[iteration].PVres, mTimeFrameGPU->getMinRs(), mTimeFrameGPU->getMaxRs(), mTimeFrameGPU->getPositionResolutions(), - mTrkParams[iteration].LayerRadii, + this->mTrkParams[iteration].LayerRadii, mTimeFrameGPU->getMSangles(), conf.nBlocks, conf.nThreads); @@ -116,7 +123,7 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) mTimeFrameGPU->createCellsLUTDevice(); auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); - for (int iLayer = 0; iLayer < mTrkParams[iteration].CellsPerRoad(); ++iLayer) { + for (int iLayer = 0; iLayer < this->mTrkParams[iteration].CellsPerRoad(); ++iLayer) { if (!mTimeFrameGPU->getNTracklets()[iLayer + 1] || !mTimeFrameGPU->getNTracklets()[iLayer]) { continue; } @@ -131,10 +138,10 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) nullptr, mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceCellLUTs()[iLayer], - mBz, - mTrkParams[iteration].MaxChi2ClusterAttachment, - mTrkParams[iteration].CellDeltaTanLambdaSigma, - mTrkParams[iteration].NSigmaCut, + this->mBz, + this->mTrkParams[iteration].MaxChi2ClusterAttachment, + this->mTrkParams[iteration].CellDeltaTanLambdaSigma, + this->mTrkParams[iteration].NSigmaCut, conf.nBlocks, conf.nThreads); mTimeFrameGPU->createCellsBuffers(iLayer); @@ -148,10 +155,10 @@ void TrackerTraitsGPU::computeLayerCells(const int iteration) mTimeFrameGPU->getDeviceCells()[iLayer], mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceCellLUTs()[iLayer], - mBz, - mTrkParams[iteration].MaxChi2ClusterAttachment, - mTrkParams[iteration].CellDeltaTanLambdaSigma, - mTrkParams[iteration].NSigmaCut, + this->mBz, + this->mTrkParams[iteration].MaxChi2ClusterAttachment, + this->mTrkParams[iteration].CellDeltaTanLambdaSigma, + this->mTrkParams[iteration].NSigmaCut, conf.nBlocks, conf.nThreads); } @@ -162,7 +169,7 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) { mTimeFrameGPU->createNeighboursIndexTablesDevice(); auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); - for (int iLayer{0}; iLayer < mTrkParams[iteration].CellsPerRoad() - 1; ++iLayer) { + for (int iLayer{0}; iLayer < this->mTrkParams[iteration].CellsPerRoad() - 1; ++iLayer) { const int nextLayerCellsNum{static_cast(mTimeFrameGPU->getNCells()[iLayer + 1])}; if (!nextLayerCellsNum) { @@ -175,8 +182,8 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceNeighbourPairs(iLayer), mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer), - mTrkParams[0].MaxChi2ClusterAttachment, - mBz, + this->mTrkParams[0].MaxChi2ClusterAttachment, + this->mBz, iLayer, mTimeFrameGPU->getNCells()[iLayer], nextLayerCellsNum, @@ -191,8 +198,8 @@ void TrackerTraitsGPU::findCellsNeighbours(const int iteration) mTimeFrameGPU->getDeviceArrayCellsLUT(), mTimeFrameGPU->getDeviceNeighbourPairs(iLayer), mTimeFrameGPU->getDeviceNeighboursIndexTables(iLayer), - mTrkParams[0].MaxChi2ClusterAttachment, - mBz, + this->mTrkParams[0].MaxChi2ClusterAttachment, + this->mBz, iLayer, mTimeFrameGPU->getNCells()[iLayer], nextLayerCellsNum, @@ -212,11 +219,11 @@ template void TrackerTraitsGPU::findRoads(const int iteration) { auto& conf = o2::its::ITSGpuTrackingParamConfig::Instance(); - for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { + for (int startLevel{this->mTrkParams[iteration].CellsPerRoad()}; startLevel >= this->mTrkParams[iteration].CellMinimumLevel(); --startLevel) { const int minimumLayer{startLevel - 1}; std::vector trackSeeds; - for (int startLayer{mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) { - if ((mTrkParams[iteration].StartLayerMask & (1 << (startLayer + 2))) == 0) { + for (int startLayer{this->mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) { + if ((this->mTrkParams[iteration].StartLayerMask & (1 << (startLayer + 2))) == 0) { continue; } processNeighboursHandler(startLayer, @@ -229,11 +236,11 @@ void TrackerTraitsGPU::findRoads(const int iteration) mTimeFrameGPU->getDeviceNeighboursLUTs(), mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), trackSeeds, - mBz, - mTrkParams[0].MaxChi2ClusterAttachment, - mTrkParams[0].MaxChi2NDF, + this->mBz, + this->mTrkParams[0].MaxChi2ClusterAttachment, + this->mTrkParams[0].MaxChi2NDF, mTimeFrameGPU->getDevicePropagator(), - mCorrType, + this->mCorrType, conf.nBlocks, conf.nThreads); } @@ -248,14 +255,14 @@ void TrackerTraitsGPU::findRoads(const int iteration) trackSeedHandler(mTimeFrameGPU->getDeviceTrackSeeds(), // CellSeed* trackSeeds mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(), // TrackingFrameInfo** foundTrackingFrameInfo mTimeFrameGPU->getDeviceTrackITSExt(), // o2::its::TrackITSExt* tracks - mTrkParams[iteration].MinPt, // std::vector& minPtsHost, + this->mTrkParams[iteration].MinPt, // std::vector& minPtsHost, trackSeeds.size(), // const size_t nSeeds - mBz, // const float Bz + this->mBz, // const float Bz startLevel, // const int startLevel, - mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment - mTrkParams[0].MaxChi2NDF, // float maxChi2NDF + this->mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment + this->mTrkParams[0].MaxChi2NDF, // float maxChi2NDF mTimeFrameGPU->getDevicePropagator(), // const o2::base::Propagator* propagator - mCorrType, // o2::base::PropagatorImpl::MatCorrType + this->mCorrType, // o2::base::PropagatorImpl::MatCorrType conf.nBlocks, conf.nThreads); @@ -269,7 +276,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) } int nShared = 0; bool isFirstShared{false}; - for (int iLayer{0}; iLayer < mTrkParams[0].NLayers; ++iLayer) { + for (int iLayer{0}; iLayer < this->mTrkParams[0].NLayers; ++iLayer) { if (track.getClusterIndex(iLayer) == UnusedIndex) { continue; } @@ -277,12 +284,12 @@ void TrackerTraitsGPU::findRoads(const int iteration) isFirstShared |= !iLayer && mTimeFrameGPU->isClusterUsed(iLayer, track.getClusterIndex(iLayer)); } - if (nShared > mTrkParams[0].ClusterSharing) { + if (nShared > this->mTrkParams[0].ClusterSharing) { continue; } std::array rofs{INT_MAX, INT_MAX, INT_MAX}; - for (int iLayer{0}; iLayer < mTrkParams[0].NLayers; ++iLayer) { + for (int iLayer{0}; iLayer < this->mTrkParams[0].NLayers; ++iLayer) { if (track.getClusterIndex(iLayer) == UnusedIndex) { continue; } @@ -307,7 +314,7 @@ void TrackerTraitsGPU::findRoads(const int iteration) } mTimeFrameGPU->loadUsedClustersDevice(); } - if (iteration == mTrkParams.size() - 1) { + if (iteration == this->mTrkParams.size() - 1) { mTimeFrameGPU->unregisterHostMemory(0); } }; @@ -333,7 +340,7 @@ int TrackerTraitsGPU::getTFNumberOfCells() const template void TrackerTraitsGPU::setBz(float bz) { - mBz = bz; + this->mBz = bz; mTimeFrameGPU->setBz(bz); } diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index bb39e9e70341b..0dc5c28b3bc65 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -864,8 +864,8 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, const float NSigmaCut, std::vector& phiCuts, const float resolutionPV, - std::vector& minRs, - std::vector& maxRs, + std::array& minRs, + std::array& maxRs, std::vector& resolutions, std::vector& radii, std::vector& mulScatAng, @@ -930,8 +930,8 @@ void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const float NSigmaCut, std::vector& phiCuts, const float resolutionPV, - std::vector& minRs, - std::vector& maxRs, + std::array& minRs, + std::array& maxRs, std::vector& resolutions, std::vector& radii, std::vector& mulScatAng, @@ -1318,8 +1318,8 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils* utils, const float NSigmaCut, std::vector& phiCuts, const float resolutionPV, - std::vector& minRs, - std::vector& maxRs, + std::array& minRs, + std::array& maxRs, std::vector& resolutions, std::vector& radii, std::vector& mulScatAng, @@ -1350,8 +1350,8 @@ template void computeTrackletsInROFsHandler<7>(const IndexTableUtils* utils, const float NSigmaCut, std::vector& phiCuts, const float resolutionPV, - std::vector& minRs, - std::vector& maxRs, + std::array& minRs, + std::array& maxRs, std::vector& resolutions, std::vector& radii, std::vector& mulScatAng, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/VertexerTraitsGPU.cxx b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/VertexerTraitsGPU.cxx index f630d2cdec76a..2a6debe8f652e 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/VertexerTraitsGPU.cxx +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/VertexerTraitsGPU.cxx @@ -28,10 +28,6 @@ namespace o2::its { -VertexerTraitsGPU::VertexerTraitsGPU() -{ - setIsGPU(true); -} void VertexerTraitsGPU::initialise(const TrackingParameters& trackingParams, const int iteration) { diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index 46c4a8e19fa47..19c24c743ff07 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -16,12 +16,9 @@ #include #include #include +#include #include -#include #include -#include -#include -#include #include "DataFormatsITS/TrackITS.h" @@ -59,14 +56,19 @@ class ROFRecord; namespace its { +namespace gpu +{ +template +class TimeFrameGPU; +} using Vertex = o2::dataformats::Vertex>; -class TimeFrame -{ - public: - friend class TimeFrameGPU; - TimeFrame(int nLayers = 7); - const Vertex& getPrimaryVertex(const int) const; +template +struct TimeFrame { + friend class gpu::TimeFrameGPU; + TimeFrame(); + virtual ~TimeFrame(); + const Vertex& getPrimaryVertex(const int ivtx) const { return mPrimaryVertices[ivtx]; } gsl::span getPrimaryVertices(int rofId) const; gsl::span getPrimaryVertices(int romin, int romax) const; gsl::span> getPrimaryVerticesMCRecInfo(const int rofId) const; @@ -91,11 +93,10 @@ class TimeFrame int getTotalClusters() const; std::vector& getTotVertIteration() { return mTotVertPerIteration; } - bool empty() const; - bool isGPU() const { return mIsGPU; } - int getSortedIndex(int rofId, int layer, int i) const; - int getSortedStartIndex(const int, const int) const; - int getNrof() const; + bool empty() const { return getTotalClusters() == 0; } + int getSortedIndex(int rofId, int layer, int idx) const { return mROFramesClusters[layer][rofId] + idx; } + int getSortedStartIndex(const int rofId, const int layer) const { return mROFramesClusters[layer][rofId]; } + int getNrof() const { return mNrof; } void resetBeamXY(const float x, const float y, const float w = 0); void setBeamPosition(const float x, const float y, const float s2, const float base = 50.f, const float systematic = 0.f) @@ -104,10 +105,10 @@ class TimeFrame resetBeamXY(x, y, s2 / o2::gpu::CAMath::Sqrt(base * base + systematic)); } - float getBeamX() const; - float getBeamY() const; - std::vector& getMinRs() { return mMinR; } - std::vector& getMaxRs() { return mMaxR; } + float getBeamX() const { return mBeamPos[0]; } + float getBeamY() const { return mBeamPos[1]; } + auto& getMinRs() { return mMinR; } + auto& getMaxRs() { return mMaxR; } float getMinR(int layer) const { return mMinR[layer]; } float getMaxR(int layer) const { return mMaxR[layer]; } float getMSangle(int layer) const { return mMSangles[layer]; } @@ -129,43 +130,43 @@ class TimeFrame gsl::span getIndexTablePerROFrange(int rofMin, int range, int layerId) const; gsl::span getIndexTable(int rofId, int layerId); std::vector& getIndexTableWhole(int layerId) { return mIndexTables[layerId]; } - const std::vector& getTrackingFrameInfoOnLayer(int layerId) const; + const std::vector& getTrackingFrameInfoOnLayer(int layerId) const { return mTrackingFrameInfo[layerId]; } const TrackingFrameInfo& getClusterTrackingFrameInfo(int layerId, const Cluster& cl) const; - const gsl::span getClusterLabels(int layerId, const Cluster& cl) const; - const gsl::span getClusterLabels(int layerId, const int clId) const; - int getClusterExternalIndex(int layerId, const int clId) const; - int getClusterSize(int clusterId) const; - void setClusterSize(const std::vector& v) { mClusterSize = v; }; + const gsl::span getClusterLabels(int layerId, const Cluster& cl) const { return getClusterLabels(layerId, cl.clusterId); } + const gsl::span getClusterLabels(int layerId, const int clId) const { return mClusterLabels->getLabels(mClusterExternalIndices[layerId][clId]); } + int getClusterExternalIndex(int layerId, const int clId) const { return mClusterExternalIndices[layerId][clId]; } + int getClusterSize(int clusterId) const { return mClusterSize[clusterId]; } + void setClusterSize(const std::vector& v) { mClusterSize = v; } std::vector& getTrackletsLabel(int layer) { return mTrackletLabels[layer]; } std::vector& getCellsLabel(int layer) { return mCellLabels[layer]; } - bool hasMCinformation() const; + bool hasMCinformation() const { return mClusterLabels; } void initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers = 7, bool resetVertices = true); void resetRofPV() { deepVectorClear(mPrimaryVertices); mROFramesPV.resize(1, 0); mTotVertPerIteration.resize(1); - }; + } - bool isClusterUsed(int layer, int clusterId) const; - void markUsedCluster(int layer, int clusterId); + bool isClusterUsed(int layer, int clusterId) const { return mUsedClusters[layer][clusterId]; } + void markUsedCluster(int layer, int clusterId) { mUsedClusters[layer][clusterId] = true; } gsl::span getUsedClusters(const int layer); - std::vector>& getTracklets(); - std::vector>& getTrackletsLookupTable(); + auto& getTracklets() { return mTracklets; } + auto& getTrackletsLookupTable() { return mTrackletsLookupTable; } - std::vector>& getClusters(); - std::vector>& getUnsortedClusters(); + auto& getClusters() { return mClusters; } + auto& getUnsortedClusters() { return mUnsortedClusters; } int getClusterROF(int iLayer, int iCluster); - std::vector>& getCells(); + auto& getCells() { return mCells; } - std::vector>& getCellsLookupTable(); - std::vector>& getCellsNeighbours(); - std::vector>& getCellsNeighboursLUT(); - std::vector>& getRoads(); + auto& getCellsLookupTable() { return mCellsLookupTable; } + auto& getCellsNeighbours() { return mCellsNeighbours; } + auto& getCellsNeighboursLUT() { return mCellsNeighboursLUT; } + std::vector>& getRoads() { return mRoads; } std::vector& getTracks(int rofId) { return mTracks[rofId]; } std::vector& getTracksLabel(const int rofId) { return mTracksLabel[rofId]; } std::vector& getLinesLabel(const int rofId) { return mLinesLabels[rofId]; } @@ -182,20 +183,20 @@ class TimeFrame bool checkMemory(unsigned long max) { return getArtefactsMemory() < max; } unsigned long getArtefactsMemory(); - int getROFCutClusterMult() const { return mCutClusterMult; }; - int getROFCutVertexMult() const { return mCutVertexMult; }; + int getROFCutClusterMult() const { return mCutClusterMult; } + int getROFCutVertexMult() const { return mCutVertexMult; } int getROFCutAllMult() const { return mCutClusterMult + mCutVertexMult; } // Vertexer void computeTrackletsPerROFScans(); void computeTracletsPerClusterScans(); - int& getNTrackletsROF(int rofId, int combId); - std::vector& getLines(int rofId); + int& getNTrackletsROF(int rofId, int combId) { return mNTrackletsPerROF[combId][rofId]; } + std::vector& getLines(int rofId) { return mLines[rofId]; } int getNLinesTotal() const { return std::accumulate(mLines.begin(), mLines.end(), 0, [](int sum, const auto& l) { return sum + l.size(); }); } - std::vector& getTrackletClusters(int rofId); + std::vector& getTrackletClusters(int rofId) { return mTrackletClusters[rofId]; } gsl::span getFoundTracklets(int rofId, int combId) const; gsl::span getFoundTracklets(int rofId, int combId); gsl::span getLabelsFoundTracklets(int rofId, int combId) const; @@ -210,8 +211,8 @@ class TimeFrame void initialiseRoadLabels(); void setRoadLabel(int i, const unsigned long long& lab, bool fake); - const unsigned long long& getRoadLabel(int i) const; - bool isRoadFake(int i) const; + const unsigned long long& getRoadLabel(int i) const { return mRoadLabels[i].first; } + bool isRoadFake(int i) const { return mRoadLabels[i].second; } void setMultiplicityCutMask(const std::vector& cutMask) { mMultiplicityCutMask = cutMask; } void setROFMask(const std::vector& rofMask) { mROFMask = rofMask; } @@ -222,33 +223,16 @@ class TimeFrame void setBz(float bz) { mBz = bz; } float getBz() const { return mBz; } - void setExternalAllocator(ExternalAllocator* allocator) - { - if (mIsGPU) { - LOGP(debug, "Setting timeFrame allocator to external"); - mAllocator = allocator; - mExtAllocator = true; // to be removed - } else { - LOGP(debug, "External allocator is currently only supported for GPU"); - } - } - - virtual void setDevicePropagator(const o2::base::PropagatorImpl*) - { - return; - }; + virtual void setDevicePropagator(const o2::base::PropagatorImpl*) { return; } const o2::base::PropagatorImpl* getDevicePropagator() const { return mPropagatorDevice; } template void addClusterToLayer(int layer, T&&... args); template void addTrackingFrameInfoToLayer(int layer, T&&... args); - void addClusterExternalIndexToLayer(int layer, const int idx); + void addClusterExternalIndexToLayer(int layer, const int idx) { mClusterExternalIndices[layer].push_back(idx); } - void resizeVectors(int nLayers); - - void setExtAllocator(bool ext) { mExtAllocator = ext; } - bool getExtAllocator() const { return mExtAllocator; } + void resetVectors(); /// Debug and printing void checkTrackletLUTs(); @@ -265,17 +249,17 @@ class TimeFrame bool mIsGPU = false; - std::vector> mClusters; - std::vector> mTrackingFrameInfo; - std::vector> mClusterExternalIndices; - std::vector> mROFramesClusters; + std::array, nLayers> mClusters; + std::array, nLayers> mTrackingFrameInfo; + std::array, nLayers> mClusterExternalIndices; + std::array, nLayers> mROFramesClusters; const dataformats::MCTruthContainer* mClusterLabels = nullptr; std::array, 2> mNTrackletsPerCluster; std::array, 2> mNTrackletsPerClusterSum; - std::vector> mNClustersPerROF; + std::array, nLayers> mNClustersPerROF; std::vector> mIndexTables; std::vector> mTrackletsLookupTable; - std::vector> mUsedClusters; + std::array, nLayers> mUsedClusters; int mNrof = 0; int mNExtendedTracks{0}; int mNExtendedUsedClusters{0}; @@ -285,12 +269,25 @@ class TimeFrame // State if memory will be externally managed. bool mExtAllocator = false; ExternalAllocator* mAllocator = nullptr; - std::vector> mUnsortedClusters; + void setExternalAllocator(ExternalAllocator* allocator) + { + if (mIsGPU) { + LOGP(debug, "Setting timeFrame allocator to external"); + mAllocator = allocator; + mExtAllocator = true; // to be removed + } else { + LOGP(fatal, "External allocator is currently only supported for GPU"); + } + } + void setExtAllocator(bool ext) { mExtAllocator = ext; } + bool getExtAllocator() const { return mExtAllocator; } + + std::array, nLayers> mUnsortedClusters; std::vector> mTracklets; std::vector> mCells; std::vector> mCellSeeds; std::vector> mCellSeedsChi2; - std::vector> mRoads; + std::vector> mRoads; std::vector> mTracks; std::vector> mCellsNeighbours; std::vector> mCellsLookupTable; @@ -318,8 +315,8 @@ class TimeFrame int mBeamPosWeight = 0; std::array mBeamPos = {0.f, 0.f}; bool isBeamPositionOverridden = false; - std::vector mMinR; - std::vector mMaxR; + std::array mMinR; + std::array mMaxR; std::vector mMSangles; std::vector mPhiCuts; std::vector mPositionResolution; @@ -341,7 +338,7 @@ class TimeFrame std::vector> mNTrackletsPerROF; std::vector> mLines; std::vector> mTrackletClusters; - std::vector> mTrackletsIndexROF; + std::array, 2> mTrackletsIndexROF; std::vector> mLinesLabels; std::vector> mVerticesMCRecInfo; std::array mTotalTracklets = {0, 0}; @@ -350,9 +347,8 @@ class TimeFrame // \Vertexer }; -inline const Vertex& TimeFrame::getPrimaryVertex(const int vertexIndex) const { return mPrimaryVertices[vertexIndex]; } - -inline gsl::span TimeFrame::getPrimaryVertices(int rofId) const +template +inline gsl::span TimeFrame::getPrimaryVertices(int rofId) const { const int start = mROFramesPV[rofId]; const int stop_idx = rofId >= mNrof - 1 ? mNrof : rofId + 1; @@ -360,7 +356,8 @@ inline gsl::span TimeFrame::getPrimaryVertices(int rofId) const return {&mPrimaryVertices[start], static_cast::size_type>(delta)}; } -inline gsl::span> TimeFrame::getPrimaryVerticesMCRecInfo(const int rofId) const +template +inline gsl::span> TimeFrame::getPrimaryVerticesMCRecInfo(const int rofId) const { const int start = mROFramesPV[rofId]; const int stop_idx = rofId >= mNrof - 1 ? mNrof : rofId + 1; @@ -368,12 +365,14 @@ inline gsl::span> TimeFrame::getPrimaryVerti return {&(mVerticesMCRecInfo[start]), static_cast>::size_type>(delta)}; } -inline gsl::span TimeFrame::getPrimaryVertices(int romin, int romax) const +template +inline gsl::span TimeFrame::getPrimaryVertices(int romin, int romax) const { return {&mPrimaryVertices[mROFramesPV[romin]], static_cast::size_type>(mROFramesPV[romax + 1] - mROFramesPV[romin])}; } -inline gsl::span> TimeFrame::getPrimaryVerticesXAlpha(int rofId) const +template +inline gsl::span> TimeFrame::getPrimaryVerticesXAlpha(int rofId) const { const int start = mROFramesPV[rofId]; const int stop_idx = rofId >= mNrof - 1 ? mNrof : rofId + 1; @@ -381,364 +380,290 @@ inline gsl::span> TimeFrame::getPrimaryVerticesXAlpha return {&(mPValphaX[start]), static_cast>::size_type>(delta)}; } -inline int TimeFrame::getPrimaryVerticesNum(int rofId) const +template +inline int TimeFrame::getPrimaryVerticesNum(int rofId) const { return rofId < 0 ? mPrimaryVertices.size() : mROFramesPV[rofId + 1] - mROFramesPV[rofId]; } -inline bool TimeFrame::empty() const { return getTotalClusters() == 0; } - -inline int TimeFrame::getSortedIndex(int rofId, int layer, int index) const { return mROFramesClusters[layer][rofId] + index; } - -inline int TimeFrame::getSortedStartIndex(const int rofId, const int layer) const { return mROFramesClusters[layer][rofId]; } - -inline int TimeFrame::getNrof() const { return mNrof; } - -inline void TimeFrame::resetBeamXY(const float x, const float y, const float w) +template +inline void TimeFrame::resetBeamXY(const float x, const float y, const float w) { mBeamPos[0] = x; mBeamPos[1] = y; mBeamPosWeight = w; } -inline float TimeFrame::getBeamX() const { return mBeamPos[0]; } - -inline float TimeFrame::getBeamY() const { return mBeamPos[1]; } - -inline gsl::span TimeFrame::getROFrameClusters(int layerId) const +template +inline gsl::span TimeFrame::getROFrameClusters(int layerId) const { return {&mROFramesClusters[layerId][0], static_cast::size_type>(mROFramesClusters[layerId].size())}; } -inline gsl::span TimeFrame::getClustersOnLayer(int rofId, int layerId) +template +inline gsl::span TimeFrame::getClustersOnLayer(int rofId, int layerId) { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } int startIdx{mROFramesClusters[layerId][rofId]}; return {&mClusters[layerId][startIdx], static_cast::size_type>(mROFramesClusters[layerId][rofId + 1] - startIdx)}; } -inline gsl::span TimeFrame::getClustersOnLayer(int rofId, int layerId) const +template +inline gsl::span TimeFrame::getClustersOnLayer(int rofId, int layerId) const { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } int startIdx{mROFramesClusters[layerId][rofId]}; return {&mClusters[layerId][startIdx], static_cast::size_type>(mROFramesClusters[layerId][rofId + 1] - startIdx)}; } -inline gsl::span TimeFrame::getUsedClustersROF(int rofId, int layerId) +template +inline gsl::span TimeFrame::getUsedClustersROF(int rofId, int layerId) { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } int startIdx{mROFramesClusters[layerId][rofId]}; - return {&mUsedClusters[layerId][startIdx], static_cast::size_type>(mROFramesClusters[layerId][rofId + 1] - startIdx)}; + return {&mUsedClusters[layerId][startIdx], static_cast::size_type>(mROFramesClusters[layerId][rofId + 1] - startIdx)}; } -inline gsl::span TimeFrame::getUsedClustersROF(int rofId, int layerId) const +template +inline gsl::span TimeFrame::getUsedClustersROF(int rofId, int layerId) const { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } int startIdx{mROFramesClusters[layerId][rofId]}; - return {&mUsedClusters[layerId][startIdx], static_cast::size_type>(mROFramesClusters[layerId][rofId + 1] - startIdx)}; + return {&mUsedClusters[layerId][startIdx], static_cast::size_type>(mROFramesClusters[layerId][rofId + 1] - startIdx)}; } -inline gsl::span TimeFrame::getClustersPerROFrange(int rofMin, int range, int layerId) const +template +inline gsl::span TimeFrame::getClustersPerROFrange(int rofMin, int range, int layerId) const { if (rofMin < 0 || rofMin >= mNrof) { - return gsl::span(); + return {}; } int startIdx{mROFramesClusters[layerId][rofMin]}; // First cluster of rofMin int endIdx{mROFramesClusters[layerId][o2::gpu::CAMath::Min(rofMin + range, mNrof)]}; return {&mClusters[layerId][startIdx], static_cast::size_type>(endIdx - startIdx)}; } -inline gsl::span TimeFrame::getROFramesClustersPerROFrange(int rofMin, int range, int layerId) const +template +inline gsl::span TimeFrame::getROFramesClustersPerROFrange(int rofMin, int range, int layerId) const { int chkdRange{o2::gpu::CAMath::Min(range, mNrof - rofMin)}; return {&mROFramesClusters[layerId][rofMin], static_cast::size_type>(chkdRange)}; } -inline gsl::span TimeFrame::getNClustersROFrange(int rofMin, int range, int layerId) const +template +inline gsl::span TimeFrame::getNClustersROFrange(int rofMin, int range, int layerId) const { int chkdRange{o2::gpu::CAMath::Min(range, mNrof - rofMin)}; return {&mNClustersPerROF[layerId][rofMin], static_cast::size_type>(chkdRange)}; } -inline int TimeFrame::getTotalClustersPerROFrange(int rofMin, int range, int layerId) const +template +inline int TimeFrame::getTotalClustersPerROFrange(int rofMin, int range, int layerId) const { int startIdx{rofMin}; // First cluster of rofMin int endIdx{o2::gpu::CAMath::Min(rofMin + range, mNrof)}; return mROFramesClusters[layerId][endIdx] - mROFramesClusters[layerId][startIdx]; } -inline gsl::span TimeFrame::getIndexTablePerROFrange(int rofMin, int range, int layerId) const +template +inline gsl::span TimeFrame::getIndexTablePerROFrange(int rofMin, int range, int layerId) const { const int iTableSize{mIndexTableUtils.getNphiBins() * mIndexTableUtils.getNzBins() + 1}; int chkdRange{o2::gpu::CAMath::Min(range, mNrof - rofMin)}; return {&mIndexTables[layerId][rofMin * iTableSize], static_cast::size_type>(chkdRange * iTableSize)}; } -inline int TimeFrame::getClusterROF(int iLayer, int iCluster) +template +inline int TimeFrame::getClusterROF(int iLayer, int iCluster) { return std::lower_bound(mROFramesClusters[iLayer].begin(), mROFramesClusters[iLayer].end(), iCluster + 1) - mROFramesClusters[iLayer].begin() - 1; } -inline gsl::span TimeFrame::getUnsortedClustersOnLayer(int rofId, int layerId) const +template +inline gsl::span TimeFrame::getUnsortedClustersOnLayer(int rofId, int layerId) const { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } int startIdx{mROFramesClusters[layerId][rofId]}; return {&mUnsortedClusters[layerId][startIdx], static_cast::size_type>(mROFramesClusters[layerId][rofId + 1] - startIdx)}; } -inline const std::vector& TimeFrame::getTrackingFrameInfoOnLayer(int layerId) const -{ - return mTrackingFrameInfo[layerId]; -} - -inline const TrackingFrameInfo& TimeFrame::getClusterTrackingFrameInfo(int layerId, const Cluster& cl) const -{ - return mTrackingFrameInfo[layerId][cl.clusterId]; -} - -inline const gsl::span TimeFrame::getClusterLabels(int layerId, const Cluster& cl) const -{ - return getClusterLabels(layerId, cl.clusterId); -} - -inline const gsl::span TimeFrame::getClusterLabels(int layerId, int clId) const -{ - return mClusterLabels->getLabels(mClusterExternalIndices[layerId][clId]); -} - -inline int TimeFrame::getClusterSize(int clusterId) const -{ - return mClusterSize[clusterId]; -} - -inline int TimeFrame::getClusterExternalIndex(int layerId, const int clId) const -{ - return mClusterExternalIndices[layerId][clId]; -} - -inline gsl::span TimeFrame::getIndexTable(int rofId, int layer) +template +inline gsl::span TimeFrame::getIndexTable(int rofId, int layer) { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } return {&mIndexTables[layer][rofId * (mIndexTableUtils.getNphiBins() * mIndexTableUtils.getNzBins() + 1)], static_cast::size_type>(mIndexTableUtils.getNphiBins() * mIndexTableUtils.getNzBins() + 1)}; } -inline std::vector& TimeFrame::getLines(int rofId) -{ - return mLines[rofId]; -} - -inline std::vector& TimeFrame::getTrackletClusters(int rofId) -{ - return mTrackletClusters[rofId]; -} - +template template -void TimeFrame::addClusterToLayer(int layer, T&&... values) +void TimeFrame::addClusterToLayer(int layer, T&&... values) { mUnsortedClusters[layer].emplace_back(std::forward(values)...); } +template template -void TimeFrame::addTrackingFrameInfoToLayer(int layer, T&&... values) +void TimeFrame::addTrackingFrameInfoToLayer(int layer, T&&... values) { mTrackingFrameInfo[layer].emplace_back(std::forward(values)...); } -inline void TimeFrame::addClusterExternalIndexToLayer(int layer, const int idx) +template +inline gsl::span TimeFrame::getUsedClusters(const int layer) { - mClusterExternalIndices[layer].push_back(idx); + return {&mUsedClusters[layer][0], static_cast::size_type>(mUsedClusters[layer].size())}; } -inline bool TimeFrame::hasMCinformation() const -{ - return mClusterLabels; -} - -inline bool TimeFrame::isClusterUsed(int layer, int clusterId) const -{ - return mUsedClusters[layer][clusterId]; -} - -inline gsl::span TimeFrame::getUsedClusters(const int layer) -{ - return {&mUsedClusters[layer][0], static_cast::size_type>(mUsedClusters[layer].size())}; -} - -inline void TimeFrame::markUsedCluster(int layer, int clusterId) { mUsedClusters[layer][clusterId] = true; } - -inline std::vector>& TimeFrame::getTracklets() -{ - return mTracklets; -} - -inline std::vector>& TimeFrame::getTrackletsLookupTable() -{ - return mTrackletsLookupTable; -} - -inline void TimeFrame::initialiseRoadLabels() +template +inline void TimeFrame::initialiseRoadLabels() { mRoadLabels.clear(); mRoadLabels.resize(mRoads.size()); } -inline void TimeFrame::setRoadLabel(int i, const unsigned long long& lab, bool fake) +template +inline void TimeFrame::setRoadLabel(int i, const unsigned long long& lab, bool fake) { mRoadLabels[i].first = lab; mRoadLabels[i].second = fake; } -inline const unsigned long long& TimeFrame::getRoadLabel(int i) const -{ - return mRoadLabels[i].first; -} - -inline gsl::span TimeFrame::getNTrackletsCluster(int rofId, int combId) +template +inline gsl::span TimeFrame::getNTrackletsCluster(int rofId, int combId) { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } auto startIdx{mROFramesClusters[1][rofId]}; return {&mNTrackletsPerCluster[combId][startIdx], static_cast::size_type>(mROFramesClusters[1][rofId + 1] - startIdx)}; } -inline gsl::span TimeFrame::getExclusiveNTrackletsCluster(int rofId, int combId) +template +inline gsl::span TimeFrame::getExclusiveNTrackletsCluster(int rofId, int combId) { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } auto clusStartIdx{mROFramesClusters[1][rofId]}; return {&mNTrackletsPerClusterSum[combId][clusStartIdx], static_cast::size_type>(mROFramesClusters[1][rofId + 1] - clusStartIdx)}; } -inline int& TimeFrame::getNTrackletsROF(int rofId, int combId) -{ - return mNTrackletsPerROF[combId][rofId]; -} - -inline bool TimeFrame::isRoadFake(int i) const -{ - return mRoadLabels[i].second; -} - -inline std::vector>& TimeFrame::getClusters() -{ - return mClusters; -} - -inline std::vector>& TimeFrame::getUnsortedClusters() -{ - return mUnsortedClusters; -} - -inline std::vector>& TimeFrame::getCells() { return mCells; } - -inline std::vector>& TimeFrame::getCellsLookupTable() -{ - return mCellsLookupTable; -} - -inline std::vector>& TimeFrame::getCellsNeighbours() { return mCellsNeighbours; } -inline std::vector>& TimeFrame::getCellsNeighboursLUT() { return mCellsNeighboursLUT; } - -inline std::vector>& TimeFrame::getRoads() { return mRoads; } - -inline gsl::span TimeFrame::getFoundTracklets(int rofId, int combId) +template +inline gsl::span TimeFrame::getFoundTracklets(int rofId, int combId) { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } auto startIdx{mNTrackletsPerROF[combId][rofId]}; return {&mTracklets[combId][startIdx], static_cast::size_type>(mNTrackletsPerROF[combId][rofId + 1] - startIdx)}; } -inline gsl::span TimeFrame::getFoundTracklets(int rofId, int combId) const +template +inline gsl::span TimeFrame::getFoundTracklets(int rofId, int combId) const { if (rofId < 0 || rofId >= mNrof) { - return gsl::span(); + return {}; } auto startIdx{mNTrackletsPerROF[combId][rofId]}; return {&mTracklets[combId][startIdx], static_cast::size_type>(mNTrackletsPerROF[combId][rofId + 1] - startIdx)}; } -inline gsl::span TimeFrame::getLabelsFoundTracklets(int rofId, int combId) const +template +inline gsl::span TimeFrame::getLabelsFoundTracklets(int rofId, int combId) const { if (rofId < 0 || rofId >= mNrof || !hasMCinformation()) { - return gsl::span(); + return {}; } auto startIdx{mNTrackletsPerROF[combId][rofId]}; return {&mTrackletLabels[combId][startIdx], static_cast::size_type>(mNTrackletsPerROF[combId][rofId + 1] - startIdx)}; } -inline int TimeFrame::getNumberOfClusters() const +template +inline int TimeFrame::getTotalClusters() const +{ + size_t totalClusters{0}; + for (const auto& clusters : mUnsortedClusters) { + totalClusters += clusters.size(); + } + return int(totalClusters); +} + +template +inline int TimeFrame::getNumberOfClusters() const { int nClusters = 0; - for (auto& layer : mClusters) { + for (const auto& layer : mClusters) { nClusters += layer.size(); } return nClusters; } -inline int TimeFrame::getNumberOfCells() const +template +inline int TimeFrame::getNumberOfCells() const { int nCells = 0; - for (auto& layer : mCells) { + for (const auto& layer : mCells) { nCells += layer.size(); } return nCells; } -inline int TimeFrame::getNumberOfTracklets() const +template +inline int TimeFrame::getNumberOfTracklets() const { int nTracklets = 0; - for (auto& layer : mTracklets) { + for (const auto& layer : mTracklets) { nTracklets += layer.size(); } return nTracklets; } -inline int TimeFrame::getNumberOfNeighbours() const +template +inline int TimeFrame::getNumberOfNeighbours() const { int n{0}; - for (auto& l : mCellsNeighbours) { + for (const auto& l : mCellsNeighbours) { n += l.size(); } return n; } -inline size_t TimeFrame::getNumberOfTracks() const +template +inline size_t TimeFrame::getNumberOfTracks() const { int nTracks = 0; - for (auto& t : mTracks) { + for (const auto& t : mTracks) { nTracks += t.size(); } return nTracks; } -inline size_t TimeFrame::getNumberOfUsedClusters() const +template +inline size_t TimeFrame::getNumberOfUsedClusters() const { size_t nClusters = 0; - for (auto& layer : mUsedClusters) { + for (const auto& layer : mUsedClusters) { nClusters += std::count(layer.begin(), layer.end(), true); } return nClusters; } -inline void TimeFrame::insertPastVertex(const Vertex& vertex, const int iteration) +template +inline void TimeFrame::insertPastVertex(const Vertex& vertex, const int iteration) { int rofId = vertex.getTimeStamp().getTimeStamp(); mPrimaryVertices.insert(mPrimaryVertices.begin() + mROFramesPV[rofId], vertex); diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h index 8f0a471b40c59..2e535e2b6a644 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Tracker.h @@ -32,6 +32,7 @@ #include "ITStracking/ROframe.h" #include "ITStracking/MathUtils.h" #include "ITStracking/TimeFrame.h" +#include "ITStracking/TrackerTraits.h" #include "ITStracking/Road.h" #include "DataFormatsITS/TrackITS.h" @@ -46,48 +47,52 @@ class GPUChainITS; } namespace its { -class TrackerTraits; class Tracker { + static constexpr int NLayers{7}; + using TrackerTraits7 = TrackerTraits; + using TimeFrame7 = TimeFrame; using LogFunc = std::function; public: - Tracker(TrackerTraits* traits); + Tracker(TrackerTraits* traits); - void adoptTimeFrame(TimeFrame& tf); + void adoptTimeFrame(TimeFrame& tf); - void clustersToTracks(LogFunc = [](std::string s) { std::cout << s << std::endl; }, LogFunc = [](std::string s) { std::cerr << s << std::endl; }); + void clustersToTracks( + LogFunc = [](const std::string& s) { std::cout << s << '\n'; }, + LogFunc = [](const std::string& s) { std::cerr << s << '\n'; }); - void setParameters(const std::vector&); + void setParameters(const std::vector& p) { mTrkParams = p; } std::vector& getParameters() { return mTrkParams; } void getGlobalConfiguration(); - void setBz(float); - void setCorrType(const o2::base::PropagatorImpl::MatCorrType type); - bool isMatLUT() const; - void setNThreads(int n); - int getNThreads() const; + void setBz(float bz) { mTraits->setBz(bz); } + void setCorrType(const o2::base::PropagatorImpl::MatCorrType type) { mTraits->setCorrType(type); } + bool isMatLUT() const { return mTraits->isMatLUT(); } + void setNThreads(int n) { mTraits->setNThreads(n); } + int getNThreads() const { return mTraits->getNThreads(); } void printSummary() const; private: - void initialiseTimeFrame(int& iteration); - void computeTracklets(int& iteration, int& iROFslice, int& iVertex); - void computeCells(int& iteration); - void findCellsNeighbours(int& iteration); - void findRoads(int& iteration); - void findShortPrimaries(); - void extendTracks(int& iteration); + void initialiseTimeFrame(int iteration) { mTraits->initialiseTimeFrame(iteration); } + void computeTracklets(int iteration, int iROFslice, int iVertex) { mTraits->computeLayerTracklets(iteration, iROFslice, iVertex); } + void computeCells(int iteration) { mTraits->computeLayerCells(iteration); } + void findCellsNeighbours(int iteration) { mTraits->findCellsNeighbours(iteration); } + void findRoads(int iteration) { mTraits->findRoads(iteration); } + void findShortPrimaries() { mTraits->findShortPrimaries(); } + void extendTracks(int iteration) { mTraits->extendTracks(iteration); } // MC interaction void computeRoadsMClabels(); void computeTracksMClabels(); void rectifyClusterIndices(); - template - float evaluateTask(void (Tracker::*)(T...), const char*, LogFunc logger, T&&... args); + template + float evaluateTask(void (Tracker::*task)(T...), const char* taskName, LogFunc logger, F&&... args); - TrackerTraits* mTraits = nullptr; /// Observer pointer, not owned by this class - TimeFrame* mTimeFrame = nullptr; /// Observer pointer, not owned by this class + TrackerTraits7* mTraits = nullptr; /// Observer pointer, not owned by this class + TimeFrame7* mTimeFrame = nullptr; /// Observer pointer, not owned by this class std::vector mTrkParams; o2::gpu::GPUChainITS* mRecoChain = nullptr; @@ -97,13 +102,8 @@ class Tracker double mTotalTime{0}; }; -inline void Tracker::setParameters(const std::vector& trkPars) -{ - mTrkParams = trkPars; -} - -template -float Tracker::evaluateTask(void (Tracker::*task)(T...), const char* taskName, LogFunc logger, T&&... args) +template +float Tracker::evaluateTask(void (Tracker::*task)(T...), const char* taskName, LogFunc logger, F&&... args) { float diff{0.f}; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h index 6b514c6e8d000..7340f247e4cf2 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -35,12 +35,13 @@ namespace its { class TrackITSExt; +template class TrackerTraits { public: virtual ~TrackerTraits() = default; - virtual void adoptTimeFrame(TimeFrame* tf); - virtual void initialiseTimeFrame(const int iteration); + virtual void adoptTimeFrame(TimeFrame* tf) { mTimeFrame = tf; } + virtual void initialiseTimeFrame(const int iteration) { mTimeFrame->initialise(iteration, mTrkParams[iteration], mTrkParams[iteration].NLayers); } virtual void computeLayerTracklets(const int iteration, int iROFslice, int iVertex); virtual void computeLayerCells(const int iteration); @@ -55,11 +56,11 @@ class TrackerTraits virtual bool trackFollowing(TrackITSExt* track, int rof, bool outward, const int iteration); virtual void processNeighbours(int iLayer, int iLevel, const std::vector& currentCellSeed, const std::vector& currentCellId, std::vector& updatedCellSeed, std::vector& updatedCellId); - void UpdateTrackingParameters(const std::vector& trkPars); - TimeFrame* getTimeFrame() { return mTimeFrame; } + void updateTrackingParameters(const std::vector& trkPars) { mTrkParams = trkPars; } + TimeFrame* getTimeFrame() { return mTimeFrame; } virtual void setBz(float bz); - float getBz() const; + float getBz() const { return mBz; } void setCorrType(const o2::base::PropagatorImpl::MatCorrType type) { mCorrType = type; } bool isMatLUT() const; virtual const char* getName() const noexcept { return "CPU"; } @@ -67,8 +68,8 @@ class TrackerTraits // Others GPUhd() static consteval int4 getEmptyBinsRect() { return int4{0, 0, 0, 0}; } - const int4 getBinsRect(const Cluster&, int layer, float z1, float z2, float maxdeltaz, float maxdeltaphi) const noexcept; - const int4 getBinsRect(int layer, float phi, float maxdeltaphi, float z, float maxdeltaz) const noexcept; + const int4 getBinsRect(int layer, float phi, float maxdeltaphi, float z, float maxdeltaz) const noexcept { return getBinsRect(layer, phi, maxdeltaphi, z, z, maxdeltaz); } + const int4 getBinsRect(const Cluster& cls, int layer, float z1, float z2, float maxdeltaz, float maxdeltaphi) const noexcept { return getBinsRect(layer, cls.phi, maxdeltaphi, z1, z2, maxdeltaz); } const int4 getBinsRect(int layer, float phi, float maxdeltaphi, float z1, float z2, float maxdeltaz) const noexcept; void SetRecoChain(o2::gpu::GPUChainITS* chain) { mChain = chain; } void setSmoothing(bool v) { mApplySmoothing = v; } @@ -79,9 +80,9 @@ class TrackerTraits o2::gpu::GPUChainITS* getChain() const { return mChain; } // TimeFrame information forwarding - virtual int getTFNumberOfClusters() const; - virtual int getTFNumberOfTracklets() const; - virtual int getTFNumberOfCells() const; + virtual int getTFNumberOfClusters() const { return mTimeFrame->getNumberOfClusters(); } + virtual int getTFNumberOfTracklets() const { return mTimeFrame->getNumberOfTracklets(); } + virtual int getTFNumberOfCells() const { return mTimeFrame->getNumberOfCells(); } float mBz = 5.f; @@ -95,36 +96,12 @@ class TrackerTraits protected: o2::base::PropagatorImpl::MatCorrType mCorrType = o2::base::PropagatorImpl::MatCorrType::USEMatCorrNONE; o2::gpu::GPUChainITS* mChain = nullptr; - TimeFrame* mTimeFrame; + TimeFrame* mTimeFrame; std::vector mTrkParams; }; -inline void TrackerTraits::initialiseTimeFrame(const int iteration) -{ - mTimeFrame->initialise(iteration, mTrkParams[iteration], mTrkParams[iteration].NLayers); -} - -inline float TrackerTraits::getBz() const -{ - return mBz; -} - -inline void TrackerTraits::UpdateTrackingParameters(const std::vector& trkPars) -{ - mTrkParams = trkPars; -} - -inline const int4 TrackerTraits::getBinsRect(const int layerIndex, float phi, float maxdeltaphi, float z, float maxdeltaz) const noexcept -{ - return getBinsRect(layerIndex, phi, maxdeltaphi, z, z, maxdeltaz); -} - -inline const int4 TrackerTraits::getBinsRect(const Cluster& currentCluster, int layerIndex, float z1, float z2, float maxdeltaz, float maxdeltaphi) const noexcept -{ - return getBinsRect(layerIndex, currentCluster.phi, maxdeltaphi, z1, z2, maxdeltaz); -} - -inline const int4 TrackerTraits::getBinsRect(const int layerIndex, float phi, float maxdeltaphi, float z1, float z2, float maxdeltaz) const noexcept +template +inline const int4 TrackerTraits::getBinsRect(const int layerIndex, float phi, float maxdeltaphi, float z1, float z2, float maxdeltaz) const noexcept { const float zRangeMin = o2::gpu::GPUCommonMath::Min(z1, z2) - maxdeltaz; const float phiRangeMin = (maxdeltaphi > constants::math::Pi) ? 0.f : phi - maxdeltaphi; @@ -142,6 +119,7 @@ inline const int4 TrackerTraits::getBinsRect(const int layerIndex, float phi, fl o2::gpu::GPUCommonMath::Min(mTrkParams[0].ZBins - 1, utils.getZBinIndex(layerIndex, zRangeMax)), // /!\ trkParams can potentially change across iterations utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))}; } + } // namespace its } // namespace o2 diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h index 6eacb94ebb1ea..89b5527f5b2c9 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h @@ -31,6 +31,10 @@ namespace o2::its { class ITSTrackingInterface { + static constexpr int NLayers{7}; + using TrackerTraits7 = TrackerTraits; + using TimeFrame7 = TimeFrame; + public: ITSTrackingInterface(bool isMC, int trgType, @@ -52,7 +56,6 @@ class ITSTrackingInterface } // Task callbacks void initialise(); - template void run(framework::ProcessingContext& pc); void printSummary() const; @@ -60,7 +63,7 @@ class ITSTrackingInterface virtual void finaliseCCDB(framework::ConcreteDataMatcher& matcher, void* obj); // Custom - void setTraitsFromProvider(VertexerTraits*, TrackerTraits*, TimeFrame*); + void setTraitsFromProvider(VertexerTraits*, TrackerTraits7*, TimeFrame7*); void setTrackingMode(TrackingMode mode = TrackingMode::Unset) { if (mode == TrackingMode::Unset) { @@ -72,7 +75,7 @@ class ITSTrackingInterface auto getTracker() const { return mTracker.get(); } auto getVertexer() const { return mVertexer.get(); } - TimeFrame* mTimeFrame = nullptr; + TimeFrame7* mTimeFrame = nullptr; protected: virtual void loadROF(gsl::span& trackROFspan, diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h index ac0cf51921176..56cd286b509ad 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h @@ -35,74 +35,62 @@ #include "GPUCommonLogger.h" -class TTree; - -namespace o2 -{ -namespace its +namespace o2::its { -using TimeFrame = o2::its::TimeFrame; + using Vertex = o2::dataformats::Vertex>; class Vertexer { + static constexpr int NLayers{7}; + using TimeFrame7 = TimeFrame; + using LogFunc = std::function; + public: Vertexer(VertexerTraits* traits); virtual ~Vertexer() = default; Vertexer(const Vertexer&) = delete; Vertexer& operator=(const Vertexer&) = delete; - void adoptTimeFrame(TimeFrame& tf); - std::vector& getVertParameters() const; - void setParameters(std::vector& vertParams); + void adoptTimeFrame(TimeFrame7& tf); + auto& getVertParameters() const { return mTraits->getVertexingParameters(); } + void setParameters(const std::vector& vertParams) { mVertParams = vertParams; } void getGlobalConfiguration(); std::vector exportVertices(); VertexerTraits* getTraits() const { return mTraits; }; - float clustersToVertices(std::function = [](std::string s) { std::cout << s << std::endl; }); - float clustersToVerticesHybrid(std::function = [](std::string s) { std::cout << s << std::endl; }); + float clustersToVertices(LogFunc = [](const std::string& s) { std::cout << s << '\n'; }); void filterMCTracklets(); template void findTracklets(T&&... args); - template - void findTrackletsHybrid(T&&... args); - void findTrivialMCTracklets(); template void validateTracklets(T&&... args); template - void validateTrackletsHybrid(T&&... args); - template void findVertices(T&&... args); - template - void findVerticesHybrid(T&&... args); void findHistVertices(); template void initialiseVertexer(T&&... args); template void initialiseTimeFrame(T&&... args); - template - void initialiseVertexerHybrid(T&&... args); - template - void initialiseTimeFrameHybrid(T&&... args); // Utils - void dumpTraits(); + void dumpTraits() { mTraits->dumpVertexerTraits(); } template - float evaluateTask(void (Vertexer::*)(T...), const char*, std::function logger, T&&... args); - void printEpilog(std::function logger, - bool isHybrid, - const unsigned int trackletN01, const unsigned int trackletN12, const unsigned selectedN, const unsigned int vertexN, - const float initT, const float trackletT, const float selecT, const float vertexT); + float evaluateTask(void (Vertexer::*)(T...), const char*, LogFunc logger, T&&... args); + void printEpilog(LogFunc& logger, + const unsigned int trackletN01, const unsigned int trackletN12, + const unsigned selectedN, const unsigned int vertexN, const float initT, + const float trackletT, const float selecT, const float vertexT); private: std::uint32_t mTimeFrameCounter = 0; VertexerTraits* mTraits = nullptr; /// Observer pointer, not owned by this class - TimeFrame* mTimeFrame = nullptr; /// Observer pointer, not owned by this class + TimeFrame7* mTimeFrame = nullptr; /// Observer pointer, not owned by this class std::vector mVertParams; }; @@ -119,21 +107,6 @@ void Vertexer::findTracklets(T&&... args) mTraits->computeTracklets(std::forward(args)...); } -inline std::vector& Vertexer::getVertParameters() const -{ - return mTraits->getVertexingParameters(); -} - -inline void Vertexer::setParameters(std::vector& vertParams) -{ - mVertParams = vertParams; -} - -inline void Vertexer::dumpTraits() -{ - mTraits->dumpVertexerTraits(); -} - template inline void Vertexer::validateTracklets(T&&... args) { @@ -147,31 +120,7 @@ inline void Vertexer::findVertices(T&&... args) } template -void Vertexer::initialiseVertexerHybrid(T&&... args) -{ - mTraits->initialiseHybrid(std::forward(args)...); -} - -template -void Vertexer::findTrackletsHybrid(T&&... args) -{ - mTraits->computeTrackletsHybrid(std::forward(args)...); -} - -template -inline void Vertexer::validateTrackletsHybrid(T&&... args) -{ - mTraits->computeTrackletMatchingHybrid(std::forward(args)...); -} - -template -inline void Vertexer::findVerticesHybrid(T&&... args) -{ - mTraits->computeVerticesHybrid(std::forward(args)...); -} - -template -float Vertexer::evaluateTask(void (Vertexer::*task)(T...), const char* taskName, std::function logger, +float Vertexer::evaluateTask(void (Vertexer::*task)(T...), const char* taskName, LogFunc logger, T&&... args) { float diff{0.f}; @@ -198,6 +147,5 @@ float Vertexer::evaluateTask(void (Vertexer::*task)(T...), const char* taskName, return diff; } -} // namespace its -} // namespace o2 +} // namespace o2::its #endif diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h index e27d74093aaca..bc7daf5e3ab24 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h @@ -47,33 +47,29 @@ enum class TrackletMode { class VertexerTraits { + static constexpr int NLayers{7}; + using TimeFrame7 = TimeFrame; + public: VertexerTraits() = default; virtual ~VertexerTraits() = default; - GPUhd() static constexpr int4 getEmptyBinsRect() + GPUhdi() static consteval int4 getEmptyBinsRect() { return int4{0, 0, 0, 0}; } GPUhd() const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi); - GPUhd() const int2 getPhiBins(float phi, float deltaPhi); - GPUhd() static const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi, const IndexTableUtils&); GPUhd() static const int2 getPhiBins(float phi, float deltaPhi, const IndexTableUtils&); + GPUhd() const int2 getPhiBins(float phi, float deltaPhi) { return getPhiBins(phi, deltaPhi, mIndexTableUtils); } // virtual vertexer interface virtual void initialise(const TrackingParameters& trackingParams, const int iteration = 0); virtual void computeTracklets(const int iteration = 0); virtual void computeTrackletMatching(const int iteration = 0); virtual void computeVertices(const int iteration = 0); - virtual void adoptTimeFrame(TimeFrame* tf); + virtual void adoptTimeFrame(TimeFrame7* tf) { mTimeFrame = tf; } virtual void updateVertexingParameters(const std::vector& vrtPar, const TimeFrameGPUParameters& gpuTfPar); - // Hybrid - virtual void initialiseHybrid(const TrackingParameters& trackingParams, const int iteration = 0) { initialise(trackingParams, iteration); }; - virtual void computeTrackletsHybrid(const int iteration = 0) { computeTracklets(iteration); }; - virtual void computeTrackletMatchingHybrid(const int iteration = 0) { computeTrackletMatching(iteration); }; - virtual void computeVerticesHybrid(const int iteration = 0) { computeVertices(iteration); }; - virtual void adoptTimeFrameHybrid(TimeFrame* tf) { adoptTimeFrame(tf); }; void computeVerticesInRof(int, gsl::span&, @@ -82,7 +78,7 @@ class VertexerTraits std::array&, std::vector&, std::vector&, - TimeFrame*, + TimeFrame7*, std::vector*, const int iteration = 0); @@ -91,14 +87,14 @@ class VertexerTraits const IndexTableUtils& utils); // utils - std::vector& getVertexingParameters() { return mVrtParams; } - std::vector getVertexingParameters() const { return mVrtParams; } - void setIsGPU(const unsigned char isgpu) { mIsGPU = isgpu; }; + auto& getVertexingParameters() { return mVrtParams; } + auto getVertexingParameters() const { return mVrtParams; } void setVertexingParameters(std::vector& vertParams) { mVrtParams = vertParams; } - unsigned char getIsGPU() const { return mIsGPU; }; void dumpVertexerTraits(); void setNThreads(int n); int getNThreads() const { return mNThreads; } + virtual bool isGPU() const noexcept { return false; } + virtual const char* getName() const noexcept { return "CPU"; } template static std::pair computeMain(const std::vector& elements) @@ -116,25 +112,18 @@ class VertexerTraits } protected: - unsigned char mIsGPU; int mNThreads = 1; std::vector mVrtParams; IndexTableUtils mIndexTableUtils; // Frame related quantities - TimeFrame* mTimeFrame = nullptr; + TimeFrame7* mTimeFrame = nullptr; // observer ptr }; inline void VertexerTraits::initialise(const TrackingParameters& trackingParams, const int iteration) { mTimeFrame->initialise(0, trackingParams, 3, (bool)(!iteration)); // iteration for initialisation must be 0 for correctly resetting the frame, we need to pass the non-reset flag for vertices as well, tho. - setIsGPU(false); -} - -GPUhdi() const int2 VertexerTraits::getPhiBins(float phi, float dPhi) -{ - return VertexerTraits::getPhiBins(phi, dPhi, mIndexTableUtils); } GPUhdi() const int2 VertexerTraits::getPhiBins(float phi, float dPhi, const IndexTableUtils& utils) @@ -169,8 +158,6 @@ GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, c return VertexerTraits::getBinsRect(currentCluster, layerIndex, directionZIntersection, maxdeltaz, maxdeltaphi, mIndexTableUtils); } -inline void VertexerTraits::adoptTimeFrame(TimeFrame* tf) { mTimeFrame = tf; } - } // namespace its } // namespace o2 #endif diff --git a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx index 506656e2777d1..80a6ef481fde6 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TimeFrame.cxx @@ -50,9 +50,7 @@ float Sq(float v) } // namespace -namespace o2 -{ -namespace its +namespace o2::its { constexpr float DefClusErrorRow = o2::itsmft::SegmentationAlpide::PitchRow * 0.5; @@ -60,21 +58,20 @@ constexpr float DefClusErrorCol = o2::itsmft::SegmentationAlpide::PitchCol * 0.5 constexpr float DefClusError2Row = DefClusErrorRow * DefClusErrorRow; constexpr float DefClusError2Col = DefClusErrorCol * DefClusErrorCol; -TimeFrame::TimeFrame(int nLayers) +template +TimeFrame::TimeFrame() +{ + resetVectors(); +} + +template +TimeFrame::~TimeFrame() { - mMinR.resize(nLayers, 10000.); - mMaxR.resize(nLayers, -1.); - mClusters.resize(nLayers); - mUnsortedClusters.resize(nLayers); - mTrackingFrameInfo.resize(nLayers); - mClusterExternalIndices.resize(nLayers); - mUsedClusters.resize(nLayers); - mROFramesClusters.resize(nLayers, {0}); /// TBC: if resetting the timeframe is required, then this has to be done - mNClustersPerROF.resize(nLayers); - mTrackletsIndexROF.resize(2, {0}); + resetVectors(); } -void TimeFrame::addPrimaryVertices(const std::vector& vertices) +template +void TimeFrame::addPrimaryVertices(const std::vector& vertices) { for (const auto& vertex : vertices) { mPrimaryVertices.emplace_back(vertex); @@ -88,17 +85,20 @@ void TimeFrame::addPrimaryVertices(const std::vector& vertices) mROFramesPV.push_back(mPrimaryVertices.size()); } -void TimeFrame::addPrimaryVertices(const std::vector& vertices, const int rofId, const int iteration) +template +void TimeFrame::addPrimaryVertices(const std::vector& vertices, const int rofId, const int iteration) { addPrimaryVertices(gsl::span(vertices), rofId, iteration); } -void TimeFrame::addPrimaryVerticesLabels(std::vector>& labels) +template +void TimeFrame::addPrimaryVerticesLabels(std::vector>& labels) { mVerticesMCRecInfo.insert(mVerticesMCRecInfo.end(), labels.begin(), labels.end()); } -void TimeFrame::addPrimaryVerticesInROF(const std::vector& vertices, const int rofId, const int iteration) +template +void TimeFrame::addPrimaryVerticesInROF(const std::vector& vertices, const int rofId, const int iteration) { mPrimaryVertices.insert(mPrimaryVertices.begin() + mROFramesPV[rofId], vertices.begin(), vertices.end()); for (int i = rofId + 1; i < mROFramesPV.size(); ++i) { @@ -107,12 +107,14 @@ void TimeFrame::addPrimaryVerticesInROF(const std::vector& vertices, con mTotVertPerIteration[iteration] += vertices.size(); } -void TimeFrame::addPrimaryVerticesLabelsInROF(const std::vector>& labels, const int rofId) +template +void TimeFrame::addPrimaryVerticesLabelsInROF(const std::vector>& labels, const int rofId) { mVerticesMCRecInfo.insert(mVerticesMCRecInfo.begin() + mROFramesPV[rofId], labels.begin(), labels.end()); } -void TimeFrame::addPrimaryVertices(const gsl::span& vertices, const int rofId, const int iteration) +template +void TimeFrame::addPrimaryVertices(const gsl::span& vertices, const int rofId, const int iteration) { std::vector futureVertices; for (const auto& vertex : vertices) { @@ -133,57 +135,18 @@ void TimeFrame::addPrimaryVertices(const gsl::span& vertices, cons } } mROFramesPV.push_back(mPrimaryVertices.size()); // current rof must have number of vertices up to present - if (futureVertices.size()) { // append future vertices. In the last rofId we cannot have ones from the next, so we are never here. - for (auto& vertex : futureVertices) { - mPrimaryVertices.emplace_back(vertex); - mTotVertPerIteration[iteration]++; - } - } -} - -int TimeFrame::loadROFrameData(const o2::itsmft::ROFRecord& rof, gsl::span clusters, - const dataformats::MCTruthContainer* mcLabels) -{ - GeometryTGeo* geom = GeometryTGeo::Instance(); - geom->fillMatrixCache(o2::math_utils::bit2Mask(o2::math_utils::TransformType::T2L, o2::math_utils::TransformType::L2G)); - int clusterId{0}; - - auto first = rof.getFirstEntry(); - auto clusters_in_frame = rof.getROFData(clusters); - for (auto& c : clusters_in_frame) { - int layer = geom->getLayer(c.getSensorID()); - - /// Clusters are stored in the tracking frame - auto xyz = c.getXYZGloRot(*geom); - addTrackingFrameInfoToLayer(layer, xyz.x(), xyz.y(), xyz.z(), c.getX(), geom->getSensorRefAlpha(c.getSensorID()), - std::array{c.getY(), c.getZ()}, - std::array{c.getSigmaY2(), c.getSigmaYZ(), c.getSigmaZ2()}); - - /// Rotate to the global frame - addClusterToLayer(layer, xyz.x(), xyz.y(), xyz.z(), mUnsortedClusters[layer].size()); - addClusterExternalIndexToLayer(layer, first + clusterId); - clusterId++; - } - - for (unsigned int iL{0}; iL < mUnsortedClusters.size(); ++iL) { - mNClustersPerROF[iL].push_back(mUnsortedClusters[iL].size() - mROFramesClusters[iL].back()); - mROFramesClusters[iL].push_back(mUnsortedClusters[iL].size()); - if (iL < 2) { - mTrackletsIndexROF[iL].push_back(mUnsortedClusters[1].size()); // Tracklets used in vertexer are always computed starting from L1 - } - } - if (mcLabels) { - mClusterLabels = mcLabels; + for (auto& vertex : futureVertices) { + mPrimaryVertices.emplace_back(vertex); + mTotVertPerIteration[iteration]++; } - mNrof++; - return clusters_in_frame.size(); } -int TimeFrame::loadROFrameData(gsl::span rofs, - gsl::span clusters, - gsl::span::iterator& pattIt, - const itsmft::TopologyDictionary* dict, - const dataformats::MCTruthContainer* mcLabels) +template +int TimeFrame::loadROFrameData(gsl::span rofs, + gsl::span clusters, + gsl::span::iterator& pattIt, + const itsmft::TopologyDictionary* dict, + const dataformats::MCTruthContainer* mcLabels) { for (int iLayer{0}; iLayer < mUnsortedClusters.size(); ++iLayer) { deepVectorClear(mUnsortedClusters[iLayer]); @@ -260,22 +223,14 @@ int TimeFrame::loadROFrameData(gsl::span rofs, mNTrackletsPerClusterSum[i].resize(mUnsortedClusters[1].size() + 1); // Exc sum "prepends" a 0 } - if (mcLabels) { + if (mcLabels != nullptr) { mClusterLabels = mcLabels; } return mNrof; } -int TimeFrame::getTotalClusters() const -{ - size_t totalClusters{0}; - for (auto& clusters : mUnsortedClusters) { - totalClusters += clusters.size(); - } - return int(totalClusters); -} - -void TimeFrame::prepareClusters(const TrackingParameters& trkParam, const int maxLayers) +template +void TimeFrame::prepareClusters(const TrackingParameters& trkParam, const int maxLayers) { std::vector cHelper; std::vector clsPerBin(trkParam.PhiBins * trkParam.ZBins, 0); @@ -342,7 +297,8 @@ void TimeFrame::prepareClusters(const TrackingParameters& trkParam, const int ma } } -void TimeFrame::initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers, bool resetVertices) +template +void TimeFrame::initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers, bool resetVertices) { if (iteration == 0) { if (maxLayers < trkParam.NLayers && resetVertices) { @@ -451,7 +407,8 @@ void TimeFrame::initialise(const int iteration, const TrackingParameters& trkPar } } -unsigned long TimeFrame::getArtefactsMemory() +template +unsigned long TimeFrame::getArtefactsMemory() { unsigned long size{0}; for (auto& trkl : mTracklets) { @@ -463,10 +420,11 @@ unsigned long TimeFrame::getArtefactsMemory() for (auto& cellsN : mCellsNeighbours) { size += sizeof(int) * cellsN.size(); } - return size + sizeof(Road<5>) * mRoads.size(); + return size + sizeof(Road) * mRoads.size(); } -void TimeFrame::fillPrimaryVerticesXandAlpha() +template +void TimeFrame::fillPrimaryVerticesXandAlpha() { if (mPValphaX.size()) { mPValphaX.clear(); @@ -477,7 +435,8 @@ void TimeFrame::fillPrimaryVerticesXandAlpha() } } -void TimeFrame::computeTrackletsPerROFScans() +template +void TimeFrame::computeTrackletsPerROFScans() { for (ushort iLayer = 0; iLayer < 2; ++iLayer) { for (unsigned int iRof{0}; iRof < mNrof; ++iRof) { @@ -490,7 +449,8 @@ void TimeFrame::computeTrackletsPerROFScans() } } -void TimeFrame::checkTrackletLUTs() +template +void TimeFrame::checkTrackletLUTs() { for (uint32_t iLayer{0}; iLayer < getTracklets().size(); ++iLayer) { int prev{-1}; @@ -522,21 +482,27 @@ void TimeFrame::checkTrackletLUTs() } } -void TimeFrame::resizeVectors(int nLayers) +template +void TimeFrame::resetVectors() { - mMinR.resize(nLayers, 10000.); - mMaxR.resize(nLayers, -1.); - mClusters.resize(nLayers); - mUnsortedClusters.resize(nLayers); - mTrackingFrameInfo.resize(nLayers); - mClusterExternalIndices.resize(nLayers); - mUsedClusters.resize(nLayers); - mROFramesClusters.resize(nLayers, {0}); - mNClustersPerROF.resize(nLayers); - mTrackletsIndexROF.resize(2, {0}); + mMinR.fill(10000.); + mMaxR.fill(-1.); + for (int iLayers{nLayers}; iLayers--;) { + mClusters[iLayers].clear(); + mUnsortedClusters[iLayers].clear(); + mTrackingFrameInfo[iLayers].clear(); + mClusterExternalIndices[iLayers].clear(); + mUsedClusters[iLayers].clear(); + mROFramesClusters[iLayers].clear(); + mNClustersPerROF[iLayers].clear(); + } + for (int i{2}; i--;) { + mTrackletsIndexROF[i].clear(); + } } -void TimeFrame::printTrackletLUTonLayer(int i) +template +void TimeFrame::printTrackletLUTonLayer(int i) { std::cout << "--------" << std::endl << "Tracklet LUT " << i << std::endl; @@ -547,7 +513,8 @@ void TimeFrame::printTrackletLUTonLayer(int i) << std::endl; } -void TimeFrame::printCellLUTonLayer(int i) +template +void TimeFrame::printCellLUTonLayer(int i) { std::cout << "--------" << std::endl << "Cell LUT " << i << std::endl; @@ -558,21 +525,24 @@ void TimeFrame::printCellLUTonLayer(int i) << std::endl; } -void TimeFrame::printTrackletLUTs() +template +void TimeFrame::printTrackletLUTs() { for (unsigned int i{0}; i < mTrackletsLookupTable.size(); ++i) { printTrackletLUTonLayer(i); } } -void TimeFrame::printCellLUTs() +template +void TimeFrame::printCellLUTs() { for (unsigned int i{0}; i < mCellsLookupTable.size(); ++i) { printCellLUTonLayer(i); } } -void TimeFrame::printVertices() +template +void TimeFrame::printVertices() { std::cout << "Vertices in ROF (nROF = " << mNrof << ", lut size = " << mROFramesPV.size() << ")" << std::endl; for (unsigned int iR{0}; iR < mROFramesPV.size(); ++iR) { @@ -585,7 +555,8 @@ void TimeFrame::printVertices() std::cout << "--------" << std::endl; } -void TimeFrame::printROFoffsets() +template +void TimeFrame::printROFoffsets() { std::cout << "--------" << std::endl; for (unsigned int iLayer{0}; iLayer < mROFramesClusters.size(); ++iLayer) { @@ -597,7 +568,8 @@ void TimeFrame::printROFoffsets() } } -void TimeFrame::printNClsPerROF() +template +void TimeFrame::printNClsPerROF() { std::cout << "--------" << std::endl; for (unsigned int iLayer{0}; iLayer < mNClustersPerROF.size(); ++iLayer) { @@ -609,7 +581,8 @@ void TimeFrame::printNClsPerROF() } } -void TimeFrame::printSliceInfo(const int startROF, const int sliceSize) +template +void TimeFrame::printSliceInfo(const int startROF, const int sliceSize) { std::cout << "Dumping slice of " << sliceSize << " rofs:" << std::endl; for (int iROF{startROF}; iROF < startROF + sliceSize; ++iROF) { @@ -625,5 +598,6 @@ void TimeFrame::printSliceInfo(const int startROF, const int sliceSize) } } -} // namespace its -} // namespace o2 +template class TimeFrame<7>; + +} // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx index c23ba0576c625..c452e20f7f707 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Tracker.cxx @@ -30,13 +30,11 @@ #include #include -namespace o2 -{ -namespace its +namespace o2::its { using o2::its::constants::GB; -Tracker::Tracker(o2::its::TrackerTraits* traits) : mTraits(traits) +Tracker::Tracker(TrackerTraits7* traits) : mTraits(traits) { /// Initialise standard configuration with 1 iteration mTrkParams.resize(1); @@ -47,7 +45,7 @@ void Tracker::clustersToTracks(LogFunc logger, LogFunc error) LogFunc evalLog = [](const std::string&) {}; double total{0}; - mTraits->UpdateTrackingParameters(mTrkParams); + mTraits->updateTrackingParameters(mTrkParams); int maxNvertices{-1}; if (mTrkParams[0].PerPrimaryVertexProcessing) { for (int iROF{0}; iROF < mTimeFrame->getNrof(); ++iROF) { @@ -149,41 +147,6 @@ void Tracker::clustersToTracks(LogFunc logger, LogFunc error) mTotalTime += total; } -void Tracker::initialiseTimeFrame(int& iteration) -{ - mTraits->initialiseTimeFrame(iteration); -} - -void Tracker::computeTracklets(int& iteration, int& iROFslice, int& iVertex) -{ - mTraits->computeLayerTracklets(iteration, iROFslice, iVertex); -} - -void Tracker::computeCells(int& iteration) -{ - mTraits->computeLayerCells(iteration); -} - -void Tracker::findCellsNeighbours(int& iteration) -{ - mTraits->findCellsNeighbours(iteration); -} - -void Tracker::findRoads(int& iteration) -{ - mTraits->findRoads(iteration); -} - -void Tracker::extendTracks(int& iteration) -{ - mTraits->extendTracks(iteration); -} - -void Tracker::findShortPrimaries() -{ - mTraits->findShortPrimaries(); -} - void Tracker::computeRoadsMClabels() { /// Moore's Voting Algorithm @@ -427,41 +390,15 @@ void Tracker::getGlobalConfiguration() } } -void Tracker::adoptTimeFrame(TimeFrame& tf) +void Tracker::adoptTimeFrame(TimeFrame7& tf) { mTimeFrame = &tf; mTraits->adoptTimeFrame(&tf); } -void Tracker::setBz(float bz) -{ - mTraits->setBz(bz); -} - -void Tracker::setCorrType(const o2::base::PropagatorImpl::MatCorrType type) -{ - mTraits->setCorrType(type); -} - -bool Tracker::isMatLUT() const -{ - return mTraits->isMatLUT(); -} - -void Tracker::setNThreads(int n) -{ - mTraits->setNThreads(n); -} - -int Tracker::getNThreads() const -{ - return mTraits->getNThreads(); -} - void Tracker::printSummary() const { LOGP(info, "Tracker summary: Processed {} TFs (dropped {}) in TOT={:.2f} s, AVG/TF={:.2f} s", mTimeFrameCounter, mNumberOfDroppedTFs, mTotalTime * 1.e-3, mTotalTime * 1.e-3 / ((mTimeFrameCounter > 0) ? (double)mTimeFrameCounter : -1.0)); } -} // namespace its -} // namespace o2 +} // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 987e8e3128fb4..da0eeb52156be 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -46,43 +46,41 @@ inline float Sq(float q) } } // namespace -namespace o2 -{ -namespace its +namespace o2::its { constexpr int debugLevel{0}; -void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, int iVertex) +template +void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, int iVertex) { - TimeFrame* tf = mTimeFrame; - #ifdef OPTIMISATION_OUTPUT static int iter{0}; std::ofstream off(std::format("tracklets{}.txt", iter++)); #endif for (int iLayer = 0; iLayer < mTrkParams[iteration].TrackletsPerRoad(); ++iLayer) { - tf->getTracklets()[iLayer].clear(); - tf->getTrackletsLabel(iLayer).clear(); + mTimeFrame->getTracklets()[iLayer].clear(); + mTimeFrame->getTrackletsLabel(iLayer).clear(); if (iLayer > 0) { - std::fill(tf->getTrackletsLookupTable()[iLayer - 1].begin(), tf->getTrackletsLookupTable()[iLayer - 1].end(), 0); + std::fill(mTimeFrame->getTrackletsLookupTable()[iLayer - 1].begin(), + mTimeFrame->getTrackletsLookupTable()[iLayer - 1].end(), 0); } } const Vertex diamondVert({mTrkParams[iteration].Diamond[0], mTrkParams[iteration].Diamond[1], mTrkParams[iteration].Diamond[2]}, {25.e-6f, 0.f, 0.f, 25.e-6f, 0.f, 36.f}, 1, 1.f); gsl::span diamondSpan(&diamondVert, 1); int startROF{mTrkParams[iteration].nROFsPerIterations > 0 ? iROFslice * mTrkParams[iteration].nROFsPerIterations : 0}; - int endROF{gpu::GPUCommonMath::Min(mTrkParams[iteration].nROFsPerIterations > 0 ? (iROFslice + 1) * mTrkParams[iteration].nROFsPerIterations + mTrkParams[iteration].DeltaROF : tf->getNrof(), tf->getNrof())}; + int endROF{o2::gpu::GPUCommonMath::Min(mTrkParams[iteration].nROFsPerIterations > 0 ? (iROFslice + 1) * mTrkParams[iteration].nROFsPerIterations + mTrkParams[iteration].DeltaROF : mTimeFrame->getNrof(), mTimeFrame->getNrof())}; for (int rof0{startROF}; rof0 < endROF; ++rof0) { - gsl::span primaryVertices = mTrkParams[iteration].UseDiamond ? diamondSpan : tf->getPrimaryVertices(rof0); + gsl::span primaryVertices = mTrkParams[iteration].UseDiamond ? diamondSpan : mTimeFrame->getPrimaryVertices(rof0); const int startVtx{iVertex >= 0 ? iVertex : 0}; const int endVtx{iVertex >= 0 ? o2::gpu::CAMath::Min(iVertex + 1, static_cast(primaryVertices.size())) : static_cast(primaryVertices.size())}; int minRof = o2::gpu::CAMath::Max(startROF, rof0 - mTrkParams[iteration].DeltaROF); int maxRof = o2::gpu::CAMath::Min(endROF - 1, rof0 + mTrkParams[iteration].DeltaROF); #pragma omp parallel for num_threads(mNThreads) for (int iLayer = 0; iLayer < mTrkParams[iteration].TrackletsPerRoad(); ++iLayer) { - gsl::span layer0 = tf->getClustersOnLayer(rof0, iLayer); + gsl::span layer0 = mTimeFrame->getClustersOnLayer(rof0, iLayer); if (layer0.empty()) { continue; } @@ -91,9 +89,9 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in const int currentLayerClustersNum{static_cast(layer0.size())}; for (int iCluster{0}; iCluster < currentLayerClustersNum; ++iCluster) { const Cluster& currentCluster{layer0[iCluster]}; - const int currentSortedIndex{tf->getSortedIndex(rof0, iLayer, iCluster)}; + const int currentSortedIndex{mTimeFrame->getSortedIndex(rof0, iLayer, iCluster)}; - if (tf->isClusterUsed(iLayer, currentCluster.clusterId)) { + if (mTimeFrame->isClusterUsed(iLayer, currentCluster.clusterId)) { continue; } const float inverseR0{1.f / currentCluster.radius}; @@ -103,18 +101,17 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in if (primaryVertex.isFlagSet(2) && iteration != 3) { continue; } - const float resolution = o2::gpu::CAMath::Sqrt(Sq(mTrkParams[iteration].PVres) / primaryVertex.getNContributors() + Sq(tf->getPositionResolution(iLayer))); + const float resolution = o2::gpu::CAMath::Sqrt(Sq(mTrkParams[iteration].PVres) / primaryVertex.getNContributors() + Sq(mTimeFrame->getPositionResolution(iLayer))); const float tanLambda{(currentCluster.zCoordinate - primaryVertex.getZ()) * inverseR0}; - const float zAtRmin{tanLambda * (tf->getMinR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate}; - const float zAtRmax{tanLambda * (tf->getMaxR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate}; + const float zAtRmin{tanLambda * (mTimeFrame->getMinR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate}; + const float zAtRmax{tanLambda * (mTimeFrame->getMaxR(iLayer + 1) - currentCluster.radius) + currentCluster.zCoordinate}; const float sqInverseDeltaZ0{1.f / (Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + 2.e-8f)}; /// protecting from overflows adding the detector resolution - const float sigmaZ{o2::gpu::CAMath::Sqrt(Sq(resolution) * Sq(tanLambda) * ((Sq(inverseR0) + sqInverseDeltaZ0) * Sq(meanDeltaR) + 1.f) + Sq(meanDeltaR * tf->getMSangle(iLayer)))}; + const float sigmaZ{o2::gpu::CAMath::Sqrt(Sq(resolution) * Sq(tanLambda) * ((Sq(inverseR0) + sqInverseDeltaZ0) * Sq(meanDeltaR) + 1.f) + Sq(meanDeltaR * mTimeFrame->getMSangle(iLayer)))}; - const int4 selectedBinsRect{getBinsRect(currentCluster, iLayer + 1, zAtRmin, zAtRmax, - sigmaZ * mTrkParams[iteration].NSigmaCut, tf->getPhiCut(iLayer))}; + const int4 selectedBinsRect{getBinsRect(currentCluster, iLayer + 1, zAtRmin, zAtRmax, sigmaZ * mTrkParams[iteration].NSigmaCut, mTimeFrame->getPhiCut(iLayer))}; if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) { continue; } @@ -126,46 +123,46 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in } for (int rof1{minRof}; rof1 <= maxRof; ++rof1) { - gsl::span layer1 = tf->getClustersOnLayer(rof1, iLayer + 1); + auto layer1 = mTimeFrame->getClustersOnLayer(rof1, iLayer + 1); if (layer1.empty()) { continue; } for (int iPhiCount{0}; iPhiCount < phiBinsNum; iPhiCount++) { int iPhiBin = (selectedBinsRect.y + iPhiCount) % mTrkParams[iteration].PhiBins; - const int firstBinIndex{tf->mIndexTableUtils.getBinIndex(selectedBinsRect.x, iPhiBin)}; + const int firstBinIndex{mTimeFrame->mIndexTableUtils.getBinIndex(selectedBinsRect.x, iPhiBin)}; const int maxBinIndex{firstBinIndex + selectedBinsRect.z - selectedBinsRect.x + 1}; if constexpr (debugLevel) { - if (firstBinIndex < 0 || firstBinIndex > tf->getIndexTable(rof1, iLayer + 1).size() || - maxBinIndex < 0 || maxBinIndex > tf->getIndexTable(rof1, iLayer + 1).size()) { - std::cout << iLayer << "\t" << iCluster << "\t" << zAtRmin << "\t" << zAtRmax << "\t" << sigmaZ * mTrkParams[iteration].NSigmaCut << "\t" << tf->getPhiCut(iLayer) << std::endl; + if (firstBinIndex < 0 || firstBinIndex > mTimeFrame->getIndexTable(rof1, iLayer + 1).size() || + maxBinIndex < 0 || maxBinIndex > mTimeFrame->getIndexTable(rof1, iLayer + 1).size()) { + std::cout << iLayer << "\t" << iCluster << "\t" << zAtRmin << "\t" << zAtRmax << "\t" << sigmaZ * mTrkParams[iteration].NSigmaCut << "\t" << mTimeFrame->getPhiCut(iLayer) << std::endl; std::cout << currentCluster.zCoordinate << "\t" << primaryVertex.getZ() << "\t" << currentCluster.radius << std::endl; - std::cout << tf->getMinR(iLayer + 1) << "\t" << currentCluster.radius << "\t" << currentCluster.zCoordinate << std::endl; + std::cout << mTimeFrame->getMinR(iLayer + 1) << "\t" << currentCluster.radius << "\t" << currentCluster.zCoordinate << std::endl; std::cout << "Illegal access to IndexTable " << firstBinIndex << "\t" << maxBinIndex << "\t" << selectedBinsRect.z << "\t" << selectedBinsRect.x << std::endl; exit(1); } } - const int firstRowClusterIndex = tf->getIndexTable(rof1, iLayer + 1)[firstBinIndex]; - const int maxRowClusterIndex = tf->getIndexTable(rof1, iLayer + 1)[maxBinIndex]; + const int firstRowClusterIndex = mTimeFrame->getIndexTable(rof1, iLayer + 1)[firstBinIndex]; + const int maxRowClusterIndex = mTimeFrame->getIndexTable(rof1, iLayer + 1)[maxBinIndex]; for (int iNextCluster{firstRowClusterIndex}; iNextCluster < maxRowClusterIndex; ++iNextCluster) { if (iNextCluster >= (int)layer1.size()) { break; } const Cluster& nextCluster{layer1[iNextCluster]}; - if (tf->isClusterUsed(iLayer + 1, nextCluster.clusterId)) { + if (mTimeFrame->isClusterUsed(iLayer + 1, nextCluster.clusterId)) { continue; } - const float deltaPhi{gpu::GPUCommonMath::Abs(currentCluster.phi - nextCluster.phi)}; - const float deltaZ{gpu::GPUCommonMath::Abs(tanLambda * (nextCluster.radius - currentCluster.radius) + - currentCluster.zCoordinate - nextCluster.zCoordinate)}; + const float deltaPhi{o2::gpu::GPUCommonMath::Abs(currentCluster.phi - nextCluster.phi)}; + const float deltaZ{o2::gpu::GPUCommonMath::Abs(tanLambda * (nextCluster.radius - currentCluster.radius) + + currentCluster.zCoordinate - nextCluster.zCoordinate)}; #ifdef OPTIMISATION_OUTPUT MCCompLabel label; int currentId{currentCluster.clusterId}; int nextId{nextCluster.clusterId}; - for (auto& lab1 : tf->getClusterLabels(iLayer, currentId)) { - for (auto& lab2 : tf->getClusterLabels(iLayer + 1, nextId)) { + for (auto& lab1 : mTimeFrame->getClusterLabels(iLayer, currentId)) { + for (auto& lab2 : mTimeFrame->getClusterLabels(iLayer + 1, nextId)) { if (lab1 == lab2 && lab1.isValid()) { label = lab1; break; @@ -179,16 +176,16 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in #endif if (deltaZ / sigmaZ < mTrkParams[iteration].NSigmaCut && - (deltaPhi < tf->getPhiCut(iLayer) || - gpu::GPUCommonMath::Abs(deltaPhi - constants::math::TwoPi) < tf->getPhiCut(iLayer))) { + (deltaPhi < mTimeFrame->getPhiCut(iLayer) || + o2::gpu::GPUCommonMath::Abs(deltaPhi - constants::math::TwoPi) < mTimeFrame->getPhiCut(iLayer))) { if (iLayer > 0) { - tf->getTrackletsLookupTable()[iLayer - 1][currentSortedIndex]++; + mTimeFrame->getTrackletsLookupTable()[iLayer - 1][currentSortedIndex]++; } const float phi{o2::gpu::GPUCommonMath::ATan2(currentCluster.yCoordinate - nextCluster.yCoordinate, currentCluster.xCoordinate - nextCluster.xCoordinate)}; const float tanL{(currentCluster.zCoordinate - nextCluster.zCoordinate) / (currentCluster.radius - nextCluster.radius)}; - tf->getTracklets()[iLayer].emplace_back(currentSortedIndex, tf->getSortedIndex(rof1, iLayer + 1, iNextCluster), tanL, phi, rof0, rof1); + mTimeFrame->getTracklets()[iLayer].emplace_back(currentSortedIndex, mTimeFrame->getSortedIndex(rof1, iLayer + 1, iNextCluster), tanL, phi, rof0, rof1); } } } @@ -197,19 +194,19 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in } } } - if (!tf->checkMemory(mTrkParams[iteration].MaxMemory)) { + if (!mTimeFrame->checkMemory(mTrkParams[iteration].MaxMemory)) { return; } #pragma omp parallel for num_threads(mNThreads) for (int iLayer = 0; iLayer < mTrkParams[iteration].CellsPerRoad(); ++iLayer) { /// Sort tracklets - auto& trkl{tf->getTracklets()[iLayer + 1]}; + auto& trkl{mTimeFrame->getTracklets()[iLayer + 1]}; std::sort(trkl.begin(), trkl.end(), [](const Tracklet& a, const Tracklet& b) { return a.firstClusterIndex < b.firstClusterIndex || (a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex < b.secondClusterIndex); }); /// Remove duplicates - auto& lut{tf->getTrackletsLookupTable()[iLayer]}; + auto& lut{mTimeFrame->getTrackletsLookupTable()[iLayer]}; int id0{-1}, id1{-1}; std::vector newTrk; newTrk.reserve(trkl.size()); @@ -229,30 +226,30 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in lut.push_back(trkl.size()); } /// Layer 0 is done outside the loop - std::sort(tf->getTracklets()[0].begin(), tf->getTracklets()[0].end(), [](const Tracklet& a, const Tracklet& b) { + std::sort(mTimeFrame->getTracklets()[0].begin(), mTimeFrame->getTracklets()[0].end(), [](const Tracklet& a, const Tracklet& b) { return a.firstClusterIndex < b.firstClusterIndex || (a.firstClusterIndex == b.firstClusterIndex && a.secondClusterIndex < b.secondClusterIndex); }); int id0{-1}, id1{-1}; std::vector newTrk; - newTrk.reserve(tf->getTracklets()[0].size()); - for (auto& trk : tf->getTracklets()[0]) { + newTrk.reserve(mTimeFrame->getTracklets()[0].size()); + for (auto& trk : mTimeFrame->getTracklets()[0]) { if (trk.firstClusterIndex != id0 || trk.secondClusterIndex != id1) { id0 = trk.firstClusterIndex; id1 = trk.secondClusterIndex; newTrk.push_back(trk); } } - tf->getTracklets()[0].swap(newTrk); + mTimeFrame->getTracklets()[0].swap(newTrk); /// Create tracklets labels - if (tf->hasMCinformation()) { + if (mTimeFrame->hasMCinformation()) { for (int iLayer{0}; iLayer < mTrkParams[iteration].TrackletsPerRoad(); ++iLayer) { - for (auto& trk : tf->getTracklets()[iLayer]) { + for (auto& trk : mTimeFrame->getTracklets()[iLayer]) { MCCompLabel label; - int currentId{tf->getClusters()[iLayer][trk.firstClusterIndex].clusterId}; - int nextId{tf->getClusters()[iLayer + 1][trk.secondClusterIndex].clusterId}; - for (auto& lab1 : tf->getClusterLabels(iLayer, currentId)) { - for (auto& lab2 : tf->getClusterLabels(iLayer + 1, nextId)) { + int currentId{mTimeFrame->getClusters()[iLayer][trk.firstClusterIndex].clusterId}; + int nextId{mTimeFrame->getClusters()[iLayer + 1][trk.secondClusterIndex].clusterId}; + for (auto& lab1 : mTimeFrame->getClusterLabels(iLayer, currentId)) { + for (auto& lab2 : mTimeFrame->getClusterLabels(iLayer + 1, nextId)) { if (lab1 == lab2 && lab1.isValid()) { label = lab1; break; @@ -262,13 +259,14 @@ void TrackerTraits::computeLayerTracklets(const int iteration, int iROFslice, in break; } } - tf->getTrackletsLabel(iLayer).emplace_back(label); + mTimeFrame->getTrackletsLabel(iLayer).emplace_back(label); } } } } -void TrackerTraits::computeLayerCells(const int iteration) +template +void TrackerTraits::computeLayerCells(const int iteration) { #ifdef OPTIMISATION_OUTPUT static int iter{0}; @@ -283,12 +281,11 @@ void TrackerTraits::computeLayerCells(const int iteration) } } - TimeFrame* tf = mTimeFrame; #pragma omp parallel for num_threads(mNThreads) for (int iLayer = 0; iLayer < mTrkParams[iteration].CellsPerRoad(); ++iLayer) { - if (tf->getTracklets()[iLayer + 1].empty() || - tf->getTracklets()[iLayer].empty()) { + if (mTimeFrame->getTracklets()[iLayer + 1].empty() || + mTimeFrame->getTracklets()[iLayer].empty()) { continue; } @@ -296,29 +293,29 @@ void TrackerTraits::computeLayerCells(const int iteration) float resolution{o2::gpu::CAMath::Sqrt(0.5f * (mTrkParams[iteration].SystErrorZ2[iLayer] + mTrkParams[iteration].SystErrorZ2[iLayer + 1] + mTrkParams[iteration].SystErrorZ2[iLayer + 2] + mTrkParams[iteration].SystErrorY2[iLayer] + mTrkParams[iteration].SystErrorY2[iLayer + 1] + mTrkParams[iteration].SystErrorY2[iLayer + 2])) / mTrkParams[iteration].LayerResolution[iLayer]}; resolution = resolution > 1.e-12 ? resolution : 1.f; #endif - const int currentLayerTrackletsNum{static_cast(tf->getTracklets()[iLayer].size())}; + const int currentLayerTrackletsNum{static_cast(mTimeFrame->getTracklets()[iLayer].size())}; for (int iTracklet{0}; iTracklet < currentLayerTrackletsNum; ++iTracklet) { - const Tracklet& currentTracklet{tf->getTracklets()[iLayer][iTracklet]}; + const Tracklet& currentTracklet{mTimeFrame->getTracklets()[iLayer][iTracklet]}; const int nextLayerClusterIndex{currentTracklet.secondClusterIndex}; const int nextLayerFirstTrackletIndex{ - tf->getTrackletsLookupTable()[iLayer][nextLayerClusterIndex]}; + mTimeFrame->getTrackletsLookupTable()[iLayer][nextLayerClusterIndex]}; const int nextLayerLastTrackletIndex{ - tf->getTrackletsLookupTable()[iLayer][nextLayerClusterIndex + 1]}; + mTimeFrame->getTrackletsLookupTable()[iLayer][nextLayerClusterIndex + 1]}; if (nextLayerFirstTrackletIndex == nextLayerLastTrackletIndex) { continue; } for (int iNextTracklet{nextLayerFirstTrackletIndex}; iNextTracklet < nextLayerLastTrackletIndex; ++iNextTracklet) { - if (tf->getTracklets()[iLayer + 1][iNextTracklet].firstClusterIndex != nextLayerClusterIndex) { + if (mTimeFrame->getTracklets()[iLayer + 1][iNextTracklet].firstClusterIndex != nextLayerClusterIndex) { break; } - const Tracklet& nextTracklet{tf->getTracklets()[iLayer + 1][iNextTracklet]}; + const Tracklet& nextTracklet{mTimeFrame->getTracklets()[iLayer + 1][iNextTracklet]}; const float deltaTanLambda{std::abs(currentTracklet.tanLambda - nextTracklet.tanLambda)}; #ifdef OPTIMISATION_OUTPUT - bool good{tf->getTrackletsLabel(iLayer)[iTracklet] == tf->getTrackletsLabel(iLayer + 1)[iNextTracklet]}; + bool good{mTimeFrame->getTrackletsLabel(iLayer)[iTracklet] == mTimeFrame->getTrackletsLabel(iLayer + 1)[iNextTracklet]}; float signedDelta{currentTracklet.tanLambda - nextTracklet.tanLambda}; off << std::format("{}\t{:d}\t{}\t{}\t{}\t{}", iLayer, good, signedDelta, signedDelta / (mTrkParams[iteration].CellDeltaTanLambdaSigma), tanLambda, resolution) << std::endl; #endif @@ -367,41 +364,42 @@ void TrackerTraits::computeLayerCells(const int iteration) if (!good) { continue; } - if (iLayer > 0 && (int)tf->getCellsLookupTable()[iLayer - 1].size() <= iTracklet) { - tf->getCellsLookupTable()[iLayer - 1].resize(iTracklet + 1, tf->getCells()[iLayer].size()); + if (iLayer > 0 && (int)mTimeFrame->getCellsLookupTable()[iLayer - 1].size() <= iTracklet) { + mTimeFrame->getCellsLookupTable()[iLayer - 1].resize(iTracklet + 1, mTimeFrame->getCells()[iLayer].size()); } - tf->getCells()[iLayer].emplace_back(iLayer, clusId[0], clusId[1], clusId[2], - iTracklet, iNextTracklet, track, chi2); + mTimeFrame->getCells()[iLayer].emplace_back(iLayer, clusId[0], clusId[1], clusId[2], + iTracklet, iNextTracklet, track, chi2); } } } if (iLayer > 0) { - tf->getCellsLookupTable()[iLayer - 1].resize(currentLayerTrackletsNum + 1, tf->getCells()[iLayer].size()); + mTimeFrame->getCellsLookupTable()[iLayer - 1].resize(currentLayerTrackletsNum + 1, mTimeFrame->getCells()[iLayer].size()); } } - if (!tf->checkMemory(mTrkParams[iteration].MaxMemory)) { + if (!mTimeFrame->checkMemory(mTrkParams[iteration].MaxMemory)) { return; } /// Create cells labels - if (tf->hasMCinformation()) { + if (mTimeFrame->hasMCinformation()) { for (int iLayer{0}; iLayer < mTrkParams[iteration].CellsPerRoad(); ++iLayer) { - for (auto& cell : tf->getCells()[iLayer]) { - MCCompLabel currentLab{tf->getTrackletsLabel(iLayer)[cell.getFirstTrackletIndex()]}; - MCCompLabel nextLab{tf->getTrackletsLabel(iLayer + 1)[cell.getSecondTrackletIndex()]}; - tf->getCellsLabel(iLayer).emplace_back(currentLab == nextLab ? currentLab : MCCompLabel()); + for (auto& cell : mTimeFrame->getCells()[iLayer]) { + MCCompLabel currentLab{mTimeFrame->getTrackletsLabel(iLayer)[cell.getFirstTrackletIndex()]}; + MCCompLabel nextLab{mTimeFrame->getTrackletsLabel(iLayer + 1)[cell.getSecondTrackletIndex()]}; + mTimeFrame->getCellsLabel(iLayer).emplace_back(currentLab == nextLab ? currentLab : MCCompLabel()); } } } if constexpr (debugLevel) { for (int iLayer{0}; iLayer < mTrkParams[iteration].CellsPerRoad(); ++iLayer) { - std::cout << "Cells on layer " << iLayer << " " << tf->getCells()[iLayer].size() << std::endl; + std::cout << "Cells on layer " << iLayer << " " << mTimeFrame->getCells()[iLayer].size() << std::endl; } } } -void TrackerTraits::findCellsNeighbours(const int iteration) +template +void TrackerTraits::findCellsNeighbours(const int iteration) { #ifdef OPTIMISATION_OUTPUT std::ofstream off(std::format("cellneighs{}.txt", iteration)); @@ -421,7 +419,6 @@ void TrackerTraits::findCellsNeighbours(const int iteration) cellsNeighbours.reserve(nextLayerCellsNum); for (int iCell{0}; iCell < layerCellsNum; ++iCell) { - const auto& currentCellSeed{mTimeFrame->getCells()[iLayer][iCell]}; const int nextLayerTrackletIndex{currentCellSeed.getSecondTrackletIndex()}; const int nextLayerFirstCellIndex{mTimeFrame->getCellsLookupTable()[iLayer][nextLayerTrackletIndex]}; @@ -469,7 +466,8 @@ void TrackerTraits::findCellsNeighbours(const int iteration) } } -void TrackerTraits::processNeighbours(int iLayer, int iLevel, const std::vector& currentCellSeed, const std::vector& currentCellId, std::vector& updatedCellSeeds, std::vector& updatedCellsIds) +template +void TrackerTraits::processNeighbours(int iLayer, int iLevel, const std::vector& currentCellSeed, const std::vector& currentCellId, std::vector& updatedCellSeeds, std::vector& updatedCellsIds) { bool print = iLayer == 3 && iLevel == 2; if (iLevel < 2 || iLayer < 1) { @@ -568,7 +566,8 @@ void TrackerTraits::processNeighbours(int iLayer, int iLevel, const std::vector< #endif } -void TrackerTraits::findRoads(const int iteration) +template +void TrackerTraits::findRoads(const int iteration) { CA_DEBUGGER(std::cout << "Finding roads, iteration " << iteration << std::endl); for (int startLevel{mTrkParams[iteration].CellsPerRoad()}; startLevel >= mTrkParams[iteration].CellMinimumLevel(); --startLevel) { @@ -676,7 +675,8 @@ void TrackerTraits::findRoads(const int iteration) } } -void TrackerTraits::extendTracks(const int iteration) +template +void TrackerTraits::extendTracks(const int iteration) { for (int rof{0}; rof < mTimeFrame->getNrof(); ++rof) { for (auto& track : mTimeFrame->getTracks(rof)) { @@ -724,7 +724,8 @@ void TrackerTraits::extendTracks(const int iteration) } } -void TrackerTraits::findShortPrimaries() +template +void TrackerTraits::findShortPrimaries() { const auto propagator = o2::base::Propagator::Instance(); mTimeFrame->fillPrimaryVerticesXandAlpha(); @@ -812,7 +813,8 @@ void TrackerTraits::findShortPrimaries() } } -bool TrackerTraits::fitTrack(TrackITSExt& track, int start, int end, int step, float chi2clcut, float chi2ndfcut, float maxQoverPt, int nCl) +template +bool TrackerTraits::fitTrack(TrackITSExt& track, int start, int end, int step, float chi2clcut, float chi2ndfcut, float maxQoverPt, int nCl) { auto propInstance = o2::base::Propagator::Instance(); @@ -851,7 +853,8 @@ bool TrackerTraits::fitTrack(TrackITSExt& track, int start, int end, int step, f return std::abs(track.getQ2Pt()) < maxQoverPt && track.getChi2() < chi2ndfcut * (nCl * 2 - 5); } -bool TrackerTraits::trackFollowing(TrackITSExt* track, int rof, bool outward, const int iteration) +template +bool TrackerTraits::trackFollowing(TrackITSExt* track, int rof, bool outward, const int iteration) { auto propInstance = o2::base::Propagator::Instance(); const int step = -1 + outward * 2; @@ -966,7 +969,8 @@ bool TrackerTraits::trackFollowing(TrackITSExt* track, int rof, bool outward, co /// Clusters are given from inside outward (cluster3 is the outermost). The outermost cluster is given in the tracking /// frame coordinates whereas the others are referred to the global frame. -track::TrackParCov TrackerTraits::buildTrackSeed(const Cluster& cluster1, const Cluster& cluster2, const TrackingFrameInfo& tf3) +template +track::TrackParCov TrackerTraits::buildTrackSeed(const Cluster& cluster1, const Cluster& cluster2, const TrackingFrameInfo& tf3) { const float ca = o2::gpu::CAMath::Cos(tf3.alphaTrackingFrame), sa = o2::gpu::CAMath::Sin(tf3.alphaTrackingFrame); const float x1 = cluster1.xCoordinate * ca + cluster1.yCoordinate * sa; @@ -997,15 +1001,21 @@ track::TrackParCov TrackerTraits::buildTrackSeed(const Cluster& cluster1, const 0.f, 0.f, 0.f, 0.f, sg2q2pt}); } -void TrackerTraits::setBz(float bz) +template +void TrackerTraits::setBz(float bz) { mBz = bz; mTimeFrame->setBz(bz); } -bool TrackerTraits::isMatLUT() const { return o2::base::Propagator::Instance()->getMatLUT() && (mCorrType == o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT); } +template +bool TrackerTraits::isMatLUT() const +{ + return o2::base::Propagator::Instance()->getMatLUT() && (mCorrType == o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT); +} -void TrackerTraits::setNThreads(int n) +template +void TrackerTraits::setNThreads(int n) { #ifdef WITH_OPENMP mNThreads = n > 0 ? n : 1; @@ -1014,25 +1024,6 @@ void TrackerTraits::setNThreads(int n) #endif } -int TrackerTraits::getTFNumberOfClusters() const -{ - return mTimeFrame->getNumberOfClusters(); -} - -int TrackerTraits::getTFNumberOfTracklets() const -{ - return mTimeFrame->getNumberOfTracklets(); -} - -int TrackerTraits::getTFNumberOfCells() const -{ - return mTimeFrame->getNumberOfCells(); -} - -void TrackerTraits::adoptTimeFrame(TimeFrame* tf) -{ - mTimeFrame = tf; -} +template class TrackerTraits<7>; -} // namespace its -} // namespace o2 +} // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx index f0dad2722a301..0df1899c9ab6a 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx @@ -139,7 +139,6 @@ void ITSTrackingInterface::initialise() mVertexer->setParameters(vertParams); } -template void ITSTrackingInterface::run(framework::ProcessingContext& pc) { auto compClusters = pc.inputs().get>("compClusters"); @@ -211,9 +210,9 @@ void ITSTrackingInterface::run(framework::ProcessingContext& pc) loadROF(trackROFspan, compClusters, pattIt, labels); pattIt = patterns.begin(); std::vector savedROF; - auto logger = [&](std::string s) { LOG(info) << s; }; - auto fatalLogger = [&](std::string s) { LOG(fatal) << s; }; - auto errorLogger = [&](std::string s) { LOG(error) << s; }; + auto logger = [&](const std::string& s) { LOG(info) << s; }; + auto fatalLogger = [&](const std::string& s) { LOG(fatal) << s; }; + auto errorLogger = [&](const std::string& s) { LOG(error) << s; }; FastMultEst multEst; // mult estimator std::vector processingMask, processUPCMask; @@ -224,11 +223,7 @@ void ITSTrackingInterface::run(framework::ProcessingContext& pc) if (mRunVertexer) { vertROFvec.reserve(trackROFvec.size()); // Run seeding vertexer - if constexpr (isGPU) { - vertexerElapsedTime = mVertexer->clustersToVerticesHybrid(logger); - } else { - vertexerElapsedTime = mVertexer->clustersToVertices(logger); - } + vertexerElapsedTime = mVertexer->clustersToVertices(logger); } else { // cosmics mTimeFrame->resetRofPV(); } @@ -436,8 +431,8 @@ void ITSTrackingInterface::printSummary() const } void ITSTrackingInterface::setTraitsFromProvider(VertexerTraits* vertexerTraits, - TrackerTraits* trackerTraits, - TimeFrame* frame) + TrackerTraits7* trackerTraits, + TimeFrame7* frame) { mVertexer = std::make_unique(vertexerTraits); mTracker = std::make_unique(trackerTraits); @@ -453,8 +448,5 @@ void ITSTrackingInterface::loadROF(gsl::span& trackROFspan, { mTimeFrame->loadROFrameData(trackROFspan, clusters, pattIt, mDict, mcLabels); } - -template void ITSTrackingInterface::run(framework::ProcessingContext& pc); -template void ITSTrackingInterface::run(framework::ProcessingContext& pc); } // namespace its } // namespace o2 diff --git a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx index 13ce03e9fba4f..0b8d59f61eb8d 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx @@ -21,76 +21,40 @@ #include "ITStracking/VertexerTraits.h" #include "ITStracking/TrackingConfigParam.h" -#include - -namespace o2 -{ -namespace its +namespace o2::its { -Vertexer::Vertexer(VertexerTraits* traits) +Vertexer::Vertexer(VertexerTraits* traits) : mTraits(traits) { - if (!traits) { + if (!mTraits) { LOG(fatal) << "nullptr passed to ITS vertexer construction."; } mVertParams.resize(1); - mTraits = traits; } -float Vertexer::clustersToVertices(std::function logger) +float Vertexer::clustersToVertices(LogFunc logger) { + LogFunc evalLog = [](const std::string&) {}; TrackingParameters trkPars; TimeFrameGPUParameters tfGPUpar; mTraits->updateVertexingParameters(mVertParams, tfGPUpar); float timeTracklet{0.f}, timeSelection{0.f}, timeVertexing{0.f}, timeInit{0.f}; for (int iteration = 0; iteration < std::min(mVertParams[0].nIterations, (int)mVertParams.size()); ++iteration) { - unsigned int nTracklets01, nTracklets12; - logger(fmt::format("ITS Seeding vertexer iteration {} summary:", iteration)); + unsigned int nTracklets01{0}, nTracklets12{0}; + logger(fmt::format("=== ITS {} Seeding vertexer iteration {} summary:", mTraits->getName(), iteration)); trkPars.PhiBins = mTraits->getVertexingParameters()[0].PhiBins; trkPars.ZBins = mTraits->getVertexingParameters()[0].ZBins; auto timeInitIteration = evaluateTask( - &Vertexer::initialiseVertexer, "Vertexer initialisation", [](std::string) {}, trkPars, iteration); + &Vertexer::initialiseVertexer, " - Vertexer initialisation", evalLog, trkPars, iteration); auto timeTrackletIteration = evaluateTask( - &Vertexer::findTracklets, "Vertexer tracklet finding", [](std::string) {}, iteration); + &Vertexer::findTracklets, " - Vertexer tracklet finding", evalLog, iteration); nTracklets01 = mTimeFrame->getTotalTrackletsTF(0); nTracklets12 = mTimeFrame->getTotalTrackletsTF(1); auto timeSelectionIteration = evaluateTask( - &Vertexer::validateTracklets, "Vertexer tracklets validation", [](std::string) {}, iteration); + &Vertexer::validateTracklets, " - Vertexer tracklets validation", evalLog, iteration); auto timeVertexingIteration = evaluateTask( - &Vertexer::findVertices, "Vertexer vertex finding", [](std::string) {}, iteration); - printEpilog(logger, false, nTracklets01, nTracklets12, mTimeFrame->getNLinesTotal(), mTimeFrame->getTotVertIteration()[iteration], timeInitIteration, timeTrackletIteration, timeSelectionIteration, timeVertexingIteration); - timeInit += timeInitIteration; - timeTracklet += timeTrackletIteration; - timeSelection += timeSelectionIteration; - timeVertexing += timeVertexingIteration; - } - - return timeInit + timeTracklet + timeSelection + timeVertexing; -} - -float Vertexer::clustersToVerticesHybrid(std::function logger) -{ - TrackingParameters trkPars; - TimeFrameGPUParameters tfGPUpar; - float timeTracklet, timeSelection, timeVertexing, timeInit; - mTraits->updateVertexingParameters(mVertParams, tfGPUpar); - for (int iteration = 0; iteration < std::min(mVertParams[0].nIterations, (int)mVertParams.size()); ++iteration) { - unsigned int nTracklets01, nTracklets12; - logger(fmt::format("ITS Hybrid seeding vertexer iteration {} summary:", iteration)); - trkPars.PhiBins = mTraits->getVertexingParameters()[0].PhiBins; - trkPars.ZBins = mTraits->getVertexingParameters()[0].ZBins; - auto timeInitIteration = evaluateTask( - &Vertexer::initialiseVertexerHybrid, "Hybrid Vertexer initialisation", [](std::string) {}, trkPars, iteration); - auto timeTrackletIteration = evaluateTask( - &Vertexer::findTrackletsHybrid, "Hybrid Vertexer tracklet finding", [](std::string) {}, iteration); - nTracklets01 = mTimeFrame->getTotalTrackletsTF(0); - nTracklets12 = mTimeFrame->getTotalTrackletsTF(1); - auto timeSelectionIteration = evaluateTask( - &Vertexer::validateTrackletsHybrid, "Hybrid Vertexer adjacent tracklets validation", [](std::string) {}, iteration); - auto timeVertexingIteration = evaluateTask( - &Vertexer::findVerticesHybrid, "Hybrid Vertexer vertex finding", [](std::string) {}, iteration); - - printEpilog(logger, true, nTracklets01, nTracklets12, mTimeFrame->getNLinesTotal(), mTimeFrame->getTotVertIteration()[iteration], timeInitIteration, timeTrackletIteration, timeSelectionIteration, timeVertexingIteration); + &Vertexer::findVertices, " - Vertexer vertex finding", evalLog, iteration); + printEpilog(logger, nTracklets01, nTracklets12, mTimeFrame->getNLinesTotal(), mTimeFrame->getTotVertIteration()[iteration], timeInitIteration, timeTrackletIteration, timeSelectionIteration, timeVertexingIteration); timeInit += timeInitIteration; timeTracklet += timeTrackletIteration; timeSelection += timeSelectionIteration; @@ -129,23 +93,22 @@ void Vertexer::getGlobalConfiguration() mVertParams[0].PhiBins = vc.PhiBins; } -void Vertexer::adoptTimeFrame(TimeFrame& tf) +void Vertexer::adoptTimeFrame(TimeFrame7& tf) { mTimeFrame = &tf; mTraits->adoptTimeFrame(&tf); } -void Vertexer::printEpilog(std::function logger, - bool isHybrid, - const unsigned int trackletN01, const unsigned int trackletN12, const unsigned selectedN, const unsigned int vertexN, - const float initT, const float trackletT, const float selecT, const float vertexT) +void Vertexer::printEpilog(LogFunc& logger, + const unsigned int trackletN01, const unsigned int trackletN12, + const unsigned selectedN, const unsigned int vertexN, const float initT, + const float trackletT, const float selecT, const float vertexT) { float total = initT + trackletT + selecT + vertexT; - logger(fmt::format(" - {}Vertexer: found {} | {} tracklets in: {} ms", isHybrid ? "Hybrid " : "", trackletN01, trackletN12, trackletT)); - logger(fmt::format(" - {}Vertexer: selected {} tracklets in: {} ms", isHybrid ? "Hybrid " : "", selectedN, selecT)); - logger(fmt::format(" - {}Vertexer: found {} vertices in: {} ms", isHybrid ? "Hybrid " : "", vertexN, vertexT)); + logger(fmt::format(" - {} Vertexer: found {} | {} tracklets in: {} ms", mTraits->getName(), trackletN01, trackletN12, trackletT)); + logger(fmt::format(" - {} Vertexer: selected {} tracklets in: {} ms", mTraits->getName(), selectedN, selecT)); + logger(fmt::format(" - {} Vertexer: found {} vertices in: {} ms", mTraits->getName(), vertexN, vertexT)); // logger(fmt::format(" - Timeframe {} vertexing completed in: {} ms, using {} thread(s).", mTimeFrameCounter++, total, mTraits->getNThreads())); } -} // namespace its -} // namespace o2 +} // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx index 3ecf48771f42e..71b54d95d3f8d 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx @@ -586,7 +586,7 @@ void VertexerTraits::computeVerticesInRof(int rofId, std::array& beamPosXY, std::vector& vertices, std::vector& verticesInRof, - TimeFrame* tf, + TimeFrame7* tf, std::vector* labels, const int iteration) { diff --git a/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/TrackerSpec.h b/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/TrackerSpec.h index be9965fc8be58..9160df6fc49fd 100644 --- a/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/TrackerSpec.h +++ b/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/TrackerSpec.h @@ -28,9 +28,7 @@ #include "TStopwatch.h" -namespace o2 -{ -namespace its +namespace o2::its { class TrackerDPL : public framework::Task @@ -41,7 +39,7 @@ class TrackerDPL : public framework::Task int trgType, const TrackingMode& trMode = TrackingMode::Unset, const bool overrBeamEst = false, - gpu::GPUDataTypes::DeviceType dType = gpu::GPUDataTypes::DeviceType::CPU); + o2::gpu::GPUDataTypes::DeviceType dType = o2::gpu::GPUDataTypes::DeviceType::CPU); ~TrackerDPL() override = default; void init(framework::InitContext& ic) final; void run(framework::ProcessingContext& pc) final; @@ -59,9 +57,8 @@ class TrackerDPL : public framework::Task }; using o2::its::TrackingMode; -framework::DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, int useTrig, const std::string& trMode, const bool overrBeamEst = false, gpu::GPUDataTypes::DeviceType dType = gpu::GPUDataTypes::DeviceType::CPU); +framework::DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, int useTrig, const std::string& trMode, const bool overrBeamEst = false, o2::gpu::GPUDataTypes::DeviceType dType = o2::gpu::GPUDataTypes::DeviceType::CPU); -} // namespace its -} // namespace o2 +} // namespace o2::its #endif /* O2_ITS_TRACKERDPL */ diff --git a/Detectors/Upgrades/ITS3/reconstruction/include/ITS3Reconstruction/IOUtils.h b/Detectors/Upgrades/ITS3/reconstruction/include/ITS3Reconstruction/IOUtils.h index b9e7fd0f6ec39..771b13539b759 100644 --- a/Detectors/Upgrades/ITS3/reconstruction/include/ITS3Reconstruction/IOUtils.h +++ b/Detectors/Upgrades/ITS3/reconstruction/include/ITS3Reconstruction/IOUtils.h @@ -71,7 +71,7 @@ void convertCompactClusters(gsl::span clusters, std::vector>& output, const its3::TopologyDictionary* dict); -int loadROFrameDataITS3(its::TimeFrame* tf, +int loadROFrameDataITS3(its::TimeFrame<7>* tf, gsl::span rofs, gsl::span clusters, gsl::span::iterator& pattIt, diff --git a/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx b/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx index 58dd56ac41f95..5c3b9670fbdb9 100644 --- a/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx +++ b/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx @@ -56,7 +56,7 @@ void convertCompactClusters(gsl::span clusters, } } -int loadROFrameDataITS3(its::TimeFrame* tf, +int loadROFrameDataITS3(its::TimeFrame<7>* tf, gsl::span rofs, gsl::span clusters, gsl::span::iterator& pattIt, diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index c76bf11c3e25d..04c9c12475fab 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -111,16 +111,16 @@ GPUReconstruction::~GPUReconstruction() } } -void GPUReconstruction::GetITSTraits(std::unique_ptr* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr* timeFrame) +void GPUReconstruction::GetITSTraits(std::unique_ptr>* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr>* timeFrame) { if (trackerTraits) { - trackerTraits->reset(new o2::its::TrackerTraits); + trackerTraits->reset(new o2::its::TrackerTraits<7>); } if (vertexerTraits) { vertexerTraits->reset(new o2::its::VertexerTraits); } if (timeFrame) { - timeFrame->reset(new o2::its::TimeFrame); + timeFrame->reset(new o2::its::TimeFrame<7>); } } diff --git a/GPU/GPUTracking/Base/GPUReconstruction.h b/GPU/GPUTracking/Base/GPUReconstruction.h index d5c0b8e828087..4c7b3da6008a0 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.h +++ b/GPU/GPUTracking/Base/GPUReconstruction.h @@ -35,8 +35,10 @@ namespace o2::its { +template class TrackerTraits; class VertexerTraits; +template class TimeFrame; } // namespace o2::its @@ -188,7 +190,7 @@ class GPUReconstruction GPUMemorySizeScalers* MemoryScalers() { return mMemoryScalers.get(); } // Helpers to fetch processors from other shared libraries - virtual void GetITSTraits(std::unique_ptr* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr* timeFrame); + virtual void GetITSTraits(std::unique_ptr>* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr>* timeFrame); bool slavesExist() { return mSlaves.size() || mMaster; } int slaveId() { return mSlaveId; } diff --git a/GPU/GPUTracking/Base/GPUReconstructionIncludesITS.h b/GPU/GPUTracking/Base/GPUReconstructionIncludesITS.h index 5891891d9da24..87829f5634375 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionIncludesITS.h +++ b/GPU/GPUTracking/Base/GPUReconstructionIncludesITS.h @@ -30,9 +30,11 @@ namespace o2::its class VertexerTraits { }; +template class TrackerTraits { }; +template class TimeFrame { }; diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu index c8e5420a8bcf3..a630cb3f541ef 100644 --- a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu +++ b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu @@ -87,13 +87,13 @@ int32_t GPUReconstructionCUDA::GPUChkErrInternal(const int64_t error, const char GPUReconstruction* GPUReconstruction_Create_CUDA(const GPUSettingsDeviceBackend& cfg) { return new GPUReconstructionCUDA(cfg); } -void GPUReconstructionCUDA::GetITSTraits(std::unique_ptr* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr* timeFrame) +void GPUReconstructionCUDA::GetITSTraits(std::unique_ptr>* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr>* timeFrame) { if (trackerTraits) { trackerTraits->reset(new o2::its::TrackerTraitsGPU); } if (vertexerTraits) { - vertexerTraits->reset(new o2::its::VertexerTraitsGPU); + vertexerTraits->reset(new o2::its::VertexerTraits); // TODO gpu-code to be implemented } if (timeFrame) { timeFrame->reset(new o2::its::gpu::TimeFrameGPU); diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h index 6c126d153d8ae..1cc7e0fc819ff 100644 --- a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h +++ b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.h @@ -76,7 +76,7 @@ class GPUReconstructionCUDA : public GPUReconstructionProcessing::KernelInterfac void RecordMarker(deviceEvent* ev, int32_t stream) override; void SetONNXGPUStream(Ort::SessionOptions& session_options, int32_t stream, int32_t* deviceId) override; - void GetITSTraits(std::unique_ptr* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr* timeFrame) override; + void GetITSTraits(std::unique_ptr>* trackerTraits, std::unique_ptr* vertexerTraits, std::unique_ptr>* timeFrame) override; #ifndef __HIPCC__ // CUDA bool CanQueryMaxMemory() override { return true; } diff --git a/GPU/GPUTracking/Global/GPUChainITS.cxx b/GPU/GPUTracking/Global/GPUChainITS.cxx index 5d36dc63ca85d..bcb99fff87a64 100644 --- a/GPU/GPUTracking/Global/GPUChainITS.cxx +++ b/GPU/GPUTracking/Global/GPUChainITS.cxx @@ -63,7 +63,7 @@ void GPUChainITS::MemorySize(size_t& gpuMem, size_t& pageLockedHostMem) int32_t GPUChainITS::Init() { return 0; } -o2::its::TrackerTraits* GPUChainITS::GetITSTrackerTraits() +o2::its::TrackerTraits<7>* GPUChainITS::GetITSTrackerTraits() { if (mITSTrackerTraits == nullptr) { mRec->GetITSTraits(&mITSTrackerTraits, nullptr, nullptr); @@ -79,7 +79,7 @@ o2::its::VertexerTraits* GPUChainITS::GetITSVertexerTraits() return mITSVertexerTraits.get(); } -o2::its::TimeFrame* GPUChainITS::GetITSTimeframe() +o2::its::TimeFrame<7>* GPUChainITS::GetITSTimeframe() { if (mITSTimeFrame == nullptr) { mRec->GetITSTraits(nullptr, nullptr, &mITSTimeFrame); diff --git a/GPU/GPUTracking/Global/GPUChainITS.h b/GPU/GPUTracking/Global/GPUChainITS.h index 37cd5acc5264c..150d66031d084 100644 --- a/GPU/GPUTracking/Global/GPUChainITS.h +++ b/GPU/GPUTracking/Global/GPUChainITS.h @@ -43,15 +43,15 @@ class GPUChainITS : public GPUChain int32_t RunChain() override; void MemorySize(size_t& gpuMem, size_t& pageLockedHostMem) override; - o2::its::TrackerTraits* GetITSTrackerTraits(); + o2::its::TrackerTraits<7>* GetITSTrackerTraits(); o2::its::VertexerTraits* GetITSVertexerTraits(); - o2::its::TimeFrame* GetITSTimeframe(); + o2::its::TimeFrame<7>* GetITSTimeframe(); protected: GPUChainITS(GPUReconstruction* rec, uint32_t maxTracks = GPUCA_MAX_ITS_FIT_TRACKS); - std::unique_ptr mITSTrackerTraits; + std::unique_ptr> mITSTrackerTraits; std::unique_ptr mITSVertexerTraits; - std::unique_ptr mITSTimeFrame; + std::unique_ptr> mITSTimeFrame; std::unique_ptr mFrameworkAllocator; uint32_t mMaxTracks; diff --git a/GPU/GPUTracking/Interface/GPUO2Interface.cxx b/GPU/GPUTracking/Interface/GPUO2Interface.cxx index 81eb2c285192b..b9921acfc2492 100644 --- a/GPU/GPUTracking/Interface/GPUO2Interface.cxx +++ b/GPU/GPUTracking/Interface/GPUO2Interface.cxx @@ -251,7 +251,7 @@ void GPUO2Interface::setErrorCodeOutput(std::vector>* v) } } -void GPUO2Interface::GetITSTraits(o2::its::TrackerTraits*& trackerTraits, o2::its::VertexerTraits*& vertexerTraits, o2::its::TimeFrame*& timeFrame) +void GPUO2Interface::GetITSTraits(o2::its::TrackerTraits<7>*& trackerTraits, o2::its::VertexerTraits*& vertexerTraits, o2::its::TimeFrame<7>*& timeFrame) { trackerTraits = mChainITS->GetITSTrackerTraits(); vertexerTraits = mChainITS->GetITSVertexerTraits(); diff --git a/GPU/GPUTracking/Interface/GPUO2Interface.h b/GPU/GPUTracking/Interface/GPUO2Interface.h index 3a819de7c7b7c..03b24c2b23877 100644 --- a/GPU/GPUTracking/Interface/GPUO2Interface.h +++ b/GPU/GPUTracking/Interface/GPUO2Interface.h @@ -43,8 +43,10 @@ struct ClusterNative; namespace o2::its { +template class TrackerTraits; class VertexerTraits; +template class TimeFrame; } // namespace o2::its @@ -77,7 +79,7 @@ class GPUO2Interface void DumpEvent(int32_t nEvent, GPUTrackingInOutPointers* data); void DumpSettings(); - void GetITSTraits(o2::its::TrackerTraits*& trackerTraits, o2::its::VertexerTraits*& vertexerTraits, o2::its::TimeFrame*& timeFrame); + void GetITSTraits(o2::its::TrackerTraits<7>*& trackerTraits, o2::its::VertexerTraits*& vertexerTraits, o2::its::TimeFrame<7>*& timeFrame); const o2::base::Propagator* GetDeviceO2Propagator(int32_t iThread = 0) const; void UseGPUPolynomialFieldInPropagator(o2::base::Propagator* prop) const; diff --git a/GPU/Workflow/include/GPUWorkflow/GPUWorkflowSpec.h b/GPU/Workflow/include/GPUWorkflow/GPUWorkflowSpec.h index 73f1f208e8889..405f4a371f0cc 100644 --- a/GPU/Workflow/include/GPUWorkflow/GPUWorkflowSpec.h +++ b/GPU/Workflow/include/GPUWorkflow/GPUWorkflowSpec.h @@ -58,6 +58,7 @@ class GeometryFlat; namespace its { +template class TimeFrame; class ITSTrackingInterface; } // namespace its @@ -214,7 +215,7 @@ class GPURecoWorkflowSpec : public o2::framework::Task std::vector mTPCSectors; std::unique_ptr mITSTrackingInterface; std::unique_ptr mPipeline; - o2::its::TimeFrame* mITSTimeFrame = nullptr; + o2::its::TimeFrame<7>* mITSTimeFrame = nullptr; std::vector mRegionInfos; const o2::itsmft::TopologyDictionary* mITSDict = nullptr; const o2::dataformats::MeanVertexObject* mMeanVertex; diff --git a/GPU/Workflow/src/GPUWorkflowITS.cxx b/GPU/Workflow/src/GPUWorkflowITS.cxx index e56958cba2c9b..31ccaad8c2783 100644 --- a/GPU/Workflow/src/GPUWorkflowITS.cxx +++ b/GPU/Workflow/src/GPUWorkflowITS.cxx @@ -31,14 +31,14 @@ int32_t GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc { mITSTimeFrame->setDevicePropagator(mGPUReco->GetDeviceO2Propagator()); LOGP(debug, "GPUChainITS is giving me device propagator: {}", (void*)mGPUReco->GetDeviceO2Propagator()); - mITSTrackingInterface->run(pc); + mITSTrackingInterface->run(pc); return 0; } void GPURecoWorkflowSpec::initFunctionITS(o2::framework::InitContext& ic) { o2::its::VertexerTraits* vtxTraits = nullptr; - o2::its::TrackerTraits* trkTraits = nullptr; + o2::its::TrackerTraits<7>* trkTraits = nullptr; #ifdef ENABLE_UPGRADES if (mSpecConfig.isITS3) { mITSTrackingInterface = std::make_unique(mSpecConfig.processMC,