diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h index fc3656aef800c..d81ba4426ca55 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Cell.h @@ -29,22 +29,6 @@ namespace o2::its class Cell final { public: - GPUhdDefault() Cell() = default; - GPUhd() Cell(const int firstClusterIndex, const int secondClusterIndex, const int thirdClusterIndex, - const int firstTrackletIndex, const int secondTrackletIndex) - : mFirstClusterIndex(firstClusterIndex), - mSecondClusterIndex(secondClusterIndex), - mThirdClusterIndex(thirdClusterIndex), - mFirstTrackletIndex(firstTrackletIndex), - mSecondTrackletIndex(secondTrackletIndex), - mLevel(1) {} - GPUhdDefault() Cell(const Cell&) = default; - GPUhdDefault() Cell(Cell&&) = default; - GPUhdDefault() ~Cell() = default; - - GPUhdDefault() Cell& operator=(const Cell&) = default; - GPUhdDefault() Cell& operator=(Cell&&) noexcept = default; - GPUhd() int getFirstClusterIndex() const { return mFirstClusterIndex; }; GPUhd() int getSecondClusterIndex() const { return mSecondClusterIndex; }; GPUhd() int getThirdClusterIndex() const { return mThirdClusterIndex; }; diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h index 48cc45e44cf1c..ab9d0c2e4d1a6 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Constants.h @@ -31,11 +31,6 @@ constexpr int UnusedIndex{-1}; constexpr float Resolution{0.0005f}; constexpr float Radl = 9.36f; // Radiation length of Si [cm] constexpr float Rho = 2.33f; // Density of Si [g/cm^3] -namespace its // to be removed -{ -constexpr int UnusedIndex{-1}; -constexpr float Resolution{0.0005f}; -} // namespace its } // namespace o2::its::constants #endif /* TRACKINGITSU_INCLUDE_CONSTANTS_H_ */ diff --git a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Road.h b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Road.h index e9cd306e63bc5..75f187f31652b 100644 --- a/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Road.h +++ b/Detectors/ITSMFT/ITS/tracking/include/ITStracking/Road.h @@ -45,14 +45,14 @@ class Road final GPUhd() void resetRoad() { for (int i = 0; i < maxRoadSize; i++) { - mCellIds[i] = constants::its::UnusedIndex; + mCellIds[i] = constants::UnusedIndex; } mRoadSize = 0; } GPUhd() void addCell(int cellLayer, int cellId) { - if (mCellIds[cellLayer] == constants::its::UnusedIndex) { + if (mCellIds[cellLayer] == constants::UnusedIndex) { ++mRoadSize; } @@ -60,8 +60,7 @@ class Road final } private: - int mCellIds[maxRoadSize]{constants::its::UnusedIndex}; - // int mLabel; + int mCellIds[maxRoadSize]{constants::UnusedIndex}; unsigned char mRoadSize{0}; bool mIsFakeRoad{false}; }; diff --git a/GPU/GPUTracking/Base/GPUConstantMem.h b/GPU/GPUTracking/Base/GPUConstantMem.h index 532c270431d99..ffb17997b9190 100644 --- a/GPU/GPUTracking/Base/GPUConstantMem.h +++ b/GPU/GPUTracking/Base/GPUConstantMem.h @@ -26,7 +26,6 @@ #include "GPUTPCConvert.h" #include "GPUTPCCompression.h" #include "GPUTPCDecompression.h" -#include "GPUITSFitter.h" #include "GPUTPCClusterFinder.h" #include "GPUTrackingRefit.h" @@ -50,7 +49,6 @@ struct GPUConstantMem { GPUTRDTrackerGPU trdTrackerGPU; GPUTRDTracker trdTrackerO2; GPUTPCClusterFinder tpcClusterer[GPUCA_NSECTORS]; - GPUITSFitter itsFitter; GPUTrackingRefitProcessor trackingRefit; GPUTrackingInOutPointers ioPtrs; GPUCalibObjectsConst calibObjects; diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index 0560ea2dd12d2..a7159549322a0 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -173,8 +173,6 @@ set(SRCS_NO_CINT ${SRCS_NO_CINT} display/GPUDisplayInterface.cxx) set(SRCS_NO_CINT ${SRCS_NO_CINT} Global/GPUChainITS.cxx - ITS/GPUITSFitter.cxx - ITS/GPUITSFitterKernels.cxx dEdx/GPUdEdx.cxx TPCConvert/GPUTPCConvert.cxx TPCConvert/GPUTPCConvertKernel.cxx @@ -220,7 +218,6 @@ set(SRCS_NO_H ${SRCS_NO_H} TPCClusterFinder/GPUTPCClusterFinderDump.cxx) set(HDRS_INSTALL ${HDRS_INSTALL} - ITS/GPUITSTrack.h TPCClusterFinder/CfArray2D.h TPCClusterFinder/CfConsts.h TPCClusterFinder/CfFragment.h @@ -298,7 +295,6 @@ set(INCDIRS ${CMAKE_CURRENT_SOURCE_DIR}/SectorTracker ${CMAKE_CURRENT_SOURCE_DIR}/TPCConvert ${CMAKE_CURRENT_SOURCE_DIR}/dEdx - ${CMAKE_CURRENT_SOURCE_DIR}/ITS ${CMAKE_CURRENT_SOURCE_DIR}/TRDTracking ${CMAKE_CURRENT_SOURCE_DIR}/qa ${CMAKE_CURRENT_SOURCE_DIR}/Global diff --git a/GPU/GPUTracking/Definitions/GPUDefParametersConstants.h b/GPU/GPUTracking/Definitions/GPUDefParametersConstants.h index 78036e47fc49d..747ee84a906a6 100644 --- a/GPU/GPUTracking/Definitions/GPUDefParametersConstants.h +++ b/GPU/GPUTracking/Definitions/GPUDefParametersConstants.h @@ -32,7 +32,6 @@ // Default maximum numbers #define GPUCA_MAX_CLUSTERS ((size_t) 1024 * 1024 * 1024) // Maximum number of TPC clusters #define GPUCA_MAX_TRD_TRACKLETS ((size_t) 128 * 1024) // Maximum number of TRD tracklets -#define GPUCA_MAX_ITS_FIT_TRACKS ((size_t) 96 * 1024) // Max number of tracks for ITS track fit #define GPUCA_MEMORY_SIZE ((size_t) 6 * 1024 * 1024 * 1024) // Size of memory allocated on Device #define GPUCA_HOST_MEMORY_SIZE ((size_t) 1 * 1024 * 1024 * 1024) // Size of memory allocated on Host #define GPUCA_GPU_STACK_SIZE ((size_t) 8 * 1024) // Stack size per GPU thread diff --git a/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h b/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h index 8fe8e8ca68e44..0591ac8c58630 100644 --- a/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h +++ b/GPU/GPUTracking/Definitions/GPUDefParametersDefaults.h @@ -438,9 +438,6 @@ #ifndef GPUCA_LB_GPUTPCGMO2Output_output #define GPUCA_LB_GPUTPCGMO2Output_output 256 #endif - #ifndef GPUCA_LB_GPUITSFitterKernels - #define GPUCA_LB_GPUITSFitterKernels 256 - #endif #ifndef GPUCA_LB_GPUTPCStartHitsFinder #define GPUCA_LB_GPUTPCStartHitsFinder 256 #endif diff --git a/GPU/GPUTracking/Global/GPUChainITS.cxx b/GPU/GPUTracking/Global/GPUChainITS.cxx index cbc19100fe4fa..18fb5ff1de939 100644 --- a/GPU/GPUTracking/Global/GPUChainITS.cxx +++ b/GPU/GPUTracking/Global/GPUChainITS.cxx @@ -44,22 +44,7 @@ GPUChainITS::~GPUChainITS() mITSVertexerTraits.reset(); } -GPUChainITS::GPUChainITS(GPUReconstruction* rec, uint32_t maxTracks) : GPUChain(rec), mMaxTracks(maxTracks) {} - -void GPUChainITS::RegisterPermanentMemoryAndProcessors() { mRec->RegisterGPUProcessor(&processors()->itsFitter, GetRecoStepsGPU() & RecoStep::ITSTracking); } - -void GPUChainITS::RegisterGPUProcessors() -{ - if (GetRecoStepsGPU() & RecoStep::ITSTracking) { - mRec->RegisterGPUDeviceProcessor(&processorsShadow()->itsFitter, &processors()->itsFitter); - } -} - -void GPUChainITS::MemorySize(size_t& gpuMem, size_t& pageLockedHostMem) -{ - gpuMem = mMaxTracks * sizeof(GPUITSTrack) + GPUCA_MEMALIGN; - pageLockedHostMem = gpuMem; -} +GPUChainITS::GPUChainITS(GPUReconstruction* rec) : GPUChain(rec) {} int32_t GPUChainITS::Init() { return 0; } diff --git a/GPU/GPUTracking/Global/GPUChainITS.h b/GPU/GPUTracking/Global/GPUChainITS.h index 150d66031d084..6821f63845b95 100644 --- a/GPU/GPUTracking/Global/GPUChainITS.h +++ b/GPU/GPUTracking/Global/GPUChainITS.h @@ -29,32 +29,31 @@ class GPUFrameworkExternalAllocator; namespace o2::gpu { -class GPUChainITS : public GPUChain +class GPUChainITS final : public GPUChain { friend class GPUReconstruction; public: ~GPUChainITS() override; - void RegisterPermanentMemoryAndProcessors() override; - void RegisterGPUProcessors() override; int32_t Init() override; int32_t PrepareEvent() override; int32_t Finalize() override; int32_t RunChain() override; - void MemorySize(size_t& gpuMem, size_t& pageLockedHostMem) override; + + void RegisterPermanentMemoryAndProcessors() final {}; + void RegisterGPUProcessors() final {}; + void MemorySize(size_t&, size_t&) final {}; o2::its::TrackerTraits<7>* GetITSTrackerTraits(); o2::its::VertexerTraits* GetITSVertexerTraits(); o2::its::TimeFrame<7>* GetITSTimeframe(); protected: - GPUChainITS(GPUReconstruction* rec, uint32_t maxTracks = GPUCA_MAX_ITS_FIT_TRACKS); + GPUChainITS(GPUReconstruction* rec); std::unique_ptr> mITSTrackerTraits; std::unique_ptr mITSVertexerTraits; std::unique_ptr> mITSTimeFrame; std::unique_ptr mFrameworkAllocator; - - uint32_t mMaxTracks; }; } // namespace o2::gpu diff --git a/GPU/GPUTracking/ITS/GPUITSFitter.cxx b/GPU/GPUTracking/ITS/GPUITSFitter.cxx deleted file mode 100644 index a954d430f7ac3..0000000000000 --- a/GPU/GPUTracking/ITS/GPUITSFitter.cxx +++ /dev/null @@ -1,64 +0,0 @@ -// Copyright 2019-2020 CERN and copyright holders of ALICE O2. -// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. -// All rights not expressly granted are reserved. -// -// This software is distributed under the terms of the GNU General Public -// License v3 (GPL Version 3), copied verbatim in the file "COPYING". -// -// In applying this license CERN does not waive the privileges and immunities -// granted to it by virtue of its status as an Intergovernmental Organization -// or submit itself to any jurisdiction. - -/// \file GPUITSFitter.cxx -/// \author David Rohr, Maximiliano Puccio - -#include "GPUITSFitter.h" - -#include "ITStracking/Road.h" -#include "ITStracking/Cluster.h" -#include "GPUITSTrack.h" -#include "GPUReconstruction.h" - -using namespace o2::gpu; - -#ifndef GPUCA_GPUCODE -void GPUITSFitter::InitializeProcessor() -{ -} - -void* GPUITSFitter::SetPointersInput(void* mem) -{ - computePointerWithAlignment(mem, mRoads, mNumberOfRoads); - for (int32_t i = 0; i < 7; i++) { - computePointerWithAlignment(mem, mTF[i], mNTF[i]); - } - return mem; -} - -void* GPUITSFitter::SetPointersTracks(void* mem) -{ - computePointerWithAlignment(mem, mTracks, mNMaxTracks); - return mem; -} - -void* GPUITSFitter::SetPointersMemory(void* mem) -{ - computePointerWithAlignment(mem, mMemory, 1); - return mem; -} - -void GPUITSFitter::RegisterMemoryAllocation() -{ - AllocateAndInitializeLate(); - mMemoryResInput = mRec->RegisterMemoryAllocation(this, &GPUITSFitter::SetPointersInput, GPUMemoryResource::MEMORY_INPUT, "ITSInput"); - mMemoryResTracks = mRec->RegisterMemoryAllocation(this, &GPUITSFitter::SetPointersTracks, GPUMemoryResource::MEMORY_OUTPUT, "ITSTracks"); - mMemoryResMemory = mRec->RegisterMemoryAllocation(this, &GPUITSFitter::SetPointersMemory, GPUMemoryResource::MEMORY_PERMANENT, "ITSMemory"); -} - -void GPUITSFitter::SetMaxData(const GPUTrackingInOutPointers& io) { mNMaxTracks = mNumberOfRoads; } -#endif - -void GPUITSFitter::clearMemory() -{ - new (mMemory) Memory; -} diff --git a/GPU/GPUTracking/ITS/GPUITSFitter.h b/GPU/GPUTracking/ITS/GPUITSFitter.h deleted file mode 100644 index 9c0995a80749f..0000000000000 --- a/GPU/GPUTracking/ITS/GPUITSFitter.h +++ /dev/null @@ -1,102 +0,0 @@ -// Copyright 2019-2020 CERN and copyright holders of ALICE O2. -// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. -// All rights not expressly granted are reserved. -// -// This software is distributed under the terms of the GNU General Public -// License v3 (GPL Version 3), copied verbatim in the file "COPYING". -// -// In applying this license CERN does not waive the privileges and immunities -// granted to it by virtue of its status as an Intergovernmental Organization -// or submit itself to any jurisdiction. - -/// \file GPUITSFitter.h -/// \author David Rohr, Maximiliano Puccio - -#ifndef GPUITSFITTER_H -#define GPUITSFITTER_H - -#include "GPUProcessor.h" -#include "GPUITSTrack.h" - -namespace o2::its -{ -template -class Road; -struct TrackingFrameInfo; -struct Cluster; -class Cell; -} // namespace o2::its - -namespace o2::gpu -{ -class GPUITSTrack; - -class GPUITSFitter : public GPUProcessor -{ - public: -#ifndef GPUCA_GPUCODE - void InitializeProcessor(); - void RegisterMemoryAllocation(); - void SetMaxData(const GPUTrackingInOutPointers& io); - - void* SetPointersInput(void* mem); - void* SetPointersTracks(void* mem); - void* SetPointersMemory(void* mem); -#endif - - GPUd() o2::its::Road<5>* roads() - { - return mRoads; - } - GPUd() void SetNumberOfRoads(int32_t v) { mNumberOfRoads = v; } - GPUd() int32_t NumberOfRoads() { return mNumberOfRoads; } - GPUd() GPUITSTrack* tracks() - { - return mTracks; - } - GPUd() GPUAtomic(uint32_t) & NumberOfTracks() - { - return mMemory->mNumberOfTracks; - } - GPUd() void SetNumberOfLayers(int32_t i) { mNumberOfLayers = i; } - GPUd() int32_t NumberOfLayers() { return mNumberOfLayers; } - GPUd() void SetNumberTF(int32_t i, int32_t v) { mNTF[i] = v; } - GPUd() o2::its::TrackingFrameInfo** trackingFrame() - { - return mTF; - } - GPUd() const o2::its::Cluster** clusters() - { - return mClusterPtrs; - } - GPUd() const o2::its::Cell** cells() - { - return mCellPtrs; - } - - void clearMemory(); - - struct Memory { - GPUAtomic(uint32_t) mNumberOfTracks = 0; - }; - - protected: - int32_t mNumberOfLayers; - int32_t mNumberOfRoads = 0; - int32_t mNMaxTracks = 0; - int32_t* mNTF = nullptr; - Memory* mMemory = nullptr; - o2::its::Road<5>* mRoads = nullptr; - o2::its::TrackingFrameInfo** mTF = {nullptr}; - GPUITSTrack* mTracks = nullptr; - - const o2::its::Cluster** mClusterPtrs; - const o2::its::Cell** mCellPtrs; - - int16_t mMemoryResInput = -1; - int16_t mMemoryResTracks = -1; - int16_t mMemoryResMemory = -1; -}; -} // namespace o2::gpu - -#endif diff --git a/GPU/GPUTracking/ITS/GPUITSFitterKernels.cxx b/GPU/GPUTracking/ITS/GPUITSFitterKernels.cxx deleted file mode 100644 index 798efb7d7964e..0000000000000 --- a/GPU/GPUTracking/ITS/GPUITSFitterKernels.cxx +++ /dev/null @@ -1,201 +0,0 @@ -// Copyright 2019-2020 CERN and copyright holders of ALICE O2. -// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. -// All rights not expressly granted are reserved. -// -// This software is distributed under the terms of the GNU General Public -// License v3 (GPL Version 3), copied verbatim in the file "COPYING". -// -// In applying this license CERN does not waive the privileges and immunities -// granted to it by virtue of its status as an Intergovernmental Organization -// or submit itself to any jurisdiction. - -/// \file GPUITSFitterKernels.cxx -/// \author David Rohr, Maximiliano Puccio - -#include "GPUITSFitterKernels.h" -#include "GPUConstantMem.h" - -#include "ITStracking/Constants.h" -#include "ITStracking/MathUtils.h" -#include "ITStracking/Road.h" -#include "ITStracking/Cluster.h" -#include "ITStracking/Cell.h" -#include "CommonConstants/MathConstants.h" - -#if defined(CA_DEBUG) && !defined(GPUCA_GPUCODE_DEVICE) -#include -#endif - -using namespace o2::gpu; -using namespace o2; -using namespace o2::its; - -GPUdii() bool GPUITSFitterKernels::fitTrack(GPUITSFitter& GPUrestrict() Fitter, GPUTPCGMPropagator& GPUrestrict() prop, GPUITSTrack& GPUrestrict() track, int32_t start, int32_t end, int32_t step) -{ - for (int32_t iLayer{start}; iLayer != end; iLayer += step) { - if (track.mClusters[iLayer] == o2::its::constants::its::UnusedIndex) { - continue; - } - const TrackingFrameInfo& GPUrestrict() trackingHit = Fitter.trackingFrame()[iLayer][track.mClusters[iLayer]]; - - if (prop.PropagateToXAlpha(trackingHit.xTrackingFrame, trackingHit.alphaTrackingFrame, step > 0)) { - return false; - } - - if (prop.Update(trackingHit.positionTrackingFrame[0], trackingHit.positionTrackingFrame[1], 0, false, trackingHit.covarianceTrackingFrame[0], trackingHit.covarianceTrackingFrame[2])) { - return false; - } - - /*const float xx0 = (iLayer > 2) ? 0.008f : 0.003f; // Rough layer thickness //FIXME - constexpr float radiationLength = 9.36f; // Radiation length of Si [cm] - constexpr float density = 2.33f; // Density of Si [g/cm^3] - if (!track.correctForMaterial(xx0, xx0 * radiationLength * density, true)) - return false;*/ - } - return true; -} - -template <> -GPUdii() void GPUITSFitterKernels::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors) -{ - GPUITSFitter& Fitter = processors.itsFitter; - - GPUTPCGMPropagator prop; - prop.SetPolynomialField(&processors.param.polynomialField); - prop.SetMaxSinPhi(GPUCA_MAX_SIN_PHI); - prop.SetFitInProjections(1); - float bz = -5.f; // FIXME - -#ifdef CA_DEBUG - int32_t roadCounters[4]{0, 0, 0, 0}; - int32_t fitCounters[4]{0, 0, 0, 0}; - int32_t backpropagatedCounters[4]{0, 0, 0, 0}; - int32_t refitCounters[4]{0, 0, 0, 0}; -#endif - for (int32_t iRoad = get_global_id(0); iRoad < Fitter.NumberOfRoads(); iRoad += get_global_size(0)) { - Road<5>& road = Fitter.roads()[iRoad]; - int32_t clusters[7] = {o2::its::constants::its::UnusedIndex, o2::its::constants::its::UnusedIndex, o2::its::constants::its::UnusedIndex, o2::its::constants::its::UnusedIndex, o2::its::constants::its::UnusedIndex, o2::its::constants::its::UnusedIndex, o2::its::constants::its::UnusedIndex}; - int32_t lastCellLevel = o2::its::constants::its::UnusedIndex; - CA_DEBUGGER(int32_t nClusters = 2); - - for (int32_t iCell{0}; iCell < Fitter.NumberOfLayers() - 2; ++iCell) { - const int32_t cellIndex = road[iCell]; - if (cellIndex == o2::its::constants::its::UnusedIndex) { - continue; - } else { - clusters[iCell] = Fitter.cells()[iCell][cellIndex].getFirstClusterIndex(); - clusters[iCell + 1] = Fitter.cells()[iCell][cellIndex].getSecondClusterIndex(); - clusters[iCell + 2] = Fitter.cells()[iCell][cellIndex].getThirdClusterIndex(); - lastCellLevel = iCell; - CA_DEBUGGER(nClusters++); - } - } - - CA_DEBUGGER(roadCounters[nClusters - 4]++); - - if (lastCellLevel == o2::its::constants::its::UnusedIndex) { - continue; - } - - /// From primary vertex context index to event index (== the one used as input of the tracking code) - for (int32_t iC{0}; iC < 7; iC++) { - if (clusters[iC] != o2::its::constants::its::UnusedIndex) { - clusters[iC] = Fitter.clusters()[iC][clusters[iC]].clusterId; - } - } - /// Track seed preparation. Clusters are numbered progressively from the outermost to the innermost. - const auto& cluster1 = Fitter.trackingFrame()[lastCellLevel + 2][clusters[lastCellLevel + 2]]; - const auto& cluster2 = Fitter.trackingFrame()[lastCellLevel + 1][clusters[lastCellLevel + 1]]; - const auto& cluster3 = Fitter.trackingFrame()[lastCellLevel][clusters[lastCellLevel]]; - - GPUITSTrack temporaryTrack; - { - const float ca = CAMath::Cos(cluster3.alphaTrackingFrame), sa = CAMath::Sin(cluster3.alphaTrackingFrame); - const float x1 = cluster1.xCoordinate * ca + cluster1.yCoordinate * sa; - const float y1 = -cluster1.xCoordinate * sa + cluster1.yCoordinate * ca; - const float z1 = cluster1.zCoordinate; - const float x2 = cluster2.xCoordinate * ca + cluster2.yCoordinate * sa; - const float y2 = -cluster2.xCoordinate * sa + cluster2.yCoordinate * ca; - const float z2 = cluster2.zCoordinate; - const float x3 = cluster3.xTrackingFrame; - const float y3 = cluster3.positionTrackingFrame[0]; - const float z3 = cluster3.positionTrackingFrame[1]; - - const float crv = o2::its::math_utils::computeCurvature(x1, y1, x2, y2, x3, y3); - const float x0 = o2::its::math_utils::computeCurvatureCentreX(x1, y1, x2, y2, x3, y3); - const float tgl12 = o2::its::math_utils::computeTanDipAngle(x1, y1, x2, y2, z1, z2); - const float tgl23 = o2::its::math_utils::computeTanDipAngle(x2, y2, x3, y3, z2, z3); - - const float r2 = CAMath::Sqrt(cluster2.xCoordinate * cluster2.xCoordinate + cluster2.yCoordinate * cluster2.yCoordinate); - const float r3 = CAMath::Sqrt(cluster3.xCoordinate * cluster3.xCoordinate + cluster3.yCoordinate * cluster3.yCoordinate); - const float fy = 1.f / (r2 - r3); - const float& tz = fy; - const float cy = (o2::its::math_utils::computeCurvature(x1, y1, x2, y2 + o2::its::constants::its::Resolution, x3, y3) - crv) / (o2::its::constants::its::Resolution * bz * constants::math::B2C) * 20.f; // FIXME: MS contribution to the cov[14] (*20 added) - constexpr float s2 = o2::its::constants::its::Resolution * o2::its::constants::its::Resolution; - - temporaryTrack.X() = cluster3.xTrackingFrame; - temporaryTrack.Y() = y3; - temporaryTrack.Z() = z3; - temporaryTrack.SinPhi() = crv * (x3 - x0); - temporaryTrack.DzDs() = 0.5f * (tgl12 + tgl23); - temporaryTrack.QPt() = CAMath::Abs(bz) < constants::math::Almost0 ? constants::math::Almost0 : crv / (bz * constants::math::B2C); - temporaryTrack.TZOffset() = 0; - temporaryTrack.Cov()[0] = s2; - temporaryTrack.Cov()[1] = 0.f; - temporaryTrack.Cov()[2] = s2; - temporaryTrack.Cov()[3] = s2 * fy; - temporaryTrack.Cov()[4] = 0.f; - temporaryTrack.Cov()[5] = s2 * fy * fy; - temporaryTrack.Cov()[6] = 0.f; - temporaryTrack.Cov()[7] = s2 * tz; - temporaryTrack.Cov()[8] = 0.f; - temporaryTrack.Cov()[9] = s2 * tz * tz; - temporaryTrack.Cov()[10] = s2 * cy; - temporaryTrack.Cov()[11] = 0.f; - temporaryTrack.Cov()[12] = s2 * fy * cy; - temporaryTrack.Cov()[13] = 0.f; - temporaryTrack.Cov()[14] = s2 * cy * cy; - temporaryTrack.SetChi2(0); - temporaryTrack.SetNDF(-5); - - prop.SetTrack(&temporaryTrack, cluster3.alphaTrackingFrame); - } - - for (size_t iC = 0; iC < 7; ++iC) { - temporaryTrack.mClusters[iC] = clusters[iC]; - } - bool fitSuccess = fitTrack(Fitter, prop, temporaryTrack, Fitter.NumberOfLayers() - 4, -1, -1); - if (!fitSuccess) { - continue; - } - CA_DEBUGGER(fitCounters[nClusters - 4]++); - temporaryTrack.ResetCovariance(); - fitSuccess = fitTrack(Fitter, prop, temporaryTrack, 0, Fitter.NumberOfLayers(), 1); - if (!fitSuccess) { - continue; - } - CA_DEBUGGER(backpropagatedCounters[nClusters - 4]++); - for (int32_t k = 0; k < 5; k++) { - temporaryTrack.mOuterParam.P[k] = temporaryTrack.Par()[k]; - } - for (int32_t k = 0; k < 15; k++) { - temporaryTrack.mOuterParam.C[k] = temporaryTrack.Cov()[k]; - } - temporaryTrack.mOuterParam.X = temporaryTrack.X(); - temporaryTrack.mOuterParam.alpha = prop.GetAlpha(); - temporaryTrack.ResetCovariance(); - fitSuccess = fitTrack(Fitter, prop, temporaryTrack, Fitter.NumberOfLayers() - 1, -1, -1); - if (!fitSuccess) { - continue; - } - CA_DEBUGGER(refitCounters[nClusters - 4]++); - int32_t trackId = CAMath::AtomicAdd(&Fitter.NumberOfTracks(), 1u); - Fitter.tracks()[trackId] = temporaryTrack; - } -#ifdef CA_DEBUG - GPUInfo("Roads: %i %i %i %i", roadCounters[0], roadCounters[1], roadCounters[2], roadCounters[3]); - GPUInfo("Fitted tracks: %i %i %i %i", fitCounters[0], fitCounters[1], fitCounters[2], fitCounters[3]); - GPUInfo("Backpropagated tracks: %i %i %i %i", backpropagatedCounters[0], backpropagatedCounters[1], backpropagatedCounters[2], backpropagatedCounters[3]); - GPUInfo("Refitted tracks: %i %i %i %i", refitCounters[0], refitCounters[1], refitCounters[2], refitCounters[3]); -#endif -} diff --git a/GPU/GPUTracking/ITS/GPUITSFitterKernels.h b/GPU/GPUTracking/ITS/GPUITSFitterKernels.h deleted file mode 100644 index 5a2a30de28823..0000000000000 --- a/GPU/GPUTracking/ITS/GPUITSFitterKernels.h +++ /dev/null @@ -1,42 +0,0 @@ -// Copyright 2019-2020 CERN and copyright holders of ALICE O2. -// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. -// All rights not expressly granted are reserved. -// -// This software is distributed under the terms of the GNU General Public -// License v3 (GPL Version 3), copied verbatim in the file "COPYING". -// -// In applying this license CERN does not waive the privileges and immunities -// granted to it by virtue of its status as an Intergovernmental Organization -// or submit itself to any jurisdiction. - -/// \file GPUITSFitterKernels.h -/// \author David Rohr, Maximiliano Puccio - -#ifndef GPUITSFITTERKERNELS_H -#define GPUITSFITTERKERNELS_H - -#include "GPUGeneralKernels.h" -namespace o2::its -{ -struct TrackingFrameInfo; -} // namespace o2::its - -namespace o2::gpu -{ -class GPUTPCGMPropagator; -class GPUITSFitter; -class GPUITSTrack; - -class GPUITSFitterKernels : public GPUKernelTemplate -{ - public: - GPUhdi() constexpr static GPUDataTypes::RecoStep GetRecoStep() { return GPUDataTypes::RecoStep::ITSTracking; } - template - GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors); - - protected: - GPUd() static bool fitTrack(GPUITSFitter& Fitter, GPUTPCGMPropagator& prop, GPUITSTrack& track, int32_t start, int32_t end, int32_t step); -}; -} // namespace o2::gpu - -#endif diff --git a/GPU/GPUTracking/ITS/GPUITSTrack.h b/GPU/GPUTracking/ITS/GPUITSTrack.h deleted file mode 100644 index 5063985692a43..0000000000000 --- a/GPU/GPUTracking/ITS/GPUITSTrack.h +++ /dev/null @@ -1,32 +0,0 @@ -// Copyright 2019-2020 CERN and copyright holders of ALICE O2. -// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders. -// All rights not expressly granted are reserved. -// -// This software is distributed under the terms of the GNU General Public -// License v3 (GPL Version 3), copied verbatim in the file "COPYING". -// -// In applying this license CERN does not waive the privileges and immunities -// granted to it by virtue of its status as an Intergovernmental Organization -// or submit itself to any jurisdiction. - -/// \file GPUITSTrack.h -/// \author David Rohr, Maximiliano Puccio - -#ifndef GPUITSTRACK_H -#define GPUITSTRACK_H - -#include "GPUTPCGMMergerTypes.h" -#include "GPUTPCGMTrackParam.h" - -namespace o2::gpu -{ -class GPUITSTrack : public GPUTPCGMTrackParam -{ - public: - gputpcgmmergertypes::GPUTPCOuterParam mOuterParam; - float mAlpha; - int32_t mClusters[7]; -}; -} // namespace o2::gpu - -#endif diff --git a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx index a624e1e55ed4b..0dbba0d8a2a99 100644 --- a/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx +++ b/GPU/GPUTracking/Standalone/Benchmark/standalone.cxx @@ -766,9 +766,9 @@ int32_t main(int argc, char** argv) chainTrackingPipeline->SetQAFromForeignChain(chainTracking); } if (!configStandalone.proc.doublePipeline) { - chainITS = rec->AddChain(0); + chainITS = rec->AddChain(); if (configStandalone.testSyncAsync) { - chainITSAsync = recAsync->AddChain(0); + chainITSAsync = recAsync->AddChain(); } } diff --git a/GPU/GPUTracking/kernels.cmake b/GPU/GPUTracking/kernels.cmake index 2a59f98a6d5b4..52cda3e8ff416 100644 --- a/GPU/GPUTracking/kernels.cmake +++ b/GPU/GPUTracking/kernels.cmake @@ -90,7 +90,6 @@ o2_gpu_add_kernel("GPUTPCGMO2Output, output" "= TPC o2_gpu_add_kernel("GPUTPCGMO2Output, mc" "= TPCMERGER") o2_gpu_add_kernel("GPUTRDTrackerKernels, gpuVersion" "= TRDTRACKER MATLUT TPCMERGER" LB GPUTRDTrackerGPU* externalInstance) o2_gpu_add_kernel("GPUTRDTrackerKernels, o2Version" "= TRDTRACKER MATLUT O2PROPAGATOR" LB GPUTRDTracker* externalInstance) -o2_gpu_add_kernel("GPUITSFitterKernels" "= TPCMERGER MATLUT" LB) o2_gpu_add_kernel("GPUTPCConvertKernel" "=" LB) o2_gpu_add_kernel("GPUTPCCompressionKernels, step0attached" "= TPCCOMPRESSION" LB) o2_gpu_add_kernel("GPUTPCCompressionKernels, step1unattached" "= ERRORS" LB)