diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h index 9731ad5c5db67..a1d52bff11f9d 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h @@ -28,6 +28,7 @@ template class TimeFrameGPU : public TimeFrame { using typename TimeFrame::CellSeedN; + using typename TimeFrame::IndexTableUtilsN; public: TimeFrameGPU(); @@ -36,8 +37,8 @@ class TimeFrameGPU : public TimeFrame /// Most relevant operations void registerHostMemory(const int); void unregisterHostMemory(const int); - void initialise(const int, const TrackingParameters&, const int, IndexTableUtils* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr); - void initDevice(IndexTableUtils*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int); + void initialise(const int, const TrackingParameters&, const int, IndexTableUtilsN* utils = nullptr, const TimeFrameGPUParameters* pars = nullptr); + void initDevice(IndexTableUtilsN*, const TrackingParameters& trkParam, const TimeFrameGPUParameters&, const int, const int); void initDeviceSAFitting(); void loadIndexTableUtils(const int); void loadTrackingFrameInfoDevice(const int, const int); @@ -98,7 +99,7 @@ class TimeFrameGPU : public TimeFrame /// interface int getNClustersInRofSpan(const int, const int, const int) const; - IndexTableUtils* getDeviceIndexTableUtils() { return mIndexTableUtilsDevice; } + IndexTableUtilsN* getDeviceIndexTableUtils() { return mIndexTableUtilsDevice; } int* getDeviceROFramesClusters(const int layer) { return mROFramesClustersDevice[layer]; } auto& getTrackITSExt() { return mTrackITSExt; } Vertex* getDeviceVertices() { return mPrimaryVerticesDevice; } @@ -165,7 +166,7 @@ class TimeFrameGPU : public TimeFrame std::array mNNeighbours; // Device pointers - IndexTableUtils* mIndexTableUtilsDevice; + IndexTableUtilsN* mIndexTableUtilsDevice; // Hybrid pref uint8_t* mMultMaskDevice; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h index f8eedb33d91eb..7d26e74692aa5 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackerTraitsGPU.h @@ -22,6 +22,8 @@ namespace o2::its template class TrackerTraitsGPU final : public TrackerTraits { + using typename TrackerTraits::IndexTableUtilsN; + public: TrackerTraitsGPU() = default; ~TrackerTraitsGPU() final = default; @@ -48,7 +50,7 @@ class TrackerTraitsGPU final : public TrackerTraits int getTFNumberOfCells() const override; private: - IndexTableUtils* mDeviceIndexTableUtils; + IndexTableUtilsN* mDeviceIndexTableUtils; gpu::TimeFrameGPU* mTimeFrameGPU; }; diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h index 53e680d474f6f..69d6799686654 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h @@ -27,13 +27,14 @@ template class CellSeed; class TrackingFrameInfo; class Tracklet; +template class IndexTableUtils; class Cluster; class TrackITSExt; class ExternalAllocator; template -void countTrackletsInROFsHandler(const IndexTableUtils* utils, +void countTrackletsInROFsHandler(const IndexTableUtils* utils, const uint8_t* multMask, const int layer, const int startROF, @@ -66,7 +67,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, gpu::Streams& streams); template -void computeTrackletsInROFsHandler(const IndexTableUtils* utils, +void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const uint8_t* multMask, const int layer, const int startROF, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt index 1f6a046a81350..3dff67dbccd80 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/CMakeLists.txt @@ -22,8 +22,8 @@ if(CUDA_ENABLED) TimeFrameGPU.cu TracerGPU.cu TrackingKernels.cu - VertexingKernels.cu - VertexerTraitsGPU.cxx + # VertexingKernels.cu + # VertexerTraitsGPU.cxx PUBLIC_INCLUDE_DIRECTORIES ../ PUBLIC_LINK_LIBRARIES O2::ITStracking O2::SimConfig diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu index d834f28e09db0..965bf27fdd12b 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TimeFrameGPU.cu @@ -61,11 +61,11 @@ void TimeFrameGPU::loadIndexTableUtils(const int iteration) { GPUTimer timer("loading indextable utils"); if (!iteration) { - GPULog("gpu-allocation: allocating IndexTableUtils buffer, for {:.2f} MB.", sizeof(IndexTableUtils) / constants::MB); - allocMem(reinterpret_cast(&mIndexTableUtilsDevice), sizeof(IndexTableUtils), this->getExtAllocator()); + GPULog("gpu-allocation: allocating IndexTableUtils buffer, for {:.2f} MB.", sizeof(IndexTableUtilsN) / constants::MB); + allocMem(reinterpret_cast(&mIndexTableUtilsDevice), sizeof(IndexTableUtilsN), this->getExtAllocator()); } - GPULog("gpu-transfer: loading IndexTableUtils object, for {:.2f} MB.", sizeof(IndexTableUtils) / constants::MB); - GPUChkErrS(cudaMemcpy(mIndexTableUtilsDevice, &(this->mIndexTableUtils), sizeof(IndexTableUtils), cudaMemcpyHostToDevice)); + GPULog("gpu-transfer: loading IndexTableUtils object, for {:.2f} MB.", sizeof(IndexTableUtilsN) / constants::MB); + GPUChkErrS(cudaMemcpy(mIndexTableUtilsDevice, &(this->mIndexTableUtils), sizeof(IndexTableUtilsN), cudaMemcpyHostToDevice)); } template @@ -547,7 +547,7 @@ template void TimeFrameGPU::initialise(const int iteration, const TrackingParameters& trkParam, const int maxLayers, - IndexTableUtils* utils, + IndexTableUtilsN* utils, const TimeFrameGPUParameters* gpuParam) { mGpuStreams.resize(nLayers); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu index 694b598334be3..94c6610ab9430 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu +++ b/Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu @@ -95,8 +95,9 @@ GPUdii() int4 getEmptyBinsRect() return int4{0, 0, 0, 0}; } +template GPUdii() const int4 getBinsRect(const Cluster& currentCluster, const int layerIndex, - const o2::its::IndexTableUtils& utils, + const IndexTableUtils& utils, const float z1, const float z2, float maxdeltaz, float maxdeltaphi) { const float zRangeMin = o2::gpu::CAMath::Min(z1, z2) - maxdeltaz; @@ -331,7 +332,7 @@ GPUg() void fitTrackSeedsKernel( temporaryTrack.resetCovariance(); temporaryTrack.setChi2(0); auto& clusters = seed.getClusters(); - for (int iL{0}; iL < 7; ++iL) { + for (int iL{0}; iL < nLayers; ++iL) { temporaryTrack.setExternalClusterIndex(iL, clusters[iL], clusters[iL] != constants::UnusedIndex); } bool fitSuccess = fitTrack(temporaryTrack, // TrackITSExt& track, @@ -523,9 +524,9 @@ GPUg() void computeLayerCellsKernel( } } -template +template GPUg() void computeLayerTrackletsMultiROFKernel( - const IndexTableUtils* utils, + const IndexTableUtils* utils, const uint8_t* multMask, const int layerIndex, const int startROF, @@ -601,7 +602,7 @@ GPUg() void computeLayerTrackletsMultiROFKernel( const float zAtRmax{tanLambda * (maxR - currentCluster.radius) + currentCluster.zCoordinate}; const float sqInverseDeltaZ0{1.f / (math_utils::Sq(currentCluster.zCoordinate - primaryVertex.getZ()) + constants::Tolerance)}; /// protecting from overflows adding the detector resolution const float sigmaZ{o2::gpu::CAMath::Sqrt(math_utils::Sq(resolution) * math_utils::Sq(tanLambda) * ((math_utils::Sq(inverseR0) + sqInverseDeltaZ0) * math_utils::Sq(meanDeltaR) + 1.f) + math_utils::Sq(meanDeltaR * MSAngle))}; - const int4 selectedBinsRect{getBinsRect(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)}; + const int4 selectedBinsRect{getBinsRect(currentCluster, layerIndex + 1, *utils, zAtRmin, zAtRmax, sigmaZ * NSigmaCut, phiCut)}; if (selectedBinsRect.x == 0 && selectedBinsRect.y == 0 && selectedBinsRect.z == 0 && selectedBinsRect.w == 0) { continue; } @@ -769,7 +770,7 @@ GPUhi() void deallocateMemory(void* p, size_t bytes, cudaStream_t stream = nullp } // namespace gpu template -void countTrackletsInROFsHandler(const IndexTableUtils* utils, +void countTrackletsInROFsHandler(const IndexTableUtils* utils, const uint8_t* multMask, const int layer, const int startROF, @@ -833,7 +834,7 @@ void countTrackletsInROFsHandler(const IndexTableUtils* utils, } template -void computeTrackletsInROFsHandler(const IndexTableUtils* utils, +void computeTrackletsInROFsHandler(const IndexTableUtils* utils, const uint8_t* multMask, const int layer, const int startROF, @@ -1241,7 +1242,7 @@ void trackSeedHandler(CellSeed* trackSeeds, } /// Explicit instantiation of ITS2 handlers -template void countTrackletsInROFsHandler<7>(const IndexTableUtils* utils, +template void countTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const uint8_t* multMask, const int layer, const int startROF, @@ -1273,7 +1274,7 @@ template void countTrackletsInROFsHandler<7>(const IndexTableUtils* utils, const int nThreads, gpu::Streams& streams); -template void computeTrackletsInROFsHandler<7>(const IndexTableUtils* utils, +template void computeTrackletsInROFsHandler<7>(const IndexTableUtils<7>* utils, const uint8_t* multMask, const int layer, const int startROF, diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt b/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt index e8e475f2232c8..dd83669311a54 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt +++ b/Detectors/ITSMFT/ITS/tracking/GPU/hip/CMakeLists.txt @@ -19,8 +19,8 @@ if(HIP_ENABLED) ../cuda/TrackerTraitsGPU.cxx ../cuda/TracerGPU.cu ../cuda/TrackingKernels.cu - ../cuda/VertexingKernels.cu - ../cuda/VertexerTraitsGPU.cxx + # ../cuda/VertexingKernels.cu + # ../cuda/VertexerTraitsGPU.cxx PUBLIC_INCLUDE_DIRECTORIES ../ PUBLIC_LINK_LIBRARIES O2::ITStracking O2::GPUTracking diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cluster.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cluster.h index eaefbee5e2aaa..dd96dc80f2926 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cluster.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cluster.h @@ -22,32 +22,23 @@ namespace o2::its { +template class IndexTableUtils; struct Cluster final { GPUhdDefault() Cluster() = default; GPUhd() Cluster(const float x, const float y, const float z, const int idx); - GPUhd() Cluster(const int, const IndexTableUtils& utils, const Cluster&); - GPUhd() Cluster(const int, const float3&, const IndexTableUtils& utils, const Cluster&); + template + GPUhd() Cluster(const int, const IndexTableUtils& utils, const Cluster&); + template + GPUhd() Cluster(const int, const float3&, const IndexTableUtils& utils, const Cluster&); GPUhdDefault() Cluster(const Cluster&) = default; GPUhdDefault() Cluster(Cluster&&) noexcept = default; GPUhdDefault() ~Cluster() = default; GPUhdDefault() Cluster& operator=(const Cluster&) = default; GPUhdDefault() Cluster& operator=(Cluster&&) noexcept = default; - - // TODO - /*GPUhdDefault() bool operator==(const Cluster&) const = default;*/ - GPUhd() bool operator==(const Cluster& other) const - { - return xCoordinate == other.xCoordinate && - yCoordinate == other.yCoordinate && - zCoordinate == other.zCoordinate && - phi == other.phi && - radius == other.radius && - clusterId == other.clusterId && - indexTableBinIndex == other.indexTableBinIndex; - } + GPUhdDefault() bool operator==(const Cluster&) const = default; GPUhd() void print() const; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h index 61072cb2410b7..118557c970c35 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/IndexTableUtils.h @@ -16,16 +16,19 @@ #ifndef TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_ #define TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_ +#include + +#include "ITStracking/Constants.h" #include "ITStracking/Configuration.h" #include "ITStracking/Definitions.h" #include "CommonConstants/MathConstants.h" #include "GPUCommonMath.h" #include "GPUCommonDef.h" -namespace o2 -{ -namespace its +namespace o2::its { + +template class IndexTableUtils { public: @@ -48,12 +51,13 @@ class IndexTableUtils int mNzBins = 0; int mNphiBins = 0; float mInversePhiBinSize = 0.f; - float mLayerZ[8] = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f}; - float mInverseZBinSize[8] = {0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f}; + std::array mLayerZ{}; + std::array mInverseZBinSize{}; }; +template template -inline void IndexTableUtils::setTrackingParameters(const T& params) +inline void IndexTableUtils::setTrackingParameters(const T& params) { mInversePhiBinSize = params.PhiBins / o2::constants::math::TwoPI; mNzBins = params.ZBins; @@ -66,28 +70,33 @@ inline void IndexTableUtils::setTrackingParameters(const T& params) } } -inline float IndexTableUtils::getInverseZCoordinate(const int layerIndex) const +template +inline float IndexTableUtils::getInverseZCoordinate(const int layerIndex) const { return 0.5f * mNzBins / mLayerZ[layerIndex]; } -GPUhdi() int IndexTableUtils::getZBinIndex(const int layerIndex, const float zCoordinate) const +template +GPUhdi() int IndexTableUtils::getZBinIndex(const int layerIndex, const float zCoordinate) const { return (zCoordinate + mLayerZ[layerIndex]) * mInverseZBinSize[layerIndex]; } -GPUhdi() int IndexTableUtils::getPhiBinIndex(const float currentPhi) const +template +GPUhdi() int IndexTableUtils::getPhiBinIndex(const float currentPhi) const { return (currentPhi * mInversePhiBinSize); } -GPUhdi() int IndexTableUtils::getBinIndex(const int zIndex, const int phiIndex) const +template +GPUhdi() int IndexTableUtils::getBinIndex(const int zIndex, const int phiIndex) const { return o2::gpu::GPUCommonMath::Min(phiIndex * mNzBins + zIndex, mNzBins * mNphiBins - 1); } -GPUhdi() int IndexTableUtils::countRowSelectedBins(const int* indexTable, const int phiBinIndex, - const int minZBinIndex, const int maxZBinIndex) const +template +GPUhdi() int IndexTableUtils::countRowSelectedBins(const int* indexTable, const int phiBinIndex, + const int minZBinIndex, const int maxZBinIndex) const { const int firstBinIndex{getBinIndex(minZBinIndex, phiBinIndex)}; const int maxBinIndex{firstBinIndex + maxZBinIndex - minZBinIndex + 1}; @@ -95,14 +104,14 @@ GPUhdi() int IndexTableUtils::countRowSelectedBins(const int* indexTable, const return indexTable[maxBinIndex] - indexTable[firstBinIndex]; } -GPUhdi() void IndexTableUtils::print() const +template +GPUhdi() void IndexTableUtils::print() const { printf("NzBins: %d, NphiBins: %d, InversePhiBinSize: %f\n", mNzBins, mNphiBins, mInversePhiBinSize); - for (int iLayer{0}; iLayer < 7; ++iLayer) { + for (int iLayer{0}; iLayer < nLayers; ++iLayer) { printf("Layer %d: Z: %f, InverseZBinSize: %f\n", iLayer, mLayerZ[iLayer], mInverseZBinSize[iLayer]); } } -} // namespace its -} // namespace o2 +} // namespace o2::its #endif /* TRACKINGITSU_INCLUDE_INDEXTABLEUTILS_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h index ebc885a3a35cf..a148049e50129 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TimeFrame.h @@ -65,6 +65,7 @@ class TimeFrameGPU; template struct TimeFrame { + using IndexTableUtilsN = IndexTableUtils; using CellSeedN = CellSeed; friend class gpu::TimeFrameGPU; @@ -273,7 +274,7 @@ struct TimeFrame { void printCellLUTs(); void printSliceInfo(const int, const int); - IndexTableUtils mIndexTableUtils; + IndexTableUtilsN mIndexTableUtils; bool mIsGPU = false; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h index 8647236b4f7e7..9d14bb91635a0 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackerTraits.h @@ -21,6 +21,7 @@ #include "DetectorsBase/Propagator.h" #include "ITStracking/Configuration.h" #include "ITStracking/MathUtils.h" +#include "ITStracking/IndexTableUtils.h" #include "ITStracking/TimeFrame.h" #include "ITStracking/Cell.h" #include "ITStracking/BoundedAllocator.h" @@ -40,9 +41,10 @@ class TrackITSExt; template class TrackerTraits { + public: + using IndexTableUtilsN = IndexTableUtils; using CellSeedN = CellSeed; - public: virtual ~TrackerTraits() = default; virtual void adoptTimeFrame(TimeFrame* tf) { mTimeFrame = tf; } virtual void initialiseTimeFrame(const int iteration) { mTimeFrame->initialise(iteration, mTrkParams[iteration], mTrkParams[iteration].NLayers); } @@ -119,7 +121,7 @@ inline const int4 TrackerTraits::getBinsRect(const int layerIndex, floa return getEmptyBinsRect(); } - const IndexTableUtils& utils{mTimeFrame->mIndexTableUtils}; + const IndexTableUtilsN& utils{mTimeFrame->mIndexTableUtils}; return int4{o2::gpu::GPUCommonMath::Max(0, utils.getZBinIndex(layerIndex, zRangeMin)), utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMin)), o2::gpu::GPUCommonMath::Min(mTrkParams[0].ZBins - 1, utils.getZBinIndex(layerIndex, zRangeMax)), // /!\ trkParams can potentially change across iterations diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h index 9d927e21202cc..787f299e15888 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackingInterface.h @@ -35,6 +35,8 @@ namespace o2::its class ITSTrackingInterface { static constexpr int NLayers{7}; + using VertexerN = Vertexer; + using VertexerTraitsN = VertexerTraits; using TrackerN = Tracker; using TrackerTraitsN = TrackerTraits; using TimeFrameN = TimeFrame; @@ -67,7 +69,7 @@ class ITSTrackingInterface virtual void finaliseCCDB(framework::ConcreteDataMatcher& matcher, void* obj); // Custom - void setTraitsFromProvider(VertexerTraits*, TrackerTraitsN*, TimeFrameN*); + void setTraitsFromProvider(VertexerTraitsN*, TrackerTraitsN*, TimeFrameN*); void setTrackingMode(TrackingMode::Type mode = TrackingMode::Unset) { mMode = mode; } auto getTracker() const { return mTracker.get(); } @@ -90,7 +92,7 @@ class ITSTrackingInterface bool mOverrideBeamEstimation = false; const o2::itsmft::TopologyDictionary* mDict = nullptr; std::unique_ptr mTracker = nullptr; - std::unique_ptr mVertexer = nullptr; + std::unique_ptr mVertexer = nullptr; const o2::dataformats::MeanVertexObject* mMeanVertex; std::shared_ptr mMemoryPool; std::shared_ptr mTaskArena; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h index 47e483c2e9f06..ab92e7c1a1523 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Vertexer.h @@ -35,47 +35,58 @@ namespace o2::its { +template class Vertexer { - static constexpr int NLayers{7}; - using TimeFrame7 = TimeFrame; + using TimeFrameN = TimeFrame; + using VertexerTraitsN = VertexerTraits; using LogFunc = std::function; public: - Vertexer(VertexerTraits* traits); + Vertexer(VertexerTraitsN* traits); virtual ~Vertexer() = default; Vertexer(const Vertexer&) = delete; Vertexer& operator=(const Vertexer&) = delete; - void adoptTimeFrame(TimeFrame7& tf); + void adoptTimeFrame(TimeFrameN& tf); auto& getVertParameters() const { return mTraits->getVertexingParameters(); } void setParameters(const std::vector& vertParams) { mVertParams = vertParams; } const auto& getParameters() const noexcept { return mVertParams; } void setMemoryPool(std::shared_ptr& pool) { mMemoryPool = pool; } std::vector exportVertices(); - VertexerTraits* getTraits() const { return mTraits; }; + VertexerTraitsN* getTraits() const { return mTraits; }; float clustersToVertices(LogFunc = [](const std::string& s) { std::cout << s << '\n'; }); void filterMCTracklets(); template - void findTracklets(T&&... args); - void findTrivialMCTracklets(); + void findTracklets(T&&... args) + { + mTraits->computeTracklets(std::forward(args)...); + } template - void validateTracklets(T&&... args); + void validateTracklets(T&&... args) + { + mTraits->computeTrackletMatching(std::forward(args)...); + } template - void findVertices(T&&... args); + void findVertices(T&&... args) + { + mTraits->computeVertices(std::forward(args)...); + } void addTruthSeeds() { mTraits->addTruthSeedingVertices(); } template - void initialiseVertexer(T&&... args); + void initialiseVertexer(T&&... args) + { + mTraits->initialise(std::forward(args)...); + } template void initialiseTimeFrame(T&&... args); // Utils - void dumpTraits() { mTraits->dumpVertexerTraits(); } template float evaluateTask(void (Vertexer::*task)(T...), std::string_view taskName, int iteration, LogFunc& logger, T&&... args); @@ -89,8 +100,8 @@ class Vertexer private: std::uint32_t mTimeFrameCounter = 0; - VertexerTraits* mTraits = nullptr; /// Observer pointer, not owned by this class - TimeFrame7* mTimeFrame = nullptr; /// Observer pointer, not owned by this class + VertexerTraitsN* mTraits = nullptr; /// Observer pointer, not owned by this class + TimeFrameN* mTimeFrame = nullptr; /// Observer pointer, not owned by this class std::vector mVertParams; std::shared_ptr mMemoryPool; @@ -107,32 +118,9 @@ class Vertexer static constexpr std::array StateNames{"Initialisation", "Tracklet finding", "Tracklet validation", "Vertex finding", "Truth seeding"}; }; +template template -void Vertexer::initialiseVertexer(T&&... args) -{ - mTraits->initialise(std::forward(args)...); -} - -template -void Vertexer::findTracklets(T&&... args) -{ - mTraits->computeTracklets(std::forward(args)...); -} - -template -inline void Vertexer::validateTracklets(T&&... args) -{ - mTraits->computeTrackletMatching(std::forward(args)...); -} - -template -inline void Vertexer::findVertices(T&&... args) -{ - mTraits->computeVertices(std::forward(args)...); -} - -template -float Vertexer::evaluateTask(void (Vertexer::*task)(T...), std::string_view taskName, int iteration, LogFunc& logger, T&&... args) +float Vertexer::evaluateTask(void (Vertexer::*task)(T...), std::string_view taskName, int iteration, LogFunc& logger, T&&... args) { float diff{0.f}; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h index 1213ad0a423b8..7da7617957179 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/VertexerTraits.h @@ -48,10 +48,11 @@ enum class TrackletMode { Layer1Layer2 = 2 }; +template class VertexerTraits { - static constexpr int NLayers{7}; - using TimeFrame7 = TimeFrame; + using IndexTableUtilsN = IndexTableUtils; + using TimeFrameN = TimeFrame; public: VertexerTraits() = default; @@ -62,8 +63,8 @@ class VertexerTraits return int4{0, 0, 0, 0}; } GPUhd() const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi); - 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() static const int4 getBinsRect(const Cluster&, const int, const float, float maxdeltaz, float maxdeltaphi, const IndexTableUtilsN&); + GPUhd() static const int2 getPhiBins(float phi, float deltaPhi, const IndexTableUtilsN&); GPUhd() const int2 getPhiBins(float phi, float deltaPhi) { return getPhiBins(phi, deltaPhi, mIndexTableUtils); } // virtual vertexer interface @@ -71,7 +72,7 @@ class VertexerTraits virtual void computeTracklets(const int iteration = 0); virtual void computeTrackletMatching(const int iteration = 0); virtual void computeVertices(const int iteration = 0); - virtual void adoptTimeFrame(TimeFrame7* tf) noexcept { mTimeFrame = tf; } + virtual void adoptTimeFrame(TimeFrameN* tf) noexcept { mTimeFrame = tf; } virtual void updateVertexingParameters(const std::vector& vrtPar, const TimeFrameGPUParameters& gpuTfPar); // truth tracking @@ -81,7 +82,6 @@ class VertexerTraits auto& getVertexingParameters() { return mVrtParams; } auto getVertexingParameters() const { return mVrtParams; } void setVertexingParameters(std::vector& vertParams) { mVrtParams = vertParams; } - void dumpVertexerTraits(); void setNThreads(int n, std::shared_ptr& arena); int getNThreads() { return mTaskArena->max_concurrency(); } virtual bool isGPU() const noexcept { return false; } @@ -112,10 +112,10 @@ class VertexerTraits protected: std::vector mVrtParams; - IndexTableUtils mIndexTableUtils; + IndexTableUtilsN mIndexTableUtils; // Frame related quantities - TimeFrame7* mTimeFrame = nullptr; // observer ptr + TimeFrameN* mTimeFrame = nullptr; // observer ptr private: std::shared_ptr mMemoryPool; std::shared_ptr mTaskArena; @@ -126,20 +126,23 @@ class VertexerTraits void debugComputeVertices(int iteration); }; -inline void VertexerTraits::initialise(const TrackingParameters& trackingParams, const int iteration) +template +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. } -GPUhdi() const int2 VertexerTraits::getPhiBins(float phi, float dPhi, const IndexTableUtils& utils) +template +GPUhdi() const int2 VertexerTraits::getPhiBins(float phi, float dPhi, const IndexTableUtilsN& utils) { return int2{utils.getPhiBinIndex(math_utils::getNormalizedPhi(phi - dPhi)), utils.getPhiBinIndex(math_utils::getNormalizedPhi(phi + dPhi))}; } -GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, const int layerIndex, - const float directionZIntersection, float maxdeltaz, float maxdeltaphi, - const IndexTableUtils& utils) +template +GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, const int layerIndex, + const float directionZIntersection, float maxdeltaz, float maxdeltaphi, + const IndexTableUtilsN& utils) { const float zRangeMin = directionZIntersection - 2 * maxdeltaz; const float phiRangeMin = currentCluster.phi - maxdeltaphi; @@ -157,8 +160,9 @@ GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, c utils.getPhiBinIndex(math_utils::getNormalizedPhi(phiRangeMax))}; } -GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, const int layerIndex, - const float directionZIntersection, float maxdeltaz, float maxdeltaphi) +template +GPUhdi() const int4 VertexerTraits::getBinsRect(const Cluster& currentCluster, const int layerIndex, + const float directionZIntersection, float maxdeltaz, float maxdeltaphi) { return VertexerTraits::getBinsRect(currentCluster, layerIndex, directionZIntersection, maxdeltaz, maxdeltaphi, mIndexTableUtils); } diff --git a/Detectors/ITSMFT/ITS/tracking/src/Cluster.cxx b/Detectors/ITSMFT/ITS/tracking/src/Cluster.cxx index 78f6683675947..c4d288bd61777 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Cluster.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Cluster.cxx @@ -37,7 +37,8 @@ Cluster::Cluster(const float x, const float y, const float z, const int index) // Nothing to do } -Cluster::Cluster(const int layerIndex, const IndexTableUtils& utils, const Cluster& other) +template +Cluster::Cluster(const int layerIndex, const IndexTableUtils& utils, const Cluster& other) : xCoordinate{other.xCoordinate}, yCoordinate{other.yCoordinate}, zCoordinate{other.zCoordinate}, @@ -51,7 +52,8 @@ Cluster::Cluster(const int layerIndex, const IndexTableUtils& utils, const Clust // Nothing to do } -Cluster::Cluster(const int layerIndex, const float3& primaryVertex, const IndexTableUtils& utils, const Cluster& other) +template +Cluster::Cluster(const int layerIndex, const float3& primaryVertex, const IndexTableUtils& utils, const Cluster& other) : xCoordinate{other.xCoordinate}, yCoordinate{other.yCoordinate}, zCoordinate{other.zCoordinate}, diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx index 280101d4dc4c7..58cfab73a7af3 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx @@ -814,7 +814,7 @@ void TrackerTraits::findRoads(const int iteration) TrackITSExt temporaryTrack{seed}; temporaryTrack.resetCovariance(); temporaryTrack.setChi2(0); - for (int iL{0}; iL < 7; ++iL) { + for (int iL{0}; iL < nLayers; ++iL) { temporaryTrack.setExternalClusterIndex(iL, seed.getCluster(iL), seed.getCluster(iL) != constants::UnusedIndex); } diff --git a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx index d2e96aad40a9e..d6ac3b5229509 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/TrackingInterface.cxx @@ -382,11 +382,11 @@ void ITSTrackingInterface::printSummary() const mTracker->printSummary(); } -void ITSTrackingInterface::setTraitsFromProvider(VertexerTraits* vertexerTraits, +void ITSTrackingInterface::setTraitsFromProvider(VertexerTraitsN* vertexerTraits, TrackerTraitsN* trackerTraits, TimeFrameN* frame) { - mVertexer = std::make_unique(vertexerTraits); + mVertexer = std::make_unique(vertexerTraits); mTracker = std::make_unique(trackerTraits); mTimeFrame = frame; mVertexer->adoptTimeFrame(*mTimeFrame); diff --git a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx index 94d69f7968ac8..69dddbf367653 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/Vertexer.cxx @@ -26,7 +26,8 @@ namespace o2::its { -Vertexer::Vertexer(VertexerTraits* traits) : mTraits(traits) +template +Vertexer::Vertexer(VertexerTraitsN* traits) : mTraits(traits) { if (!mTraits) { LOG(fatal) << "nullptr passed to ITS vertexer construction."; @@ -34,7 +35,8 @@ Vertexer::Vertexer(VertexerTraits* traits) : mTraits(traits) mVertParams.resize(1); } -float Vertexer::clustersToVertices(LogFunc logger) +template +float Vertexer::clustersToVertices(LogFunc logger) { LogFunc evalLog = [](const std::string&) {}; @@ -87,16 +89,18 @@ float Vertexer::clustersToVertices(LogFunc logger) return timeInit + timeTracklet + timeSelection + timeVertexing; } -void Vertexer::adoptTimeFrame(TimeFrame7& tf) +template +void Vertexer::adoptTimeFrame(TimeFrameN& tf) { mTimeFrame = &tf; mTraits->adoptTimeFrame(&tf); } -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) +template +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) { 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)); @@ -107,4 +111,6 @@ void Vertexer::printEpilog(LogFunc& logger, } } +template class Vertexer<7>; + } // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx index bcafa98972d78..153d7b6faa358 100644 --- a/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx +++ b/Detectors/ITSMFT/ITS/tracking/src/VertexerTraits.cxx @@ -28,9 +28,10 @@ #include "DetectorsRaw/HBFUtils.h" #include "CommonUtils/TreeStreamRedirector.h" -using namespace o2::its; +namespace o2::its +{ -template +template static void trackleterKernelHost( const gsl::span& clustersNextLayer, // 0 2 const gsl::span& clustersCurrentLayer, // 1 1 @@ -39,7 +40,7 @@ static void trackleterKernelHost( const float phiCut, bounded_vector& tracklets, gsl::span foundTracklets, - const IndexTableUtils& utils, + const IndexTableUtils& utils, const short pivotRof, const short targetRof, gsl::span rofFoundTrackletsOffsets, // we want to change those, to keep track of the offset in deltaRof>0 @@ -51,7 +52,7 @@ static void trackleterKernelHost( for (int iCurrentLayerClusterIndex = 0; iCurrentLayerClusterIndex < clustersCurrentLayer.size(); ++iCurrentLayerClusterIndex) { int storedTracklets{0}; const Cluster& currentCluster{clustersCurrentLayer[iCurrentLayerClusterIndex]}; - const int4 selectedBinsRect{VertexerTraits::getBinsRect(currentCluster, (int)Mode, 0.f, 50.f, phiCut / 2, utils)}; + const int4 selectedBinsRect{VertexerTraits::getBinsRect(currentCluster, (int)Mode, 0.f, 50.f, phiCut / 2, utils)}; if (selectedBinsRect.x != 0 || selectedBinsRect.y != 0 || selectedBinsRect.z != 0 || selectedBinsRect.w != 0) { int phiBinsNum{selectedBinsRect.w - selectedBinsRect.y + 1}; if (phiBinsNum < 0) { @@ -151,7 +152,8 @@ static void trackletSelectionKernelHost( } } -void VertexerTraits::updateVertexingParameters(const std::vector& vrtPar, const TimeFrameGPUParameters& tfPar) +template +void VertexerTraits::updateVertexingParameters(const std::vector& vrtPar, const TimeFrameGPUParameters& tfPar) { mVrtParams = vrtPar; mIndexTableUtils.setTrackingParameters(vrtPar[0]); @@ -162,7 +164,8 @@ void VertexerTraits::updateVertexingParameters(const std::vector +void VertexerTraits::computeTracklets(const int iteration) { mTaskArena->execute([&] { tbb::parallel_for( @@ -283,7 +286,8 @@ void VertexerTraits::computeTracklets(const int iteration) #endif } -void VertexerTraits::computeTrackletMatching(const int iteration) +template +void VertexerTraits::computeTrackletMatching(const int iteration) { mTaskArena->execute([&] { tbb::parallel_for( @@ -341,7 +345,8 @@ void VertexerTraits::computeTrackletMatching(const int iteration) deepVectorClear(mTimeFrame->getTracklets()[1]); } -void VertexerTraits::computeVertices(const int iteration) +template +void VertexerTraits::computeVertices(const int iteration) { auto nsigmaCut{std::min(mVrtParams[iteration].vertNsigmaCut * mVrtParams[iteration].vertNsigmaCut * (mVrtParams[iteration].vertRadiusSigma * mVrtParams[iteration].vertRadiusSigma + mVrtParams[iteration].trackletSigma * mVrtParams[iteration].trackletSigma), 1.98f)}; bounded_vector vertices(mMemoryPool.get()); @@ -497,7 +502,8 @@ void VertexerTraits::computeVertices(const int iteration) #endif } -void VertexerTraits::addTruthSeedingVertices() +template +void VertexerTraits::addTruthSeedingVertices() { LOGP(info, "Using truth seeds as vertices; will skip computations"); mTimeFrame->resetRofPV(); @@ -562,7 +568,8 @@ void VertexerTraits::addTruthSeedingVertices() LOGP(info, "Found {}/{} ROFs with {} vertices -> ={:.2f}", vertices.size(), mTimeFrame->getNrof(), nVerts, (float)nVerts / (float)vertices.size()); } -void VertexerTraits::setNThreads(int n, std::shared_ptr& arena) +template +void VertexerTraits::setNThreads(int n, std::shared_ptr& arena) { #if defined(VTX_DEBUG) LOGP(info, "Vertexer with debug output forcing single thread"); @@ -578,7 +585,8 @@ void VertexerTraits::setNThreads(int n, std::shared_ptr& arena) #endif } -void VertexerTraits::debugComputeTracklets(int iteration) +template +void VertexerTraits::debugComputeTracklets(int iteration) { auto stream = new utils::TreeStreamRedirector("artefacts_tf.root", "recreate"); LOGP(info, "writing debug output for computeTracklets"); @@ -597,7 +605,8 @@ void VertexerTraits::debugComputeTracklets(int iteration) delete stream; } -void VertexerTraits::debugComputeTrackletMatching(int iteration) +template +void VertexerTraits::debugComputeTrackletMatching(int iteration) { auto stream = new utils::TreeStreamRedirector("artefacts_tf.root", "update"); LOGP(info, "writing debug output for computeTrackletMatching"); @@ -718,7 +727,8 @@ void VertexerTraits::debugComputeTrackletMatching(int iteration) delete stream; } -void VertexerTraits::debugComputeVertices(int iteration) +template +void VertexerTraits::debugComputeVertices(int iteration) { auto stream = new utils::TreeStreamRedirector("artefacts_tf.root", "update"); LOGP(info, "writing debug output for computeVertices"); @@ -831,3 +841,6 @@ void VertexerTraits::debugComputeVertices(int iteration) stream->Close(); delete stream; } + +template class VertexerTraits<7>; +} // namespace o2::its diff --git a/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/CookedTrackerSpec.h b/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/CookedTrackerSpec.h index eb55081889c12..4ecc98eed9cfb 100644 --- a/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/CookedTrackerSpec.h +++ b/Detectors/ITSMFT/ITS/workflow/include/ITSWorkflow/CookedTrackerSpec.h @@ -58,8 +58,8 @@ class CookedTrackerDPL : public Task const o2::itsmft::TopologyDictionary* mDict = nullptr; std::unique_ptr mGRP = nullptr; o2::its::CookedTracker mTracker; - std::unique_ptr mVertexerTraitsPtr = nullptr; - std::unique_ptr mVertexerPtr = nullptr; + std::unique_ptr> mVertexerTraitsPtr = nullptr; + std::unique_ptr> mVertexerPtr = nullptr; std::shared_ptr mMemoryPool; std::shared_ptr mTaskArena; TStopwatch mTimer; diff --git a/Detectors/ITSMFT/ITS/workflow/src/CookedTrackerSpec.cxx b/Detectors/ITSMFT/ITS/workflow/src/CookedTrackerSpec.cxx index 98662e46269d1..b989a78e59b7c 100644 --- a/Detectors/ITSMFT/ITS/workflow/src/CookedTrackerSpec.cxx +++ b/Detectors/ITSMFT/ITS/workflow/src/CookedTrackerSpec.cxx @@ -53,8 +53,8 @@ using Vertex = o2::dataformats::Vertex>; CookedTrackerDPL::CookedTrackerDPL(std::shared_ptr gr, bool useMC, int trgType, TrackingMode::Type trMode) : mGGCCDBRequest(gr), mUseMC(useMC), mUseTriggers{trgType}, mMode(trMode) { - mVertexerTraitsPtr = std::make_unique(); - mVertexerPtr = std::make_unique(mVertexerTraitsPtr.get()); + mVertexerTraitsPtr = std::make_unique>(); + mVertexerPtr = std::make_unique>(mVertexerTraitsPtr.get()); } void CookedTrackerDPL::init(InitContext& ic) diff --git a/GPU/GPUTracking/Base/GPUReconstruction.cxx b/GPU/GPUTracking/Base/GPUReconstruction.cxx index c7b61a976021a..09aae2aacf16d 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.cxx +++ b/GPU/GPUTracking/Base/GPUReconstruction.cxx @@ -111,13 +111,13 @@ 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<7>); } if (vertexerTraits) { - vertexerTraits->reset(new o2::its::VertexerTraits); + vertexerTraits->reset(new o2::its::VertexerTraits<7>); } if (timeFrame) { timeFrame->reset(new o2::its::TimeFrame<7>); diff --git a/GPU/GPUTracking/Base/GPUReconstruction.h b/GPU/GPUTracking/Base/GPUReconstruction.h index 70c504cdee6e1..420e602e61352 100644 --- a/GPU/GPUTracking/Base/GPUReconstruction.h +++ b/GPU/GPUTracking/Base/GPUReconstruction.h @@ -38,6 +38,7 @@ namespace o2::its { template class TrackerTraits; +template class VertexerTraits; template class TimeFrame; @@ -191,7 +192,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 f4ef3b464c24f..c4202e9980d24 100644 --- a/GPU/GPUTracking/Base/GPUReconstructionIncludesITS.h +++ b/GPU/GPUTracking/Base/GPUReconstructionIncludesITS.h @@ -21,26 +21,28 @@ #include "ITStracking/TimeFrame.h" #if defined(__CUDACC__) || defined(__HIPCC__) #include "ITStrackingGPU/TrackerTraitsGPU.h" -#include "ITStrackingGPU/VertexerTraitsGPU.h" +// #include "ITStrackingGPU/VertexerTraitsGPU.h" #include "ITStrackingGPU/TimeFrameGPU.h" #endif #else namespace o2::its { +template class VertexerTraits { }; -template +template class TrackerTraits { }; -template +template class TimeFrame { }; -class VertexerTraitsGPU : public VertexerTraits -{ -}; +// template +// class VertexerTraitsGPU : public VertexerTraits +// { +// }; template class TrackerTraitsGPU : public TrackerTraits { diff --git a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu index dc904fa96fa2d..0ee91dd2eaf21 100644 --- a/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu +++ b/GPU/GPUTracking/Base/cuda/GPUReconstructionCUDA.cu @@ -91,13 +91,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::VertexerTraits); // TODO gpu-code to be implemented + vertexerTraits->reset(new o2::its::VertexerTraits<7>); // 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 1cc7e0fc819ff..b3562eff4096d 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 18fb5ff1de939..e53f5db3a2549 100644 --- a/GPU/GPUTracking/Global/GPUChainITS.cxx +++ b/GPU/GPUTracking/Global/GPUChainITS.cxx @@ -56,7 +56,7 @@ o2::its::TrackerTraits<7>* GPUChainITS::GetITSTrackerTraits() return mITSTrackerTraits.get(); } -o2::its::VertexerTraits* GPUChainITS::GetITSVertexerTraits() +o2::its::VertexerTraits<7>* GPUChainITS::GetITSVertexerTraits() { if (mITSVertexerTraits == nullptr) { mRec->GetITSTraits(nullptr, &mITSVertexerTraits, nullptr); diff --git a/GPU/GPUTracking/Global/GPUChainITS.h b/GPU/GPUTracking/Global/GPUChainITS.h index 6821f63845b95..a607f66322bab 100644 --- a/GPU/GPUTracking/Global/GPUChainITS.h +++ b/GPU/GPUTracking/Global/GPUChainITS.h @@ -45,13 +45,13 @@ class GPUChainITS final : public GPUChain void MemorySize(size_t&, size_t&) final {}; o2::its::TrackerTraits<7>* GetITSTrackerTraits(); - o2::its::VertexerTraits* GetITSVertexerTraits(); + o2::its::VertexerTraits<7>* GetITSVertexerTraits(); o2::its::TimeFrame<7>* GetITSTimeframe(); protected: GPUChainITS(GPUReconstruction* rec); std::unique_ptr> mITSTrackerTraits; - std::unique_ptr mITSVertexerTraits; + std::unique_ptr> mITSVertexerTraits; std::unique_ptr> mITSTimeFrame; std::unique_ptr mFrameworkAllocator; }; diff --git a/GPU/GPUTracking/Interface/GPUO2Interface.cxx b/GPU/GPUTracking/Interface/GPUO2Interface.cxx index f1504a430db36..65907528a3dba 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<7>*& trackerTraits, o2::its::VertexerTraits*& vertexerTraits, o2::its::TimeFrame<7>*& timeFrame) +void GPUO2Interface::GetITSTraits(o2::its::TrackerTraits<7>*& trackerTraits, o2::its::VertexerTraits<7>*& 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 03b24c2b23877..9b7390f2ed663 100644 --- a/GPU/GPUTracking/Interface/GPUO2Interface.h +++ b/GPU/GPUTracking/Interface/GPUO2Interface.h @@ -45,6 +45,7 @@ namespace o2::its { template class TrackerTraits; +template class VertexerTraits; template class TimeFrame; @@ -79,7 +80,7 @@ class GPUO2Interface void DumpEvent(int32_t nEvent, GPUTrackingInOutPointers* data); void DumpSettings(); - void GetITSTraits(o2::its::TrackerTraits<7>*& trackerTraits, o2::its::VertexerTraits*& vertexerTraits, o2::its::TimeFrame<7>*& timeFrame); + void GetITSTraits(o2::its::TrackerTraits<7>*& trackerTraits, o2::its::VertexerTraits<7>*& 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/src/GPUWorkflowITS.cxx b/GPU/Workflow/src/GPUWorkflowITS.cxx index 31ccaad8c2783..fe55b83633972 100644 --- a/GPU/Workflow/src/GPUWorkflowITS.cxx +++ b/GPU/Workflow/src/GPUWorkflowITS.cxx @@ -37,7 +37,7 @@ int32_t GPURecoWorkflowSpec::runITSTracking(o2::framework::ProcessingContext& pc void GPURecoWorkflowSpec::initFunctionITS(o2::framework::InitContext& ic) { - o2::its::VertexerTraits* vtxTraits = nullptr; + o2::its::VertexerTraits<7>* vtxTraits = nullptr; o2::its::TrackerTraits<7>* trkTraits = nullptr; #ifdef ENABLE_UPGRADES if (mSpecConfig.isITS3) {