diff --git a/Common/DCAFitter/include/DCAFitter/DCAFitterN.h b/Common/DCAFitter/include/DCAFitter/DCAFitterN.h index 226f008f5fce2..340a0be199c68 100644 --- a/Common/DCAFitter/include/DCAFitter/DCAFitterN.h +++ b/Common/DCAFitter/include/DCAFitter/DCAFitterN.h @@ -753,7 +753,7 @@ GPUd() void DCAFitterN::calcTrackResiduals() //___________________________________________________________________ template -GPUd() inline void DCAFitterN::calcTrackDerivatives() +GPUdi() void DCAFitterN::calcTrackDerivatives() { // calculate track derivatives over X param for (int i = N; i--;) { @@ -763,7 +763,7 @@ GPUd() inline void DCAFitterN::calcTrackDerivatives() //___________________________________________________________________ template -GPUd() inline double DCAFitterN::calcChi2() const +GPUdi() double DCAFitterN::calcChi2() const { // calculate current chi2 double chi2 = 0; @@ -777,7 +777,7 @@ GPUd() inline double DCAFitterN::calcChi2() const //___________________________________________________________________ template -GPUd() inline double DCAFitterN::calcChi2NoErr() const +GPUdi() double DCAFitterN::calcChi2NoErr() const { // calculate current chi2 of abs. distance minimization double chi2 = 0; @@ -842,7 +842,7 @@ GPUd() bool DCAFitterN::propagateTracksToVertex(int icand) //___________________________________________________________________ template -GPUd() inline o2::track::TrackPar DCAFitterN::getTrackParamAtPCA(int i, int icand) +GPUdi() o2::track::TrackPar DCAFitterN::getTrackParamAtPCA(int i, int icand) { // propagate tracks param only to current vertex (if not already done) int ord = mOrder[icand]; @@ -858,7 +858,7 @@ GPUd() inline o2::track::TrackPar DCAFitterN::getTrackParamAtPCA(int //___________________________________________________________________ template -GPUd() inline double DCAFitterN::getAbsMax(const VecND& v) +GPUdi() double DCAFitterN::getAbsMax(const VecND& v) { double mx = -1; for (int i = N; i--;) { @@ -1075,7 +1075,7 @@ GPUd() o2::track::TrackPar DCAFitterN::createParentTrackPar(int cand //___________________________________________________________________ template -GPUd() inline bool DCAFitterN::propagateParamToX(o2::track::TrackPar& t, float x) +GPUdi() bool DCAFitterN::propagateParamToX(o2::track::TrackPar& t, float x) { bool res = true; if (mUsePropagator || mMatCorr != o2::base::Propagator::MatCorrType::USEMatCorrNONE) { @@ -1091,7 +1091,7 @@ GPUd() inline bool DCAFitterN::propagateParamToX(o2::track::TrackPar //___________________________________________________________________ template -GPUd() inline bool DCAFitterN::propagateToX(o2::track::TrackParCov& t, float x) +GPUdi() bool DCAFitterN::propagateToX(o2::track::TrackParCov& t, float x) { bool res = true; if (mUsePropagator || mMatCorr != o2::base::Propagator::MatCorrType::USEMatCorrNONE) { diff --git a/DataFormats/Detectors/TPC/include/DataFormatsTPC/TrackTPC.h b/DataFormats/Detectors/TPC/include/DataFormatsTPC/TrackTPC.h index 42c098b7bdd54..0b40090de1c2e 100644 --- a/DataFormats/Detectors/TPC/include/DataFormatsTPC/TrackTPC.h +++ b/DataFormats/Detectors/TPC/include/DataFormatsTPC/TrackTPC.h @@ -79,8 +79,8 @@ class TrackTPC : public o2::track::TrackParCov GPUd() void setClusterRef(uint32_t entry, uint16_t ncl) { mClustersReference.set(entry, ncl); } template - GPUd() static inline void getClusterReference(T& clinfo, int nCluster, - uint8_t& sectorIndex, uint8_t& rowIndex, uint32_t& clusterIndex, const ClusRef& ref) + GPUdi() static void getClusterReference(T& clinfo, int nCluster, + uint8_t& sectorIndex, uint8_t& rowIndex, uint32_t& clusterIndex, const ClusRef& ref) { // data for given tracks starts at clinfo[ ref.getFirstEntry() ], // 1st ref.getEntries() cluster indices are stored as uint32_t @@ -95,15 +95,15 @@ class TrackTPC : public o2::track::TrackParCov } template - GPUd() inline void getClusterReference(T& clinfo, int nCluster, - uint8_t& sectorIndex, uint8_t& rowIndex, uint32_t& clusterIndex) const + GPUdi() void getClusterReference(T& clinfo, int nCluster, + uint8_t& sectorIndex, uint8_t& rowIndex, uint32_t& clusterIndex) const { getClusterReference(clinfo, nCluster, sectorIndex, rowIndex, clusterIndex, mClustersReference); } template - GPUd() static inline const o2::tpc::ClusterNative& getCluster(T& clinfo, int nCluster, - const o2::tpc::ClusterNativeAccess& clusters, uint8_t& sectorIndex, uint8_t& rowIndex, const ClusRef& ref) + GPUdi() static const o2::tpc::ClusterNative& getCluster(T& clinfo, int nCluster, + const o2::tpc::ClusterNativeAccess& clusters, uint8_t& sectorIndex, uint8_t& rowIndex, const ClusRef& ref) { uint32_t clusterIndex; getClusterReference(clinfo, nCluster, sectorIndex, rowIndex, clusterIndex, ref); @@ -111,15 +111,15 @@ class TrackTPC : public o2::track::TrackParCov } template - GPUd() inline const o2::tpc::ClusterNative& getCluster(T& clinfo, int nCluster, - const o2::tpc::ClusterNativeAccess& clusters, uint8_t& sectorIndex, uint8_t& rowIndex) const + GPUdi() const o2::tpc::ClusterNative& getCluster(T& clinfo, int nCluster, + const o2::tpc::ClusterNativeAccess& clusters, uint8_t& sectorIndex, uint8_t& rowIndex) const { return getCluster(clinfo, nCluster, clusters, sectorIndex, rowIndex, mClustersReference); } template - GPUd() inline const o2::tpc::ClusterNative& getCluster(T& clinfo, int nCluster, - const o2::tpc::ClusterNativeAccess& clusters) const + GPUdi() const o2::tpc::ClusterNative& getCluster(T& clinfo, int nCluster, + const o2::tpc::ClusterNativeAccess& clusters) const { uint8_t sectorIndex, rowIndex; return (getCluster(clinfo, nCluster, clusters, sectorIndex, rowIndex)); diff --git a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h b/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h index 2ef5577ec6bf2..a988c96168170 100644 --- a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h +++ b/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackParametrization.h @@ -199,9 +199,9 @@ class TrackParametrization GPUd() value_t getPt() const; GPUd() value_t getE2() const; GPUd() value_t getE() const; - GPUd() static inline value_t getdEdxBB(value_t betagamma) { return BetheBlochSolid(betagamma); } - GPUd() static inline value_t getdEdxBBOpt(value_t betagamma) { return BetheBlochSolidOpt(betagamma); } - GPUd() static inline value_t getBetheBlochSolidDerivativeApprox(value_T dedx, value_T bg) { return BetheBlochSolidDerivative(dedx, bg); } + GPUdi() static value_t getdEdxBB(value_t betagamma) { return BetheBlochSolid(betagamma); } + GPUdi() static value_t getdEdxBBOpt(value_t betagamma) { return BetheBlochSolidOpt(betagamma); } + GPUdi() static value_t getBetheBlochSolidDerivativeApprox(value_T dedx, value_T bg) { return BetheBlochSolidDerivative(dedx, bg); } GPUd() value_t getTheta() const; GPUd() value_t getEta() const; diff --git a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackUtils.h b/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackUtils.h index 96ab53fde5e5c..0ee0ca4461ab0 100644 --- a/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackUtils.h +++ b/DataFormats/Reconstruction/include/ReconstructionDataFormats/TrackUtils.h @@ -190,7 +190,7 @@ GPUd() value_T BetheBlochSolidOpt(value_T bg) //____________________________________________________ template -GPUd() value_T inline BetheBlochSolidDerivative(value_T dedx, value_T bg) +GPUdi() value_T BetheBlochSolidDerivative(value_T dedx, value_T bg) { // // This is approximate derivative of the BB over betagamm, NO check for the consistency of the provided dedx and bg is done diff --git a/DataFormats/simulation/include/SimulationDataFormat/MCTrack.h b/DataFormats/simulation/include/SimulationDataFormat/MCTrack.h index a5322b2f53dbf..b4cf26b11f82e 100644 --- a/DataFormats/simulation/include/SimulationDataFormat/MCTrack.h +++ b/DataFormats/simulation/include/SimulationDataFormat/MCTrack.h @@ -139,6 +139,12 @@ class MCTrackT } } + Double_t GetTgl() const + { + auto pT = GetPt(); + return pT > 1e-6 ? mStartVertexMomentumZ / pT : (GetStartVertexMomentumZ() > 0 ? 999. : -999.); + } + Double_t GetTheta() const { double mz(mStartVertexMomentumZ); diff --git a/Detectors/Base/include/DetectorsBase/MatLayerCyl.h b/Detectors/Base/include/DetectorsBase/MatLayerCyl.h index 229bc327d3039..869234e03f6c1 100644 --- a/Detectors/Base/include/DetectorsBase/MatLayerCyl.h +++ b/Detectors/Base/include/DetectorsBase/MatLayerCyl.h @@ -54,9 +54,11 @@ class MatLayerCyl : public o2::gpu::FlatObject Within = 0, Above = 1 }; +#ifndef GPUCA_GPUCODE MatLayerCyl(); MatLayerCyl(const MatLayerCyl& src) CON_DELETE; ~MatLayerCyl() CON_DEFAULT; +#endif #ifndef GPUCA_ALIGPUCODE // this part is unvisible on GPU version MatLayerCyl(float rMin, float rMax, float zHalfSpan, float dzMin, float drphiMin); diff --git a/Detectors/Base/include/DetectorsBase/MatLayerCylSet.h b/Detectors/Base/include/DetectorsBase/MatLayerCylSet.h index 3ffa7424ee61e..83fed8caf42eb 100644 --- a/Detectors/Base/include/DetectorsBase/MatLayerCylSet.h +++ b/Detectors/Base/include/DetectorsBase/MatLayerCylSet.h @@ -51,9 +51,11 @@ class MatLayerCylSet : public o2::gpu::FlatObject { public: +#ifndef GPUCA_GPUCODE MatLayerCylSet() CON_DEFAULT; ~MatLayerCylSet() CON_DEFAULT; MatLayerCylSet(const MatLayerCylSet& src) CON_DELETE; +#endif GPUd() const MatLayerCylSetLayout* get() const { return reinterpret_cast(mFlatBufferPtr); } GPUd() MatLayerCylSetLayout* get() { return reinterpret_cast(mFlatBufferPtr); } diff --git a/Detectors/Base/src/MatLayerCyl.cxx b/Detectors/Base/src/MatLayerCyl.cxx index 04f68fb81866f..2346946ea6a8a 100644 --- a/Detectors/Base/src/MatLayerCyl.cxx +++ b/Detectors/Base/src/MatLayerCyl.cxx @@ -23,10 +23,12 @@ using namespace o2::base; using flatObject = o2::gpu::FlatObject; +#ifndef GPUCA_GPUCODE //________________________________________________________________________________ MatLayerCyl::MatLayerCyl() : mNZBins(0), mNPhiBins(0), mNPhiSlices(0), mZHalf(0.f), mRMin2(0.f), mRMax2(0.f), mDZ(0.f), mDZInv(0.f), mDPhi(0.f), mDPhiInv(0.f), mPhiBin2Slice(nullptr), mSliceCos(nullptr), mSliceSin(nullptr), mCells(nullptr) { } +#endif #ifndef GPUCA_ALIGPUCODE // this part is unvisible on GPU version //________________________________________________________________________________ diff --git a/Detectors/Base/src/Propagator.cxx b/Detectors/Base/src/Propagator.cxx index b55ef22b857f0..c7c7b461034e5 100644 --- a/Detectors/Base/src/Propagator.cxx +++ b/Detectors/Base/src/Propagator.cxx @@ -788,16 +788,12 @@ namespace o2::base { #if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE) // FIXME: DR: WORKAROUND to avoid CUDA bug creating host symbols for device code. template class PropagatorImpl; +template bool GPUdni() PropagatorImpl::propagateToAlphaX::TrackPar_t>(PropagatorImpl::TrackPar_t&, float, float, bool, float, float, int, PropagatorImpl::MatCorrType matCorr, track::TrackLTIntegral*, int) const; +template bool GPUdni() PropagatorImpl::propagateToAlphaX::TrackParCov_t>(PropagatorImpl::TrackParCov_t&, float, float, bool, float, float, int, PropagatorImpl::MatCorrType matCorr, track::TrackLTIntegral*, int) const; #endif #ifndef GPUCA_GPUCODE template class PropagatorImpl; -#endif -#ifndef __HIPCC__ // TODO: Fixme: must prevent HIP from compiling this, should file bug report -template bool GPUd() PropagatorImpl::propagateToAlphaX::TrackPar_t>(PropagatorImpl::TrackPar_t&, float, float, bool, float, float, int, PropagatorImpl::MatCorrType matCorr, track::TrackLTIntegral*, int) const; -template bool GPUd() PropagatorImpl::propagateToAlphaX::TrackParCov_t>(PropagatorImpl::TrackParCov_t&, float, float, bool, float, float, int, PropagatorImpl::MatCorrType matCorr, track::TrackLTIntegral*, int) const; -#ifndef GPUCA_GPUCODE_DEVICE template bool PropagatorImpl::propagateToAlphaX::TrackPar_t>(PropagatorImpl::TrackPar_t&, double, double, bool, double, double, int, PropagatorImpl::MatCorrType matCorr, track::TrackLTIntegral*, int) const; template bool PropagatorImpl::propagateToAlphaX::TrackParCov_t>(PropagatorImpl::TrackParCov_t&, double, double, bool, double, double, int, PropagatorImpl::MatCorrType matCorr, track::TrackLTIntegral*, int) const; #endif -#endif } // namespace o2::base diff --git a/Detectors/FIT/macros/CMakeLists.txt b/Detectors/FIT/macros/CMakeLists.txt index 320cc19b00b8c..81f2cc05e0b25 100644 --- a/Detectors/FIT/macros/CMakeLists.txt +++ b/Detectors/FIT/macros/CMakeLists.txt @@ -40,5 +40,5 @@ o2_add_test_root_macro(readFITDCSdata.C O2::CCDB LABELS fit) -o2_data_file(COPY readFITDCSdata.C DESTINATION Detectors/FIT/macros/readFITDCSdata.C) -o2_data_file(COPY readFITDeadChannelMap.C DESTINATION Detectors/FIT/macros/readFITDeadChannelMap.C) \ No newline at end of file +o2_data_file(COPY readFITDCSdata.C DESTINATION Detectors/FIT/macros/) +o2_data_file(COPY readFITDeadChannelMap.C DESTINATION Detectors/FIT/macros/) \ No newline at end of file diff --git a/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwd.h b/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwd.h index 164bdb8d9f1f4..7aaadaad8ffff 100644 --- a/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwd.h +++ b/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwd.h @@ -333,11 +333,14 @@ class MatchGlobalFwd std::vector mMFTMatchPlaneParams; ///< MFT track parameters at matching plane std::vector mMCHMatchPlaneParams; ///< MCH track parameters at matching plane + std::map>> mCandidates; ///< map each MCH track id to vector of best match candidates + const o2::itsmft::TopologyDictionary* mMFTDict{nullptr}; // cluster patterns dictionary o2::itsmft::ChipMappingMFT mMFTMapping; bool mMCTruthON = false; ///< Flag availability of MC truth bool mUseMIDMCHMatch = false; ///< Flag for using MCHMID matches (TrackMCHMID) - int mSaveMode = 0; ///< Output mode [0 = SaveBestMatch; 1 = SaveAllMatches; 2 = SaveTrainingData] + int mSaveMode = 0; ///< Output mode [0 = SaveBestMatch; 1 = SaveAllMatches; 2 = SaveTrainingData; 3 = SaveNCandidates] + int mNCandidates = 5; ///< Numbers of matching candidates to save in savemode=3 MatchingType mMatchingType = MATCHINGUNDEFINED; TGeoManager* mGeoManager; }; diff --git a/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwdParam.h b/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwdParam.h index 9ba6563cec334..5633decb80985 100644 --- a/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwdParam.h +++ b/Detectors/GlobalTracking/include/GlobalTracking/MatchGlobalFwdParam.h @@ -26,7 +26,8 @@ namespace globaltracking enum SaveMode { kBestMatch = 0, kSaveAll, - kSaveTrainingData }; + kSaveTrainingData, + kSaveNCandidates }; struct GlobalFwdMatchingParam : public o2::conf::ConfigurableParamHelper { @@ -41,6 +42,7 @@ struct GlobalFwdMatchingParam : public o2::conf::ConfigurableParamHelper using namespace o2::globaltracking; @@ -66,6 +67,8 @@ void MatchGlobalFwd::init() mSaveMode = matchingParam.saveMode; LOG(info) << "Save mode MFTMCH candidates = " << mSaveMode; + + mNCandidates = matchingParam.nCandidates; } //_________________________________________________________ @@ -98,6 +101,9 @@ void MatchGlobalFwd::run(const o2::globaltracking::RecoContainer& inp) case kSaveTrainingData: doMatching(); break; + case kSaveNCandidates: + doMatching(); + break; default: LOG(fatal) << "Invalid MFTMCH save mode"; } @@ -132,6 +138,7 @@ void MatchGlobalFwd::clear() mMatchLabels.clear(); mMFTTrackROFContMapping.clear(); mMatchingInfo.clear(); + mCandidates.clear(); } //_________________________________________________________ @@ -400,6 +407,25 @@ void MatchGlobalFwd::doMatching() if (mMCTruthON) { LOG(info) << " MFT-MCH Matching: nFakes = " << nFakes << " nTrue = " << nTrue; } + } else if constexpr (saveAllMode == SaveMode::kSaveNCandidates) { + int nFakes = 0, nTrue = 0; + auto& matchAllChi2 = mMatchingFunctionMap["matchALL"]; + for (auto MCHId = 0; MCHId < mMCHWork.size(); MCHId++) { + auto& thisMCHTrack = mMCHWork[MCHId]; + for (auto& pairCandidate : mCandidates[MCHId]) { + thisMCHTrack.setMFTTrackID(pairCandidate.second); + auto& thisMFTTrack = mMFTWork[pairCandidate.second]; + auto chi2 = matchAllChi2(thisMCHTrack, thisMFTTrack); // Matching chi2 is stored independently + thisMCHTrack.setMFTMCHMatchingScore(pairCandidate.first); + thisMCHTrack.setMFTMCHMatchingChi2(chi2); + mMatchedTracks.emplace_back(thisMCHTrack); + mMatchingInfo.emplace_back(thisMCHTrack); + if (mMCTruthON) { + mMatchLabels.push_back(computeLabel(MCHId, pairCandidate.second)); + mMatchLabels.back().isFake() ? nFakes++ : nTrue++; + } + } + } } } @@ -413,6 +439,10 @@ void MatchGlobalFwd::ROFMatch(int MFTROFId, int firstMCHROFId, int lastMCHROFId) const auto& lastMCHROF = mMCHTrackROFRec[lastMCHROFId]; int nFakes = 0, nTrue = 0; + auto compare = [](const std::pair& a, const std::pair& b) { + return a.first < b.first; + }; + auto firstMFTTrackID = thisMFTROF.getFirstEntry(); auto lastMFTTrackID = firstMFTTrackID + thisMFTROF.getNEntries() - 1; @@ -464,12 +494,21 @@ void MatchGlobalFwd::ROFMatch(int MFTROFId, int firstMCHROFId, int lastMCHROFId) } } + if constexpr (saveAllMode == SaveMode::kSaveNCandidates) { // In saveAllmode save all pairs to output container + auto score = mMatchFunc(thisMCHTrack, thisMFTTrack); + std::pair scoreID = {score, MFTId}; + mCandidates[MCHId].push_back(scoreID); + std::sort(mCandidates[MCHId].begin(), mCandidates[MCHId].end(), compare); + if (mCandidates[MCHId].size() > mNCandidates) { + mCandidates[MCHId].pop_back(); + } + } + if constexpr (saveAllMode == SaveMode::kSaveTrainingData) { // In save training data mode store track parameters at matching plane thisMCHTrack.setMFTTrackID(MFTId); mMatchingInfo.emplace_back(thisMCHTrack); mMCHMatchPlaneParams.emplace_back(thisMCHTrack); mMFTMatchPlaneParams.emplace_back(static_cast(thisMFTTrack)); - if (mMCTruthON) { mMatchLabels.push_back(matchLabel); mMatchLabels.back().isFake() ? nFakes++ : nTrue++; diff --git a/Detectors/GlobalTrackingWorkflow/study/CMakeLists.txt b/Detectors/GlobalTrackingWorkflow/study/CMakeLists.txt index 17bad37b3c14e..adae8dcfb8b10 100644 --- a/Detectors/GlobalTrackingWorkflow/study/CMakeLists.txt +++ b/Detectors/GlobalTrackingWorkflow/study/CMakeLists.txt @@ -9,7 +9,7 @@ # granted to it by virtue of its status as an Intergovernmental Organization # or submit itself to any jurisdiction. -# add_compile_options(-O0 -g -fPIC) +add_compile_options(-O0 -g -fPIC) o2_add_library(GlobalTrackingStudy SOURCES src/TPCTrackStudy.cxx @@ -21,6 +21,8 @@ o2_add_library(GlobalTrackingStudy src/DumpTracks.cxx src/V0Ext.cxx src/TrackInfoExt.cxx + src/TrackMCStudyConfig.cxx + src/TrackMCStudyTypes.cxx PUBLIC_LINK_LIBRARIES O2::GlobalTracking O2::GlobalTrackingWorkflowReaders O2::GlobalTrackingWorkflowHelpers @@ -33,6 +35,8 @@ o2_target_root_dictionary( GlobalTrackingStudy HEADERS include/GlobalTrackingStudy/V0Ext.h include/GlobalTrackingStudy/TrackInfoExt.h + include/GlobalTrackingStudy/TrackMCStudyConfig.h + include/GlobalTrackingStudy/TrackMCStudyTypes.h ) o2_add_executable(study-workflow diff --git a/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudy.h b/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudy.h index 3dd315909ed1e..f7d402577a18c 100644 --- a/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudy.h +++ b/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudy.h @@ -12,17 +12,15 @@ #ifndef O2_TRACKING_STUDY_H #define O2_TRACKING_STUDY_H -#include "ReconstructionDataFormats/GlobalTrackID.h" -#include "Framework/Task.h" #include "Framework/DataProcessorSpec.h" -#include "ReconstructionDataFormats/Track.h" -#include "MathUtils/detail/Bracket.h" -#include "DataFormatsTPC/ClusterNative.h" +#include "Framework/Task.h" +#include "ReconstructionDataFormats/GlobalTrackID.h" namespace o2::trackstudy { + /// create a processor spec -o2::framework::DataProcessorSpec getTrackMCStudySpec(o2::dataformats::GlobalTrackID::mask_t srcTracks, o2::dataformats::GlobalTrackID::mask_t srcClus, bool checkMatching); +o2::framework::DataProcessorSpec getTrackMCStudySpec(o2::dataformats::GlobalTrackID::mask_t srcTracks, o2::dataformats::GlobalTrackID::mask_t srcClus); } // namespace o2::trackstudy diff --git a/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudyConfig.h b/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudyConfig.h new file mode 100644 index 0000000000000..8ba48ea2010a4 --- /dev/null +++ b/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudyConfig.h @@ -0,0 +1,34 @@ +// 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. + +#ifndef O2_TRACKING_STUDY_CONFIG_H +#define O2_TRACKING_STUDY_CONFIG_H +#include "CommonUtils/ConfigurableParam.h" +#include "CommonUtils/ConfigurableParamHelper.h" + +namespace o2::trackstudy +{ +struct TrackMCStudyConfig : o2::conf::ConfigurableParamHelper { + float minPt = 0.05; + float maxTgl = 1.2; + float minPtMC = 0.05; + float maxTglMC = 1.2; + float minRMC = 33.; + float maxPosTglMC = 2.; + float maxPVZOffset = 15.; + float decayMotherMaxT = 0.1; // max TOF in ns for mother particles to study + bool requireITSorTPCTrackRefs = true; + int decayPDG[5] = {310, 3122, -1, -1, -1}; // decays to study, must end by -1 + O2ParamDef(TrackMCStudyConfig, "trmcconf"); +}; +} // namespace o2::trackstudy + +#endif diff --git a/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudyTypes.h b/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudyTypes.h new file mode 100644 index 0000000000000..0334af3448fba --- /dev/null +++ b/Detectors/GlobalTrackingWorkflow/study/include/GlobalTrackingStudy/TrackMCStudyTypes.h @@ -0,0 +1,68 @@ +// 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. + +#ifndef O2_TRACKING_STUDY_TYPES_H +#define O2_TRACKING_STUDY_TYPES_H +#include "ReconstructionDataFormats/GlobalTrackID.h" +#include "ReconstructionDataFormats/VtxTrackIndex.h" +#include "ReconstructionDataFormats/Track.h" +#include "SimulationDataFormat/MCCompLabel.h" +#include "Framework/Task.h" + +namespace o2::trackstudy +{ +struct MCTrackInfo { + o2::track::TrackPar track{}; + o2::MCCompLabel label{}; + float occTPC = -1.f; + int bcInTF = -1; + int pdg = 0; + int pdgParent = 0; + int16_t nTPCCl = 0; + int16_t nTPCClShared = 0; + uint8_t minTPCRow = -1; + uint8_t maxTPCRow = 0; + int8_t nITSCl = 0; + int8_t pattITSCl = 0; + ClassDefNV(MCTrackInfo, 1); +}; + +struct RecTrack { + o2::track::TrackParCov track{}; + o2::dataformats::VtxTrackIndex gid{}; + uint8_t nClITS = 0; + uint8_t nClTPC = 0; + uint8_t pattITS = 0; + int8_t lowestPadRow = -1; + bool isFake = false; + ClassDefNV(RecTrack, 1); +}; + +struct TrackFamily { // set of tracks related to the same MC label + MCTrackInfo mcTrackInfo{}; + std::vector recTracks{}; + int8_t entITS = -1; + int8_t entTPC = -1; + int8_t entITSTPC = -1; + bool contains(const o2::dataformats::VtxTrackIndex& ref) const + { + for (const auto& tr : recTracks) { + if (ref == tr.gid) { + return true; + } + } + return false; + } + + ClassDefNV(TrackFamily, 1); +}; +} // namespace o2::trackstudy +#endif diff --git a/Detectors/GlobalTrackingWorkflow/study/src/GlobalTrackingStudyLinkDef.h b/Detectors/GlobalTrackingWorkflow/study/src/GlobalTrackingStudyLinkDef.h index 9b14e24d03cb4..a0cfc30726e9a 100644 --- a/Detectors/GlobalTrackingWorkflow/study/src/GlobalTrackingStudyLinkDef.h +++ b/Detectors/GlobalTrackingWorkflow/study/src/GlobalTrackingStudyLinkDef.h @@ -21,5 +21,12 @@ #pragma link C++ class std::vector < o2::dataformats::TrackInfoExt> + ; #pragma link C++ class std::vector < o2::dataformats::ProngInfoExt> + ; #pragma link C++ class std::vector < o2::dataformats::V0Ext> + ; - +#pragma link C++ class o2::trackstudy::TrackMCStudyConfig + ; +#pragma link C++ class o2::conf::ConfigurableParamHelper < o2::trackstudy::TrackMCStudyConfig> + ; +#pragma link C++ class o2::trackstudy::RecTrack + ; +#pragma link C++ class std::vector < o2::trackstudy::RecTrack> + ; +#pragma link C++ class o2::trackstudy::TrackFamily + ; +#pragma link C++ class std::vector < o2::trackstudy::TrackFamily> + ; +#pragma link C++ class o2::trackstudy::MCTrackInfo + ; +#pragma link C++ class std::vector < o2::trackstudy::MCTrackInfo> + ; #endif diff --git a/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudy.cxx b/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudy.cxx index 4a20e3d2e022d..bd259ff499ae2 100644 --- a/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudy.cxx +++ b/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudy.cxx @@ -16,11 +16,13 @@ #include "ReconstructionDataFormats/TrackTPCITS.h" #include "ReconstructionDataFormats/GlobalTrackID.h" #include "TPCCalibration/VDriftHelper.h" -#include "ITSMFTBase/DPLAlpideParam.h" +#include "TPCCalibration/CorrectionMapsLoader.h" +#include "ITSMFTReconstruction/ChipMappingITS.h" #include "DetectorsBase/Propagator.h" #include "DetectorsBase/GeometryManager.h" #include "SimulationDataFormat/MCEventLabel.h" #include "SimulationDataFormat/MCUtils.h" +#include "SimulationDataFormat/O2DatabasePDG.h" #include "CommonDataFormat/BunchFilling.h" #include "CommonUtils/NameConf.h" #include "DataFormatsFT0/RecPoints.h" @@ -32,6 +34,8 @@ #include "DetectorsCommonDataFormats/DetID.h" #include "DetectorsBase/GRPGeomHelper.h" #include "GlobalTrackingStudy/TrackMCStudy.h" +#include "GlobalTrackingStudy/TrackMCStudyConfig.h" +#include "GlobalTrackingStudy/TrackMCStudyTypes.h" #include "GlobalTracking/MatchTPCITSParams.h" #include "TPCBase/ParameterElectronics.h" #include "ReconstructionDataFormats/PrimaryVertex.h" @@ -41,10 +45,17 @@ #include "ReconstructionDataFormats/VtxTrackRef.h" #include "ReconstructionDataFormats/DCA.h" #include "Steer/MCKinematicsReader.h" +#include "CommonUtils/ConfigurableParam.h" +#include "CommonUtils/ConfigurableParamHelper.h" +#include "GPUO2InterfaceRefit.h" +#include "GPUParam.h" +#include "GPUParam.inc" #include "MathUtils/fit.h" #include +#include #include #include +#include // workflow to study relation of reco tracks to MCTruth // o2-trackmc-study-workflow --device-verbosity 3 -b --run @@ -68,27 +79,38 @@ using timeEst = o2::dataformats::TimeStampWithError; class TrackMCStudy : public Task { public: - TrackMCStudy(std::shared_ptr dr, std::shared_ptr gr, GTrackID::mask_t src, bool checkMatching) - : mDataRequest(dr), mGGCCDBRequest(gr), mTracksSrc(src), mCheckMatching(checkMatching) {} + TrackMCStudy(std::shared_ptr dr, std::shared_ptr gr, GTrackID::mask_t src) + : mDataRequest(dr), mGGCCDBRequest(gr), mTracksSrc(src) {} ~TrackMCStudy() final = default; void init(InitContext& ic) final; void run(ProcessingContext& pc) final; void endOfStream(EndOfStreamContext& ec) final; void finaliseCCDB(ConcreteDataMatcher& matcher, void* obj) final; - void process(o2::globaltracking::RecoContainer& recoData); + void process(const o2::globaltracking::RecoContainer& recoData); private: - void prepareITSData(o2::globaltracking::RecoContainer& recoData); + void loadTPCOccMap(const o2::globaltracking::RecoContainer& recoData); + void fillMCClusterInfo(const o2::globaltracking::RecoContainer& recoData); + void prepareITSData(const o2::globaltracking::RecoContainer& recoData); + bool processMCParticle(int src, int ev, int trid); + bool addMCParticle(const MCTrack& mctr, const o2::MCCompLabel& lb, TParticlePDG* pPDG = nullptr); + bool acceptMCCharged(const MCTrack& tr, const o2::MCCompLabel& lb, int followDec = -1); bool propagateToRefX(o2::track::TrackParCov& trcTPC, o2::track::TrackParCov& trcITS); void updateTimeDependentParams(ProcessingContext& pc); float getDCAYCut(float pt) const; + gsl::span mCurrMCTracks; + TVector3 mCurrMCVertex; o2::tpc::VDriftHelper mTPCVDriftHelper{}; std::shared_ptr mDataRequest; std::shared_ptr mGGCCDBRequest; std::unique_ptr mDBGOut; + std::vector mTBinClOcc; ///< TPC occupancy histo: i-th entry is the integrated occupancy for ~1 orbit starting from the TB = i*mNTPCOccBinLength + std::vector mIntBC; ///< interaction global BC wrt TF start + std::vector mTPCOcc; ///< TPC occupancy for this interaction time + int mNTPCOccBinLength = 0; ///< TPC occ. histo bin length in TBs + float mNTPCOccBinLengthInv; int mVerbose = 0; - bool mCheckMatching = false; float mITSTimeBiasMUS = 0.f; float mITSROFrameLengthMUS = 0.f; ///< ITS RO frame in mus float mTPCTBinMUS = 0.; ///< TPC time bin duration in microseconds @@ -96,15 +118,24 @@ class TrackMCStudy : public Task float mTPCDCAYCut = 2.; float mTPCDCAZCut = 2.; float mMinX = 6.; - float mMaxEta = 0.8; - float mMinPt = 0.03; int mMinTPCClusters = 10; + int mNCheckDecays = 0; std::string mDCAYFormula = "0.0105 + 0.0350 / pow(x, 1.1)"; GTrackID::mask_t mTracksSrc{}; o2::steer::MCKinematicsReader mcReader; // reader of MC information std::vector mITSROF; std::vector mITSROFBracket; + std::vector mDecProdLblPool; // labels of decay products to watch, added to MC map + struct DecayRef { + o2::MCCompLabel mother{}; + o2::track::TrackPar parent{}; + int pdg = 0; + int daughterFirst = -1; + int daughterLast = -1; + }; + std::vector> mDecaysMaps; // for every parent particle to watch store its label and entries of 1st/last decay product labels in mDecProdLblPool + std::unordered_map mSelMCTracks; static constexpr float MaxSnp = 0.9; // max snp of ITS or TPC track at xRef to be matched }; @@ -119,15 +150,28 @@ void TrackMCStudy::init(InitContext& ic) mTPCDCAYCut = ic.options().get("max-tpc-dcay"); mTPCDCAZCut = ic.options().get("max-tpc-dcaz"); mMinX = ic.options().get("min-x-prop"); - mMaxEta = ic.options().get("max-eta"); - mMinPt = ic.options().get("min-pt"); mMinTPCClusters = ic.options().get("min-tpc-clusters"); mDCAYFormula = ic.options().get("dcay-vs-pt"); + + const auto& params = o2::trackstudy::TrackMCStudyConfig::Instance(); + for (int id = 0; id < sizeof(params.decayPDG) / sizeof(int); id++) { + if (params.decayPDG[id] < 0) { + break; + } + mNCheckDecays++; + } + mDecaysMaps.resize(mNCheckDecays); } void TrackMCStudy::run(ProcessingContext& pc) { o2::globaltracking::RecoContainer recoData; + for (int i = 0; i < mNCheckDecays; i++) { + mDecaysMaps[i].clear(); + } + mDecProdLblPool.clear(); + mCurrMCTracks = {}; + recoData.collectData(pc, *mDataRequest.get()); // select tracks of needed type, with minimal cuts, the real selected will be done in the vertexer updateTimeDependentParams(pc); // Make sure this is called after recoData.collectData, which may load some conditions process(recoData); @@ -156,8 +200,9 @@ void TrackMCStudy::updateTimeDependentParams(ProcessingContext& pc) } } -void TrackMCStudy::process(o2::globaltracking::RecoContainer& recoData) +void TrackMCStudy::process(const o2::globaltracking::RecoContainer& recoData) { + const auto& params = o2::trackstudy::TrackMCStudyConfig::Instance(); auto pvvec = recoData.getPrimaryVertices(); auto trackIndex = recoData.getPrimaryVertexMatchedTracks(); // Global ID's for associated tracks auto vtxRefs = recoData.getPrimaryVertexMatchedTrackRefs(); // references from vertex to these track IDs @@ -166,10 +211,8 @@ void TrackMCStudy::process(o2::globaltracking::RecoContainer& recoData) float vdriftTB = mTPCVDriftHelper.getVDriftObject().getVDrift() * o2::tpc::ParameterElectronics::Instance().ZbinWidth; // VDrift expressed in cm/TimeBin float itsBias = 0.5 * mITSROFrameLengthMUS + o2::itsmft::DPLAlpideParam::Instance().roFrameBiasInBC * o2::constants::lhc::LHCBunchSpacingNS * 1e-3; // ITS time is supplied in \mus as beginning of ROF - if (mCheckMatching) { - prepareITSData(recoData); - } - + prepareITSData(recoData); + loadTPCOccMap(recoData); auto getITSPatt = [&](GTrackID gid, uint8_t& ncl) { int16_t patt = 0; if (gid.getSource() == VTIndex::ITSAB) { @@ -204,10 +247,42 @@ void TrackMCStudy::process(o2::globaltracking::RecoContainer& recoData) return -1; }; - std::map> MCTRMap; + const auto* digconst = mcReader.getDigitizationContext(); + const auto& mcEvRecords = digconst->getEventRecords(false); + for (const auto& mcIR : mcEvRecords) { + long tbc = mcIR.differenceInBC(recoData.startIR); + mIntBC.push_back(tbc); + int occBin = tbc / 8 * mNTPCOccBinLengthInv; + mTPCOcc.push_back(occBin < 0 ? mTBinClOcc[0] : (occBin >= mTBinClOcc.size() ? mTBinClOcc.back() : mTBinClOcc[occBin])); + } + // collect interesting MC particle (tracks and parents) + int curSrcMC = 0, curEvMC = 0; + for (curSrcMC = 0; curSrcMC < (int)mcReader.getNSources(); curSrcMC++) { + if (mVerbose > 1) { + LOGP(info, "Source {}", curSrcMC); + } + for (curEvMC = 0; curEvMC < (int)mcReader.getNEvents(curSrcMC); curEvMC++) { + if (mVerbose > 1) { + LOGP(info, "Event {}", curEvMC); + } + const auto& mt = mcReader.getTracks(curSrcMC, curEvMC); + mCurrMCTracks = gsl::span(mt.data(), mt.size()); + const_cast(mcReader.getMCEventHeader(curSrcMC, curEvMC)).GetVertex(mCurrMCVertex); + for (int itr = 0; itr < mCurrMCTracks.size(); itr++) { + processMCParticle(curSrcMC, curEvMC, itr); + } + } + } + if (mVerbose > 0) { + for (int id = 0; id < mNCheckDecays; id++) { + LOGP(info, "Decay PDG={} : {} entries", params.decayPDG[id], mDecaysMaps[id].size()); + } + } + + // add reconstruction info to MC particles. If MC particle was not selected before but was reconstrected, account MC info for (int iv = 0; iv < nv; iv++) { - if (mVerbose > 0) { + if (mVerbose > 1) { LOGP(info, "processing PV {} of {}", iv, nv); } const auto& vtref = vtxRefs[iv]; @@ -220,222 +295,171 @@ void TrackMCStudy::process(o2::globaltracking::RecoContainer& recoData) for (int i = idMin; i < idMax; i++) { auto vid = trackIndex[i]; const auto& trc = recoData.getTrackParam(vid); - if (trc.getPt() < mMinPt) { + if (trc.getPt() < params.minPt || std::abs(trc.getTgl()) > params.maxTgl) { continue; } auto lbl = recoData.getTrackMCLabel(vid); if (lbl.isValid()) { lbl.setFakeFlag(false); - auto& vvids = MCTRMap[lbl]; - if (vid.isAmbiguous() || vvids.empty()) { // do not repeat ambiguous tracks - bool skip = false; - for (const auto& va : vvids) { - if (va.second == vid) { - skip = true; - break; - } + auto entry = mSelMCTracks.find(lbl); + if (entry == mSelMCTracks.end()) { // add the track which was not added during MC scan + if (lbl.getSourceID() != curSrcMC || lbl.getEventID() != curEvMC) { + curSrcMC = lbl.getSourceID(); + curEvMC = lbl.getEventID(); + const auto& mt = mcReader.getTracks(curSrcMC, curEvMC); + mCurrMCTracks = gsl::span(mt.data(), mt.size()); + const_cast(mcReader.getMCEventHeader(curSrcMC, curEvMC)).GetVertex(mCurrMCVertex); } - if (skip) { + if (!acceptMCCharged(mCurrMCTracks[lbl.getTrackID()], lbl)) { continue; } + entry = mSelMCTracks.find(lbl); } - vvids.emplace_back(iv, vid); + auto& trackFamily = entry->second; + if (vid.isAmbiguous()) { // do not repeat ambiguous tracks + if (trackFamily.contains(vid)) { + continue; + } + } + auto& trf = trackFamily.recTracks.emplace_back(); + trf.gid = vid; // account(iv, vid); + if (mVerbose > 1) { + LOGP(info, "Matched rec track {} to MC track {}", vid.asString(), entry->first.asString()); + } + } else { + continue; } } } } - o2::track::TrackParCov dummyTrackParCov; - o2::track::TrackPar dummyTrackPar; - dummyTrackParCov.invalidate(); - dummyTrackPar.invalidate(); - - const std::vector* mcTracks = nullptr; - o2::MCCompLabel prevLbl; - std::vector recTracks; - std::vector recGIDs; - std::vector recFakes; - std::vector lowestPadrows; - std::vector itsPatterns; - std::vector tpcNcls; - std::vector itsNcls; - LOGP(info, "Recorded {} reconstructed tracks", MCTRMap.size()); - size_t count = 0; - for (auto ent : MCTRMap) { - count++; - auto lbl = ent.first; - if (lbl.getEventID() != prevLbl.getEventID() || lbl.getSourceID() != prevLbl.getSourceID()) { - if (mVerbose > 0) { - LOGP(info, "Loading MC Event={} / Src={}", lbl.getEventID(), lbl.getSourceID()); - } - mcTracks = &mcReader.getTracks(lbl.getSourceID(), lbl.getEventID()); - prevLbl = lbl; - } - const auto& mcPart = (*mcTracks)[lbl.getTrackID()]; - int pdg = mcPart.GetPdgCode(), pdgParent = 0; - std::array xyz{(float)mcPart.GetStartVertexCoordinatesX(), (float)mcPart.GetStartVertexCoordinatesY(), (float)mcPart.GetStartVertexCoordinatesZ()}; - std::array pxyz{(float)mcPart.GetStartVertexMomentumX(), (float)mcPart.GetStartVertexMomentumY(), (float)mcPart.GetStartVertexMomentumZ()}; - TParticlePDG* pPDG = TDatabasePDG::Instance()->GetParticle(pdg); - if (!pPDG) { - LOGP(error, "Unknown particle {}, skip. Was at {} of {}", pdg, count, MCTRMap.size()); - continue; - } - o2::track::TrackPar mctrO2(xyz, pxyz, TMath::Nint(pPDG->Charge() / 3), false); - bool primary = mcPart.isPrimary(); - auto parID = primary ? -1 : mcPart.getMotherTrackId(); - if (parID >= 0) { - const auto& mcPartPar = (*mcTracks)[parID]; - pdgParent = mcPartPar.GetPdgCode(); - } - auto& vgids = ent.second; - // make sure the more global tracks come 1st - if (vgids.size() > 1) { - std::sort(vgids.begin(), vgids.end(), [](VTIndexV& lhs, VTIndexV& rhs) { return lhs.second.getSource() > rhs.second.getSource(); }); + // collect ITS/TPC cluster info for selected MC particles + fillMCClusterInfo(recoData); + + LOGP(info, "collected {} MC tracks", mSelMCTracks.size()); + int mcnt = 0; + for (auto& entry : mSelMCTracks) { + auto& trackFam = entry.second; + auto& tracks = trackFam.recTracks; + mcnt++; + if (tracks.empty()) { + continue; } - recTracks.clear(); - recGIDs.clear(); - recFakes.clear(); - lowestPadrows.clear(); - itsPatterns.clear(); - itsNcls.clear(); - tpcNcls.clear(); if (mVerbose > 1) { - LOGP(info, "[{}] Lbl:{} PDG:{:+5d} (par: {:+5d}) | MC: {}", vgids.size(), lbl.asString(), pdg, pdgParent, mctrO2.asString()); + LOGP(info, "Processing MC track#{} {} -> {} reconstructed tracks", mcnt - 1, entry.first.asString(), tracks.size()); } - int entITS = -1, entTPC = -1, entITSTPC = -1; - for (size_t i = 0; i < vgids.size(); i++) { - auto vid = vgids[i].second; - auto lbl = recoData.getTrackMCLabel(vid); - const auto& trc = recoData.getTrackParam(vid); - int16_t itsPatt = 0; - uint8_t nclITS = 0; - uint8_t nclTPC = 0; - if (mVerbose > 1) { - LOGP(info, " :{} {:22} | [{}] {}", lbl.asString(), vid.asString(), i, ((const o2::track::TrackPar&)trc).asString()); - } - recTracks.push_back(trc); - recGIDs.push_back(vid); - recFakes.push_back(recoData.getTrackMCLabel(vid).isFake()); - auto msk = vid.getSourceDetectorsMask(); - if (mCheckMatching) { - lowestPadrows.push_back(-1); - } - if (msk[DetID::ITS]) { - auto gidITS = recoData.getITSContributorGID(vid); - itsPatt = getITSPatt(gidITS, nclITS); - if (vid.getSource() == VTIndex::ITS) { - entITS = i; - } - if (msk[DetID::TPC]) { - entITSTPC = i; - } - } - if (msk[DetID::TPC]) { - if (vid.getSource() == VTIndex::TPC) { - entTPC = i; + // sort according to the gid complexity (in principle, should be already sorted due to the backwards loop over NSources above + std::sort(tracks.begin(), tracks.end(), [](RecTrack& lhs, RecTrack& rhs) { return lhs.gid.getSource() > rhs.gid.getSource(); }); + // fill track params + int tcnt = 0; + for (auto& tref : tracks) { + if (tref.gid.isSourceSet()) { + tref.track = recoData.getTrackParam(tref.gid); + tref.isFake = recoData.getTrackMCLabel(tref.gid).isFake(); + auto msk = tref.gid.getSourceDetectorsMask(); + if (msk[DetID::ITS]) { + auto gidITS = recoData.getITSContributorGID(tref.gid); + tref.pattITS = getITSPatt(gidITS, tref.nClITS); + if (tref.gid.getSource() == VTIndex::ITS && trackFam.entITS < 0) { + trackFam.entITS = tcnt; + } + if (msk[DetID::TPC] && trackFam.entITSTPC < 0) { + trackFam.entITSTPC = tcnt; + } } - auto gidTPC = recoData.getTPCContributorGID(vid); - const auto& trtpc = recoData.getTPCTrack(gidTPC); - nclTPC = trtpc.getNClusters(); - if (mCheckMatching) { - auto& lr = lowestPadrows.back(); - lr = getLowestPadrow(recoData.getTPCTrack(recoData.getTPCContributorGID(vid))); + if (msk[DetID::TPC] && trackFam.entTPC < 0) { + if (tref.gid.getSource() == VTIndex::TPC) { + trackFam.entTPC = tcnt; + } + auto gidTPC = recoData.getTPCContributorGID(tref.gid); + const auto& trtpc = recoData.getTPCTrack(gidTPC); + tref.nClTPC = trtpc.getNClusters(); + tref.lowestPadRow = getLowestPadrow(recoData.getTPCTrack(gidTPC)); } + } else { + LOGP(info, "Invalid entry {} of {} getTrackMCLabel {}", tcnt, tracks.size(), tref.gid.asString()); } - tpcNcls.push_back(nclTPC); - itsNcls.push_back(nclITS); - itsPatterns.push_back(itsPatt); + tcnt++; } - (*mDBGOut) << "tracks" - << "lbl=" << lbl - << "mcTr=" << mctrO2 - << "pdg=" << pdg - << "pdgPar=" << pdgParent - << "recTr=" << recTracks - << "recGID=" << recGIDs - << "recFake=" << recFakes - << "itsPatt=" << itsPatterns - << "nClITS=" << itsNcls - << "nClTPC=" << tpcNcls; - if (mCheckMatching) { - (*mDBGOut) << "tracks" - << "lowestPadRow=" << lowestPadrows; + if (trackFam.entITSTPC < 0 && trackFam.entITS > -1 && trackFam.entTPC > -1) { // ITS and TPC were found but matching failed + auto vidITS = tracks[trackFam.entITS].gid; + auto vidTPC = tracks[trackFam.entTPC].gid; + auto trcTPC = recoData.getTrackParam(vidTPC); + auto trcITS = recoData.getTrackParamOut(vidITS); + if (!propagateToRefX(trcTPC, trcITS)) { + // break; + } } - (*mDBGOut) << "tracks" - << "\n"; - - // special ITS-TPC matching failure output - while (mCheckMatching) { - if (entITSTPC < 0 && entITS > -1 && entTPC > -1) { // ITS and TPC were found but matching failed - auto vidITS = vgids[entITS].second; - auto vidTPC = recoData.getTPCContributorGID(vgids[entTPC].second); // might be TPC match to outer detector, extract TPC - auto trcTPC = recoData.getTrackParam(vidTPC); - auto trcITS = recoData.getTrackParam(vidITS); - if (!propagateToRefX(trcTPC, trcITS)) { - break; + } + + // single tracks + for (auto& entry : mSelMCTracks) { + auto& trackFam = entry.second; + (*mDBGOut) << "tracks" << "tr=" << trackFam << "\n"; + } + // decays + std::vector decFam; + for (int id = 0; id < mNCheckDecays; id++) { + std::string decTreeName = fmt::format("dec{}", params.decayPDG[id]); + for (const auto& dec : mDecaysMaps[id]) { + decFam.clear(); + for (int idd = dec.daughterFirst; idd <= dec.daughterLast; idd++) { + auto dtLbl = mDecProdLblPool[idd]; // daughter MC label + const auto& dtFamily = mSelMCTracks[dtLbl]; + if (dtFamily.mcTrackInfo.pdgParent != dec.pdg) { + LOGP(error, "{}-th decay (pdg={}): {} in {}:{} range refers to MC track with pdgParent = {}", id, params.decayPDG[id], idd, dec.daughterFirst, dec.daughterLast, dtFamily.mcTrackInfo.pdgParent); + continue; } - const auto& trcTPCOrig = recoData.getTPCTrack(vidTPC); - const auto& trcITSOrig = recoData.getITSTrack(vidITS); - int lowestTPCRow = lowestPadrows[entTPC]; - float tpcT0 = trcTPCOrig.getTime0(), tF = trcTPCOrig.getDeltaTFwd(), tB = trcTPCOrig.getDeltaTBwd(); - TBracket tpcBr((tpcT0 - tB) * mTPCTBinMUS, (tpcT0 + tF) * mTPCTBinMUS); - - (*mDBGOut) << "failMatch" - << "mcTr=" << mctrO2 - << "pdg=" << pdg - << "pdgPar=" << pdgParent - << "labelITS=" << recoData.getTrackMCLabel(vidITS) - << "labelTPC=" << recoData.getTrackMCLabel(vidTPC) - << "gidITS=" << vidITS - << "gidTPC=" << vidTPC - << "itsBracket=" << mITSROFBracket[mITSROF[vidITS.getIndex()]] - << "tpcBracket=" << tpcBr - << "itsRef=" << trcITS - << "tpcRef=" << trcTPC - << "itsOrig=" << trcITSOrig - << "tpcOrig=" << trcTPCOrig - << "itsPatt=" << itsPatterns[entITS] - << "tpcLowestRow=" << lowestTPCRow - << "\n"; - } else if (entITSTPC > -1) { // match was found - auto contribIDs = recoData.getSingleDetectorRefs(vgids[entITSTPC].second); - auto vidMatch = contribIDs[VTIndex::ITSTPC]; - auto vidTPC = contribIDs[VTIndex::TPC]; - auto vidITS = contribIDs[VTIndex::ITSAB].isSourceSet() ? contribIDs[VTIndex::ITSAB] : contribIDs[VTIndex::ITS]; - const auto& trcTPCOrig = recoData.getTPCTrack(vidTPC); - o2::MCCompLabel itsLb; - int nITScl = 0; - if (vidITS.getSource() == VTIndex::ITS) { - itsLb = recoData.getTrackMCLabel(vidITS); - nITScl = recoData.getITSTrack(vidITS).getNClusters(); - } else { - itsLb = recoData.getITSABMCLabels()[vidITS]; - nITScl = recoData.getITSABRefs()[vidITS].getNClusters(); + decFam.push_back(dtFamily); + } + (*mDBGOut) << decTreeName.c_str() << "pdgPar=" << dec.pdg << "trPar=" << dec.parent << "prod=" << decFam << "\n"; + } + } +} + +void TrackMCStudy::fillMCClusterInfo(const o2::globaltracking::RecoContainer& recoData) +{ + // TPC clusters info + const auto& TPCClusterIdxStruct = recoData.inputsTPCclusters->clusterIndex; + const auto* TPCClMClab = recoData.inputsTPCclusters->clusterIndex.clustersMCTruth; + for (int sector = 0; sector < 36; sector++) { + for (int row = 0; row < 152; row++) { + unsigned int offs = TPCClusterIdxStruct.clusterOffset[sector][row]; + for (unsigned int icl0 = 0; icl0 < TPCClusterIdxStruct.nClusters[sector][row]; icl0++) { + const auto labels = TPCClMClab->getLabels(icl0 + offs); + for (const auto& lbl : labels) { + auto entry = mSelMCTracks.find(lbl); + if (entry == mSelMCTracks.end()) { // not selected + continue; + } + auto& mctr = entry->second.mcTrackInfo; + mctr.nTPCCl++; + if (mctr.maxTPCRow < row) { + mctr.maxTPCRow = row; + } + if (mctr.minTPCRow > row) { + mctr.minTPCRow = row; + } + if (labels.size() > 1) { + mctr.nTPCClShared++; + } } - int lowestTPCRow = lowestPadrows[entITSTPC]; - const auto& trackITSTPC = recoData.getTPCITSTrack(vidMatch); - float timeTB = trackITSTPC.getTimeMUS().getTimeStamp() / o2::constants::lhc::LHCBunchSpacingMUS / 8; // ITS-TPC time in TPC timebins - - (*mDBGOut) << "match" - << "mcTr=" << mctrO2 - << "pdg=" << pdg - << "pdgPar=" << pdgParent - << "labelMatch=" << recoData.getTrackMCLabel(vidMatch) - << "labelTPC=" << recoData.getTrackMCLabel(vidTPC) - << "labelITS=" << itsLb - << "gidTPC=" << vidTPC - << "gidITS=" << vidITS - << "tpcOrig=" << trcTPCOrig - << "nClITS=" << itsNcls[entITSTPC] - << "itsPatt=" << itsPatterns[entITSTPC] - << "itstpc=" << ((o2::track::TrackParCov&)trackITSTPC) - << "matchChi2=" << trackITSTPC.getChi2Match() - << "refitChi2=" << trackITSTPC.getChi2Refit() - << "timeTB=" << timeTB - << "tpcLowestRow=" << lowestTPCRow - << "\n"; } - break; + } + } + // fill ITS cluster info + const auto* mcITSClusters = recoData.getITSClustersMCLabels(); + const auto& ITSClusters = recoData.getITSClusters(); + for (unsigned int icl = 0; icl < ITSClusters.size(); icl++) { + const auto labels = mcITSClusters->getLabels(icl); + for (const auto& lbl : labels) { + auto entry = mSelMCTracks.find(lbl); + if (entry == mSelMCTracks.end()) { // not selected + continue; + } + auto& mctr = entry->second.mcTrackInfo; + mctr.nITSCl++; + mctr.pattITSCl |= 0x1 << o2::itsmft::ChipMappingITS::getLayer(ITSClusters[icl].getChipID()); } } } @@ -498,7 +522,7 @@ void TrackMCStudy::finaliseCCDB(ConcreteDataMatcher& matcher, void* obj) } //_____________________________________________________ -void TrackMCStudy::prepareITSData(o2::globaltracking::RecoContainer& recoData) +void TrackMCStudy::prepareITSData(const o2::globaltracking::RecoContainer& recoData) { auto ITSTracksArray = recoData.getITSTracks(); auto ITSTrackROFRec = recoData.getITSTracksROFRecords(); @@ -525,7 +549,167 @@ float TrackMCStudy::getDCAYCut(float pt) const return fun.Eval(pt); } -DataProcessorSpec getTrackMCStudySpec(GTrackID::mask_t srcTracks, GTrackID::mask_t srcClusters, bool checkMatching) +bool TrackMCStudy::processMCParticle(int src, int ev, int trid) +{ + const auto& mcPart = mCurrMCTracks[trid]; + int pdg = mcPart.GetPdgCode(); + bool res = false; + while (true) { + auto lbl = o2::MCCompLabel(trid, ev, src); + int decay = -1; // is this decay to watch? + const auto& params = o2::trackstudy::TrackMCStudyConfig::Instance(); + if (mcPart.T() < params.decayMotherMaxT) { + for (int id = 0; id < mNCheckDecays; id++) { + if (params.decayPDG[id] == std::abs(pdg)) { + decay = id; + break; + } + } + if (decay >= 0) { // check if decay and kinematics is acceptable + int idd0 = mcPart.getFirstDaughterTrackId(), idd1 = mcPart.getLastDaughterTrackId(); // we want only charged and trackable daughters + int dtStart = mDecProdLblPool.size(), dtEnd = -1; + if (idd0 < 0) { + break; + } + for (int idd = idd0; idd <= idd1; idd++) { + const auto& product = mCurrMCTracks[idd]; + auto lbld = o2::MCCompLabel(idd, ev, src); + if (!acceptMCCharged(product, lbld, decay)) { + decay = -1; // discard decay + mDecProdLblPool.resize(dtStart); + break; + } + mDecProdLblPool.push_back(lbld); // register prong entry and label + } + if (decay >= 0) { + // account decay + dtEnd = mDecProdLblPool.size() - 1; + std::array xyz{(float)mcPart.GetStartVertexCoordinatesX(), (float)mcPart.GetStartVertexCoordinatesY(), (float)mcPart.GetStartVertexCoordinatesZ()}; + std::array pxyz{(float)mcPart.GetStartVertexMomentumX(), (float)mcPart.GetStartVertexMomentumY(), (float)mcPart.GetStartVertexMomentumZ()}; + mDecaysMaps[decay].emplace_back(DecayRef{lbl, + o2::track::TrackPar(xyz, pxyz, TMath::Nint(O2DatabasePDG::Instance()->GetParticle(mcPart.GetPdgCode())->Charge() / 3), false), + mcPart.GetPdgCode(), dtStart, dtEnd}); + if (mVerbose > 1) { + LOGP(info, "Adding MC parent pdg={} {}, with prongs in {}:{} range", pdg, lbl.asString(), dtStart, dtEnd); + } + res = true; // Accept! + } + break; + } + } + // check if this is a charged which should be processed but was not accounted as a decay product + if (mSelMCTracks.find(lbl) == mSelMCTracks.end()) { + res = acceptMCCharged(mcPart, lbl); + } + break; + } + return res; +} + +bool TrackMCStudy::acceptMCCharged(const MCTrack& tr, const o2::MCCompLabel& lb, int followDecay) +{ + const auto& params = o2::trackstudy::TrackMCStudyConfig::Instance(); + if (tr.GetPt() < params.minPtMC || + std::abs(tr.GetTgl()) > params.maxTglMC || + tr.R2() > params.minRMC * params.minRMC) { + if (mVerbose > 1 && followDecay > -1) { + LOGP(info, "rejecting decay {} prong : pdg={}, pT={}, tgL={}, r={}", followDecay, tr.GetPdgCode(), tr.GetPt(), tr.GetTgl(), std::sqrt(tr.R2())); + } + return false; + } + float dx = tr.GetStartVertexCoordinatesX() - mCurrMCVertex.X(), dy = tr.GetStartVertexCoordinatesY() - mCurrMCVertex.Y(), dz = tr.GetStartVertexCoordinatesZ() - mCurrMCVertex.Z(); + float r2 = dx * dx + dy * dy; + float posTgl2 = r2 > 1 && std::abs(dz) < 20 ? dz * dz / r2 : 0; + if (posTgl2 > params.maxPosTglMC * params.maxPosTglMC) { + if (mVerbose > 1 && followDecay > -1) { + LOGP(info, "rejecting decay {} prong : pdg={}, pT={}, tgL={}, dr={}, dz={} r={}, z={}, posTgl={}", followDecay, tr.GetPdgCode(), tr.GetPt(), tr.GetTgl(), std::sqrt(r2), dz, std::sqrt(tr.R2()), tr.GetStartVertexCoordinatesZ(), std::sqrt(posTgl2)); + } + return false; + } + if (params.requireITSorTPCTrackRefs) { + auto trspan = mcReader.getTrackRefs(lb.getSourceID(), lb.getEventID(), lb.getTrackID()); + bool ok = false; + for (const auto& trf : trspan) { + if (trf.getDetectorId() == DetID::ITS || trf.getDetectorId() == DetID::TPC) { + ok = true; + break; + } + } + if (!ok) { + return false; + } + } + TParticlePDG* pPDG = O2DatabasePDG::Instance()->GetParticle(tr.GetPdgCode()); + if (!pPDG) { + LOGP(debug, "Unknown particle {}", tr.GetPdgCode()); + return false; + } + if (pPDG->Charge() == 0.) { + return false; + } + return addMCParticle(tr, lb, pPDG); +} + +bool TrackMCStudy::addMCParticle(const MCTrack& mcPart, const o2::MCCompLabel& lb, TParticlePDG* pPDG) +{ + std::array xyz{(float)mcPart.GetStartVertexCoordinatesX(), (float)mcPart.GetStartVertexCoordinatesY(), (float)mcPart.GetStartVertexCoordinatesZ()}; + std::array pxyz{(float)mcPart.GetStartVertexMomentumX(), (float)mcPart.GetStartVertexMomentumY(), (float)mcPart.GetStartVertexMomentumZ()}; + if (!pPDG && !(pPDG = O2DatabasePDG::Instance()->GetParticle(mcPart.GetPdgCode()))) { + LOGP(debug, "Unknown particle {}", mcPart.GetPdgCode()); + return false; + } + auto& mcEntry = mSelMCTracks[lb]; + mcEntry.mcTrackInfo.pdg = mcPart.GetPdgCode(); + mcEntry.mcTrackInfo.track = o2::track::TrackPar(xyz, pxyz, TMath::Nint(pPDG->Charge() / 3), false); + mcEntry.mcTrackInfo.label = lb; + mcEntry.mcTrackInfo.bcInTF = mIntBC[lb.getEventID()]; + mcEntry.mcTrackInfo.occTPC = mTPCOcc[lb.getEventID()]; + int moth = -1; + o2::MCCompLabel mclbPar; + if (!mcPart.isPrimary() && (moth = mcPart.getMotherTrackId()) >= 0) { + const auto& mcPartPar = mCurrMCTracks[moth]; + mcEntry.mcTrackInfo.pdgParent = mcPartPar.GetPdgCode(); + } + if (mVerbose > 1) { + LOGP(info, "Adding charged MC pdg={} {} ", mcPart.GetPdgCode(), lb.asString()); + } + return true; +} + +void TrackMCStudy::loadTPCOccMap(const o2::globaltracking::RecoContainer& recoData) +{ + auto NHBPerTF = o2::base::GRPGeomHelper::instance().getGRPECS()->getNHBFPerTF(); + o2::tpc::CorrectionMapsLoader TPCCorrMapsLoader{}; + const auto& TPCOccMap = recoData.occupancyMapTPC; + auto prop = o2::base::Propagator::Instance(); + auto TPCRefitter = std::make_unique(&recoData.inputsTPCclusters->clusterIndex, &TPCCorrMapsLoader, prop->getNominalBz(), + recoData.getTPCTracksClusterRefs().data(), 0, recoData.clusterShMapTPC.data(), TPCOccMap.data(), TPCOccMap.size(), nullptr, prop); + mNTPCOccBinLength = TPCRefitter->getParam()->rec.tpc.occupancyMapTimeBins; + mTBinClOcc.clear(); + if (mNTPCOccBinLength > 1 && TPCOccMap.size()) { + mNTPCOccBinLengthInv = 1. / mNTPCOccBinLength; + int nTPCBins = NHBPerTF * o2::constants::lhc::LHCMaxBunches / 8, ninteg = 0; + int nTPCOccBins = nTPCBins * mNTPCOccBinLengthInv, sumBins = std::max(1, int(o2::constants::lhc::LHCMaxBunches / 8 * mNTPCOccBinLengthInv)); + mTBinClOcc.resize(nTPCOccBins); + std::vector mltHistTB(nTPCOccBins); + float sm = 0., tb = 0.5 * mNTPCOccBinLength; + for (int i = 0; i < nTPCOccBins; i++) { + mltHistTB[i] = TPCRefitter->getParam()->GetUnscaledMult(tb); + tb += mNTPCOccBinLength; + } + for (int i = nTPCOccBins; i--;) { + sm += mltHistTB[i]; + if (i + sumBins < nTPCOccBins) { + sm -= mltHistTB[i + sumBins]; + } + mTBinClOcc[i] = sm; + } + } else { + mTBinClOcc.resize(1); + } +} + +DataProcessorSpec getTrackMCStudySpec(GTrackID::mask_t srcTracks, GTrackID::mask_t srcClusters) { std::vector outputs; auto dataRequest = std::make_shared(); @@ -548,15 +732,13 @@ DataProcessorSpec getTrackMCStudySpec(GTrackID::mask_t srcTracks, GTrackID::mask "track-mc-study", dataRequest->inputs, outputs, - AlgorithmSpec{adaptFromTask(dataRequest, ggRequest, srcTracks, checkMatching)}, + AlgorithmSpec{adaptFromTask(dataRequest, ggRequest, srcTracks)}, Options{ {"device-verbosity", VariantType::Int, 0, {"Verbosity level"}}, {"dcay-vs-pt", VariantType::String, "0.0105 + 0.0350 / pow(x, 1.1)", {"Formula for global tracks DCAy vs pT cut"}}, {"min-tpc-clusters", VariantType::Int, 60, {"Cut on TPC clusters"}}, {"max-tpc-dcay", VariantType::Float, 2.f, {"Cut on TPC dcaY"}}, {"max-tpc-dcaz", VariantType::Float, 2.f, {"Cut on TPC dcaZ"}}, - {"max-eta", VariantType::Float, 1.5f, {"Cut on track eta"}}, - {"min-pt", VariantType::Float, 0.02f, {"Cut on track pT"}}, {"min-x-prop", VariantType::Float, 6.f, {"track should be propagated to this X at least"}}, }}; } diff --git a/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudyConfig.cxx b/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudyConfig.cxx new file mode 100644 index 0000000000000..300e69c514ae3 --- /dev/null +++ b/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudyConfig.cxx @@ -0,0 +1,14 @@ +// 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. + +#include "GlobalTrackingStudy/TrackMCStudyConfig.h" + +O2ParamImpl(o2::trackstudy::TrackMCStudyConfig); diff --git a/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudyTypes.cxx b/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudyTypes.cxx new file mode 100644 index 0000000000000..b95fda9699420 --- /dev/null +++ b/Detectors/GlobalTrackingWorkflow/study/src/TrackMCStudyTypes.cxx @@ -0,0 +1,12 @@ +// 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. + +#include "GlobalTrackingStudy/TrackMCStudyTypes.h" diff --git a/Detectors/GlobalTrackingWorkflow/study/src/trackMCStudy-workflow.cxx b/Detectors/GlobalTrackingWorkflow/study/src/trackMCStudy-workflow.cxx index 93e549dcc2ef3..3674fd5f3e3a8 100644 --- a/Detectors/GlobalTrackingWorkflow/study/src/trackMCStudy-workflow.cxx +++ b/Detectors/GlobalTrackingWorkflow/study/src/trackMCStudy-workflow.cxx @@ -40,7 +40,6 @@ void customize(std::vector& workflowOptions) {"cluster-sources", VariantType::String, std::string{GID::ALL}, {"comma-separated list of cluster sources to use"}}, {"disable-root-input", VariantType::Bool, false, {"disable root-files input reader"}}, {"disable-mc", o2::framework::VariantType::Bool, false, {"disable MC propagation, never use it"}}, - {"check-its-tpc", VariantType::Bool, false, {"Special output for failed ITS-TPC matches"}}, {"configKeyValues", VariantType::String, "", {"Semicolon separated key=value strings ..."}}}; o2::raw::HBFUtilsInitializer::addConfigOption(options); std::swap(workflowOptions, options); @@ -57,22 +56,19 @@ WorkflowSpec defineDataProcessing(ConfigContext const& configcontext) if (!useMC) { throw std::runtime_error("MC cannot be disabled for this workflow"); } - bool checkMatching = configcontext.options().get("check-its-tpc"); GID::mask_t allowedSourcesTrc = GID::getSourcesMask("ITS,TPC,ITS-TPC,TPC-TOF,TPC-TRD,ITS-TPC-TRD,TPC-TRD-TOF,ITS-TPC-TOF,ITS-TPC-TRD-TOF"); - GID::mask_t allowedSourcesClus = GID::getSourcesMask(checkMatching ? "TPC" : "none"); + GID::mask_t allowedSourcesClus = GID::getSourcesMask("ITS,TPC"); // Update the (declared) parameters if changed from the command line o2::conf::ConfigurableParam::updateFromString(configcontext.options().get("configKeyValues")); GID::mask_t srcTrc = allowedSourcesTrc & GID::getSourcesMask(configcontext.options().get("track-sources")); GID::mask_t srcCls = allowedSourcesClus & GID::getSourcesMask(configcontext.options().get("cluster-sources")); - if (checkMatching) { - srcCls |= GID::getSourcesMask("TPC"); - } + srcCls |= GID::getSourcesMask("ITS,TPC"); o2::globaltracking::InputHelper::addInputSpecs(configcontext, specs, srcCls, srcTrc, srcTrc, true); o2::globaltracking::InputHelper::addInputSpecsPVertex(configcontext, specs, true); // P-vertex is always needed - specs.emplace_back(o2::trackstudy::getTrackMCStudySpec(srcTrc, srcCls, checkMatching)); + specs.emplace_back(o2::trackstudy::getTrackMCStudySpec(srcTrc, srcCls)); // configure dpl timer to inject correct firstTForbit: start from the 1st orbit of TF containing 1st sampled orbit o2::raw::HBFUtilsInitializer hbfIni(configcontext, specs); diff --git a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/ClusterLinesGPU.h b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/ClusterLinesGPU.h index ecb25b7bc9d71..75d75e0f67700 100644 --- a/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/ClusterLinesGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/ClusterLinesGPU.h @@ -57,7 +57,7 @@ class ClusterLinesGPU final public: GPUd() ClusterLinesGPU(const Line& firstLine, const Line& secondLine); // poor man solution to calculate duplets' centroid GPUd() void computeClusterCentroid(); - GPUd() inline float* getVertex() { return mVertex; } + GPUdi() float* getVertex() { return mVertex; } private: float mAMatrix[6]; // AX=B diff --git a/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/RawPixelDecoder.h b/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/RawPixelDecoder.h index 4846eab140773..810bff1037513 100644 --- a/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/RawPixelDecoder.h +++ b/Detectors/ITSMFT/common/reconstruction/include/ITSMFTReconstruction/RawPixelDecoder.h @@ -95,7 +95,7 @@ class RawPixelDecoder final : public PixelReader bool getAlwaysParseTrigger() const { return mAlwaysParseTrigger; } void printReport(bool decstat = true, bool skipNoErr = true) const; - void produceRawDataDumps(int dump, const o2::framework::TimingInfo& tinfo); + size_t produceRawDataDumps(int dump, const o2::framework::TimingInfo& tinfo); void clearStat(bool resetRaw = false); diff --git a/Detectors/ITSMFT/common/reconstruction/src/RawPixelDecoder.cxx b/Detectors/ITSMFT/common/reconstruction/src/RawPixelDecoder.cxx index 209a9fde82632..dc61bea9f406e 100644 --- a/Detectors/ITSMFT/common/reconstruction/src/RawPixelDecoder.cxx +++ b/Detectors/ITSMFT/common/reconstruction/src/RawPixelDecoder.cxx @@ -515,8 +515,9 @@ void RawPixelDecoder::clearStat(bool resetRaw) ///______________________________________________________________________ template -void RawPixelDecoder::produceRawDataDumps(int dump, const o2::framework::TimingInfo& tinfo) +size_t RawPixelDecoder::produceRawDataDumps(int dump, const o2::framework::TimingInfo& tinfo) { + size_t outSize = 0; bool dumpFullTF = false; for (auto& ru : mRUDecodeVec) { if (ru.linkHBFToDump.size()) { @@ -550,6 +551,7 @@ void RawPixelDecoder::produceRawDataDumps(int dump, const o2::framework break; } ostrm.write(reinterpret_cast(piece->data), piece->size); + outSize += piece->size; entry++; } LOG(info) << "produced " << std::filesystem::current_path().c_str() << '/' << fnm; @@ -569,11 +571,13 @@ void RawPixelDecoder::produceRawDataDumps(int dump, const o2::framework for (size_t i = 0; i < lnk.rawData.getNPieces(); i++) { const auto* piece = lnk.rawData.getPiece(i); ostrm.write(reinterpret_cast(piece->data), piece->size); + outSize += piece->size; } } LOG(info) << "produced " << std::filesystem::current_path().c_str() << '/' << fnm; break; } + return outSize; } ///______________________________________________________________________ diff --git a/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/DeadMapBuilderSpec.h b/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/DeadMapBuilderSpec.h index 00f9fc931e679..7bce60d172222 100644 --- a/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/DeadMapBuilderSpec.h +++ b/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/DeadMapBuilderSpec.h @@ -102,6 +102,7 @@ class ITSMFTDeadMapBuilder : public Task std::vector mDeadMapTF{}; + int mRunNumber = -1; unsigned long mFirstOrbitTF = 0x0; unsigned long mFirstOrbitRun = 0x0; diff --git a/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/STFDecoderSpec.h b/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/STFDecoderSpec.h index 297cfb446f8ef..a6876c456842d 100644 --- a/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/STFDecoderSpec.h +++ b/Detectors/ITSMFT/common/workflow/include/ITSMFTWorkflow/STFDecoderSpec.h @@ -81,6 +81,7 @@ class STFDecoder : public Task bool mApplyNoiseMap = true; bool mUseClusterDictionary = true; bool mVerifyDecoder = false; + bool mDumpFrom1stPipeline = false; int mDumpOnError = 0; int mNThreads = 1; int mVerbosity = 0; @@ -91,6 +92,8 @@ class STFDecoder : public Task size_t mEstNClusPatt = 0; size_t mEstNCalib = 0; size_t mEstNROF = 0; + size_t mMaxRawDumpsSize = 0; + size_t mRawDumpedSize = 0; std::string mInputSpec; std::string mSelfName; std::unique_ptr> mDecoder; diff --git a/Detectors/ITSMFT/common/workflow/src/DeadMapBuilderSpec.cxx b/Detectors/ITSMFT/common/workflow/src/DeadMapBuilderSpec.cxx index 39d1bb7df7728..c97c3440afcc3 100644 --- a/Detectors/ITSMFT/common/workflow/src/DeadMapBuilderSpec.cxx +++ b/Detectors/ITSMFT/common/workflow/src/DeadMapBuilderSpec.cxx @@ -64,6 +64,7 @@ void ITSMFTDeadMapBuilder::init(InitContext& ic) mLocalOutputDir = ic.options().get("output-dir"); mSkipStaticMap = ic.options().get("skip-static-map"); + isEnded = false; mTimeStart = o2::ccdb::getCurrentTimestamp(); if (mRunMFT) { @@ -89,7 +90,7 @@ void ITSMFTDeadMapBuilder::init(InitContext& ic) } /////////////////////////////////////////////////////////////////// -// TODO: can ChipMappingITS help here? +// TODO: can ChipMappingITS help here? std::vector ITSMFTDeadMapBuilder::getChipIDsOnSameCable(uint16_t chip) { if (mRunMFT || chip < N_CHIPS_ITSIB) { @@ -178,19 +179,32 @@ void ITSMFTDeadMapBuilder::finalizeOutput() void ITSMFTDeadMapBuilder::run(ProcessingContext& pc) { + // Skip everything in case of garbage (potentially at EoS) + if (pc.services().get().firstTForbit == -1U) { + LOG(info) << "Skipping the processing of inputs for timeslice " << pc.services().get().timeslice << " (firstTForbit is " << pc.services().get().firstTForbit << ")"; + return; + } + std::chrono::time_point start; std::chrono::time_point end; start = std::chrono::high_resolution_clock::now(); - mTFCounter++; - - mFirstOrbitTF = pc.services().get().firstTForbit; + const auto& tinfo = pc.services().get(); - if (mFirstOrbitRun == 0x0) { + if (tinfo.globalRunNumberChanged || mFirstOrbitRun == 0x0) { // new run is starting + mRunNumber = tinfo.runNumber; mFirstOrbitRun = mFirstOrbitTF; + mTFCounter = 0; + isEnded = false; } + if (isEnded) { + return; + } + mFirstOrbitTF = tinfo.firstTForbit; + mTFCounter++; + long sampled_orbit = mFirstOrbitTF - mFirstOrbitRun; if (!acceptTF(sampled_orbit)) { @@ -198,7 +212,7 @@ void ITSMFTDeadMapBuilder::run(ProcessingContext& pc) } mStepCounter++; - LOG(info) << "Processing step #" << mStepCounter << " out of " << mTFCounter << " TF received. First orbit " << mFirstOrbitTF; + LOG(info) << "Processing step #" << mStepCounter << " out of " << mTFCounter << " good TF received. First orbit " << mFirstOrbitTF; mDeadMapTF.clear(); @@ -280,6 +294,12 @@ void ITSMFTDeadMapBuilder::run(ProcessingContext& pc) LOG(info) << "Elapsed time in TF processing: " << difference / 1000. << " ms"; + if (pc.transitionState() == TransitionHandlingState::Requested && !isEnded) { + std::string detname = mRunMFT ? "MFT" : "ITS"; + LOG(warning) << "Transition state requested for " << detname << " process, calling stop() and stopping the process of new data."; + stop(); + } + return; } @@ -291,8 +311,7 @@ void ITSMFTDeadMapBuilder::PrepareOutputCcdb(EndOfStreamContext* ec, std::string long tend = o2::ccdb::getCurrentTimestamp(); - std::map md = { - {"map_version", MAP_VERSION}}; + std::map md = {{"map_version", MAP_VERSION}, {"runNumber", std::to_string(mRunNumber)}}; std::string path = mRunMFT ? "MFT/Calib/" : "ITS/Calib/"; std::string name_str = "TimeDeadMap"; @@ -355,6 +374,7 @@ void ITSMFTDeadMapBuilder::endOfStream(EndOfStreamContext& ec) } else { LOG(warning) << "Time-dependent dead map is empty and will not be forwarded as output"; } + LOG(info) << "Stop process of new data because of endOfStream"; isEnded = true; } return; @@ -374,6 +394,7 @@ void ITSMFTDeadMapBuilder::stop() } else { LOG(alarm) << "endOfStream not processed. Nothing forwarded as output."; } + LOG(info) << "Stop process of new data because of stop() call."; isEnded = true; } return; diff --git a/Detectors/ITSMFT/common/workflow/src/STFDecoderSpec.cxx b/Detectors/ITSMFT/common/workflow/src/STFDecoderSpec.cxx index bbd1a52a919e5..d798bbc62204d 100644 --- a/Detectors/ITSMFT/common/workflow/src/STFDecoderSpec.cxx +++ b/Detectors/ITSMFT/common/workflow/src/STFDecoderSpec.cxx @@ -87,8 +87,16 @@ void STFDecoder::init(InitContext& ic) mDecoder->setNThreads(mNThreads); mUnmutExtraLanes = ic.options().get("unmute-extra-lanes"); mVerbosity = ic.options().get("decoder-verbosity"); + auto dmpSz = ic.options().get("stop-raw-data-dumps-after-size"); + if (dmpSz > 0) { + mMaxRawDumpsSize = size_t(dmpSz) * 1024 * 1024; + } mDumpOnError = ic.options().get("raw-data-dumps"); - if (mDumpOnError < 0 || mDumpOnError >= int(GBTLink::RawDataDumps::DUMP_NTYPES)) { + if (mDumpOnError < 0) { + mDumpOnError = -mDumpOnError; + mDumpFrom1stPipeline = true; + } + if (mDumpOnError >= int(GBTLink::RawDataDumps::DUMP_NTYPES)) { throw std::runtime_error(fmt::format("unknown raw data dump level {} requested", mDumpOnError)); } auto dumpDir = ic.options().get("raw-data-dumps-directory"); @@ -243,8 +251,13 @@ void STFDecoder::run(ProcessingContext& pc) pc.outputs().snapshot(Output{orig, "PHYSTRIG", 0}, mDecoder->getExternalTriggers()); - if (mDumpOnError != int(GBTLink::RawDataDumps::DUMP_NONE)) { - mDecoder->produceRawDataDumps(mDumpOnError, pc.services().get()); + if (mDumpOnError != int(GBTLink::RawDataDumps::DUMP_NONE) && + (!mDumpFrom1stPipeline || pc.services().get().inputTimesliceId == 0)) { + mRawDumpedSize += mDecoder->produceRawDataDumps(mDumpOnError, pc.services().get()); + if (mRawDumpedSize > mMaxRawDumpsSize && mMaxRawDumpsSize > 0) { + LOGP(info, "Max total dumped size {} MB exceeded allowed limit, disabling further dumping", mRawDumpedSize / (1024 * 1024)); + mDumpOnError = int(GBTLink::RawDataDumps::DUMP_NONE); + } } if (mDoClusters) { @@ -417,8 +430,9 @@ DataProcessorSpec getSTFDecoderSpec(const STFDecoderInp& inp) {"nthreads", VariantType::Int, 1, {"Number of decoding/clustering threads"}}, {"decoder-verbosity", VariantType::Int, 0, {"Verbosity level (-1: silent, 0: errors, 1: headers, 2: data, 3: raw data dump) of 1st lane"}}, {"always-parse-trigger", VariantType::Bool, false, {"parse trigger word even if flags continuation of old trigger"}}, - {"raw-data-dumps", VariantType::Int, int(GBTLink::RawDataDumps::DUMP_NONE), {"Raw data dumps on error (0: none, 1: HBF for link, 2: whole TF for all links"}}, + {"raw-data-dumps", VariantType::Int, int(GBTLink::RawDataDumps::DUMP_NONE), {"Raw data dumps on error (0: none, 1: HBF for link, 2: whole TF for all links. If negative, dump only on from 1st pipeline."}}, {"raw-data-dumps-directory", VariantType::String, "", {"Destination directory for the raw data dumps"}}, + {"stop-raw-data-dumps-after-size", VariantType::Int, 1024, {"Stop dumping once this size in MB is accumulated. 0: no limit"}}, {"unmute-extra-lanes", VariantType::Bool, false, {"allow extra lanes to be as verbose as 1st one"}}, {"allow-empty-rofs", VariantType::Bool, false, {"record ROFs w/o any hit"}}, {"ignore-noise-map", VariantType::Bool, false, {"do not mask pixels flagged in the noise map"}}, diff --git a/Detectors/TRD/base/include/TRDBase/GeometryFlat.h b/Detectors/TRD/base/include/TRDBase/GeometryFlat.h index 73dd11247dd66..6474c73e55756 100644 --- a/Detectors/TRD/base/include/TRDBase/GeometryFlat.h +++ b/Detectors/TRD/base/include/TRDBase/GeometryFlat.h @@ -34,7 +34,7 @@ class Geometry; class GeometryFlat : public o2::gpu::FlatObject, public GeometryBase { public: -#ifndef GPUCA_GPUCODE_DEVICE +#ifndef GPUCA_GPUCODE GeometryFlat() = default; GeometryFlat(const GeometryFlat& v) : FlatObject(), GeometryBase() { diff --git a/Detectors/Upgrades/ITS3/reconstruction/CMakeLists.txt b/Detectors/Upgrades/ITS3/reconstruction/CMakeLists.txt index 128e145b31317..a666e15cc21e4 100644 --- a/Detectors/Upgrades/ITS3/reconstruction/CMakeLists.txt +++ b/Detectors/Upgrades/ITS3/reconstruction/CMakeLists.txt @@ -19,7 +19,6 @@ o2_add_library(ITS3Reconstruction src/BuildTopologyDictionary.cxx src/LookUp.cxx src/IOUtils.cxx - # src/FastMultEst.cxx PUBLIC_LINK_LIBRARIES O2::ITSMFTBase O2::ITSMFTReconstruction O2::ITS3Base @@ -41,6 +40,15 @@ o2_target_root_dictionary( include/ITS3Reconstruction/IOUtils.h ) +o2_add_library(ITS3TrackingInterface + TARGETVARNAME targetName + SOURCES src/TrackingInterface.cxx + PRIVATE_LINK_LIBRARIES + O2::ITS3Reconstruction + O2::ITSTrackingInterface + O2::Framework + O2::GPUTracking) + if (OpenMP_CXX_FOUND) target_compile_definitions(${targetName} PRIVATE WITH_OPENMP) target_link_libraries(${targetName} PRIVATE OpenMP::OpenMP_CXX) diff --git a/Detectors/Upgrades/ITS3/reconstruction/include/ITS3Reconstruction/TrackingInterface.h b/Detectors/Upgrades/ITS3/reconstruction/include/ITS3Reconstruction/TrackingInterface.h new file mode 100644 index 0000000000000..ab2ff0086200b --- /dev/null +++ b/Detectors/Upgrades/ITS3/reconstruction/include/ITS3Reconstruction/TrackingInterface.h @@ -0,0 +1,42 @@ +// 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. + +#ifndef O2_ITS3_TRACKINGINTERFACE +#define O2_ITS3_TRACKINGINTERFACE + +#include "ITStracking/TrackingInterface.h" +#include "ITS3Reconstruction/TopologyDictionary.h" + +namespace o2::its3 +{ + +class ITS3TrackingInterface final : public its::ITSTrackingInterface +{ + public: + using its::ITSTrackingInterface::ITSTrackingInterface; + + void setClusterDictionary(const o2::its3::TopologyDictionary* d) { mDict = d; } + void updateTimeDependentParams(framework::ProcessingContext& pc) final; + void finaliseCCDB(framework::ConcreteDataMatcher& matcher, void* obj) final; + + protected: + void loadROF(gsl::span& trackROFspan, + gsl::span clusters, + gsl::span::iterator& pattIt, + const dataformats::MCTruthContainer* mcLabels) final; + + private: + const o2::its3::TopologyDictionary* mDict{nullptr}; +}; + +} // namespace o2::its3 + +#endif diff --git a/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx b/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx index f4cf31deb1ef8..50e651f7f5675 100644 --- a/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx +++ b/Detectors/Upgrades/ITS3/reconstruction/src/IOUtils.cxx @@ -119,6 +119,9 @@ int loadROFrameDataITS3(its::TimeFrame* tf, for (auto& v : tf->mNTrackletsPerCluster) { v.resize(tf->getUnsortedClusters()[1].size()); } + for (auto& v : tf->mNTrackletsPerClusterSum) { + v.resize(tf->getUnsortedClusters()[1].size() + 1); + } if (mcLabels != nullptr) { tf->mClusterLabels = mcLabels; diff --git a/Detectors/Upgrades/ITS3/reconstruction/src/TrackingInterface.cxx b/Detectors/Upgrades/ITS3/reconstruction/src/TrackingInterface.cxx new file mode 100644 index 0000000000000..10c6b9265a8bb --- /dev/null +++ b/Detectors/Upgrades/ITS3/reconstruction/src/TrackingInterface.cxx @@ -0,0 +1,74 @@ +// 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. + +#include "ITS3Reconstruction/TrackingInterface.h" +#include "ITS3Reconstruction/IOUtils.h" +#include "ITSBase/GeometryTGeo.h" +#include "ITSMFTBase/DPLAlpideParam.h" +#include "DetectorsBase/GRPGeomHelper.h" + +namespace o2::its3 +{ + +void ITS3TrackingInterface::updateTimeDependentParams(framework::ProcessingContext& pc) +{ + o2::base::GRPGeomHelper::instance().checkUpdates(pc); + static bool initOnceDone = false; + if (!initOnceDone) { // this params need to be queried only once + initOnceDone = true; + pc.inputs().get("cldict"); // just to trigger the finaliseCCDB + pc.inputs().get*>("alppar"); + if (pc.inputs().getPos("itsTGeo") >= 0) { + pc.inputs().get("itsTGeo"); + } + auto geom = its::GeometryTGeo::Instance(); + geom->fillMatrixCache(o2::math_utils::bit2Mask(o2::math_utils::TransformType::T2L, o2::math_utils::TransformType::T2GRot, o2::math_utils::TransformType::T2G)); + getConfiguration(pc); + } +} + +void ITS3TrackingInterface::finaliseCCDB(framework::ConcreteDataMatcher& matcher, void* obj) +{ + if (o2::base::GRPGeomHelper::instance().finaliseCCDB(matcher, obj)) { + return; + } + if (matcher == framework::ConcreteDataMatcher("IT3", "CLUSDICT", 0)) { + LOG(info) << "cluster dictionary updated"; + setClusterDictionary((const o2::its3::TopologyDictionary*)obj); + return; + } + if (matcher == framework::ConcreteDataMatcher("ITS", "ALPIDEPARAM", 0)) { + LOG(info) << "Alpide param updated"; + const auto& par = o2::itsmft::DPLAlpideParam::Instance(); + par.printKeyValues(); + return; + } + if (matcher == framework::ConcreteDataMatcher("GLO", "MEANVERTEX", 0)) { + LOGP(info, "Mean vertex acquired"); + setMeanVertex((const o2::dataformats::MeanVertexObject*)obj); + return; + } + if (matcher == framework::ConcreteDataMatcher("ITS", "GEOMTGEO", 0)) { + LOG(info) << "ITS GeometryTGeo loaded from ccdb"; + o2::its::GeometryTGeo::adopt((o2::its::GeometryTGeo*)obj); + return; + } +} + +void ITS3TrackingInterface::loadROF(gsl::span& trackROFspan, + gsl::span clusters, + gsl::span::iterator& pattIt, + const dataformats::MCTruthContainer* mcLabels) +{ + ioutils::loadROFrameDataITS3(mTimeFrame, trackROFspan, clusters, pattIt, mDict, mcLabels); +} + +} // namespace o2::its3 diff --git a/Detectors/Upgrades/ITS3/workflow/CMakeLists.txt b/Detectors/Upgrades/ITS3/workflow/CMakeLists.txt index c4867f0e94ec1..649e4d737d42c 100644 --- a/Detectors/Upgrades/ITS3/workflow/CMakeLists.txt +++ b/Detectors/Upgrades/ITS3/workflow/CMakeLists.txt @@ -19,7 +19,6 @@ o2_add_library(ITS3Workflow src/ClustererSpec.cxx src/ClusterWriterSpec.cxx src/TrackerSpec.cxx - # src/CookedTrackerSpec.cxx src/TrackWriterSpec.cxx src/TrackReaderSpec.cxx src/VertexReaderSpec.cxx @@ -31,6 +30,7 @@ o2_add_library(ITS3Workflow O2::ITStracking O2::ITSMFTReconstruction O2::ITS3Reconstruction + O2::ITS3TrackingInterface O2::ITSWorkflow O2::GPUTracking O2::ITSBase) diff --git a/Detectors/Upgrades/ITS3/workflow/include/ITS3Workflow/TrackerSpec.h b/Detectors/Upgrades/ITS3/workflow/include/ITS3Workflow/TrackerSpec.h index f350c09b0192c..f5c1d7bf0e947 100644 --- a/Detectors/Upgrades/ITS3/workflow/include/ITS3Workflow/TrackerSpec.h +++ b/Detectors/Upgrades/ITS3/workflow/include/ITS3Workflow/TrackerSpec.h @@ -15,57 +15,55 @@ #define O2_ITS3_TRACKERDPL #include "DataFormatsParameters/GRPObject.h" -#include "ITS3Reconstruction/TopologyDictionary.h" +#include "DataFormatsITSMFT/TopologyDictionary.h" +#include "DataFormatsCalibration/MeanVertexObject.h" #include "Framework/DataProcessorSpec.h" #include "Framework/Task.h" -#include "ITStracking/TimeFrame.h" -#include "ITStracking/Tracker.h" -#include "ITStracking/TrackerTraits.h" -#include "ITStracking/Vertexer.h" -#include "ITStracking/VertexerTraits.h" +#include "ITS3Reconstruction/TrackingInterface.h" -#include "GPUO2Interface.h" -#include "GPUReconstruction.h" -#include "GPUChainITS.h" -#include "CommonUtils/StringUtils.h" -#include "TStopwatch.h" +#include "GPUDataTypes.h" #include "DetectorsBase/GRPGeomHelper.h" +#include "TStopwatch.h" + namespace o2::its3 { class TrackerDPL : public framework::Task { public: - TrackerDPL(std::shared_ptr gr, bool isMC, int trgType, const std::string& trModeS, o2::gpu::GPUDataTypes::DeviceType dType = o2::gpu::GPUDataTypes::DeviceType::CPU); + TrackerDPL(std::shared_ptr gr, + bool isMC, + int trgType, + const its::TrackingMode& trMode = its::TrackingMode::Unset, + const bool overrBeamEst = false, + gpu::GPUDataTypes::DeviceType dType = gpu::GPUDataTypes::DeviceType::CPU); + ~TrackerDPL() override = default; + TrackerDPL(const TrackerDPL&) = delete; + TrackerDPL(TrackerDPL&&) = delete; + TrackerDPL& operator=(const TrackerDPL&) = delete; + TrackerDPL& operator=(TrackerDPL&&) = delete; + void init(framework::InitContext& ic) final; void run(framework::ProcessingContext& pc) final; void endOfStream(framework::EndOfStreamContext& ec) final; void finaliseCCDB(framework::ConcreteDataMatcher& matcher, void* obj) final; - void setClusterDictionary(const o2::its3::TopologyDictionary* d) { mDict = d; } + void stop() final; private: void updateTimeDependentParams(framework::ProcessingContext& pc); - - std::shared_ptr mGGCCDBRequest{}; - bool mIsMC{false}; - int mUseTriggers{0}; - std::string mMode{"sync"}; - std::unique_ptr mRecChain{}; - bool mRunVertexer{true}; - bool mCosmicsProcessing{false}; - const o2::its3::TopologyDictionary* mDict{}; - std::unique_ptr mChainITS{}; - std::unique_ptr mTracker{}; - std::unique_ptr mVertexer{}; - TStopwatch mTimer{}; + std::unique_ptr mRecChain = nullptr; + std::unique_ptr mChainITS = nullptr; + std::shared_ptr mGGCCDBRequest; + ITS3TrackingInterface mITS3TrackingInterface; + TStopwatch mTimer; }; /// create a processor spec /// run ITS CA tracker -framework::DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, int useTrig, const std::string& trModeS, o2::gpu::GPUDataTypes::DeviceType dType); +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); } // namespace o2::its3 diff --git a/Detectors/Upgrades/ITS3/workflow/src/TrackerSpec.cxx b/Detectors/Upgrades/ITS3/workflow/src/TrackerSpec.cxx index cc0cbcf333dd3..90f94e625d6ea 100644 --- a/Detectors/Upgrades/ITS3/workflow/src/TrackerSpec.cxx +++ b/Detectors/Upgrades/ITS3/workflow/src/TrackerSpec.cxx @@ -46,346 +46,57 @@ namespace its3 { using Vertex = o2::dataformats::Vertex>; -TrackerDPL::TrackerDPL(std::shared_ptr gr, bool isMC, int trgType, const std::string& trModeS, o2::gpu::GPUDataTypes::DeviceType dType) : mGGCCDBRequest(gr), mIsMC{isMC}, mUseTriggers{trgType}, mMode{trModeS}, mRecChain{o2::gpu::GPUReconstruction::CreateInstance(dType, true)} +TrackerDPL::TrackerDPL(std::shared_ptr gr, + bool isMC, + int trgType, + const its::TrackingMode& trMode, + const bool overrBeamEst, + o2::gpu::GPUDataTypes::DeviceType dType) : mGGCCDBRequest(gr), + mRecChain{o2::gpu::GPUReconstruction::CreateInstance(dType, true)}, + mITS3TrackingInterface{isMC, trgType, overrBeamEst} { - std::transform(mMode.begin(), mMode.end(), mMode.begin(), [](unsigned char c) { return std::tolower(c); }); + mITS3TrackingInterface.setTrackingMode(trMode); } -void TrackerDPL::init(InitContext& /*ic*/) +void TrackerDPL::init(InitContext& ic) { mTimer.Stop(); mTimer.Reset(); - o2::base::GRPGeomHelper::instance().setRequest(mGGCCDBRequest); - mChainITS.reset(mRecChain->AddChain()); - mVertexer = std::make_unique(mChainITS->GetITSVertexerTraits()); - mTracker = std::make_unique(mChainITS->GetITSTrackerTraits()); - mRunVertexer = true; - mCosmicsProcessing = false; - std::vector trackParams; - - if (mMode == "async") { - - trackParams.resize(3); - trackParams[1].TrackletMinPt = 0.2f; - trackParams[1].CellDeltaTanLambdaSigma *= 2.; - trackParams[2].TrackletMinPt = 0.1f; - trackParams[2].CellDeltaTanLambdaSigma *= 4.; - trackParams[2].MinTrackLength = 4; - - LOG(info) << "Initializing tracker in async. phase reconstruction with " << trackParams.size() << " passes"; - - } else if (mMode == "async_misaligned") { - - trackParams.resize(3); - trackParams[0].PhiBins = 32; - trackParams[0].ZBins = 64; - trackParams[0].CellDeltaTanLambdaSigma *= 3.; - trackParams[0].SystErrorZ2[0] = 1.e-4; - trackParams[0].SystErrorZ2[1] = 1.e-4; - trackParams[0].SystErrorZ2[2] = 1.e-4; - std::copy(trackParams[0].SystErrorZ2.begin(), trackParams[0].SystErrorZ2.end(), trackParams[0].SystErrorY2.begin()); - trackParams[0].MaxChi2ClusterAttachment = 60.; - trackParams[0].MaxChi2NDF = 40.; - trackParams[1] = trackParams[0]; - trackParams[2] = trackParams[0]; - trackParams[1].MinTrackLength = 6; - trackParams[2].MinTrackLength = 4; - LOG(info) << "Initializing tracker in misaligned async. phase reconstruction with " << trackParams.size() << " passes"; - - } else if (mMode == "sync_misaligned") { - - trackParams.resize(3); - trackParams[0].PhiBins = 32; - trackParams[0].ZBins = 64; - trackParams[0].CellDeltaTanLambdaSigma *= 3.; - trackParams[0].SystErrorZ2[0] = 1.e-4; - trackParams[0].SystErrorZ2[1] = 1.e-4; - trackParams[0].SystErrorZ2[2] = 1.e-4; - trackParams[0].SystErrorZ2[3] = 9.e-4; - trackParams[0].SystErrorZ2[4] = 9.e-4; - trackParams[0].SystErrorZ2[5] = 9.e-4; - trackParams[0].SystErrorZ2[6] = 9.e-4; - std::copy(trackParams[0].SystErrorZ2.begin(), trackParams[0].SystErrorZ2.end(), trackParams[0].SystErrorY2.begin()); - trackParams[0].MaxChi2ClusterAttachment = 60.; - trackParams[0].MaxChi2NDF = 40.; - trackParams[1] = trackParams[0]; - trackParams[2] = trackParams[0]; - trackParams[1].MinTrackLength = 6; - trackParams[2].MinTrackLength = 4; - LOG(info) << "Initializing tracker in misaligned sync. phase reconstruction with " << trackParams.size() << " passes"; - - } else if (mMode == "sync") { - trackParams.resize(1); - LOG(info) << "Initializing tracker in sync. phase reconstruction with " << trackParams.size() << " passes"; - } else if (mMode == "cosmics") { - mCosmicsProcessing = true; - mRunVertexer = false; - trackParams.resize(1); - trackParams[0].MinTrackLength = 4; - trackParams[0].CellDeltaTanLambdaSigma *= 10; - trackParams[0].PhiBins = 4; - trackParams[0].ZBins = 16; - trackParams[0].PVres = 1.e5f; - trackParams[0].MaxChi2ClusterAttachment = 60.; - trackParams[0].MaxChi2NDF = 40.; - trackParams[0].TrackletsPerClusterLimit = 100.; - trackParams[0].CellsPerClusterLimit = 100.; - LOG(info) << "Initializing tracker in reconstruction for cosmics with " << trackParams.size() << " passes"; - - } else { - throw std::runtime_error(fmt::format("Unsupported ITS tracking mode {:s} ", mMode)); - } - - for (auto& params : trackParams) { - params.CorrType = o2::base::PropagatorImpl::MatCorrType::USEMatCorrLUT; - for (int iLayer{0}; iLayer < params.NLayers - 4; ++iLayer) { // initialise ITS3 radii and lengths - params.LayerZ[iLayer] = constants::segment::lengthSensitive; - params.LayerRadii[iLayer] = constants::radii[iLayer]; - } - } + mITS3TrackingInterface.setTraitsFromProvider(mChainITS->GetITSVertexerTraits(), + mChainITS->GetITSTrackerTraits(), + mChainITS->GetITSTimeframe()); + mITS3TrackingInterface.initialise(); +} - mTracker->setParameters(trackParams); +void TrackerDPL::stop() +{ + LOGF(info, "CPU Reconstruction total timing: Cpu: %.3e Real: %.3e s in %d slots", mTimer.CpuTime(), mTimer.RealTime(), mTimer.Counter() - 1); } void TrackerDPL::run(ProcessingContext& pc) { + auto cput = mTimer.CpuTime(); + auto realt = mTimer.RealTime(); mTimer.Start(false); - updateTimeDependentParams(pc); - auto compClusters = pc.inputs().get>("compClusters"); - gsl::span patterns = pc.inputs().get>("patterns"); - gsl::span physTriggers; - std::vector fromTRD; - if (mUseTriggers == 2) { // use TRD triggers - o2::InteractionRecord ir{0, pc.services().get().firstTForbit}; - auto trdTriggers = pc.inputs().get>("phystrig"); - for (const auto& trig : trdTriggers) { - if (trig.getBCData() >= ir && trig.getNumberOfTracklets()) { - ir = trig.getBCData(); - fromTRD.emplace_back(o2::itsmft::PhysTrigger{ir, 0}); - } - } - physTriggers = gsl::span(fromTRD.data(), fromTRD.size()); - } else if (mUseTriggers == 1) { // use Phys triggers from ITS stream - physTriggers = pc.inputs().get>("phystrig"); - } - - // code further down does assignment to the rofs and the altered object is used for output - // we therefore need a copy of the vector rather than an object created directly on the input data, - // the output vector however is created directly inside the message memory thus avoiding copy by - // snapshot - auto orig = o2::header::gDataOriginITS; - auto rofsinput = pc.inputs().get>("ROframes"); - auto& rofs = pc.outputs().make>(Output{orig, "ITSTrackROF", 0}, rofsinput.begin(), rofsinput.end()); - - auto& irFrames = pc.outputs().make>(Output{orig, "IRFRAMES", 0}); - - const auto& alpParams = o2::itsmft::DPLAlpideParam::Instance(); // RS: this should come from CCDB - int nBCPerTF = alpParams.roFrameLengthInBC; - - LOG(info) << "ITS3Tracker pulled " << compClusters.size() << " clusters, " << rofs.size() << " RO frames"; - - const dataformats::MCTruthContainer* labels = nullptr; - gsl::span mc2rofs; - if (mIsMC) { - labels = pc.inputs().get*>("labels").release(); - // get the array as read-only span, a snapshot is send forward - mc2rofs = pc.inputs().get>("MC2ROframes"); - LOG(info) << labels->getIndexedSize() << " MC label objects , in " << mc2rofs.size() << " MC events"; - } - - auto& allClusIdx = pc.outputs().make>(Output{orig, "TRACKCLSID", 0}); - std::vector trackLabels; - std::vector verticesLabels; - auto& allTracks = pc.outputs().make>(Output{orig, "TRACKS", 0}); - std::vector allTrackLabels; - std::vector allVerticesLabels; - std::vector allVerticesPurities; - - auto& vertROFvec = pc.outputs().make>(Output{orig, "VERTICESROF", 0}); - auto& vertices = pc.outputs().make>(Output{orig, "VERTICES", 0}); - - TimeFrame* timeFrame = mChainITS->GetITSTimeframe(); - timeFrame->resizeVectors(7); - mTracker->adoptTimeFrame(*timeFrame); - - mTracker->setBz(o2::base::Propagator::Instance()->getNominalBz()); - mVertexer->adoptTimeFrame(*timeFrame); - auto pattIt = patterns.begin(); - - const gsl::span rofspan(rofs); - ioutils::loadROFrameDataITS3(timeFrame, rofspan, compClusters, pattIt, mDict, labels); - std::vector savedROF; - auto logger = [&](const std::string& s) { LOG(info) << s; }; - auto errorLogger = [&](const std::string& s) { LOG(error) << s; }; - - // o2::its3::FastMultEst multEst; // mult estimator - std::vector processingMask(rofs.size(), true); // Override mult estimator - int cutVertexMult{0}; - // int cutRandomMult = int(rofs.size()) - multEst.selectROFs(rofs, compClusters, physTriggers, processingMask); - timeFrame->setMultiplicityCutMask(processingMask); - float vertexerElapsedTime{0.f}; - if (mRunVertexer) { - // Run seeding vertexer - vertexerElapsedTime = mVertexer->clustersToVertices(logger); - } - const auto& multEstConf = FastMultEstConfig::Instance(); // parameters for mult estimation and cuts - gsl::span> vMCRecInfo; - for (size_t iRof{0}; iRof < rofspan.size(); ++iRof) { - std::vector vtxVecLoc; - auto& vtxROF = vertROFvec.emplace_back(rofspan[iRof]); - vtxROF.setFirstEntry(vertices.size()); - if (mRunVertexer) { - auto vtxSpan = timeFrame->getPrimaryVertices(iRof); - if (mIsMC) { - vMCRecInfo = timeFrame->getPrimaryVerticesMCRecInfo(iRof); - } - vtxROF.setNEntries(vtxSpan.size()); - bool selROF = vtxSpan.size() == 0; - for (size_t iV{0}; iV < vtxSpan.size(); ++iV) { - auto& v = vtxSpan[iV]; - if (multEstConf.isVtxMultCutRequested() && !multEstConf.isPassingVtxMultCut(v.getNContributors())) { - continue; // skip vertex of unwanted multiplicity - } - selROF = true; - vertices.push_back(v); - if (mIsMC) { - allVerticesLabels.push_back(vMCRecInfo[iV].first); - allVerticesPurities.push_back(vMCRecInfo[iV].second); - } - } - if (processingMask[iRof] && !selROF) { // passed selection in clusters and not in vertex multiplicity - // LOG(debug) << fmt::format("ROF {} rejected by the vertex multiplicity selection [{},{}]", - // iRof, - // multEstConf.cutMultVtxLow, - // multEstConf.cutMultVtxHigh); - processingMask[iRof] = selROF; - cutVertexMult++; - } - } else { // cosmics - vtxVecLoc.emplace_back(Vertex()); - vtxVecLoc.back().setNContributors(1); - vtxROF.setNEntries(vtxVecLoc.size()); - for (auto& v : vtxVecLoc) { - vertices.push_back(v); - } - timeFrame->addPrimaryVertices(vtxVecLoc, iRof, 0); - } - } - // LOG(info) << fmt::format(" - rejected {}/{} ROFs: random/mult.sel:{} (seed {}), vtx.sel:{}", cutRandomMult + cutVertexMult, rofspan.size(), cutRandomMult, multEst.lastRandomSeed, cutVertexMult); - LOG(info) << fmt::format(" - Vertex seeding total elapsed time: {} ms for {} vertices found in {} ROFs", vertexerElapsedTime, timeFrame->getPrimaryVerticesNum(), rofspan.size()); - LOG(info) << fmt::format(" - Beam position computed for the TF: {}, {}", timeFrame->getBeamX(), timeFrame->getBeamY()); - - if (mCosmicsProcessing && compClusters.size() > 1500 * rofspan.size()) { - LOG(error) << "Cosmics processing was requested with an average detector occupancy exceeding 1.e-7, skipping TF processing."; - } else { - - timeFrame->setMultiplicityCutMask(processingMask); - // Run CA tracker - mTracker->clustersToTracks(logger, errorLogger); - if (timeFrame->hasBogusClusters() != 0) { - LOG(warning) << fmt::format(" - The processed timeframe had {} clusters with wild z coordinates, check the dictionaries", timeFrame->hasBogusClusters()); - } - - for (unsigned int iROF{0}; iROF < rofs.size(); ++iROF) { - auto& rof{rofs[iROF]}; - auto& tracks = timeFrame->getTracks(iROF); - trackLabels = timeFrame->getTracksLabel(iROF); - auto number{tracks.size()}; - auto first{allTracks.size()}; - rof.setFirstEntry(first); - rof.setNEntries(number); - if (processingMask[iROF]) { - irFrames.emplace_back(rof.getBCData(), rof.getBCData() + nBCPerTF - 1).info = tracks.size(); - } - - std::copy(trackLabels.begin(), trackLabels.end(), std::back_inserter(allTrackLabels)); - // Some conversions that needs to be moved in the tracker internals - for (unsigned int iTrk{0}; iTrk < tracks.size(); ++iTrk) { - auto& trc{tracks[iTrk]}; - trc.setFirstClusterEntry(allClusIdx.size()); // before adding tracks, create final cluster indices - int nclf = 0, ncl = (int)allClusIdx.size(); - for (int ic = TrackITSExt::MaxClusters; (ic--) != 0;) { // track internally keeps in->out cluster indices, but we want to store the references as out->in!!! - auto clid = trc.getClusterIndex(ic); - if (clid >= 0) { - trc.setClusterSize(ic, timeFrame->getClusterSize(clid)); - allClusIdx.push_back(clid); - nclf++; - } - } - assert(ncl == nclf); - allTracks.emplace_back(trc); - } - } - - LOGP(info, "ITS3Tracker pushed {} tracks and {} vertices", allTracks.size(), vertices.size()); - if (mIsMC) { - LOGP(info, "ITS3Tracker pushed {} track labels", allTrackLabels.size()); - LOGP(info, "ITS3Tracker pushed {} vertex labels", allVerticesLabels.size()); - - pc.outputs().snapshot(Output{orig, "TRACKSMCTR", 0}, allTrackLabels); - pc.outputs().snapshot(Output{orig, "VERTICESMCTR", 0}, allVerticesLabels); - pc.outputs().snapshot(Output{orig, "VERTICESMCPUR", 0}, allVerticesPurities); - pc.outputs().snapshot(Output{orig, "ITSTrackMC2ROF", 0}, mc2rofs); - } - } + mITS3TrackingInterface.updateTimeDependentParams(pc); + mITS3TrackingInterface.run(pc); mTimer.Stop(); + LOGP(info, "CPU Reconstruction time for this TF {} s (cpu), {} s (wall)", mTimer.CpuTime() - cput, mTimer.RealTime() - realt); } -///_______________________________________ -void TrackerDPL::updateTimeDependentParams(ProcessingContext& pc) -{ - o2::base::GRPGeomHelper::instance().checkUpdates(pc); - static bool initOnceDone = false; - if (!initOnceDone) { // this params need to be queried only once - initOnceDone = true; - pc.inputs().get("cldict"); // just to trigger the finaliseCCDB - pc.inputs().get*>("alppar"); - - // Check if lightweight geometry was requested, otherwise full geometry is loaded - if (pc.inputs().getPos("itsTGeo") >= 0) { - pc.inputs().get("itsTGeo"); - } - o2::its::GeometryTGeo* geom = o2::its::GeometryTGeo::Instance(); - geom->fillMatrixCache(o2::math_utils::bit2Mask(o2::math_utils::TransformType::T2L, o2::math_utils::TransformType::T2GRot, o2::math_utils::TransformType::T2G)); - mVertexer->getGlobalConfiguration(); - mTracker->getGlobalConfiguration(); - } -} - -///_______________________________________ void TrackerDPL::finaliseCCDB(ConcreteDataMatcher& matcher, void* obj) { - if (o2::base::GRPGeomHelper::instance().finaliseCCDB(matcher, obj)) { - return; - } - if (matcher == ConcreteDataMatcher("IT3", "CLUSDICT", 0)) { - LOG(info) << "cluster dictionary updated"; - setClusterDictionary((o2::its3::TopologyDictionary*)obj); - return; - } - if (matcher == ConcreteDataMatcher("ITS", "GEOMTGEO", 0)) { - LOG(info) << "IT3 GeometryTGeo loaded from ccdb"; - o2::its::GeometryTGeo::adopt((o2::its::GeometryTGeo*)obj); - return; - } - // Note: strictly speaking, for Configurable params we don't need finaliseCCDB check, the singletons are updated at the CCDB fetcher level - if (matcher == ConcreteDataMatcher("ITS", "ALPIDEPARAM", 0)) { - LOG(info) << "Alpide param updated"; - const auto& par = o2::itsmft::DPLAlpideParam::Instance(); - par.printKeyValues(); - return; - } + mITS3TrackingInterface.finaliseCCDB(matcher, obj); } void TrackerDPL::endOfStream(EndOfStreamContext& ec) { - LOGF(info, "ITS3 CA-Tracker total timing: Cpu: %.3e Real: %.3e s in %d slots", - mTimer.CpuTime(), mTimer.RealTime(), mTimer.Counter() - 1); + LOGF(info, "ITS3 CA-Tracker total timing: Cpu: %.3e Real: %.3e s in %d slots", mTimer.CpuTime(), mTimer.RealTime(), mTimer.Counter() - 1); } -DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, const int trgType, const std::string& trModeS, o2::gpu::GPUDataTypes::DeviceType dType) +DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, int trgType, const std::string& trModeS, const bool overrBeamEst, o2::gpu::GPUDataTypes::DeviceType dType) { std::vector inputs; inputs.emplace_back("compClusters", "ITS", "COMPCLUSTERS", 0, Lifetime::Timeframe); @@ -399,7 +110,7 @@ DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, const int trgType, co inputs.emplace_back("cldict", "IT3", "CLUSDICT", 0, Lifetime::Condition, ccdbParamSpec("IT3/Calib/ClusterDictionary")); inputs.emplace_back("alppar", "ITS", "ALPIDEPARAM", 0, Lifetime::Condition, ccdbParamSpec("ITS/Config/AlpideParam")); auto ggRequest = std::make_shared(false, // orbitResetTime - false, // GRPECS + true, // GRPECS false, // GRPLHCIF true, // GRPMagField true, // askMatLUT @@ -410,6 +121,9 @@ DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, const int trgType, co if (!useGeom) { // load light-weight geometry inputs.emplace_back("itsTGeo", "ITS", "GEOMTGEO", 0, Lifetime::Condition, ccdbParamSpec("ITS/Config/Geometry")); } + if (overrBeamEst) { + inputs.emplace_back("meanvtx", "GLO", "MEANVERTEX", 0, Lifetime::Condition, ccdbParamSpec("GLO/Calib/MeanVertex", {}, 1)); + } std::vector outputs; outputs.emplace_back("ITS", "TRACKS", 0, Lifetime::Timeframe); @@ -420,8 +134,8 @@ DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, const int trgType, co outputs.emplace_back("ITS", "IRFRAMES", 0, Lifetime::Timeframe); if (useMC) { - inputs.emplace_back("labels", "ITS", "CLUSTERSMCTR", 0, Lifetime::Timeframe); - inputs.emplace_back("MC2ROframes", "ITS", "CLUSTERSMC2ROF", 0, Lifetime::Timeframe); + inputs.emplace_back("itsmclabels", "ITS", "CLUSTERSMCTR", 0, Lifetime::Timeframe); + inputs.emplace_back("ITSMC2ROframes", "ITS", "CLUSTERSMC2ROF", 0, Lifetime::Timeframe); outputs.emplace_back("ITS", "VERTICESMCTR", 0, Lifetime::Timeframe); outputs.emplace_back("ITS", "VERTICESMCPUR", 0, Lifetime::Timeframe); outputs.emplace_back("ITS", "TRACKSMCTR", 0, Lifetime::Timeframe); @@ -432,7 +146,10 @@ DataProcessorSpec getTrackerSpec(bool useMC, bool useGeom, const int trgType, co "its3-tracker", inputs, outputs, - AlgorithmSpec{adaptFromTask(ggRequest, useMC, trgType, trModeS, dType)}, + AlgorithmSpec{adaptFromTask(ggRequest, useMC, trgType, + trModeS == "sync" ? o2::its::TrackingMode::Sync : trModeS == "async" ? o2::its::TrackingMode::Async + : o2::its::TrackingMode::Cosmics, + overrBeamEst, dType)}, Options{}}; } diff --git a/Framework/Foundation/3rdparty/include/Framework/SHA1.h b/Framework/Foundation/3rdparty/include/Framework/SHA1.h index cdf6bf18b17d9..3fd2a037302fb 100644 --- a/Framework/Foundation/3rdparty/include/Framework/SHA1.h +++ b/Framework/Foundation/3rdparty/include/Framework/SHA1.h @@ -263,7 +263,7 @@ static void SHA1Final( memset(&finalcount, '\0', sizeof(finalcount)); } -void SHA1( +static void SHA1( char* hash_out, const char* str, unsigned int len) diff --git a/GPU/Common/GPUCommonDefAPI.h b/GPU/Common/GPUCommonDefAPI.h index 52691e094590a..5a152de8a2216 100644 --- a/GPU/Common/GPUCommonDefAPI.h +++ b/GPU/Common/GPUCommonDefAPI.h @@ -152,7 +152,11 @@ #define GPUbarrierWarp() #define GPUAtomic(type) type #elif defined(__CUDACC__) //Defines for CUDA - #define GPUd() __device__ + #ifndef GPUCA_GPUCODE_DEVICE + #define GPUd() __device__ inline // FIXME: DR: Workaround: mark device function as inline such that nvcc does not create bogus host symbols + #else + #define GPUd() __device__ + #endif #define GPUdDefault() #define GPUhdDefault() #define GPUdi() __device__ inline diff --git a/GPU/GPUTracking/Base/GPUConstantMem.h b/GPU/GPUTracking/Base/GPUConstantMem.h index 17e6268764606..7445233f02e8c 100644 --- a/GPU/GPUTracking/Base/GPUConstantMem.h +++ b/GPU/GPUTracking/Base/GPUConstantMem.h @@ -122,14 +122,16 @@ GPUdi() auto& GPUConstantMem::getTRDTracker<1>() #ifdef GPUCA_NOCOMPAT union GPUConstantMemCopyable { - GPUConstantMemCopyable() {} // NOLINT: We want an empty constructor, not a default one - ~GPUConstantMemCopyable() {} // NOLINT: We want an empty destructor, not a default one - GPUConstantMemCopyable(const GPUConstantMemCopyable& o) +#if !defined(__OPENCL__) || defined(__OPENCL_HOST__) + GPUh() GPUConstantMemCopyable() {} // NOLINT: We want an empty constructor, not a default one + GPUh() ~GPUConstantMemCopyable() {} // NOLINT: We want an empty destructor, not a default one + GPUh() GPUConstantMemCopyable(const GPUConstantMemCopyable& o) { for (unsigned int k = 0; k < sizeof(GPUConstantMem) / sizeof(int); k++) { ((int*)&v)[k] = ((int*)&o.v)[k]; } } +#endif GPUConstantMem v; }; #endif diff --git a/GPU/GPUTracking/CMakeLists.txt b/GPU/GPUTracking/CMakeLists.txt index 23fb231634a13..c4ad13f2a888e 100644 --- a/GPU/GPUTracking/CMakeLists.txt +++ b/GPU/GPUTracking/CMakeLists.txt @@ -295,6 +295,8 @@ string(REPLACE ".cxx" ".h" HDRS_TMP "${SRCS_NO_CINT}") set(HDRS_INSTALL ${HDRS_INSTALL} ${HDRS_TMP}) string(REPLACE ".cxx" ".h" HDRS_TMP "${SRCS_DATATYPES}") set(HDRS_CINT_DATATYPES ${HDRS_CINT_DATATYPES} ${HDRS_TMP}) +string(REPLACE ".cxx" ".h" HDRS_TMP "${SRCS_DATATYPE_HEADERS}") +set(HDRS_INSTALL ${HDRS_INSTALL} ${HDRS_TMP}) unset(HDRS_TMP) # Main CMake part for O2 diff --git a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx index 3fcd4853573e0..8d68509e42473 100644 --- a/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx +++ b/GPU/GPUTracking/DataCompression/GPUTPCCompressionKernels.cxx @@ -33,7 +33,7 @@ GPUdii() void GPUTPCCompressionKernels::Thread -GPUd() GPUTRDTrack_t::GPUTRDTrack_t() -{ - // default constructor - initialize(); -} - -template -GPUd() void GPUTRDTrack_t::initialize() -{ - // set all members to their default values (needed since in-class initialization not possible with AliRoot) - mChi2 = 0.f; - mSignal = -1.f; - mRefGlobalTrackId = 0; - mCollisionId = -1; - mFlags = 0; - mIsCrossingNeighbor = 0; - for (int i = 0; i < kNLayers; ++i) { - mAttachedTracklets[i] = -1; - } -} - -#ifdef GPUCA_ALIROOT_LIB -#include "AliHLTExternalTrackParam.h" -#include "GPUTRDTrackData.h" - -template -GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const AliHLTExternalTrackParam& t) : T(t) -{ - initialize(); -} - -template -GPUd() void GPUTRDTrack_t::ConvertTo(GPUTRDTrackDataRecord& t) const -{ - //------------------------------------------------------------------ - // convert to GPU structure - //------------------------------------------------------------------ - t.mAlpha = T::getAlpha(); - t.fX = T::getX(); - t.fY = T::getY(); - t.fZ = T::getZ(); - t.fq1Pt = T::getQ2Pt(); - t.mSinPhi = T::getSnp(); - t.fTgl = T::getTgl(); - for (int i = 0; i < 15; i++) { - t.fC[i] = T::getCov()[i]; - } - t.fTPCTrackID = getRefGlobalTrackIdRaw(); - for (int i = 0; i < kNLayers; i++) { - t.fAttachedTracklets[i] = getTrackletIndex(i); - } -} - -template -GPUd() void GPUTRDTrack_t::ConvertFrom(const GPUTRDTrackDataRecord& t) -{ - //------------------------------------------------------------------ - // convert from GPU structure - //------------------------------------------------------------------ - T::set(t.fX, t.mAlpha, &(t.fY), t.fC); - setRefGlobalTrackIdRaw(t.fTPCTrackID); - mChi2 = 0.f; - mSignal = -1.f; - mFlags = 0; - mIsCrossingNeighbor = 0; - mCollisionId = -1; - for (int iLayer = 0; iLayer < kNLayers; iLayer++) { - mAttachedTracklets[iLayer] = t.fAttachedTracklets[iLayer]; - } -} - -#endif - -#if defined(GPUCA_HAVE_O2HEADERS) -#include "ReconstructionDataFormats/TrackTPCITS.h" -#include "DataFormatsTPC/TrackTPC.h" - -template -GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const o2::dataformats::TrackTPCITS& t) : T(t) -{ - initialize(); -} - -template -GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const o2::tpc::TrackTPC& t) : T(t) -{ - initialize(); -} - -#endif - -template -GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const GPUTRDTrack_t& t) - : T(t), mChi2(t.mChi2), mSignal(t.mSignal), mRefGlobalTrackId(t.mRefGlobalTrackId), mCollisionId(t.mCollisionId), mFlags(t.mFlags), mIsCrossingNeighbor(t.mIsCrossingNeighbor) -{ - // copy constructor - for (int i = 0; i < kNLayers; ++i) { - mAttachedTracklets[i] = t.mAttachedTracklets[i]; - } -} - -template -GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const T& t) : T(t) -{ - // copy constructor from anything - initialize(); -} - -template -GPUd() GPUTRDTrack_t& GPUTRDTrack_t::operator=(const GPUTRDTrack_t& t) -{ - // assignment operator - if (&t == this) { - return *this; - } - *(T*)this = t; - mChi2 = t.mChi2; - mSignal = t.mSignal; - mRefGlobalTrackId = t.mRefGlobalTrackId; - mCollisionId = t.mCollisionId; - mFlags = t.mFlags; - mIsCrossingNeighbor = t.mIsCrossingNeighbor; - for (int i = 0; i < kNLayers; ++i) { - mAttachedTracklets[i] = t.mAttachedTracklets[i]; - } - return *this; -} - -template -GPUd() int GPUTRDTrack_t::getNlayersFindable() const -{ - // returns number of layers in which the track is in active area of TRD - int retVal = 0; - for (int iLy = 0; iLy < kNLayers; iLy++) { - if ((mFlags >> iLy) & 0x1) { - ++retVal; - } - } - return retVal; -} - -template -GPUd() int GPUTRDTrack_t::getNmissingConsecLayers(int iLayer) const -{ - // returns number of consecutive layers in which the track was - // inside the deadzone up to (and including) the given layer - int retVal = 0; - while (!getIsFindable(iLayer)) { - ++retVal; - --iLayer; - if (iLayer < 0) { - break; - } - } - return retVal; -} - -#if !defined(GPUCA_GPUCODE) && !defined(GPU_TRD_TRACK_O2) +#if !defined(GPUCA_GPUCODE) namespace GPUCA_NAMESPACE { namespace gpu diff --git a/GPU/GPUTracking/DataTypes/GPUTRDTrack.inc b/GPU/GPUTracking/DataTypes/GPUTRDTrack.inc new file mode 100644 index 0000000000000..bba48e6cfc7a1 --- /dev/null +++ b/GPU/GPUTracking/DataTypes/GPUTRDTrack.inc @@ -0,0 +1,178 @@ +// 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 GPUTRDTrack.inc +/// \author Ole Schmidt, Sergey Gorbunov + +#ifndef GPUTRDTRACK_INC_H +#define GPUTRDTRACK_INC_H + +#include "GPUTRDTrack.h" + +template +GPUd() GPUTRDTrack_t::GPUTRDTrack_t() +{ + // default constructor + initialize(); +} + +template +GPUd() void GPUTRDTrack_t::initialize() +{ + // set all members to their default values (needed since in-class initialization not possible with AliRoot) + mChi2 = 0.f; + mSignal = -1.f; + mRefGlobalTrackId = 0; + mCollisionId = -1; + mFlags = 0; + mIsCrossingNeighbor = 0; + for (int i = 0; i < kNLayers; ++i) { + mAttachedTracklets[i] = -1; + } +} + +#ifdef GPUCA_ALIROOT_LIB +#include "AliHLTExternalTrackParam.h" +#include "GPUTRDTrackData.h" + +template +GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const AliHLTExternalTrackParam& t) : T(t) +{ + initialize(); +} + +template +GPUd() void GPUTRDTrack_t::ConvertTo(GPUTRDTrackDataRecord& t) const +{ + //------------------------------------------------------------------ + // convert to GPU structure + //------------------------------------------------------------------ + t.mAlpha = T::getAlpha(); + t.fX = T::getX(); + t.fY = T::getY(); + t.fZ = T::getZ(); + t.fq1Pt = T::getQ2Pt(); + t.mSinPhi = T::getSnp(); + t.fTgl = T::getTgl(); + for (int i = 0; i < 15; i++) { + t.fC[i] = T::getCov()[i]; + } + t.fTPCTrackID = getRefGlobalTrackIdRaw(); + for (int i = 0; i < kNLayers; i++) { + t.fAttachedTracklets[i] = getTrackletIndex(i); + } +} + +template +GPUd() void GPUTRDTrack_t::ConvertFrom(const GPUTRDTrackDataRecord& t) +{ + //------------------------------------------------------------------ + // convert from GPU structure + //------------------------------------------------------------------ + T::set(t.fX, t.mAlpha, &(t.fY), t.fC); + setRefGlobalTrackIdRaw(t.fTPCTrackID); + mChi2 = 0.f; + mSignal = -1.f; + mFlags = 0; + mIsCrossingNeighbor = 0; + mCollisionId = -1; + for (int iLayer = 0; iLayer < kNLayers; iLayer++) { + mAttachedTracklets[iLayer] = t.fAttachedTracklets[iLayer]; + } +} + +#endif + +#if defined(GPUCA_HAVE_O2HEADERS) +#include "ReconstructionDataFormats/TrackTPCITS.h" +#include "DataFormatsTPC/TrackTPC.h" + +template +GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const o2::dataformats::TrackTPCITS& t) : T(t) +{ + initialize(); +} + +template +GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const o2::tpc::TrackTPC& t) : T(t) +{ + initialize(); +} + +#endif + +template +GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const GPUTRDTrack_t& t) + : T(t), mChi2(t.mChi2), mSignal(t.mSignal), mRefGlobalTrackId(t.mRefGlobalTrackId), mCollisionId(t.mCollisionId), mFlags(t.mFlags), mIsCrossingNeighbor(t.mIsCrossingNeighbor) +{ + // copy constructor + for (int i = 0; i < kNLayers; ++i) { + mAttachedTracklets[i] = t.mAttachedTracklets[i]; + } +} + +template +GPUd() GPUTRDTrack_t::GPUTRDTrack_t(const T& t) : T(t) +{ + // copy constructor from anything + initialize(); +} + +template +GPUd() GPUTRDTrack_t& GPUTRDTrack_t::operator=(const GPUTRDTrack_t& t) +{ + // assignment operator + if (&t == this) { + return *this; + } + *(T*)this = t; + mChi2 = t.mChi2; + mSignal = t.mSignal; + mRefGlobalTrackId = t.mRefGlobalTrackId; + mCollisionId = t.mCollisionId; + mFlags = t.mFlags; + mIsCrossingNeighbor = t.mIsCrossingNeighbor; + for (int i = 0; i < kNLayers; ++i) { + mAttachedTracklets[i] = t.mAttachedTracklets[i]; + } + return *this; +} + +template +GPUd() int GPUTRDTrack_t::getNlayersFindable() const +{ + // returns number of layers in which the track is in active area of TRD + int retVal = 0; + for (int iLy = 0; iLy < kNLayers; iLy++) { + if ((mFlags >> iLy) & 0x1) { + ++retVal; + } + } + return retVal; +} + +template +GPUd() int GPUTRDTrack_t::getNmissingConsecLayers(int iLayer) const +{ + // returns number of consecutive layers in which the track was + // inside the deadzone up to (and including) the given layer + int retVal = 0; + while (!getIsFindable(iLayer)) { + ++retVal; + --iLayer; + if (iLayer < 0) { + break; + } + } + return retVal; +} + +#endif // GPUTRDTRACK_INC_H diff --git a/GPU/GPUTracking/DataTypes/GPUTRDTrackO2.cxx b/GPU/GPUTracking/DataTypes/GPUTRDTrackO2.cxx index 6141a8bd9d958..d2404f9d3b74b 100644 --- a/GPU/GPUTracking/DataTypes/GPUTRDTrackO2.cxx +++ b/GPU/GPUTracking/DataTypes/GPUTRDTrackO2.cxx @@ -12,8 +12,10 @@ /// \file GPUTRDTrackO2.cxx /// \author David Rohr -#define GPU_TRD_TRACK_O2 -#include "GPUTRDTrack.cxx" +#include "GPUTRDTrackO2.h" +using namespace GPUCA_NAMESPACE::gpu; + +#include "GPUTRDTrack.inc" #include "ReconstructionDataFormats/GlobalTrackID.h" #include "ReconstructionDataFormats/Track.h" diff --git a/GPU/GPUTracking/DataTypes/GPUTRDTrackO2.h b/GPU/GPUTracking/DataTypes/GPUTRDTrackO2.h new file mode 100644 index 0000000000000..d7f811b004e02 --- /dev/null +++ b/GPU/GPUTracking/DataTypes/GPUTRDTrackO2.h @@ -0,0 +1,20 @@ +// 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 GPUTRDTrackO2.h +/// \author David Rohr + +#ifndef GPUTRDTRACKO2_H +#define GPUTRDTRACKO2_H + +#include "GPUTRDTrack.h" + +#endif // GPUTRDTRACKO2_H diff --git a/GPU/GPUTracking/Definitions/GPUSettingsList.h b/GPU/GPUTracking/Definitions/GPUSettingsList.h index 06bf4b5817f3e..610cfe1627b2f 100644 --- a/GPU/GPUTracking/Definitions/GPUSettingsList.h +++ b/GPU/GPUTracking/Definitions/GPUSettingsList.h @@ -128,7 +128,7 @@ AddOptionRTC(cfInnerThreshold, unsigned char, 0, "", 0, "Cluster Finder extends AddOptionRTC(cfMinSplitNum, unsigned char, 1, "", 0, "Minimum number of split charges in a cluster for the cluster to be marked as split") AddOptionRTC(cfNoiseSuppressionEpsilon, unsigned char, 10, "", 0, "Cluster Finder: Difference between peak and charge for the charge to count as a minima during noise suppression") AddOptionRTC(cfNoiseSuppressionEpsilonRelative, unsigned char, 76, "", 0, "Cluster Finder: Difference between peak and charge for the charge to count as a minima during noise suppression, relative as fraction of 255") -AddOptionRTC(nWays, signed char, 3, "", 0, "Do N fit passes in final fit of merger") +AddOptionRTC(nWays, unsigned char, 3, "", 0, "Do N fit passes in final fit of merger") AddOptionRTC(nWaysOuter, signed char, 0, "", 0, "Store outer param") AddOptionRTC(trackFitRejectMode, signed char, 5, "", 0, "0: no limit on rejection or missed hits, >0: break after n rejected hits, <0: reject at max -n hits") AddOptionRTC(dEdxTruncLow, unsigned char, 2, "", 0, "Low truncation threshold, fraction of 128") diff --git a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx index da167ca9731bf..0737456df2302 100644 --- a/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx +++ b/GPU/GPUTracking/Global/GPUChainTrackingDebugAndProfiling.cxx @@ -244,7 +244,7 @@ void GPUChainTracking::PrintOutputStat() } snprintf(trdText, 1024, " - TRD Tracker reconstructed %d tracks (%d tracklets)", nTRDTracks, nTRDTracklets); } - GPUInfo("Output Tracks: %d (%d / %d / %d / %d clusters (fitted / attached / adjacent / total))%s", nTracks, nAttachedClustersFitted, nAttachedClusters, nAdjacentClusters, nCls, trdText); + GPUInfo("Output Tracks: %d (%d / %d / %d / %d clusters (fitted / attached / adjacent / total) - %s format)%s", nTracks, nAttachedClustersFitted, nAttachedClusters, nAdjacentClusters, nCls, ProcessingSettings().createO2Output > 1 ? "O2" : "GPU", trdText); } void GPUChainTracking::SanityCheck() diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx index 9f7e359d1630b..671ec9a7794a0 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMerger.cxx @@ -1007,9 +1007,11 @@ GPUd() void GPUTPCGMMerger::MergeBorderTracks(int nBlocks, int nThreads, int iBl MergeBorderTracks(nBlocks, nThreads, iBlock, iThread, iSlice, b1, n1, jSlice, b2, n2, mergeMode); } -template GPUd() void GPUTPCGMMerger::MergeBorderTracks<0>(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice, char withinSlice, char mergeMode); -template GPUd() void GPUTPCGMMerger::MergeBorderTracks<1>(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice, char withinSlice, char mergeMode); -template GPUd() void GPUTPCGMMerger::MergeBorderTracks<2>(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice, char withinSlice, char mergeMode); +#if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE) // FIXME: DR: WORKAROUND to avoid CUDA bug creating host symbols for device code. +template GPUdni() void GPUTPCGMMerger::MergeBorderTracks<0>(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice, char withinSlice, char mergeMode); +template GPUdni() void GPUTPCGMMerger::MergeBorderTracks<1>(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice, char withinSlice, char mergeMode); +template GPUdni() void GPUTPCGMMerger::MergeBorderTracks<2>(int nBlocks, int nThreads, int iBlock, int iThread, int iSlice, char withinSlice, char mergeMode); +#endif GPUd() void GPUTPCGMMerger::MergeWithinSlicesPrepare(int nBlocks, int nThreads, int iBlock, int iThread) { @@ -2087,7 +2089,7 @@ GPUd() void GPUTPCGMMerger::Finalize1(int nBlocks, int nThreads, int iBlock, int if (!trk.OK() || trk.NClusters() == 0) { continue; } - char goodLeg = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].leg; + unsigned char goodLeg = mClusters[trk.FirstClusterRef() + trk.NClusters() - 1].leg; for (unsigned int j = 0; j < trk.NClusters(); j++) { int id = mClusters[trk.FirstClusterRef() + j].num; unsigned int weight = mTrackOrderAttach[i] | attachAttached; diff --git a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx index 1a9f9c536565e..e7f9e7dbf3385 100644 --- a/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx +++ b/GPU/GPUTracking/Merger/GPUTPCGMMergerGPU.cxx @@ -120,9 +120,11 @@ GPUdii() void GPUTPCGMMergerMergeBorders::Thread(int nBlocks, int nThreads, int { merger.MergeBorderTracks(nBlocks, nThreads, iBlock, iThread, args...); } -template GPUd() void GPUTPCGMMergerMergeBorders::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int iSlice, char withinSlice, char mergeMode); -template GPUd() void GPUTPCGMMergerMergeBorders::Thread<2>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int iSlice, char withinSlice, char mergeMode); -template GPUd() void GPUTPCGMMergerMergeBorders::Thread<3>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, gputpcgmmergertypes::GPUTPCGMBorderRange* range, int N, int cmpMax); +#if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE) // FIXME: DR: WORKAROUND to avoid CUDA bug creating host symbols for device code. +template GPUdni() void GPUTPCGMMergerMergeBorders::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int iSlice, char withinSlice, char mergeMode); +template GPUdni() void GPUTPCGMMergerMergeBorders::Thread<2>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int iSlice, char withinSlice, char mergeMode); +template GPUdni() void GPUTPCGMMergerMergeBorders::Thread<3>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, gputpcgmmergertypes::GPUTPCGMBorderRange* range, int N, int cmpMax); +#endif template <> GPUdii() void GPUTPCGMMergerMergeBorders::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() merger, int iSlice, char withinSlice, char mergeMode) { diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx index 5b42ad462cc84..a78565f020ab1 100644 --- a/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx +++ b/GPU/GPUTracking/Refit/GPUTrackingRefit.cxx @@ -421,12 +421,14 @@ GPUd() int GPUTrackingRefit::RefitTrack(T& trkX, bool outward, bool resetCov) return nFitted; } -template GPUd() int GPUTrackingRefit::RefitTrack(GPUTPCGMMergedTrack& trk, bool outward, bool resetCov); -template GPUd() int GPUTrackingRefit::RefitTrack(GPUTPCGMMergedTrack& trk, bool outward, bool resetCov); -template GPUd() int GPUTrackingRefit::RefitTrack(TrackTPC& trk, bool outward, bool resetCov); -template GPUd() int GPUTrackingRefit::RefitTrack(TrackTPC& trk, bool outward, bool resetCov); -template GPUd() int GPUTrackingRefit::RefitTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, bool outward, bool resetCov); -template GPUd() int GPUTrackingRefit::RefitTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, bool outward, bool resetCov); +#if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE) // FIXME: DR: WORKAROUND to avoid CUDA bug creating host symbols for device code. +template GPUdni() int GPUTrackingRefit::RefitTrack(GPUTPCGMMergedTrack& trk, bool outward, bool resetCov); +template GPUdni() int GPUTrackingRefit::RefitTrack(GPUTPCGMMergedTrack& trk, bool outward, bool resetCov); +template GPUdni() int GPUTrackingRefit::RefitTrack(TrackTPC& trk, bool outward, bool resetCov); +template GPUdni() int GPUTrackingRefit::RefitTrack(TrackTPC& trk, bool outward, bool resetCov); +template GPUdni() int GPUTrackingRefit::RefitTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, bool outward, bool resetCov); +template GPUdni() int GPUTrackingRefit::RefitTrack(GPUTrackingRefit::TrackParCovWithArgs& trk, bool outward, bool resetCov); +#endif #ifndef GPUCA_GPUCODE void GPUTrackingRefit::SetPtrsFromGPUConstantMem(const GPUConstantMem* v, MEM_CONSTANT(GPUParam) * p) diff --git a/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx index 5d1aba3c0d07d..87005299b9484 100644 --- a/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx +++ b/GPU/GPUTracking/Refit/GPUTrackingRefitKernel.cxx @@ -44,5 +44,7 @@ GPUdii() void GPUTrackingRefitKernel::Thread(int nBlocks, int nThreads, int iBlo } } } -template GPUd() void GPUTrackingRefitKernel::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors); -template GPUd() void GPUTrackingRefitKernel::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors); +#if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE) // FIXME: DR: WORKAROUND to avoid CUDA bug creating host symbols for device code. +template GPUdni() void GPUTrackingRefitKernel::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors); +template GPUdni() void GPUTrackingRefitKernel::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors); +#endif diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCTrackLinearisation.h b/GPU/GPUTracking/SliceTracker/GPUTPCTrackLinearisation.h index 5b94992b0d708..d9f332beabd7d 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCTrackLinearisation.h +++ b/GPU/GPUTracking/SliceTracker/GPUTPCTrackLinearisation.h @@ -67,7 +67,8 @@ class GPUTPCTrackLinearisation float mQPt; // QPt }; -GPUd() MEM_CLASS_PRE2() inline GPUTPCTrackLinearisation::GPUTPCTrackLinearisation(const MEM_LG2(GPUTPCTrackParam) & GPUrestrict() t) : mSinPhi(t.SinPhi()), mCosPhi(0), mDzDs(t.DzDs()), mQPt(t.QPt()) +MEM_CLASS_PRE2() +GPUdi() GPUTPCTrackLinearisation::GPUTPCTrackLinearisation(const MEM_LG2(GPUTPCTrackParam) & GPUrestrict() t) : mSinPhi(t.SinPhi()), mCosPhi(0), mDzDs(t.DzDs()), mQPt(t.QPt()) { if (mSinPhi > GPUCA_MAX_SIN_PHI) { mSinPhi = GPUCA_MAX_SIN_PHI; diff --git a/GPU/GPUTracking/SliceTracker/GPUTPCTrackParam.h b/GPU/GPUTracking/SliceTracker/GPUTPCTrackParam.h index 4b71a28401787..38d81c6e594b5 100644 --- a/GPU/GPUTracking/SliceTracker/GPUTPCTrackParam.h +++ b/GPU/GPUTracking/SliceTracker/GPUTPCTrackParam.h @@ -157,7 +157,8 @@ class GPUTPCTrackParam int mNDF; // the Number of Degrees of Freedom }; -GPUd() MEM_CLASS_PRE() inline void MEM_LG(GPUTPCTrackParam)::InitParam() +MEM_CLASS_PRE() +GPUdi() void MEM_LG(GPUTPCTrackParam)::InitParam() { // Initialize Tracklet Parameters using default values SetSinPhi(0); diff --git a/GPU/GPUTracking/TRDTracking/GPUTRDTrackerDebug.h b/GPU/GPUTracking/TRDTracking/GPUTRDTrackerDebug.h index 911efb04a7b19..e05b2e719c3d2 100644 --- a/GPU/GPUTracking/TRDTracking/GPUTRDTrackerDebug.h +++ b/GPU/GPUTracking/TRDTracking/GPUTRDTrackerDebug.h @@ -345,12 +345,14 @@ class GPUTRDTrackerDebug GPUd() void SetFindable(bool* findable) {} GPUd() void Output() {} }; +#if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE) // FIXME: DR: WORKAROUND to avoid CUDA bug creating host symbols for device code. #ifndef GPUCA_ALIROOT_LIB template class GPUTRDTrackerDebug; #endif #if !defined(GPUCA_STANDALONE) && !defined(GPUCA_GPUCODE) template class GPUTRDTrackerDebug; #endif +#endif } // namespace gpu } // namespace GPUCA_NAMESPACE diff --git a/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx b/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx index c221050816815..20db05948c598 100644 --- a/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx +++ b/GPU/GPUTracking/TRDTracking/GPUTRDTrackerKernels.cxx @@ -42,7 +42,9 @@ GPUdii() void GPUTRDTrackerKernels::Thread(int nBlocks, int nThreads, int iBlock } } -template GPUd() void GPUTRDTrackerKernels::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTrackerGPU* externalInstance); +#if !defined(GPUCA_GPUCODE) || defined(GPUCA_GPUCODE_DEVICE) // FIXME: DR: WORKAROUND to avoid CUDA bug creating host symbols for device code. +template GPUdni() void GPUTRDTrackerKernels::Thread<0>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTrackerGPU* externalInstance); #ifdef GPUCA_HAVE_O2HEADERS -template GPUd() void GPUTRDTrackerKernels::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTracker* externalInstance); +template GPUdni() void GPUTRDTrackerKernels::Thread<1>(int nBlocks, int nThreads, int iBlock, int iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUTRDTracker* externalInstance); #endif // GPUCA_HAVE_O2HEADERS +#endif diff --git a/GPU/GPUTracking/display/GPUDisplayFrontendGlfw.cxx b/GPU/GPUTracking/display/GPUDisplayFrontendGlfw.cxx index 8be4e31fa22e3..54f6591d2286c 100644 --- a/GPU/GPUTracking/display/GPUDisplayFrontendGlfw.cxx +++ b/GPU/GPUTracking/display/GPUDisplayFrontendGlfw.cxx @@ -167,7 +167,7 @@ void GPUDisplayFrontendGlfw::GetKey(int key, int scancode, int mods, int& keyOut if (specialKey) { keyOut = keyPressOut = specialKey; } else { - keyOut = keyPressOut = localeKey; + keyOut = keyPressOut = (unsigned char)localeKey; if (keyPressOut >= 'a' && keyPressOut <= 'z') { keyPressOut += 'A' - 'a'; } @@ -193,7 +193,7 @@ void GPUDisplayFrontendGlfw::key_callback(GLFWwindow* window, int key, int scanc if (action == GLFW_PRESS) { me->mLastKeyDown = handleKey; } else if (action == GLFW_RELEASE) { - keyPress = me->mKeyDownMap[handleKey]; + keyPress = (unsigned char)me->mKeyDownMap[handleKey]; me->mKeys[keyPress] = false; me->mKeysShift[keyPress] = false; } diff --git a/GPU/GPUTracking/display/GPUDisplayFrontendX11.cxx b/GPU/GPUTracking/display/GPUDisplayFrontendX11.cxx index 2be2e243f83c1..7e397813c12f2 100644 --- a/GPU/GPUTracking/display/GPUDisplayFrontendX11.cxx +++ b/GPU/GPUTracking/display/GPUDisplayFrontendX11.cxx @@ -147,7 +147,7 @@ void GPUDisplayFrontendX11::GetKey(XEvent& event, int& keyOut, int& keyPressOut) tmpString[0] = 0; } int specialKey = GetKey(sym); - int localeKey = tmpString[0]; + int localeKey = (unsigned char)tmpString[0]; // GPUInfo("Key: keycode %d -> sym %d (%c) key %d (%c) special %d (%c)", (int)event.xkey.keycode, (int)sym, (char)sym, (int)localeKey, (char)localeKey, (int)specialKey, (char)specialKey); if (specialKey) { diff --git a/GPU/Utils/FlatObject.h b/GPU/Utils/FlatObject.h index 59bbde3ebeb42..ac06d7646abd7 100644 --- a/GPU/Utils/FlatObject.h +++ b/GPU/Utils/FlatObject.h @@ -178,10 +178,14 @@ class FlatObject /// _____________ Constructors / destructors __________________________ /// Default constructor / destructor - FlatObject() CON_DEFAULT; +#ifndef GPUCA_GPUCODE + FlatObject() CON_DEFAULT; // No object derrived from FlatObject should be created on the GPU ~FlatObject(); FlatObject(const FlatObject&) CON_DELETE; FlatObject& operator=(const FlatObject&) CON_DELETE; +#else + FlatObject() CON_DELETE; +#endif protected: /// _____________ Memory alignment __________________________ diff --git a/dependencies/FindO2GPU.cmake b/dependencies/FindO2GPU.cmake index 771e326c21f53..fa82e2b6128d6 100644 --- a/dependencies/FindO2GPU.cmake +++ b/dependencies/FindO2GPU.cmake @@ -105,8 +105,8 @@ if(ENABLE_CUDA) message(${FAILURE_SEVERITY} "CUDA found but thrust not available") set(CMAKE_CUDA_COMPILER OFF) endif() - if (NOT CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "11.4") - message(${FAILURE_SEVERITY} "CUDA Version too old: ${CMAKE_CUDA_COMPILER_VERSION}, 11.4 required") + if (NOT CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "12.6") + message(${FAILURE_SEVERITY} "CUDA Version too old: ${CMAKE_CUDA_COMPILER_VERSION}, 12.6 required") set(CMAKE_CUDA_COMPILER OFF) endif() endif() @@ -115,9 +115,7 @@ if(ENABLE_CUDA) if(GPUCA_KERNEL_RESOURCE_USAGE_VERBOSE) string(APPEND CMAKE_CUDA_FLAGS " -Xptxas -v") endif() - if(CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL "12.3") - string(APPEND CMAKE_CUDA_FLAGS " -Xcudafe --diag_suppress=20257") # TODO: Cleanup - endif() + string(APPEND CMAKE_CUDA_FLAGS " -Xcudafe --diag_suppress=114") if (NOT ENABLE_CUDA STREQUAL "AUTO") string(APPEND CMAKE_CUDA_FLAGS " --allow-unsupported-compiler") endif()