From 034022106e1499c0e7264b919b0bf396946ffc5a Mon Sep 17 00:00:00 2001 From: Andres Rios Tascon Date: Fri, 30 Aug 2024 10:42:33 -0700 Subject: [PATCH] Cleaned up namespaces --- .../plugins/alpaka/LSTModulesDevESProducer.cc | 12 +- RecoTracker/LST/plugins/alpaka/LSTProducer.cc | 6 +- .../LSTCore/interface/alpaka/Constants.h | 171 ++++---- RecoTracker/LSTCore/interface/alpaka/LST.h | 182 +++++---- RecoTracker/LSTCore/src/alpaka/Event.dev.cc | 48 +-- RecoTracker/LSTCore/src/alpaka/Event.h | 364 +++++++++--------- RecoTracker/LSTCore/src/alpaka/Hit.h | 2 - RecoTracker/LSTCore/src/alpaka/LST.dev.cc | 110 +++--- RecoTracker/LSTCore/src/alpaka/MiniDoublet.h | 74 ++-- .../LSTCore/src/alpaka/PixelQuintuplet.h | 125 +++--- RecoTracker/LSTCore/src/alpaka/PixelTriplet.h | 113 +++--- RecoTracker/LSTCore/src/alpaka/Quintuplet.h | 132 +++---- RecoTracker/LSTCore/src/alpaka/Segment.h | 62 ++- RecoTracker/LSTCore/src/alpaka/Triplet.h | 32 +- RecoTracker/LSTCore/standalone/bin/lst.cc | 7 +- RecoTracker/LSTCore/standalone/bin/lst.h | 6 +- .../standalone/code/core/AccessHelper.cc | 6 +- .../standalone/code/core/AnalysisConfig.h | 2 +- .../LSTCore/standalone/code/core/trkCore.h | 3 +- .../standalone/code/core/write_lst_ntuple.cc | 22 +- 20 files changed, 702 insertions(+), 777 deletions(-) diff --git a/RecoTracker/LST/plugins/alpaka/LSTModulesDevESProducer.cc b/RecoTracker/LST/plugins/alpaka/LSTModulesDevESProducer.cc index 46c99993c5ed9..0f0c53344de18 100644 --- a/RecoTracker/LST/plugins/alpaka/LSTModulesDevESProducer.cc +++ b/RecoTracker/LST/plugins/alpaka/LSTModulesDevESProducer.cc @@ -1,3 +1,7 @@ +// LST includes +#include "RecoTracker/LSTCore/interface/Module.h" +#include "RecoTracker/LSTCore/interface/alpaka/LST.h" + #include "FWCore/ParameterSet/interface/ParameterSet.h" #include "HeterogeneousCore/AlpakaCore/interface/alpaka/ESProducer.h" @@ -7,10 +11,6 @@ #include "RecoTracker/Record/interface/TrackerRecoGeometryRecord.h" -// LST includes -#include "RecoTracker/LSTCore/interface/Module.h" -#include "RecoTracker/LSTCore/interface/alpaka/LST.h" - namespace ALPAKA_ACCELERATOR_NAMESPACE { class LSTModulesDevESProducer : public ESProducer { @@ -22,8 +22,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { descriptions.addWithDefaultLabel(desc); } - std::unique_ptr<::lst::LSTESData> produce(TrackerRecoGeometryRecord const& iRecord) { - return ::lst::loadAndFillESHost(); + std::unique_ptr> produce(TrackerRecoGeometryRecord const& iRecord) { + return lst::loadAndFillESHost(); } }; diff --git a/RecoTracker/LST/plugins/alpaka/LSTProducer.cc b/RecoTracker/LST/plugins/alpaka/LSTProducer.cc index 21a3ad93eb1e3..7eb6c57ade05c 100644 --- a/RecoTracker/LST/plugins/alpaka/LSTProducer.cc +++ b/RecoTracker/LST/plugins/alpaka/LSTProducer.cc @@ -1,5 +1,7 @@ #include +#include "RecoTracker/LSTCore/interface/alpaka/LST.h" + #include "FWCore/MessageLogger/interface/MessageLogger.h" #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" #include "FWCore/ParameterSet/interface/ParameterSet.h" @@ -19,8 +21,6 @@ #include "RecoTracker/Record/interface/TrackerRecoGeometryRecord.h" -#include "RecoTracker/LSTCore/interface/alpaka/LST.h" - namespace ALPAKA_ACCELERATOR_NAMESPACE { class LSTProducer : public stream::SynchronizingEDProducer<> { @@ -86,7 +86,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE { private: edm::EDGetTokenT lstPixelSeedInputToken_; edm::EDGetTokenT lstPhase2OTHitsInputToken_; - device::ESGetToken<::lst::LSTESData, TrackerRecoGeometryRecord> lstESToken_; + device::ESGetToken, TrackerRecoGeometryRecord> lstESToken_; const bool verbose_, nopLSDupClean_, tcpLSTriplets_; edm::EDPutTokenT lstOutputToken_; diff --git a/RecoTracker/LSTCore/interface/alpaka/Constants.h b/RecoTracker/LSTCore/interface/alpaka/Constants.h index 9fed7760c721a..1a16dad68420e 100644 --- a/RecoTracker/LSTCore/interface/alpaka/Constants.h +++ b/RecoTracker/LSTCore/interface/alpaka/Constants.h @@ -9,118 +9,99 @@ #include #endif -namespace ALPAKA_ACCELERATOR_NAMESPACE { - namespace lst { - - // Re-export some useful things from the main namespace - using ::lst::allocBufWrapper; - using ::lst::Buf; - using ::lst::max_blocks; - using ::lst::max_connected_modules; - using ::lst::n_max_nonpixel_track_candidates; - using ::lst::n_max_pixel_md_per_modules; - using ::lst::n_max_pixel_quintuplets; - using ::lst::n_max_pixel_segments_per_module; - using ::lst::n_max_pixel_track_candidates; - using ::lst::n_max_pixel_triplets; - using ::lst::Params_LS; - using ::lst::Params_pLS; - using ::lst::Params_pT3; - using ::lst::Params_pT5; - using ::lst::Params_T3; - using ::lst::Params_T5; - using ::lst::size_superbins; +namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { + + using namespace ::lst; // Half precision wrapper functions. #if defined(FP16_Base) #define __F2H __float2half #define __H2F __half2float - typedef __half float FPX; + typedef __half float FPX; #else #define __F2H #define __H2F - typedef float FPX; + typedef float FPX; #endif - Vec3D constexpr elementsPerThread(Vec3D::all(static_cast(1))); + Vec3D constexpr elementsPerThread(Vec3D::all(static_cast(1))); // Needed for files that are compiled by g++ to not throw an error. // uint4 is defined only for CUDA, so we will have to revisit this soon when running on other backends. #if !defined(ALPAKA_ACC_GPU_CUDA_ENABLED) && !defined(ALPAKA_ACC_GPU_HIP_ENABLED) - struct uint4 { - unsigned int x; - unsigned int y; - unsigned int z; - unsigned int w; - }; + struct uint4 { + unsigned int x; + unsigned int y; + unsigned int z; + unsigned int w; + }; #endif - // Adjust grid and block sizes based on backend configuration - template > - ALPAKA_FN_HOST ALPAKA_FN_INLINE WorkDiv createWorkDiv(const Vec& blocksPerGrid, - const Vec& threadsPerBlock, - const Vec& elementsPerThreadArg) { - Vec adjustedBlocks = blocksPerGrid; - Vec adjustedThreads = threadsPerBlock; - - // special overrides for CPU/host cases - if constexpr (std::is_same_v) { - adjustedBlocks = Vec::all(static_cast(1)); - - if constexpr (alpaka::accMatchesTags) { - // Serial execution, set threads to 1 as well - adjustedThreads = Vec::all(static_cast(1)); // probably redundant - } + // Adjust grid and block sizes based on backend configuration + template > + ALPAKA_FN_HOST ALPAKA_FN_INLINE WorkDiv createWorkDiv(const Vec& blocksPerGrid, + const Vec& threadsPerBlock, + const Vec& elementsPerThreadArg) { + Vec adjustedBlocks = blocksPerGrid; + Vec adjustedThreads = threadsPerBlock; + + // special overrides for CPU/host cases + if constexpr (std::is_same_v) { + adjustedBlocks = Vec::all(static_cast(1)); + + if constexpr (alpaka::accMatchesTags) { + // Serial execution, set threads to 1 as well + adjustedThreads = Vec::all(static_cast(1)); // probably redundant } - - return WorkDiv(adjustedBlocks, adjustedThreads, elementsPerThreadArg); } - // The constants below are usually used in functions like alpaka::math::min(), - // expecting a reference (T const&) in the arguments. Hence, - // ALPAKA_STATIC_ACC_MEM_GLOBAL needs to be used in addition to constexpr. - - // 15 MeV constant from the approximate Bethe-Bloch formula - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMulsInGeV = 0.015; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleBarrel[6] = { - 0.0052, 0.0038, 0.0034, 0.0034, 0.0032, 0.0034}; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleEndcap[5] = {0.006, 0.006, 0.006, 0.006, 0.006}; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanBarrel[6] = { - 25.007152356, 37.2186993757, 52.3104270826, 68.6658656666, 85.9770373007, 108.301772384}; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanEndcap[5] = { - 130.992832231, 154.813883559, 185.352604327, 221.635123002, 265.022076742}; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float k2Rinv1GeVf = (2.99792458e-3 * 3.8) / 2; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kR1GeVf = 1. / (2.99792458e-3 * 3.8); - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kSinAlphaMax = 0.95; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float ptCut = PT_CUT; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kDeltaZLum = 15.0; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPixelPSZpitch = 0.15; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStripPSZpitch = 2.4; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStrip2SZpitch = 5.0; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidth2S = 0.009; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidthPS = 0.01; - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPt_betaMax = 7.0; - // Since C++ can't represent infinity, lst_INF = 123456789 was used to represent infinity in the data table - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float lst_INF = 123456789.0; - - namespace t5dnn { - - // Working points matching LST fake rate (43.9%) or signal acceptance (82.0%) - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp1 = 0.3418833f; // 94.0% TPR, 43.9% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp2 = 0.6177366f; // 82.0% TPR, 20.0% FPR - // Other working points - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp70 = 0.7776195f; // 70.0% TPR, 10.0% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp75 = 0.7181118f; // 75.0% TPR, 13.5% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp80 = 0.6492643f; // 80.0% TPR, 17.9% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp85 = 0.5655319f; // 85.0% TPR, 23.8% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp90 = 0.4592205f; // 90.0% TPR, 32.6% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp95 = 0.3073708f; // 95.0% TPR, 47.7% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp97p5 = 0.2001348f; // 97.5% TPR, 61.2% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99 = 0.1120605f; // 99.0% TPR, 75.9% FPR - ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99p9 = 0.0218196f; // 99.9% TPR, 95.4% FPR - - } // namespace t5dnn - - } //namespace lst -} //namespace ALPAKA_ACCELERATOR_NAMESPACE + return WorkDiv(adjustedBlocks, adjustedThreads, elementsPerThreadArg); + } + + // The constants below are usually used in functions like alpaka::math::min(), + // expecting a reference (T const&) in the arguments. Hence, + // ALPAKA_STATIC_ACC_MEM_GLOBAL needs to be used in addition to constexpr. + + // 15 MeV constant from the approximate Bethe-Bloch formula + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMulsInGeV = 0.015; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleBarrel[6] = { + 0.0052, 0.0038, 0.0034, 0.0034, 0.0032, 0.0034}; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleEndcap[5] = {0.006, 0.006, 0.006, 0.006, 0.006}; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanBarrel[6] = { + 25.007152356, 37.2186993757, 52.3104270826, 68.6658656666, 85.9770373007, 108.301772384}; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanEndcap[5] = { + 130.992832231, 154.813883559, 185.352604327, 221.635123002, 265.022076742}; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float k2Rinv1GeVf = (2.99792458e-3 * 3.8) / 2; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kR1GeVf = 1. / (2.99792458e-3 * 3.8); + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kSinAlphaMax = 0.95; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float ptCut = PT_CUT; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kDeltaZLum = 15.0; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPixelPSZpitch = 0.15; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStripPSZpitch = 2.4; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStrip2SZpitch = 5.0; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidth2S = 0.009; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidthPS = 0.01; + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPt_betaMax = 7.0; + // Since C++ can't represent infinity, lst_INF = 123456789 was used to represent infinity in the data table + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float lst_INF = 123456789.0; + + namespace t5dnn { + + // Working points matching LST fake rate (43.9%) or signal acceptance (82.0%) + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp1 = 0.3418833f; // 94.0% TPR, 43.9% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp2 = 0.6177366f; // 82.0% TPR, 20.0% FPR + // Other working points + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp70 = 0.7776195f; // 70.0% TPR, 10.0% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp75 = 0.7181118f; // 75.0% TPR, 13.5% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp80 = 0.6492643f; // 80.0% TPR, 17.9% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp85 = 0.5655319f; // 85.0% TPR, 23.8% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp90 = 0.4592205f; // 90.0% TPR, 32.6% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp95 = 0.3073708f; // 95.0% TPR, 47.7% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp97p5 = 0.2001348f; // 97.5% TPR, 61.2% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99 = 0.1120605f; // 99.0% TPR, 75.9% FPR + ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99p9 = 0.0218196f; // 99.9% TPR, 95.4% FPR + + } // namespace t5dnn + +} // namespace ALPAKA_ACCELERATOR_NAMESPACE::lst #endif diff --git a/RecoTracker/LSTCore/interface/alpaka/LST.h b/RecoTracker/LSTCore/interface/alpaka/LST.h index b7ec844154575..1f3c08804540f 100644 --- a/RecoTracker/LSTCore/interface/alpaka/LST.h +++ b/RecoTracker/LSTCore/interface/alpaka/LST.h @@ -1,110 +1,106 @@ #ifndef RecoTracker_LSTCore_interface_alpaka_LST_h #define RecoTracker_LSTCore_interface_alpaka_LST_h -#include "RecoTracker/LSTCore/interface/Constants.h" +#include "RecoTracker/LSTCore/interface/alpaka/Constants.h" #include "RecoTracker/LSTCore/interface/LSTESData.h" #include #include #include -using ::lst::LSTESData; +namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { + class Event; -namespace ALPAKA_ACCELERATOR_NAMESPACE { - namespace lst { - class Event; + class LST { + public: + LST() = default; - class LST { - public: - LST() = default; + void run(Queue& queue, + bool verbose, + LSTESData const* deviceESData, + std::vector const& see_px, + std::vector const& see_py, + std::vector const& see_pz, + std::vector const& see_dxy, + std::vector const& see_dz, + std::vector const& see_ptErr, + std::vector const& see_etaErr, + std::vector const& see_stateTrajGlbX, + std::vector const& see_stateTrajGlbY, + std::vector const& see_stateTrajGlbZ, + std::vector const& see_stateTrajGlbPx, + std::vector const& see_stateTrajGlbPy, + std::vector const& see_stateTrajGlbPz, + std::vector const& see_q, + std::vector> const& see_hitIdx, + std::vector const& ph2_detId, + std::vector const& ph2_x, + std::vector const& ph2_y, + std::vector const& ph2_z, + bool no_pls_dupclean, + bool tc_pls_triplets); + std::vector> const& hits() const { return out_tc_hitIdxs_; } + std::vector const& len() const { return out_tc_len_; } + std::vector const& seedIdx() const { return out_tc_seedIdx_; } + std::vector const& trackCandidateType() const { return out_tc_trackCandidateType_; } - void run(Queue& queue, - bool verbose, - LSTESData const* deviceESData, - std::vector const& see_px, - std::vector const& see_py, - std::vector const& see_pz, - std::vector const& see_dxy, - std::vector const& see_dz, - std::vector const& see_ptErr, - std::vector const& see_etaErr, - std::vector const& see_stateTrajGlbX, - std::vector const& see_stateTrajGlbY, - std::vector const& see_stateTrajGlbZ, - std::vector const& see_stateTrajGlbPx, - std::vector const& see_stateTrajGlbPy, - std::vector const& see_stateTrajGlbPz, - std::vector const& see_q, - std::vector> const& see_hitIdx, - std::vector const& ph2_detId, - std::vector const& ph2_x, - std::vector const& ph2_y, - std::vector const& ph2_z, - bool no_pls_dupclean, - bool tc_pls_triplets); - std::vector> const& hits() const { return out_tc_hitIdxs_; } - std::vector const& len() const { return out_tc_len_; } - std::vector const& seedIdx() const { return out_tc_seedIdx_; } - std::vector const& trackCandidateType() const { return out_tc_trackCandidateType_; } + private: + void prepareInput(std::vector const& see_px, + std::vector const& see_py, + std::vector const& see_pz, + std::vector const& see_dxy, + std::vector const& see_dz, + std::vector const& see_ptErr, + std::vector const& see_etaErr, + std::vector const& see_stateTrajGlbX, + std::vector const& see_stateTrajGlbY, + std::vector const& see_stateTrajGlbZ, + std::vector const& see_stateTrajGlbPx, + std::vector const& see_stateTrajGlbPy, + std::vector const& see_stateTrajGlbPz, + std::vector const& see_q, + std::vector> const& see_hitIdx, + std::vector const& ph2_detId, + std::vector const& ph2_x, + std::vector const& ph2_y, + std::vector const& ph2_z); - private: - void prepareInput(std::vector const& see_px, - std::vector const& see_py, - std::vector const& see_pz, - std::vector const& see_dxy, - std::vector const& see_dz, - std::vector const& see_ptErr, - std::vector const& see_etaErr, - std::vector const& see_stateTrajGlbX, - std::vector const& see_stateTrajGlbY, - std::vector const& see_stateTrajGlbZ, - std::vector const& see_stateTrajGlbPx, - std::vector const& see_stateTrajGlbPy, - std::vector const& see_stateTrajGlbPz, - std::vector const& see_q, - std::vector> const& see_hitIdx, - std::vector const& ph2_detId, - std::vector const& ph2_x, - std::vector const& ph2_y, - std::vector const& ph2_z); + void getOutput(Event& event); + std::vector getHitIdxs(short trackCandidateType, + unsigned int TCIdx, + unsigned int const* TCHitIndices, + unsigned int const* hitIndices); - void getOutput(Event& event); - std::vector getHitIdxs(short trackCandidateType, - unsigned int TCIdx, - unsigned int const* TCHitIndices, - unsigned int const* hitIndices); + // Input and output vectors + std::vector in_trkX_; + std::vector in_trkY_; + std::vector in_trkZ_; + std::vector in_hitId_; + std::vector in_hitIdxs_; + std::vector in_hitIndices_vec0_; + std::vector in_hitIndices_vec1_; + std::vector in_hitIndices_vec2_; + std::vector in_hitIndices_vec3_; + std::vector in_deltaPhi_vec_; + std::vector in_ptIn_vec_; + std::vector in_ptErr_vec_; + std::vector in_px_vec_; + std::vector in_py_vec_; + std::vector in_pz_vec_; + std::vector in_eta_vec_; + std::vector in_etaErr_vec_; + std::vector in_phi_vec_; + std::vector in_charge_vec_; + std::vector in_seedIdx_vec_; + std::vector in_superbin_vec_; + std::vector in_pixelType_vec_; + std::vector in_isQuad_vec_; + std::vector> out_tc_hitIdxs_; + std::vector out_tc_len_; + std::vector out_tc_seedIdx_; + std::vector out_tc_trackCandidateType_; + }; - // Input and output vectors - std::vector in_trkX_; - std::vector in_trkY_; - std::vector in_trkZ_; - std::vector in_hitId_; - std::vector in_hitIdxs_; - std::vector in_hitIndices_vec0_; - std::vector in_hitIndices_vec1_; - std::vector in_hitIndices_vec2_; - std::vector in_hitIndices_vec3_; - std::vector in_deltaPhi_vec_; - std::vector in_ptIn_vec_; - std::vector in_ptErr_vec_; - std::vector in_px_vec_; - std::vector in_py_vec_; - std::vector in_pz_vec_; - std::vector in_eta_vec_; - std::vector in_etaErr_vec_; - std::vector in_phi_vec_; - std::vector in_charge_vec_; - std::vector in_seedIdx_vec_; - std::vector in_superbin_vec_; - std::vector<::lst::PixelType> in_pixelType_vec_; - std::vector in_isQuad_vec_; - std::vector> out_tc_hitIdxs_; - std::vector out_tc_len_; - std::vector out_tc_seedIdx_; - std::vector out_tc_trackCandidateType_; - }; - - } // namespace lst -} // namespace ALPAKA_ACCELERATOR_NAMESPACE +} // namespace ALPAKA_ACCELERATOR_NAMESPACE::lst #endif diff --git a/RecoTracker/LSTCore/src/alpaka/Event.dev.cc b/RecoTracker/LSTCore/src/alpaka/Event.dev.cc index 7badc77baa94d..fdd7d32f91757 100644 --- a/RecoTracker/LSTCore/src/alpaka/Event.dev.cc +++ b/RecoTracker/LSTCore/src/alpaka/Event.dev.cc @@ -180,8 +180,8 @@ void Event::addHitToEvent(std::vector const& x, alpaka::exec(queue, hit_loop_workdiv, HitLoopKernel{}, - ::lst::Endcap, - ::lst::TwoS, + Endcap, + TwoS, nModules_, nEndCapMap_, endcapGeometryBuffers_.geoMapDetId_buf.data(), @@ -214,7 +214,7 @@ void Event::addPixelSegmentToEvent(std::vector const& hitIndices0, std::vector const& charge, std::vector const& seedIdx, std::vector const& superbin, - std::vector<::lst::PixelType> const& pixelType, + std::vector const& pixelType, std::vector const& isQuad) { unsigned int size = ptIn.size(); @@ -685,7 +685,7 @@ void Event::createPixelTriplets() { } auto superbins_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); - auto pixelTypes_buf = allocBufWrapper<::lst::PixelType>(devHost, n_max_pixel_segments_per_module, queue); + auto pixelTypes_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); alpaka::memcpy(queue, superbins_buf, segmentsBuffers->superbin_buf); alpaka::memcpy(queue, pixelTypes_buf, segmentsBuffers->pixelType_buf); @@ -717,11 +717,11 @@ void Event::createPixelTriplets() { // TODO: check if a map/reduction to just eligible pLSs would speed up the kernel // the current selection still leaves a significant fraction of unmatchable pLSs for (unsigned int i = 0; i < nInnerSegments; i++) { // loop over # pLS - ::lst::PixelType pixelType = pixelTypes[i]; // Get pixel type for this pLS + PixelType pixelType = pixelTypes[i]; // Get pixel type for this pLS int superbin = superbins[i]; // Get superbin for this pixel if ((superbin < 0) or (superbin >= (int)size_superbins) or - ((pixelType != ::lst::PixelType::kHighPt) and (pixelType != ::lst::PixelType::kLowPtPosCurv) and - (pixelType != ::lst::PixelType::kLowPtNegCurv))) { + ((pixelType != PixelType::kHighPt) and (pixelType != PixelType::kLowPtPosCurv) and + (pixelType != PixelType::kLowPtNegCurv))) { connectedPixelSize_host[i] = 0; connectedPixelIndex_host[i] = 0; continue; @@ -729,21 +729,21 @@ void Event::createPixelTriplets() { // Used pixel type to select correct size-index arrays switch (pixelType) { - case ::lst::PixelType::kInvalid: + case PixelType::kInvalid: break; - case ::lst::PixelType::kHighPt: + case PixelType::kHighPt: // number of connected modules to this pixel connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizes[superbin]; // index to get start of connected modules for this superbin in map connectedPixelIndex_host[i] = pixelMapping_.connectedPixelsIndex[superbin]; break; - case ::lst::PixelType::kLowPtPosCurv: + case PixelType::kLowPtPosCurv: // number of connected modules to this pixel connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizesPos[superbin]; // index to get start of connected modules for this superbin in map connectedPixelIndex_host[i] = pixelMapping_.connectedPixelsIndexPos[superbin] + pixelIndexOffsetPos; break; - case ::lst::PixelType::kLowPtNegCurv: + case PixelType::kLowPtNegCurv: // number of connected modules to this pixel connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizesNeg[superbin]; // index to get start of connected modules for this superbin in map @@ -888,7 +888,7 @@ void Event::createPixelQuintuplets() { } auto superbins_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); - auto pixelTypes_buf = allocBufWrapper<::lst::PixelType>(devHost, n_max_pixel_segments_per_module, queue); + auto pixelTypes_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); alpaka::memcpy(queue, superbins_buf, segmentsBuffers->superbin_buf); alpaka::memcpy(queue, pixelTypes_buf, segmentsBuffers->pixelType_buf); @@ -919,11 +919,11 @@ void Event::createPixelQuintuplets() { // Loop over # pLS for (unsigned int i = 0; i < nInnerSegments; i++) { - ::lst::PixelType pixelType = pixelTypes[i]; // Get pixel type for this pLS - int superbin = superbins[i]; // Get superbin for this pixel + PixelType pixelType = pixelTypes[i]; // Get pixel type for this pLS + int superbin = superbins[i]; // Get superbin for this pixel if ((superbin < 0) or (superbin >= (int)size_superbins) or - ((pixelType != ::lst::PixelType::kHighPt) and (pixelType != ::lst::PixelType::kLowPtPosCurv) and - (pixelType != ::lst::PixelType::kLowPtNegCurv))) { + ((pixelType != PixelType::kHighPt) and (pixelType != PixelType::kLowPtPosCurv) and + (pixelType != PixelType::kLowPtNegCurv))) { connectedPixelSize_host[i] = 0; connectedPixelIndex_host[i] = 0; continue; @@ -931,21 +931,21 @@ void Event::createPixelQuintuplets() { // Used pixel type to select correct size-index arrays switch (pixelType) { - case ::lst::PixelType::kInvalid: + case PixelType::kInvalid: break; - case ::lst::PixelType::kHighPt: + case PixelType::kHighPt: // number of connected modules to this pixel connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizes[superbin]; // index to get start of connected modules for this superbin in map connectedPixelIndex_host[i] = pixelMapping_.connectedPixelsIndex[superbin]; break; - case ::lst::PixelType::kLowPtPosCurv: + case PixelType::kLowPtPosCurv: // number of connected modules to this pixel connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizesPos[superbin]; // index to get start of connected modules for this superbin in map connectedPixelIndex_host[i] = pixelMapping_.connectedPixelsIndexPos[superbin] + pixelIndexOffsetPos; break; - case ::lst::PixelType::kLowPtNegCurv: + case PixelType::kLowPtNegCurv: // number of connected modules to this pixel connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizesNeg[superbin]; // index to get start of connected modules for this superbin in map @@ -1030,7 +1030,7 @@ void Event::addMiniDoubletsToEventExplicit() { for (unsigned int i = 0; i < nLowerModules_; i++) { if (!(nMDsCPU[i] == 0 or module_hitRanges[i * 2] == -1)) { - if (module_subdets[i] == ::lst::Barrel) { + if (module_subdets[i] == Barrel) { n_minidoublets_by_layer_barrel_[module_layers[i] - 1] += nMDsCPU[i]; } else { n_minidoublets_by_layer_endcap_[module_layers[i] - 1] += nMDsCPU[i]; @@ -1058,7 +1058,7 @@ void Event::addSegmentsToEventExplicit() { for (unsigned int i = 0; i < nLowerModules_; i++) { if (!(nSegmentsCPU[i] == 0)) { - if (module_subdets[i] == ::lst::Barrel) { + if (module_subdets[i] == Barrel) { n_segments_by_layer_barrel_[module_layers[i] - 1] += nSegmentsCPU[i]; } else { n_segments_by_layer_endcap_[module_layers[i] - 1] += nSegmentsCPU[i]; @@ -1090,7 +1090,7 @@ void Event::addQuintupletsToEventExplicit() { for (uint16_t i = 0; i < nLowerModules_; i++) { if (!(nQuintupletsCPU[i] == 0 or module_quintupletModuleIndices[i] == -1)) { - if (module_subdets[i] == ::lst::Barrel) { + if (module_subdets[i] == Barrel) { n_quintuplets_by_layer_barrel_[module_layers[i] - 1] += nQuintupletsCPU[i]; } else { n_quintuplets_by_layer_endcap_[module_layers[i] - 1] += nQuintupletsCPU[i]; @@ -1118,7 +1118,7 @@ void Event::addTripletsToEventExplicit() { for (uint16_t i = 0; i < nLowerModules_; i++) { if (nTripletsCPU[i] != 0) { - if (module_subdets[i] == ::lst::Barrel) { + if (module_subdets[i] == Barrel) { n_triplets_by_layer_barrel_[module_layers[i] - 1] += nTripletsCPU[i]; } else { n_triplets_by_layer_endcap_[module_layers[i] - 1] += nTripletsCPU[i]; diff --git a/RecoTracker/LSTCore/src/alpaka/Event.h b/RecoTracker/LSTCore/src/alpaka/Event.h index d3f6525b0333a..ceaeea1de2ddc 100644 --- a/RecoTracker/LSTCore/src/alpaka/Event.h +++ b/RecoTracker/LSTCore/src/alpaka/Event.h @@ -19,190 +19,182 @@ #include "HeterogeneousCore/AlpakaInterface/interface/host.h" -using ::lst::EndcapGeometryBuffer; -using ::lst::LSTESData; -using ::lst::ModulesBuffer; -using ::lst::PixelMap; - -namespace ALPAKA_ACCELERATOR_NAMESPACE { - namespace lst { - - class Event { - private: - Queue queue; - Device devAcc; - DevHost devHost; - bool addObjects; - - std::array n_hits_by_layer_barrel_; - std::array n_hits_by_layer_endcap_; - std::array n_minidoublets_by_layer_barrel_; - std::array n_minidoublets_by_layer_endcap_; - std::array n_segments_by_layer_barrel_; - std::array n_segments_by_layer_endcap_; - std::array n_triplets_by_layer_barrel_; - std::array n_triplets_by_layer_endcap_; - std::array n_trackCandidates_by_layer_barrel_; - std::array n_trackCandidates_by_layer_endcap_; - std::array n_quintuplets_by_layer_barrel_; - std::array n_quintuplets_by_layer_endcap_; - unsigned int nTotalSegments_; - - //Device stuff - std::optional rangesInGPU; - std::optional> rangesBuffers; - std::optional hitsInGPU; - std::optional> hitsBuffers; - std::optional mdsInGPU; - std::optional> miniDoubletsBuffers; - std::optional segmentsInGPU; - std::optional> segmentsBuffers; - std::optional tripletsInGPU; - std::optional> tripletsBuffers; - std::optional quintupletsInGPU; - std::optional> quintupletsBuffers; - std::optional trackCandidatesInGPU; - std::optional> trackCandidatesBuffers; - std::optional pixelTripletsInGPU; - std::optional> pixelTripletsBuffers; - std::optional pixelQuintupletsInGPU; - std::optional> pixelQuintupletsBuffers; - - //CPU interface stuff - std::optional> rangesInCPU; - std::optional> hitsInCPU; - std::optional> mdsInCPU; - std::optional> segmentsInCPU; - std::optional> tripletsInCPU; - std::optional> trackCandidatesInCPU; - std::optional> modulesInCPU; - std::optional> quintupletsInCPU; - std::optional> pixelTripletsInCPU; - std::optional> pixelQuintupletsInCPU; - - void initSync(bool verbose); - - const uint16_t nModules_; - const uint16_t nLowerModules_; - const unsigned int nPixels_; - const unsigned int nEndCapMap_; - ModulesBuffer const& modulesBuffers_; - PixelMap const& pixelMapping_; - EndcapGeometryBuffer const& endcapGeometryBuffers_; - - public: - // Constructor used for CMSSW integration. Uses an external queue. - Event(bool verbose, Queue const& q, const LSTESData* deviceESData) - : queue(q), - devAcc(alpaka::getDev(q)), - devHost(cms::alpakatools::host()), - nModules_(deviceESData->nModules), - nLowerModules_(deviceESData->nLowerModules), - nPixels_(deviceESData->nPixels), - nEndCapMap_(deviceESData->nEndCapMap), - modulesBuffers_(deviceESData->modulesBuffers), - pixelMapping_(*deviceESData->pixelMapping), - endcapGeometryBuffers_(deviceESData->endcapGeometryBuffers) { - initSync(verbose); - } - void resetEventSync(); // synchronizes - void wait() const { alpaka::wait(queue); } - - // Calls the appropriate hit function, then increments the counter - void addHitToEvent(std::vector const& x, - std::vector const& y, - std::vector const& z, - std::vector const& detId, - std::vector const& idxInNtuple); - void addPixelSegmentToEvent(std::vector const& hitIndices0, - std::vector const& hitIndices1, - std::vector const& hitIndices2, - std::vector const& hitIndices3, - std::vector const& dPhiChange, - std::vector const& ptIn, - std::vector const& ptErr, - std::vector const& px, - std::vector const& py, - std::vector const& pz, - std::vector const& eta, - std::vector const& etaErr, - std::vector const& phi, - std::vector const& charge, - std::vector const& seedIdx, - std::vector const& superbin, - std::vector<::lst::PixelType> const& pixelType, - std::vector const& isQuad); - - void createMiniDoublets(); - void createSegmentsWithModuleMap(); - void createTriplets(); - void createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets); - void createPixelTriplets(); - void createQuintuplets(); - void pixelLineSegmentCleaning(bool no_pls_dupclean); - void createPixelQuintuplets(); - - // functions that map the objects to the appropriate modules - void addMiniDoubletsToEventExplicit(); - void addSegmentsToEventExplicit(); - void addQuintupletsToEventExplicit(); - void addTripletsToEventExplicit(); - void resetObjectsInModule(); - - unsigned int getNumberOfHits(); - unsigned int getNumberOfHitsByLayer(unsigned int layer); - unsigned int getNumberOfHitsByLayerBarrel(unsigned int layer); - unsigned int getNumberOfHitsByLayerEndcap(unsigned int layer); - - unsigned int getNumberOfMiniDoublets(); - unsigned int getNumberOfMiniDoubletsByLayer(unsigned int layer); - unsigned int getNumberOfMiniDoubletsByLayerBarrel(unsigned int layer); - unsigned int getNumberOfMiniDoubletsByLayerEndcap(unsigned int layer); - - unsigned int getNumberOfSegments(); - unsigned int getNumberOfSegmentsByLayer(unsigned int layer); - unsigned int getNumberOfSegmentsByLayerBarrel(unsigned int layer); - unsigned int getNumberOfSegmentsByLayerEndcap(unsigned int layer); - - unsigned int getNumberOfTriplets(); - unsigned int getNumberOfTripletsByLayer(unsigned int layer); - unsigned int getNumberOfTripletsByLayerBarrel(unsigned int layer); - unsigned int getNumberOfTripletsByLayerEndcap(unsigned int layer); - - int getNumberOfPixelTriplets(); - int getNumberOfPixelQuintuplets(); - - unsigned int getNumberOfQuintuplets(); - unsigned int getNumberOfQuintupletsByLayer(unsigned int layer); - unsigned int getNumberOfQuintupletsByLayerBarrel(unsigned int layer); - unsigned int getNumberOfQuintupletsByLayerEndcap(unsigned int layer); - - int getNumberOfTrackCandidates(); - int getNumberOfPT5TrackCandidates(); - int getNumberOfPT3TrackCandidates(); - int getNumberOfPLSTrackCandidates(); - int getNumberOfPixelTrackCandidates(); - int getNumberOfT5TrackCandidates(); - - // sync adds alpaka::wait at the end of filling a buffer during lazy fill - // (has no effect on repeated calls) - // set to false may allow faster operation with concurrent calls of get* - // HANDLE WITH CARE - HitsBuffer& getHits(bool sync = true); - HitsBuffer& getHitsInCMSSW(bool sync = true); - ObjectRangesBuffer& getRanges(bool sync = true); - MiniDoubletsBuffer& getMiniDoublets(bool sync = true); - SegmentsBuffer& getSegments(bool sync = true); - TripletsBuffer& getTriplets(bool sync = true); - QuintupletsBuffer& getQuintuplets(bool sync = true); - PixelTripletsBuffer& getPixelTriplets(bool sync = true); - PixelQuintupletsBuffer& getPixelQuintuplets(bool sync = true); - TrackCandidatesBuffer& getTrackCandidates(bool sync = true); - TrackCandidatesBuffer& getTrackCandidatesInCMSSW(bool sync = true); - ModulesBuffer& getModules(bool isFull = false, bool sync = true); - }; - - } // namespace lst - -} // namespace ALPAKA_ACCELERATOR_NAMESPACE +namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { + + class Event { + private: + Queue queue; + Device devAcc; + DevHost devHost; + bool addObjects; + + std::array n_hits_by_layer_barrel_; + std::array n_hits_by_layer_endcap_; + std::array n_minidoublets_by_layer_barrel_; + std::array n_minidoublets_by_layer_endcap_; + std::array n_segments_by_layer_barrel_; + std::array n_segments_by_layer_endcap_; + std::array n_triplets_by_layer_barrel_; + std::array n_triplets_by_layer_endcap_; + std::array n_trackCandidates_by_layer_barrel_; + std::array n_trackCandidates_by_layer_endcap_; + std::array n_quintuplets_by_layer_barrel_; + std::array n_quintuplets_by_layer_endcap_; + unsigned int nTotalSegments_; + + //Device stuff + std::optional rangesInGPU; + std::optional> rangesBuffers; + std::optional hitsInGPU; + std::optional> hitsBuffers; + std::optional mdsInGPU; + std::optional> miniDoubletsBuffers; + std::optional segmentsInGPU; + std::optional> segmentsBuffers; + std::optional tripletsInGPU; + std::optional> tripletsBuffers; + std::optional quintupletsInGPU; + std::optional> quintupletsBuffers; + std::optional trackCandidatesInGPU; + std::optional> trackCandidatesBuffers; + std::optional pixelTripletsInGPU; + std::optional> pixelTripletsBuffers; + std::optional pixelQuintupletsInGPU; + std::optional> pixelQuintupletsBuffers; + + //CPU interface stuff + std::optional> rangesInCPU; + std::optional> hitsInCPU; + std::optional> mdsInCPU; + std::optional> segmentsInCPU; + std::optional> tripletsInCPU; + std::optional> trackCandidatesInCPU; + std::optional> modulesInCPU; + std::optional> quintupletsInCPU; + std::optional> pixelTripletsInCPU; + std::optional> pixelQuintupletsInCPU; + + void initSync(bool verbose); + + const uint16_t nModules_; + const uint16_t nLowerModules_; + const unsigned int nPixels_; + const unsigned int nEndCapMap_; + ModulesBuffer const& modulesBuffers_; + PixelMap const& pixelMapping_; + EndcapGeometryBuffer const& endcapGeometryBuffers_; + + public: + // Constructor used for CMSSW integration. Uses an external queue. + Event(bool verbose, Queue const& q, const LSTESData* deviceESData) + : queue(q), + devAcc(alpaka::getDev(q)), + devHost(cms::alpakatools::host()), + nModules_(deviceESData->nModules), + nLowerModules_(deviceESData->nLowerModules), + nPixels_(deviceESData->nPixels), + nEndCapMap_(deviceESData->nEndCapMap), + modulesBuffers_(deviceESData->modulesBuffers), + pixelMapping_(*deviceESData->pixelMapping), + endcapGeometryBuffers_(deviceESData->endcapGeometryBuffers) { + initSync(verbose); + } + void resetEventSync(); // synchronizes + void wait() const { alpaka::wait(queue); } + + // Calls the appropriate hit function, then increments the counter + void addHitToEvent(std::vector const& x, + std::vector const& y, + std::vector const& z, + std::vector const& detId, + std::vector const& idxInNtuple); + void addPixelSegmentToEvent(std::vector const& hitIndices0, + std::vector const& hitIndices1, + std::vector const& hitIndices2, + std::vector const& hitIndices3, + std::vector const& dPhiChange, + std::vector const& ptIn, + std::vector const& ptErr, + std::vector const& px, + std::vector const& py, + std::vector const& pz, + std::vector const& eta, + std::vector const& etaErr, + std::vector const& phi, + std::vector const& charge, + std::vector const& seedIdx, + std::vector const& superbin, + std::vector const& pixelType, + std::vector const& isQuad); + + void createMiniDoublets(); + void createSegmentsWithModuleMap(); + void createTriplets(); + void createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets); + void createPixelTriplets(); + void createQuintuplets(); + void pixelLineSegmentCleaning(bool no_pls_dupclean); + void createPixelQuintuplets(); + + // functions that map the objects to the appropriate modules + void addMiniDoubletsToEventExplicit(); + void addSegmentsToEventExplicit(); + void addQuintupletsToEventExplicit(); + void addTripletsToEventExplicit(); + void resetObjectsInModule(); + + unsigned int getNumberOfHits(); + unsigned int getNumberOfHitsByLayer(unsigned int layer); + unsigned int getNumberOfHitsByLayerBarrel(unsigned int layer); + unsigned int getNumberOfHitsByLayerEndcap(unsigned int layer); + + unsigned int getNumberOfMiniDoublets(); + unsigned int getNumberOfMiniDoubletsByLayer(unsigned int layer); + unsigned int getNumberOfMiniDoubletsByLayerBarrel(unsigned int layer); + unsigned int getNumberOfMiniDoubletsByLayerEndcap(unsigned int layer); + + unsigned int getNumberOfSegments(); + unsigned int getNumberOfSegmentsByLayer(unsigned int layer); + unsigned int getNumberOfSegmentsByLayerBarrel(unsigned int layer); + unsigned int getNumberOfSegmentsByLayerEndcap(unsigned int layer); + + unsigned int getNumberOfTriplets(); + unsigned int getNumberOfTripletsByLayer(unsigned int layer); + unsigned int getNumberOfTripletsByLayerBarrel(unsigned int layer); + unsigned int getNumberOfTripletsByLayerEndcap(unsigned int layer); + + int getNumberOfPixelTriplets(); + int getNumberOfPixelQuintuplets(); + + unsigned int getNumberOfQuintuplets(); + unsigned int getNumberOfQuintupletsByLayer(unsigned int layer); + unsigned int getNumberOfQuintupletsByLayerBarrel(unsigned int layer); + unsigned int getNumberOfQuintupletsByLayerEndcap(unsigned int layer); + + int getNumberOfTrackCandidates(); + int getNumberOfPT5TrackCandidates(); + int getNumberOfPT3TrackCandidates(); + int getNumberOfPLSTrackCandidates(); + int getNumberOfPixelTrackCandidates(); + int getNumberOfT5TrackCandidates(); + + // sync adds alpaka::wait at the end of filling a buffer during lazy fill + // (has no effect on repeated calls) + // set to false may allow faster operation with concurrent calls of get* + // HANDLE WITH CARE + HitsBuffer& getHits(bool sync = true); + HitsBuffer& getHitsInCMSSW(bool sync = true); + ObjectRangesBuffer& getRanges(bool sync = true); + MiniDoubletsBuffer& getMiniDoublets(bool sync = true); + SegmentsBuffer& getSegments(bool sync = true); + TripletsBuffer& getTriplets(bool sync = true); + QuintupletsBuffer& getQuintuplets(bool sync = true); + PixelTripletsBuffer& getPixelTriplets(bool sync = true); + PixelQuintupletsBuffer& getPixelQuintuplets(bool sync = true); + TrackCandidatesBuffer& getTrackCandidates(bool sync = true); + TrackCandidatesBuffer& getTrackCandidatesInCMSSW(bool sync = true); + ModulesBuffer& getModules(bool isFull = false, bool sync = true); + }; + +} // namespace ALPAKA_ACCELERATOR_NAMESPACE::lst #endif diff --git a/RecoTracker/LSTCore/src/alpaka/Hit.h b/RecoTracker/LSTCore/src/alpaka/Hit.h index 1a54008d4331c..3f559f4492df7 100644 --- a/RecoTracker/LSTCore/src/alpaka/Hit.h +++ b/RecoTracker/LSTCore/src/alpaka/Hit.h @@ -4,8 +4,6 @@ #include "RecoTracker/LSTCore/interface/alpaka/Constants.h" #include "RecoTracker/LSTCore/interface/Module.h" -using ::lst::Modules; - namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { struct Hits { unsigned int* nHits; diff --git a/RecoTracker/LSTCore/src/alpaka/LST.dev.cc b/RecoTracker/LSTCore/src/alpaka/LST.dev.cc index 402d8ce62132e..65543720a1d34 100644 --- a/RecoTracker/LSTCore/src/alpaka/LST.dev.cc +++ b/RecoTracker/LSTCore/src/alpaka/LST.dev.cc @@ -2,7 +2,7 @@ #include "Event.h" -using namespace ALPAKA_ACCELERATOR_NAMESPACE; +using namespace ALPAKA_ACCELERATOR_NAMESPACE::lst; #include "Math/Vector3D.h" #include "Math/VectorUtil.h" @@ -20,25 +20,25 @@ namespace { } } // namespace -void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::prepareInput(std::vector const& see_px, - std::vector const& see_py, - std::vector const& see_pz, - std::vector const& see_dxy, - std::vector const& see_dz, - std::vector const& see_ptErr, - std::vector const& see_etaErr, - std::vector const& see_stateTrajGlbX, - std::vector const& see_stateTrajGlbY, - std::vector const& see_stateTrajGlbZ, - std::vector const& see_stateTrajGlbPx, - std::vector const& see_stateTrajGlbPy, - std::vector const& see_stateTrajGlbPz, - std::vector const& see_q, - std::vector> const& see_hitIdx, - std::vector const& ph2_detId, - std::vector const& ph2_x, - std::vector const& ph2_y, - std::vector const& ph2_z) { +void LST::prepareInput(std::vector const& see_px, + std::vector const& see_py, + std::vector const& see_pz, + std::vector const& see_dxy, + std::vector const& see_dz, + std::vector const& see_ptErr, + std::vector const& see_etaErr, + std::vector const& see_stateTrajGlbX, + std::vector const& see_stateTrajGlbY, + std::vector const& see_stateTrajGlbZ, + std::vector const& see_stateTrajGlbPx, + std::vector const& see_stateTrajGlbPy, + std::vector const& see_stateTrajGlbPz, + std::vector const& see_q, + std::vector> const& see_hitIdx, + std::vector const& ph2_detId, + std::vector const& ph2_x, + std::vector const& ph2_y, + std::vector const& ph2_z) { unsigned int count = 0; auto n_see = see_stateTrajGlbPx.size(); std::vector px_vec; @@ -78,7 +78,7 @@ void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::prepareInput(std::vector con std::vector hitIdxs(ph2_detId.size()); std::vector superbin_vec; - std::vector<::lst::PixelType> pixelType_vec; + std::vector pixelType_vec; std::vector isQuad_vec; std::iota(hitIdxs.begin(), hitIdxs.end(), 0); const int hit_size = trkX.size(); @@ -102,15 +102,15 @@ void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::prepareInput(std::vector con float pz = p3LH.z(); int charge = see_q[iSeed]; - ::lst::PixelType pixtype = ::lst::PixelType::kInvalid; + PixelType pixtype = PixelType::kInvalid; if (ptIn >= 2.0) - pixtype = ::lst::PixelType::kHighPt; + pixtype = PixelType::kHighPt; else if (ptIn >= (0.8 - 2 * ptErr) and ptIn < 2.0) { if (pixelSegmentDeltaPhiChange >= 0) - pixtype = ::lst::PixelType::kLowPtPosCurv; + pixtype = PixelType::kLowPtPosCurv; else - pixtype = ::lst::PixelType::kLowPtNegCurv; + pixtype = PixelType::kLowPtNegCurv; } else continue; @@ -212,10 +212,10 @@ void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::prepareInput(std::vector con in_isQuad_vec_ = isQuad_vec; } -std::vector ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::getHitIdxs(short trackCandidateType, - unsigned int TCIdx, - unsigned int const* TCHitIndices, - unsigned int const* hitIndices) { +std::vector LST::getHitIdxs(short trackCandidateType, + unsigned int TCIdx, + unsigned int const* TCHitIndices, + unsigned int const* hitIndices) { std::vector hits; unsigned int maxNHits = 0; @@ -247,7 +247,7 @@ std::vector ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::getHitIdxs(sho return hits; } -void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::getOutput(ALPAKA_ACCELERATOR_NAMESPACE::lst::Event& event) { +void LST::getOutput(Event& event) { std::vector> tc_hitIdxs; std::vector tc_len; std::vector tc_seedIdx; @@ -275,31 +275,31 @@ void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::getOutput(ALPAKA_ACCELERATOR_NAMESP out_tc_trackCandidateType_ = tc_trackCandidateType; } -void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::run(Queue& queue, - bool verbose, - LSTESData const* deviceESData, - std::vector const& see_px, - std::vector const& see_py, - std::vector const& see_pz, - std::vector const& see_dxy, - std::vector const& see_dz, - std::vector const& see_ptErr, - std::vector const& see_etaErr, - std::vector const& see_stateTrajGlbX, - std::vector const& see_stateTrajGlbY, - std::vector const& see_stateTrajGlbZ, - std::vector const& see_stateTrajGlbPx, - std::vector const& see_stateTrajGlbPy, - std::vector const& see_stateTrajGlbPz, - std::vector const& see_q, - std::vector> const& see_hitIdx, - std::vector const& ph2_detId, - std::vector const& ph2_x, - std::vector const& ph2_y, - std::vector const& ph2_z, - bool no_pls_dupclean, - bool tc_pls_triplets) { - auto event = ALPAKA_ACCELERATOR_NAMESPACE::lst::Event(verbose, queue, deviceESData); +void LST::run(Queue& queue, + bool verbose, + LSTESData const* deviceESData, + std::vector const& see_px, + std::vector const& see_py, + std::vector const& see_pz, + std::vector const& see_dxy, + std::vector const& see_dz, + std::vector const& see_ptErr, + std::vector const& see_etaErr, + std::vector const& see_stateTrajGlbX, + std::vector const& see_stateTrajGlbY, + std::vector const& see_stateTrajGlbZ, + std::vector const& see_stateTrajGlbPx, + std::vector const& see_stateTrajGlbPy, + std::vector const& see_stateTrajGlbPz, + std::vector const& see_q, + std::vector> const& see_hitIdx, + std::vector const& ph2_detId, + std::vector const& ph2_x, + std::vector const& ph2_y, + std::vector const& ph2_z, + bool no_pls_dupclean, + bool tc_pls_triplets) { + auto event = Event(verbose, queue, deviceESData); prepareInput(see_px, see_py, see_pz, diff --git a/RecoTracker/LSTCore/src/alpaka/MiniDoublet.h b/RecoTracker/LSTCore/src/alpaka/MiniDoublet.h index 335ceeea2ab79..27ce7b97bffdd 100644 --- a/RecoTracker/LSTCore/src/alpaka/MiniDoublet.h +++ b/RecoTracker/LSTCore/src/alpaka/MiniDoublet.h @@ -209,8 +209,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { mdsInGPU.moduleIndices[idx] = lowerModuleIdx; unsigned int anchorHitIndex, outerHitIndex; - if (modulesInGPU.moduleType[lowerModuleIdx] == ::lst::PS and - modulesInGPU.moduleLayerType[lowerModuleIdx] == ::lst::Strip) { + if (modulesInGPU.moduleType[lowerModuleIdx] == PS and modulesInGPU.moduleLayerType[lowerModuleIdx] == Strip) { mdsInGPU.anchorHitIndices[idx] = upperHitIdx; mdsInGPU.outerHitIndices[idx] = lowerHitIdx; @@ -271,10 +270,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { short side = modulesInGPU.sides[moduleIndex]; short rod = modulesInGPU.rods[moduleIndex]; - if (subdet == ::lst::Barrel) { - if ((side != ::lst::Center and layer == 3) or (side == ::lst::NegZ and layer == 2 and rod > 5) or - (side == ::lst::PosZ and layer == 2 and rod < 8) or (side == ::lst::NegZ and layer == 1 and rod > 9) or - (side == ::lst::PosZ and layer == 1 and rod < 4)) + if (subdet == Barrel) { + if ((side != Center and layer == 3) or (side == NegZ and layer == 2 and rod > 5) or + (side == PosZ and layer == 2 and rod < 8) or (side == NegZ and layer == 1 and rod > 9) or + (side == PosZ and layer == 1 and rod < 4)) return true; else return false; @@ -319,11 +318,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float moduleSeparation = 0; - if (subdet == ::lst::Barrel and side == ::lst::Center) { + if (subdet == Barrel and side == Center) { moduleSeparation = miniDeltaFlat[iL]; } else if (isTighterTiltedModules(modulesInGPU, moduleIndex)) { moduleSeparation = miniDeltaTilted[iL]; - } else if (subdet == ::lst::Endcap) { + } else if (subdet == Endcap) { moduleSeparation = miniDeltaEndcap[iL][iR]; } else //Loose tilted modules { @@ -348,19 +347,16 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int iL = modulesInGPU.layers[moduleIndex] - 1; const float miniSlope = alpaka::math::asin(acc, alpaka::math::min(acc, rt * k2Rinv1GeVf / ptCut, kSinAlphaMax)); const float rLayNominal = - ((modulesInGPU.subdets[moduleIndex] == ::lst::Barrel) ? kMiniRminMeanBarrel[iL] : kMiniRminMeanEndcap[iL]); + ((modulesInGPU.subdets[moduleIndex] == Barrel) ? kMiniRminMeanBarrel[iL] : kMiniRminMeanEndcap[iL]); const float miniPVoff = 0.1f / rLayNominal; - const float miniMuls = - ((modulesInGPU.subdets[moduleIndex] == ::lst::Barrel) ? kMiniMulsPtScaleBarrel[iL] * 3.f / ptCut - : kMiniMulsPtScaleEndcap[iL] * 3.f / ptCut); - const bool isTilted = - modulesInGPU.subdets[moduleIndex] == ::lst::Barrel and modulesInGPU.sides[moduleIndex] != ::lst::Center; + const float miniMuls = ((modulesInGPU.subdets[moduleIndex] == Barrel) ? kMiniMulsPtScaleBarrel[iL] * 3.f / ptCut + : kMiniMulsPtScaleEndcap[iL] * 3.f / ptCut); + const bool isTilted = modulesInGPU.subdets[moduleIndex] == Barrel and modulesInGPU.sides[moduleIndex] != Center; //the lower module is sent in irrespective of its layer type. We need to fetch the drdz properly float drdz; if (isTilted) { - if (modulesInGPU.moduleType[moduleIndex] == ::lst::PS and - modulesInGPU.moduleLayerType[moduleIndex] == ::lst::Strip) { + if (modulesInGPU.moduleType[moduleIndex] == PS and modulesInGPU.moduleLayerType[moduleIndex] == Strip) { drdz = modulesInGPU.drdzs[moduleIndex]; } else { drdz = modulesInGPU.drdzs[modulesInGPU.partnerModuleIndices[moduleIndex]]; @@ -379,12 +375,12 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { // Return the threshold value // ================================================================= // Following condition is met if the module is central and flatly lying - if (modulesInGPU.subdets[moduleIndex] == ::lst::Barrel and modulesInGPU.sides[moduleIndex] == ::lst::Center) { + if (modulesInGPU.subdets[moduleIndex] == Barrel and modulesInGPU.sides[moduleIndex] == Center) { return miniSlope + alpaka::math::sqrt(acc, miniMuls * miniMuls + miniPVoff * miniPVoff); } // Following condition is met if the module is central and tilted - else if (modulesInGPU.subdets[moduleIndex] == ::lst::Barrel and - modulesInGPU.sides[moduleIndex] != ::lst::Center) //all types of tilted modules + else if (modulesInGPU.subdets[moduleIndex] == Barrel and + modulesInGPU.sides[moduleIndex] != Center) //all types of tilted modules { return miniSlope + alpaka::math::sqrt(acc, miniMuls * miniMuls + miniPVoff * miniPVoff + miniTilt2 * miniSlope * miniSlope); @@ -453,11 +449,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float absdzprime; // The distance between the two points after shifting const float& drdz_ = modulesInGPU.drdzs[lowerModuleIndex]; // Assign hit pointers based on their hit type - if (modulesInGPU.moduleType[lowerModuleIndex] == ::lst::PS) { + if (modulesInGPU.moduleType[lowerModuleIndex] == PS) { // TODO: This is somewhat of an mystery.... somewhat confused why this is the case - if (modulesInGPU.subdets[lowerModuleIndex] == ::lst::Barrel - ? modulesInGPU.moduleLayerType[lowerModuleIndex] != ::lst::Pixel - : modulesInGPU.moduleLayerType[lowerModuleIndex] == ::lst::Pixel) { + if (modulesInGPU.subdets[lowerModuleIndex] == Barrel ? modulesInGPU.moduleLayerType[lowerModuleIndex] != Pixel + : modulesInGPU.moduleLayerType[lowerModuleIndex] == Pixel) { xo = xUpper; yo = yUpper; xp = xLower; @@ -482,7 +477,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } // If it is endcap some of the math gets simplified (and also computers don't like infinities) - isEndcap = modulesInGPU.subdets[lowerModuleIndex] == ::lst::Endcap; + isEndcap = modulesInGPU.subdets[lowerModuleIndex] == Endcap; // NOTE: TODO: Keep in mind that the sin(atan) function can be simplified to something like x / sqrt(1 + x^2) and similar for cos // I am not sure how slow sin, atan, cos, functions are in c++. If x / sqrt(1 + x^2) are faster change this later to reduce arithmetic computation time @@ -497,8 +492,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { moduleSeparation = moduleGapSize(modulesInGPU, lowerModuleIndex); // Sign flips if the pixel is later layer - if (modulesInGPU.moduleType[lowerModuleIndex] == ::lst::PS and - modulesInGPU.moduleLayerType[lowerModuleIndex] != ::lst::Pixel) { + if (modulesInGPU.moduleType[lowerModuleIndex] == PS and modulesInGPU.moduleLayerType[lowerModuleIndex] != Pixel) { moduleSeparation *= -1; } @@ -550,7 +544,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { angleA)); // module separation sign is for shifting in radial direction for z-axis direction take care of the sign later // Depending on which one as closer to the interactin point compute the new z wrt to the pixel properly - if (modulesInGPU.moduleLayerType[lowerModuleIndex] == ::lst::Pixel) { + if (modulesInGPU.moduleLayerType[lowerModuleIndex] == Pixel) { abszn = alpaka::math::abs(acc, zp) + absdzprime; } else { abszn = alpaka::math::abs(acc, zp) - absdzprime; @@ -587,7 +581,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float zUpper, float rtUpper) { dz = zLower - zUpper; - const float dzCut = modulesInGPU.moduleType[lowerModuleIndex] == ::lst::PS ? 2.f : 10.f; + const float dzCut = modulesInGPU.moduleType[lowerModuleIndex] == PS ? 2.f : 10.f; const float sign = ((dz > 0) - (dz < 0)) * ((zLower > 0) - (zLower < 0)); const float invertedcrossercut = (alpaka::math::abs(acc, dz) > 2) * sign; @@ -596,7 +590,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float miniCut = 0; - miniCut = modulesInGPU.moduleLayerType[lowerModuleIndex] == ::lst::Pixel + miniCut = modulesInGPU.moduleLayerType[lowerModuleIndex] == Pixel ? dPhiThreshold(acc, rtLower, modulesInGPU, lowerModuleIndex) : dPhiThreshold(acc, rtUpper, modulesInGPU, lowerModuleIndex); @@ -604,7 +598,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { // Ref to original code: https://github.com/slava77/cms-tkph2-ntuple/blob/184d2325147e6930030d3d1f780136bc2dd29ce6/doubletAnalysis.C#L3085 float xn = 0.f, yn = 0.f; // , zn = 0; float shiftedRt2; - if (modulesInGPU.sides[lowerModuleIndex] != ::lst::Center) // If barrel and not center it is tilted + if (modulesInGPU.sides[lowerModuleIndex] != Center) // If barrel and not center it is tilted { // Shift the hits and calculate new xn, yn position float shiftedCoords[3]; @@ -627,7 +621,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { yn = shiftedCoords[1]; // Lower or the upper hit needs to be modified depending on which one was actually shifted - if (modulesInGPU.moduleLayerType[lowerModuleIndex] == ::lst::Pixel) { + if (modulesInGPU.moduleLayerType[lowerModuleIndex] == Pixel) { shiftedX = xn; shiftedY = yn; shiftedZ = zUpper; @@ -656,10 +650,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { // Cut #3: The dphi change going from lower Hit to upper Hit // Ref to original code: https://github.com/slava77/cms-tkph2-ntuple/blob/184d2325147e6930030d3d1f780136bc2dd29ce6/doubletAnalysis.C#L3076 - if (modulesInGPU.sides[lowerModuleIndex] != ::lst::Center) { + if (modulesInGPU.sides[lowerModuleIndex] != Center) { // When it is tilted, use the new shifted positions // TODO: This is somewhat of an mystery.... somewhat confused why this is the case - if (modulesInGPU.moduleLayerType[lowerModuleIndex] != ::lst::Pixel) { + if (modulesInGPU.moduleLayerType[lowerModuleIndex] != Pixel) { // dPhi Change should be calculated so that the upper hit has higher rt. // In principle, this kind of check rt_lower < rt_upper should not be necessary because the hit shifting should have taken care of this. // (i.e. the strip hit is shifted to be aligned in the line of sight from interaction point to pixel hit of PS module guaranteeing rt ordering) @@ -726,7 +720,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { return false; // Cut #2 : drt cut. The dz difference can't be larger than 1cm. (max separation is 4mm for modules in the endcap) // Ref to original code: https://github.com/slava77/cms-tkph2-ntuple/blob/184d2325147e6930030d3d1f780136bc2dd29ce6/doubletAnalysis.C#L3100 - const float drtCut = modulesInGPU.moduleType[lowerModuleIndex] == ::lst::PS ? 2.f : 10.f; + const float drtCut = modulesInGPU.moduleType[lowerModuleIndex] == PS ? 2.f : 10.f; drt = rtLower - rtUpper; if (alpaka::math::abs(acc, drt) >= drtCut) return false; @@ -754,9 +748,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { yn = shiftedCoords[1]; zn = shiftedCoords[2]; - if (modulesInGPU.moduleType[lowerModuleIndex] == ::lst::PS) { + if (modulesInGPU.moduleType[lowerModuleIndex] == PS) { // Appropriate lower or upper hit is modified after checking which one was actually shifted - if (modulesInGPU.moduleLayerType[lowerModuleIndex] == ::lst::Pixel) { + if (modulesInGPU.moduleLayerType[lowerModuleIndex] == Pixel) { shiftedX = xn; shiftedY = yn; shiftedZ = zUpper; @@ -779,12 +773,12 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { // dz needs to change if it is a PS module where the strip hits are shifted in order to properly account for the case when a tilted module falls under "endcap logic" // if it was an endcap it will have zero effect - if (modulesInGPU.moduleType[lowerModuleIndex] == ::lst::PS) { - dz = modulesInGPU.moduleLayerType[lowerModuleIndex] == ::lst::Pixel ? zLower - zn : zUpper - zn; + if (modulesInGPU.moduleType[lowerModuleIndex] == PS) { + dz = modulesInGPU.moduleLayerType[lowerModuleIndex] == Pixel ? zLower - zn : zUpper - zn; } float miniCut = 0; - miniCut = modulesInGPU.moduleLayerType[lowerModuleIndex] == ::lst::Pixel + miniCut = modulesInGPU.moduleLayerType[lowerModuleIndex] == Pixel ? dPhiThreshold(acc, rtLower, modulesInGPU, lowerModuleIndex, dPhi, dz) : dPhiThreshold(acc, rtUpper, modulesInGPU, lowerModuleIndex, dPhi, dz); @@ -824,7 +818,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float yUpper, float zUpper, float rtUpper) { - if (modulesInGPU.subdets[lowerModuleIndex] == ::lst::Barrel) { + if (modulesInGPU.subdets[lowerModuleIndex] == Barrel) { return runMiniDoubletDefaultAlgoBarrel(acc, modulesInGPU, lowerModuleIndex, diff --git a/RecoTracker/LSTCore/src/alpaka/PixelQuintuplet.h b/RecoTracker/LSTCore/src/alpaka/PixelQuintuplet.h index 1ecc256887c77..180d8acf88bae 100644 --- a/RecoTracker/LSTCore/src/alpaka/PixelQuintuplet.h +++ b/RecoTracker/LSTCore/src/alpaka/PixelQuintuplet.h @@ -209,26 +209,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t lowerModuleIndex4, uint16_t lowerModuleIndex5, float rzChiSquared) { - const int layer1 = modulesInGPU.layers[lowerModuleIndex1] + - 6 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex1] == ::lst::TwoS); - const int layer2 = modulesInGPU.layers[lowerModuleIndex2] + - 6 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex2] == ::lst::TwoS); - const int layer3 = modulesInGPU.layers[lowerModuleIndex3] + - 6 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex3] == ::lst::TwoS); - const int layer4 = modulesInGPU.layers[lowerModuleIndex4] + - 6 * (modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex4] == ::lst::TwoS); - const int layer5 = modulesInGPU.layers[lowerModuleIndex5] + - 6 * (modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex5] == ::lst::TwoS); + const int layer1 = + modulesInGPU.layers[lowerModuleIndex1] + 6 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap and modulesInGPU.moduleType[lowerModuleIndex1] == TwoS); + const int layer2 = + modulesInGPU.layers[lowerModuleIndex2] + 6 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap and modulesInGPU.moduleType[lowerModuleIndex2] == TwoS); + const int layer3 = + modulesInGPU.layers[lowerModuleIndex3] + 6 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.moduleType[lowerModuleIndex3] == TwoS); + const int layer4 = + modulesInGPU.layers[lowerModuleIndex4] + 6 * (modulesInGPU.subdets[lowerModuleIndex4] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex4] == Endcap and modulesInGPU.moduleType[lowerModuleIndex4] == TwoS); + const int layer5 = + modulesInGPU.layers[lowerModuleIndex5] + 6 * (modulesInGPU.subdets[lowerModuleIndex5] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex5] == Endcap and modulesInGPU.moduleType[lowerModuleIndex5] == TwoS); if (layer1 == 1 and layer2 == 2 and layer3 == 3) { if (layer4 == 12 and layer5 == 13) { @@ -299,26 +294,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t lowerModuleIndex4, uint16_t lowerModuleIndex5, float rPhiChiSquared) { - const int layer1 = modulesInGPU.layers[lowerModuleIndex1] + - 6 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex1] == ::lst::TwoS); - const int layer2 = modulesInGPU.layers[lowerModuleIndex2] + - 6 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex2] == ::lst::TwoS); - const int layer3 = modulesInGPU.layers[lowerModuleIndex3] + - 6 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex3] == ::lst::TwoS); - const int layer4 = modulesInGPU.layers[lowerModuleIndex4] + - 6 * (modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex4] == ::lst::TwoS); - const int layer5 = modulesInGPU.layers[lowerModuleIndex5] + - 6 * (modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex5] == ::lst::TwoS); + const int layer1 = + modulesInGPU.layers[lowerModuleIndex1] + 6 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap and modulesInGPU.moduleType[lowerModuleIndex1] == TwoS); + const int layer2 = + modulesInGPU.layers[lowerModuleIndex2] + 6 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap and modulesInGPU.moduleType[lowerModuleIndex2] == TwoS); + const int layer3 = + modulesInGPU.layers[lowerModuleIndex3] + 6 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.moduleType[lowerModuleIndex3] == TwoS); + const int layer4 = + modulesInGPU.layers[lowerModuleIndex4] + 6 * (modulesInGPU.subdets[lowerModuleIndex4] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex4] == Endcap and modulesInGPU.moduleType[lowerModuleIndex4] == TwoS); + const int layer5 = + modulesInGPU.layers[lowerModuleIndex5] + 6 * (modulesInGPU.subdets[lowerModuleIndex5] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex5] == Endcap and modulesInGPU.moduleType[lowerModuleIndex5] == TwoS); if (layer1 == 1 and layer2 == 2 and layer3 == 3) { if (layer4 == 12 and layer5 == 13) { @@ -446,7 +436,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { need not always be a PS strip module, but all non-anchor hits sit on strip modules. */ - ::lst::ModuleType moduleType; + ModuleType moduleType; short moduleSubdet, moduleSide; float inv1 = kWidthPS / kWidth2S; float inv2 = kPixelPSZpitch / kWidth2S; @@ -458,21 +448,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { const float& drdz = modulesInGPU.drdzs[lowerModuleIndices[i]]; slopes[i] = modulesInGPU.dxdys[lowerModuleIndices[i]]; //category 1 - barrel PS flat - if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::PS and moduleSide == ::lst::Center) { + if (moduleSubdet == Barrel and moduleType == PS and moduleSide == Center) { delta1[i] = inv1; delta2[i] = inv1; slopes[i] = -999.f; isFlat[i] = true; } //category 2 - barrel 2S - else if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::TwoS) { + else if (moduleSubdet == Barrel and moduleType == TwoS) { delta1[i] = 1.f; delta2[i] = 1.f; slopes[i] = -999.f; isFlat[i] = true; } //category 3 - barrel PS tilted - else if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::PS and moduleSide != ::lst::Center) { + else if (moduleSubdet == Barrel and moduleType == PS and moduleSide != Center) { delta1[i] = inv1; isFlat[i] = false; @@ -483,7 +473,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } } //category 4 - endcap PS - else if (moduleSubdet == ::lst::Endcap and moduleType == ::lst::PS) { + else if (moduleSubdet == Endcap and moduleType == PS) { delta1[i] = inv1; isFlat[i] = false; /* @@ -498,7 +488,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } } //category 5 - endcap 2S - else if (moduleSubdet == ::lst::Endcap and moduleType == ::lst::TwoS) { + else if (moduleSubdet == Endcap and moduleType == TwoS) { delta1[i] = 1.f; delta2[i] = 500.f * inv1; isFlat[i] = false; @@ -559,26 +549,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t lowerModuleIndex4, uint16_t lowerModuleIndex5, float rPhiChiSquared) { - const int layer1 = modulesInGPU.layers[lowerModuleIndex1] + - 6 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex1] == ::lst::TwoS); - const int layer2 = modulesInGPU.layers[lowerModuleIndex2] + - 6 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex2] == ::lst::TwoS); - const int layer3 = modulesInGPU.layers[lowerModuleIndex3] + - 6 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex3] == ::lst::TwoS); - const int layer4 = modulesInGPU.layers[lowerModuleIndex4] + - 6 * (modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex4] == ::lst::TwoS); - const int layer5 = modulesInGPU.layers[lowerModuleIndex5] + - 6 * (modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex5] == ::lst::TwoS); + const int layer1 = + modulesInGPU.layers[lowerModuleIndex1] + 6 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap and modulesInGPU.moduleType[lowerModuleIndex1] == TwoS); + const int layer2 = + modulesInGPU.layers[lowerModuleIndex2] + 6 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap and modulesInGPU.moduleType[lowerModuleIndex2] == TwoS); + const int layer3 = + modulesInGPU.layers[lowerModuleIndex3] + 6 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.moduleType[lowerModuleIndex3] == TwoS); + const int layer4 = + modulesInGPU.layers[lowerModuleIndex4] + 6 * (modulesInGPU.subdets[lowerModuleIndex4] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex4] == Endcap and modulesInGPU.moduleType[lowerModuleIndex4] == TwoS); + const int layer5 = + modulesInGPU.layers[lowerModuleIndex5] + 6 * (modulesInGPU.subdets[lowerModuleIndex5] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex5] == Endcap and modulesInGPU.moduleType[lowerModuleIndex5] == TwoS); if (layer1 == 1 and layer2 == 2 and layer3 == 3) { if (layer4 == 12 and layer5 == 13) { @@ -664,8 +649,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { const int moduleSide = modulesInGPU.sides[lowerModuleIndex]; const int moduleSubdet = modulesInGPU.subdets[lowerModuleIndex]; - residual = (moduleSubdet == ::lst::Barrel) ? (zs[i] - zPix[0]) - slope * (rts[i] - rtPix[0]) - : (rts[i] - rtPix[0]) - (zs[i] - zPix[0]) / slope; + residual = (moduleSubdet == Barrel) ? (zs[i] - zPix[0]) - slope * (rts[i] - rtPix[0]) + : (rts[i] - rtPix[0]) - (zs[i] - zPix[0]) / slope; const float& drdz = modulesInGPU.drdzs[lowerModuleIndex]; //PS Modules if (moduleType == 0) { @@ -676,7 +661,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } //special dispensation to tilted PS modules! - if (moduleType == 0 and moduleSubdet == ::lst::Barrel and moduleSide != ::lst::Center) { + if (moduleType == 0 and moduleSubdet == Barrel and moduleSide != Center) { error2 /= (1.f + drdz * drdz); } RMSE += (residual * residual) / error2; @@ -858,7 +843,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t quintupletLowerModuleIndex = modulesInGPU.connectedPixels[iLSModule]; if (quintupletLowerModuleIndex >= *modulesInGPU.nLowerModules) continue; - if (modulesInGPU.moduleType[quintupletLowerModuleIndex] == ::lst::TwoS) + if (modulesInGPU.moduleType[quintupletLowerModuleIndex] == TwoS) continue; uint16_t pixelModuleIndex = *modulesInGPU.nLowerModules; if (segmentsInGPU.isDup[i_pLS]) diff --git a/RecoTracker/LSTCore/src/alpaka/PixelTriplet.h b/RecoTracker/LSTCore/src/alpaka/PixelTriplet.h index 710c760fb809f..f7f7c4da72a51 100644 --- a/RecoTracker/LSTCore/src/alpaka/PixelTriplet.h +++ b/RecoTracker/LSTCore/src/alpaka/PixelTriplet.h @@ -228,8 +228,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int thirdMDIndex = segmentsInGPU.mdIndices[Params_LS::kLayers * outerSegmentIndex]; unsigned int fourthMDIndex = segmentsInGPU.mdIndices[Params_LS::kLayers * outerSegmentIndex + 1]; - if (outerInnerLowerModuleSubdet == ::lst::Barrel and - (outerOuterLowerModuleSubdet == ::lst::Barrel or outerOuterLowerModuleSubdet == ::lst::Endcap)) { + if (outerInnerLowerModuleSubdet == Barrel and + (outerOuterLowerModuleSubdet == Barrel or outerOuterLowerModuleSubdet == Endcap)) { return runTripletDefaultAlgoPPBB(acc, modulesInGPU, rangesInGPU, @@ -244,7 +244,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { secondMDIndex, thirdMDIndex, fourthMDIndex); - } else if (outerInnerLowerModuleSubdet == ::lst::Endcap and outerOuterLowerModuleSubdet == ::lst::Endcap) { + } else if (outerInnerLowerModuleSubdet == Endcap and outerOuterLowerModuleSubdet == Endcap) { return runTripletDefaultAlgoPPEE(acc, modulesInGPU, rangesInGPU, @@ -268,18 +268,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t lowerModuleIndex2, uint16_t lowerModuleIndex3, float rzChiSquared) { - const int layer1 = modulesInGPU.layers[lowerModuleIndex1] + - 6 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex1] == ::lst::TwoS); - const int layer2 = modulesInGPU.layers[lowerModuleIndex2] + - 6 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex2] == ::lst::TwoS); - const int layer3 = modulesInGPU.layers[lowerModuleIndex3] + - 6 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex3] == ::lst::TwoS); + const int layer1 = + modulesInGPU.layers[lowerModuleIndex1] + 6 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap and modulesInGPU.moduleType[lowerModuleIndex1] == TwoS); + const int layer2 = + modulesInGPU.layers[lowerModuleIndex2] + 6 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap and modulesInGPU.moduleType[lowerModuleIndex2] == TwoS); + const int layer3 = + modulesInGPU.layers[lowerModuleIndex3] + 6 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.moduleType[lowerModuleIndex3] == TwoS); if (layer1 == 8 and layer2 == 9 and layer3 == 10) { return rzChiSquared < 13.6067f; @@ -379,33 +376,33 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float inv1 = kWidthPS / kWidth2S; float inv2 = kPixelPSZpitch / kWidth2S; for (size_t i = 0; i < 3; i++) { - ::lst::ModuleType moduleType = modulesInGPU.moduleType[lowerModuleIndices[i]]; + ModuleType moduleType = modulesInGPU.moduleType[lowerModuleIndices[i]]; short moduleSubdet = modulesInGPU.subdets[lowerModuleIndices[i]]; short moduleSide = modulesInGPU.sides[lowerModuleIndices[i]]; float drdz = modulesInGPU.drdzs[lowerModuleIndices[i]]; slopes[i] = modulesInGPU.dxdys[lowerModuleIndices[i]]; //category 1 - barrel PS flat - if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::PS and moduleSide == ::lst::Center) { + if (moduleSubdet == Barrel and moduleType == PS and moduleSide == Center) { delta1[i] = inv1; delta2[i] = inv1; slopes[i] = -999; isFlat[i] = true; } //category 2 - barrel 2S - else if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::TwoS) { + else if (moduleSubdet == Barrel and moduleType == TwoS) { delta1[i] = 1; delta2[i] = 1; slopes[i] = -999; isFlat[i] = true; } //category 3 - barrel PS tilted - else if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::PS and moduleSide != ::lst::Center) { + else if (moduleSubdet == Barrel and moduleType == PS and moduleSide != Center) { delta1[i] = inv1; isFlat[i] = false; delta2[i] = (inv2 * drdz / alpaka::math::sqrt(acc, 1 + drdz * drdz)); } //category 4 - endcap PS - else if (moduleSubdet == ::lst::Endcap and moduleType == ::lst::PS) { + else if (moduleSubdet == Endcap and moduleType == PS) { delta1[i] = inv1; isFlat[i] = false; @@ -416,7 +413,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { delta2[i] = inv2; } //category 5 - endcap 2S - else if (moduleSubdet == ::lst::Endcap and moduleType == ::lst::TwoS) { + else if (moduleSubdet == Endcap and moduleType == TwoS) { delta1[i] = 1; delta2[i] = 500 * inv1; isFlat[i] = false; @@ -452,18 +449,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t lowerModuleIndex2, uint16_t lowerModuleIndex3, float chiSquared) { - const int layer1 = modulesInGPU.layers[lowerModuleIndex1] + - 6 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex1] == ::lst::TwoS); - const int layer2 = modulesInGPU.layers[lowerModuleIndex2] + - 6 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex2] == ::lst::TwoS); - const int layer3 = modulesInGPU.layers[lowerModuleIndex3] + - 6 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex3] == ::lst::TwoS); + const int layer1 = + modulesInGPU.layers[lowerModuleIndex1] + 6 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap and modulesInGPU.moduleType[lowerModuleIndex1] == TwoS); + const int layer2 = + modulesInGPU.layers[lowerModuleIndex2] + 6 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap and modulesInGPU.moduleType[lowerModuleIndex2] == TwoS); + const int layer3 = + modulesInGPU.layers[lowerModuleIndex3] + 6 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.moduleType[lowerModuleIndex3] == TwoS); if (layer1 == 8 and layer2 == 9 and layer3 == 10) { return chiSquared < 7.003f; @@ -499,18 +493,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t lowerModuleIndex2, uint16_t lowerModuleIndex3, float chiSquared) { - const int layer1 = modulesInGPU.layers[lowerModuleIndex1] + - 6 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex1] == ::lst::TwoS); - const int layer2 = modulesInGPU.layers[lowerModuleIndex2] + - 6 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex2] == ::lst::TwoS); - const int layer3 = modulesInGPU.layers[lowerModuleIndex3] + - 6 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap) + - 5 * (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex3] == ::lst::TwoS); + const int layer1 = + modulesInGPU.layers[lowerModuleIndex1] + 6 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex1] == Endcap and modulesInGPU.moduleType[lowerModuleIndex1] == TwoS); + const int layer2 = + modulesInGPU.layers[lowerModuleIndex2] + 6 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex2] == Endcap and modulesInGPU.moduleType[lowerModuleIndex2] == TwoS); + const int layer3 = + modulesInGPU.layers[lowerModuleIndex3] + 6 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap) + + 5 * (modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.moduleType[lowerModuleIndex3] == TwoS); if (layer1 == 7 and layer2 == 8 and layer3 == 9) // endcap layer 1,2,3, ps { @@ -670,11 +661,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { int16_t lowerModuleIndex, uint16_t middleModuleIndex, uint16_t upperModuleIndex) { - if (modulesInGPU.subdets[lowerModuleIndex] == ::lst::Endcap) { + if (modulesInGPU.subdets[lowerModuleIndex] == Endcap) { return passRadiusCriterionEEE(acc, pixelRadius, pixelRadiusError, tripletRadius); - } else if (modulesInGPU.subdets[middleModuleIndex] == ::lst::Endcap) { + } else if (modulesInGPU.subdets[middleModuleIndex] == Endcap) { return passRadiusCriterionBEE(acc, pixelRadius, pixelRadiusError, tripletRadius); - } else if (modulesInGPU.subdets[upperModuleIndex] == ::lst::Endcap) { + } else if (modulesInGPU.subdets[upperModuleIndex] == Endcap) { return passRadiusCriterionBBE(acc, pixelRadius, pixelRadiusError, tripletRadius); } else { return passRadiusCriterionBBB(acc, pixelRadius, pixelRadiusError, tripletRadius); @@ -724,14 +715,14 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float p = alpaka::math::sqrt(acc, Px * Px + Py * Py + Pz * Pz); float rou = a / p; - if (moduleSubdet == ::lst::Endcap) { + if (moduleSubdet == Endcap) { float s = (zsi - z1) * p / Pz; float x = x1 + Px / a * alpaka::math::sin(acc, rou * s) - Py / a * (1 - alpaka::math::cos(acc, rou * s)); float y = y1 + Py / a * alpaka::math::sin(acc, rou * s) + Px / a * (1 - alpaka::math::cos(acc, rou * s)); diffr = alpaka::math::abs(acc, rtsi - alpaka::math::sqrt(acc, x * x + y * y)) * 100; } - if (moduleSubdet == ::lst::Barrel) { + if (moduleSubdet == Barrel) { float paraA = r1 * r1 + 2 * (Px * Px + Py * Py) / (a * a) + 2 * (y1 * Px - x1 * Py) / a - rtsi * rtsi; float paraB = 2 * (x1 * Px + y1 * Py) / a; float paraC = 2 * (y1 * Px - x1 * Py) / a + 2 * (Px * Px + Py * Py) / (a * a); @@ -747,7 +738,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { diffz = alpaka::math::min(acc, diffz1, diffz2); } - residual = moduleSubdet == ::lst::Barrel ? diffz : diffr; + residual = moduleSubdet == Barrel ? diffz : diffr; //PS Modules if (moduleType == 0) { @@ -758,7 +749,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } //special dispensation to tilted PS modules! - if (moduleType == 0 and moduleSubdet == ::lst::Barrel and moduleSide != ::lst::Center) { + if (moduleType == 0 and moduleSubdet == Barrel and moduleSide != Center) { float drdz = modulesInGPU.drdzs[lowerModuleIndex]; error2 /= (1 + drdz * drdz); } @@ -959,7 +950,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } #endif //Removes 2S-2S :FIXME: filter these out in the pixel map - if (modulesInGPU.moduleType[tripletLowerModuleIndex] == ::lst::TwoS) + if (modulesInGPU.moduleType[tripletLowerModuleIndex] == TwoS) continue; uint16_t pixelModuleIndex = *modulesInGPU.nLowerModules; @@ -990,7 +981,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { outerTripletArrayIndex += gridThreadExtent[2]) { unsigned int outerTripletIndex = rangesInGPU.tripletModuleIndices[tripletLowerModuleIndex] + outerTripletArrayIndex; - if (modulesInGPU.moduleType[tripletsInGPU.lowerModuleIndices[3 * outerTripletIndex + 1]] == ::lst::TwoS) + if (modulesInGPU.moduleType[tripletsInGPU.lowerModuleIndices[3 * outerTripletIndex + 1]] == TwoS) continue; //REMOVES PS-2S if (tripletsInGPU.partOfPT5[outerTripletIndex]) @@ -1178,7 +1169,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int fourthMDIndex) { float dPhi, betaIn, betaOut, pt_beta, zLo, zHi, zLoPointed, zHiPointed, dPhiCut, betaOutCut; - bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == ::lst::PS); + bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == PS); float rt_InLo = mdsInGPU.anchorRt[firstMDIndex]; float rt_InUp = mdsInGPU.anchorRt[secondMDIndex]; @@ -1284,8 +1275,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float alpha_InLo = __H2F(segmentsInGPU.dPhiChanges[innerSegmentIndex]); float alpha_OutLo = __H2F(segmentsInGPU.dPhiChanges[outerSegmentIndex]); - bool isEC_lastLayer = modulesInGPU.subdets[outerOuterLowerModuleIndex] == ::lst::Endcap and - modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::TwoS; + bool isEC_lastLayer = modulesInGPU.subdets[outerOuterLowerModuleIndex] == Endcap and + modulesInGPU.moduleType[outerOuterLowerModuleIndex] == TwoS; float alpha_OutUp, alpha_OutUp_highEdge, alpha_OutUp_lowEdge; alpha_OutUp = deltaPhi(acc, x_OutUp, y_OutUp, x_OutUp - x_OutLo, y_OutUp - y_OutLo); @@ -1436,7 +1427,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int fourthMDIndex) { float dPhi, betaIn, betaOut, pt_beta, rtLo, rtHi, dPhiCut, betaOutCut; - bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == ::lst::PS); + bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == PS); float z_InUp = mdsInGPU.anchorZ[secondMDIndex]; float z_OutLo = mdsInGPU.anchorZ[thirdMDIndex]; @@ -1480,7 +1471,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { const float dzDrtScale = alpaka::math::tan(acc, slope) / slope; //FIXME: need approximate value const float dLum = alpaka::math::copysign(acc, kDeltaZLum, z_InUp); - bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerInnerLowerModuleIndex] == ::lst::PS; + bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerInnerLowerModuleIndex] == PS; const float rtGeom1 = isOutSgInnerMDPS ? kPixelPSZpitch @@ -1546,8 +1537,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float alpha_InLo = __H2F(segmentsInGPU.dPhiChanges[innerSegmentIndex]); float alpha_OutLo = __H2F(segmentsInGPU.dPhiChanges[outerSegmentIndex]); - bool isEC_lastLayer = modulesInGPU.subdets[outerOuterLowerModuleIndex] == ::lst::Endcap and - modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::TwoS; + bool isEC_lastLayer = modulesInGPU.subdets[outerOuterLowerModuleIndex] == Endcap and + modulesInGPU.moduleType[outerOuterLowerModuleIndex] == TwoS; float alpha_OutUp, alpha_OutUp_highEdge, alpha_OutUp_lowEdge; diff --git a/RecoTracker/LSTCore/src/alpaka/Quintuplet.h b/RecoTracker/LSTCore/src/alpaka/Quintuplet.h index 4ff67d66d2844..1b75100c874e8 100644 --- a/RecoTracker/LSTCore/src/alpaka/Quintuplet.h +++ b/RecoTracker/LSTCore/src/alpaka/Quintuplet.h @@ -586,14 +586,13 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { subdets = modulesInGPU.subdets[lowerModuleIndex3]; } if (i == 2 || i == 3) { - residual = (layeri <= 6 && ((side == ::lst::Center) or (drdz < 1))) ? diffz : diffr; + residual = (layeri <= 6 && ((side == Center) or (drdz < 1))) ? diffz : diffr; float projection_missing2 = 1.f; if (drdz < 1) - projection_missing2 = ((subdets == ::lst::Endcap) or (side == ::lst::Center)) - ? 1.f - : 1.f / (1 + drdz * drdz); // cos(atan(drdz)), if dr/dz<1 + projection_missing2 = + ((subdets == Endcap) or (side == Center)) ? 1.f : 1.f / (1 + drdz * drdz); // cos(atan(drdz)), if dr/dz<1 if (drdz > 1) - projection_missing2 = ((subdets == ::lst::Endcap) or (side == ::lst::Center)) + projection_missing2 = ((subdets == Endcap) or (side == Center)) ? 1.f : (drdz * drdz) / (1 + drdz * drdz); //sin(atan(drdz)), if dr/dz>1 error2 = error2 * projection_missing2; @@ -1029,7 +1028,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { modules. */ - ::lst::ModuleType moduleType; + ModuleType moduleType; short moduleSubdet, moduleSide; float inv1 = kWidthPS / kWidth2S; float inv2 = kPixelPSZpitch / kWidth2S; @@ -1041,21 +1040,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { const float& drdz = modulesInGPU.drdzs[lowerModuleIndices[i]]; slopes[i] = modulesInGPU.dxdys[lowerModuleIndices[i]]; //category 1 - barrel PS flat - if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::PS and moduleSide == ::lst::Center) { + if (moduleSubdet == Barrel and moduleType == PS and moduleSide == Center) { delta1[i] = inv1; delta2[i] = inv1; slopes[i] = -999.f; isFlat[i] = true; } //category 2 - barrel 2S - else if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::TwoS) { + else if (moduleSubdet == Barrel and moduleType == TwoS) { delta1[i] = 1.f; delta2[i] = 1.f; slopes[i] = -999.f; isFlat[i] = true; } //category 3 - barrel PS tilted - else if (moduleSubdet == ::lst::Barrel and moduleType == ::lst::PS and moduleSide != ::lst::Center) { + else if (moduleSubdet == Barrel and moduleType == PS and moduleSide != Center) { delta1[i] = inv1; isFlat[i] = false; @@ -1066,7 +1065,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } } //category 4 - endcap PS - else if (moduleSubdet == ::lst::Endcap and moduleType == ::lst::PS) { + else if (moduleSubdet == Endcap and moduleType == PS) { delta1[i] = inv1; isFlat[i] = false; @@ -1082,7 +1081,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } } //category 5 - endcap 2S - else if (moduleSubdet == ::lst::Endcap and moduleType == ::lst::TwoS) { + else if (moduleSubdet == Endcap and moduleType == TwoS) { delta1[i] = 1.f; delta2[i] = 500.f * inv1; isFlat[i] = false; @@ -1355,8 +1354,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int secondMDIndex, unsigned int thirdMDIndex, unsigned int fourthMDIndex) { - bool isPS_InLo = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == ::lst::PS); - bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == ::lst::PS); + bool isPS_InLo = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == PS); + bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == PS); float rt_InLo = mdsInGPU.anchorRt[firstMDIndex]; float rt_InOut = mdsInGPU.anchorRt[secondMDIndex]; @@ -1435,8 +1434,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float alpha_InLo = __H2F(segmentsInGPU.dPhiChanges[innerSegmentIndex]); float alpha_OutLo = __H2F(segmentsInGPU.dPhiChanges[outerSegmentIndex]); - bool isEC_lastLayer = modulesInGPU.subdets[outerOuterLowerModuleIndex] == ::lst::Endcap and - modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::TwoS; + bool isEC_lastLayer = modulesInGPU.subdets[outerOuterLowerModuleIndex] == Endcap and + modulesInGPU.moduleType[outerOuterLowerModuleIndex] == TwoS; float alpha_OutUp, alpha_OutUp_highEdge, alpha_OutUp_lowEdge; @@ -1600,8 +1599,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int secondMDIndex, unsigned int thirdMDIndex, unsigned int fourthMDIndex) { - bool isPS_InLo = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == ::lst::PS); - bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == ::lst::PS); + bool isPS_InLo = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == PS); + bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == PS); float rt_InLo = mdsInGPU.anchorRt[firstMDIndex]; float rt_InOut = mdsInGPU.anchorRt[secondMDIndex]; @@ -1625,7 +1624,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { return false; float dLum = alpaka::math::copysign(acc, kDeltaZLum, z_InLo); - bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerInnerLowerModuleIndex] == ::lst::PS; + bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerInnerLowerModuleIndex] == PS; float rtGeom1 = isOutSgInnerMDPS ? kPixelPSZpitch : kStrip2SZpitch; float zGeom1 = alpaka::math::copysign(acc, zGeom, z_InLo); float rtLo = rt_InLo * (1.f + (z_OutLo - z_InLo - zGeom1) / (z_InLo + zGeom1 + dLum) / dzDrtScale) - @@ -1716,8 +1715,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float betaOutRHmin = betaOut; float betaOutRHmax = betaOut; - bool isEC_secondLayer = (modulesInGPU.subdets[innerOuterLowerModuleIndex] == ::lst::Endcap) and - (modulesInGPU.moduleType[innerOuterLowerModuleIndex] == ::lst::TwoS); + bool isEC_secondLayer = (modulesInGPU.subdets[innerOuterLowerModuleIndex] == Endcap) and + (modulesInGPU.moduleType[innerOuterLowerModuleIndex] == TwoS); if (isEC_secondLayer) { betaInRHmin = betaIn - sdIn_alpha_min + sdIn_alpha; @@ -1802,7 +1801,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { const float dBetaRIn2 = 0; // TODO-RH float dBetaROut = 0; - if (modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::TwoS) { + if (modulesInGPU.moduleType[outerOuterLowerModuleIndex] == TwoS) { dBetaROut = (alpaka::math::sqrt(acc, mdsInGPU.anchorHighEdgeX[fourthMDIndex] * mdsInGPU.anchorHighEdgeX[fourthMDIndex] + @@ -1866,8 +1865,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { return false; float dLum = alpaka::math::copysign(acc, kDeltaZLum, z_InLo); - bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerInnerLowerModuleIndex] == ::lst::PS; - bool isInSgInnerMDPS = modulesInGPU.moduleType[innerInnerLowerModuleIndex] == ::lst::PS; + bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerInnerLowerModuleIndex] == PS; + bool isInSgInnerMDPS = modulesInGPU.moduleType[innerInnerLowerModuleIndex] == PS; float rtGeom = (isInSgInnerMDPS and isOutSgInnerMDPS) ? 2.f * kPixelPSZpitch : (isInSgInnerMDPS or isOutSgInnerMDPS) ? kPixelPSZpitch + kStrip2SZpitch @@ -1885,7 +1884,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { if ((rtOut < rtLo) || (rtOut > rtHi)) return false; - bool isInSgOuterMDPS = modulesInGPU.moduleType[innerOuterLowerModuleIndex] == ::lst::PS; + bool isInSgOuterMDPS = modulesInGPU.moduleType[innerOuterLowerModuleIndex] == PS; const float drtSDIn = rt_InOut - rt_InLo; const float dzSDIn = z_InOut - z_InLo; @@ -2078,8 +2077,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { short outerInnerLowerModuleSubdet = modulesInGPU.subdets[outerInnerLowerModuleIndex]; short outerOuterLowerModuleSubdet = modulesInGPU.subdets[outerOuterLowerModuleIndex]; - if (innerInnerLowerModuleSubdet == ::lst::Barrel and innerOuterLowerModuleSubdet == ::lst::Barrel and - outerInnerLowerModuleSubdet == ::lst::Barrel and outerOuterLowerModuleSubdet == ::lst::Barrel) { + if (innerInnerLowerModuleSubdet == Barrel and innerOuterLowerModuleSubdet == Barrel and + outerInnerLowerModuleSubdet == Barrel and outerOuterLowerModuleSubdet == Barrel) { return runQuintupletDefaultAlgoBBBB(acc, modulesInGPU, mdsInGPU, @@ -2094,8 +2093,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { secondMDIndex, thirdMDIndex, fourthMDIndex); - } else if (innerInnerLowerModuleSubdet == ::lst::Barrel and innerOuterLowerModuleSubdet == ::lst::Barrel and - outerInnerLowerModuleSubdet == ::lst::Endcap and outerOuterLowerModuleSubdet == ::lst::Endcap) { + } else if (innerInnerLowerModuleSubdet == Barrel and innerOuterLowerModuleSubdet == Barrel and + outerInnerLowerModuleSubdet == Endcap and outerOuterLowerModuleSubdet == Endcap) { return runQuintupletDefaultAlgoBBEE(acc, modulesInGPU, mdsInGPU, @@ -2110,8 +2109,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { secondMDIndex, thirdMDIndex, fourthMDIndex); - } else if (innerInnerLowerModuleSubdet == ::lst::Barrel and innerOuterLowerModuleSubdet == ::lst::Barrel and - outerInnerLowerModuleSubdet == ::lst::Barrel and outerOuterLowerModuleSubdet == ::lst::Endcap) { + } else if (innerInnerLowerModuleSubdet == Barrel and innerOuterLowerModuleSubdet == Barrel and + outerInnerLowerModuleSubdet == Barrel and outerOuterLowerModuleSubdet == Endcap) { return runQuintupletDefaultAlgoBBBB(acc, modulesInGPU, mdsInGPU, @@ -2126,8 +2125,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { secondMDIndex, thirdMDIndex, fourthMDIndex); - } else if (innerInnerLowerModuleSubdet == ::lst::Barrel and innerOuterLowerModuleSubdet == ::lst::Endcap and - outerInnerLowerModuleSubdet == ::lst::Endcap and outerOuterLowerModuleSubdet == ::lst::Endcap) { + } else if (innerInnerLowerModuleSubdet == Barrel and innerOuterLowerModuleSubdet == Endcap and + outerInnerLowerModuleSubdet == Endcap and outerOuterLowerModuleSubdet == Endcap) { return runQuintupletDefaultAlgoBBEE(acc, modulesInGPU, mdsInGPU, @@ -2142,8 +2141,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { secondMDIndex, thirdMDIndex, fourthMDIndex); - } else if (innerInnerLowerModuleSubdet == ::lst::Endcap and innerOuterLowerModuleSubdet == ::lst::Endcap and - outerInnerLowerModuleSubdet == ::lst::Endcap and outerOuterLowerModuleSubdet == ::lst::Endcap) { + } else if (innerInnerLowerModuleSubdet == Endcap and innerOuterLowerModuleSubdet == Endcap and + outerInnerLowerModuleSubdet == Endcap and outerOuterLowerModuleSubdet == Endcap) { return runQuintupletDefaultAlgoEEEE(acc, modulesInGPU, mdsInGPU, @@ -2258,24 +2257,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float x3Vec[] = {x3, x3, x3}; float y3Vec[] = {y3, y3, y3}; - if (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex1] == ::lst::TwoS) { + if (modulesInGPU.subdets[lowerModuleIndex1] == Endcap and modulesInGPU.moduleType[lowerModuleIndex1] == TwoS) { x1Vec[1] = mdsInGPU.anchorLowEdgeX[firstMDIndex]; x1Vec[2] = mdsInGPU.anchorHighEdgeX[firstMDIndex]; y1Vec[1] = mdsInGPU.anchorLowEdgeY[firstMDIndex]; y1Vec[2] = mdsInGPU.anchorHighEdgeY[firstMDIndex]; } - if (modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex2] == ::lst::TwoS) { + if (modulesInGPU.subdets[lowerModuleIndex2] == Endcap and modulesInGPU.moduleType[lowerModuleIndex2] == TwoS) { x2Vec[1] = mdsInGPU.anchorLowEdgeX[secondMDIndex]; x2Vec[2] = mdsInGPU.anchorHighEdgeX[secondMDIndex]; y2Vec[1] = mdsInGPU.anchorLowEdgeY[secondMDIndex]; y2Vec[2] = mdsInGPU.anchorHighEdgeY[secondMDIndex]; } - if (modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex3] == ::lst::TwoS) { + if (modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.moduleType[lowerModuleIndex3] == TwoS) { x3Vec[1] = mdsInGPU.anchorLowEdgeX[thirdMDIndex]; x3Vec[2] = mdsInGPU.anchorHighEdgeX[thirdMDIndex]; @@ -2290,8 +2286,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { x1Vec[i] = x4; y1Vec[i] = y4; } - if (modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex4] == ::lst::TwoS) { + if (modulesInGPU.subdets[lowerModuleIndex4] == Endcap and modulesInGPU.moduleType[lowerModuleIndex4] == TwoS) { x1Vec[1] = mdsInGPU.anchorLowEdgeX[fourthMDIndex]; x1Vec[2] = mdsInGPU.anchorHighEdgeX[fourthMDIndex]; @@ -2306,8 +2301,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { x2Vec[i] = x5; y2Vec[i] = y5; } - if (modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap and - modulesInGPU.moduleType[lowerModuleIndex5] == ::lst::TwoS) { + if (modulesInGPU.subdets[lowerModuleIndex5] == Endcap and modulesInGPU.moduleType[lowerModuleIndex5] == TwoS) { x2Vec[1] = mdsInGPU.anchorLowEdgeX[fifthMDIndex]; x2Vec[2] = mdsInGPU.anchorHighEdgeX[fifthMDIndex]; @@ -2356,23 +2350,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { //split by category bool matchedRadii; - if (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Barrel) { + if (modulesInGPU.subdets[lowerModuleIndex1] == Barrel and modulesInGPU.subdets[lowerModuleIndex2] == Barrel and + modulesInGPU.subdets[lowerModuleIndex3] == Barrel and modulesInGPU.subdets[lowerModuleIndex4] == Barrel and + modulesInGPU.subdets[lowerModuleIndex5] == Barrel) { matchedRadii = matchRadiiBBBBB(acc, innerRadius, bridgeRadius, outerRadius); - } else if (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap) { + } else if (modulesInGPU.subdets[lowerModuleIndex1] == Barrel and + modulesInGPU.subdets[lowerModuleIndex2] == Barrel and + modulesInGPU.subdets[lowerModuleIndex3] == Barrel and + modulesInGPU.subdets[lowerModuleIndex4] == Barrel and + modulesInGPU.subdets[lowerModuleIndex5] == Endcap) { matchedRadii = matchRadiiBBBBE(acc, innerRadius, bridgeRadius, outerRadius); - } else if (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap and - modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap) { + } else if (modulesInGPU.subdets[lowerModuleIndex1] == Barrel and + modulesInGPU.subdets[lowerModuleIndex2] == Barrel and + modulesInGPU.subdets[lowerModuleIndex3] == Barrel and + modulesInGPU.subdets[lowerModuleIndex4] == Endcap and + modulesInGPU.subdets[lowerModuleIndex5] == Endcap) { if (modulesInGPU.layers[lowerModuleIndex1] == 1) { matchedRadii = matchRadiiBBBEE12378(acc, innerRadius, bridgeRadius, outerRadius, bridgeRadiusMin2S, bridgeRadiusMax2S); @@ -2385,17 +2377,15 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } } - else if (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap and - modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap) { + else if (modulesInGPU.subdets[lowerModuleIndex1] == Barrel and modulesInGPU.subdets[lowerModuleIndex2] == Barrel and + modulesInGPU.subdets[lowerModuleIndex3] == Endcap and modulesInGPU.subdets[lowerModuleIndex4] == Endcap and + modulesInGPU.subdets[lowerModuleIndex5] == Endcap) { matchedRadii = matchRadiiBBEEE(acc, innerRadius, bridgeRadius, outerRadius, bridgeRadiusMin2S, bridgeRadiusMax2S); - } else if (modulesInGPU.subdets[lowerModuleIndex1] == ::lst::Barrel and - modulesInGPU.subdets[lowerModuleIndex2] == ::lst::Endcap and - modulesInGPU.subdets[lowerModuleIndex3] == ::lst::Endcap and - modulesInGPU.subdets[lowerModuleIndex4] == ::lst::Endcap and - modulesInGPU.subdets[lowerModuleIndex5] == ::lst::Endcap) { + } else if (modulesInGPU.subdets[lowerModuleIndex1] == Barrel and + modulesInGPU.subdets[lowerModuleIndex2] == Endcap and + modulesInGPU.subdets[lowerModuleIndex3] == Endcap and + modulesInGPU.subdets[lowerModuleIndex4] == Endcap and + modulesInGPU.subdets[lowerModuleIndex5] == Endcap) { matchedRadii = matchRadiiBEEEE(acc, innerRadius, bridgeRadius, @@ -2678,9 +2668,9 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { if (tripletsInGPU.nTriplets[i] == 0) continue; - if (module_subdets == ::lst::Barrel and module_layers >= 3) + if (module_subdets == Barrel and module_layers >= 3) continue; - if (module_subdets == ::lst::Endcap and module_layers > 1) + if (module_subdets == Endcap and module_layers > 1) continue; int nEligibleT5Modules = alpaka::atomicAdd(acc, &nEligibleT5Modulesx, 1, alpaka::hierarchy::Threads{}); diff --git a/RecoTracker/LSTCore/src/alpaka/Segment.h b/RecoTracker/LSTCore/src/alpaka/Segment.h index b6cb96a4dc6b7..bc2d1d82a5fc9 100644 --- a/RecoTracker/LSTCore/src/alpaka/Segment.h +++ b/RecoTracker/LSTCore/src/alpaka/Segment.h @@ -31,7 +31,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int* nSegments; //number of segments per inner lower module unsigned int* totOccupancySegments; //number of segments per inner lower module uint4* pLSHitsIdxs; - ::lst::PixelType* pixelType; + PixelType* pixelType; char* isQuad; char* isDup; bool* partOfPT5; @@ -107,7 +107,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { Buf nSegments_buf; Buf totOccupancySegments_buf; Buf pLSHitsIdxs_buf; - Buf pixelType_buf; + Buf pixelType_buf; Buf isQuad_buf; Buf isDup_buf; Buf partOfPT5_buf; @@ -150,7 +150,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { nSegments_buf(allocBufWrapper(devAccIn, nLowerModules + 1, queue)), totOccupancySegments_buf(allocBufWrapper(devAccIn, nLowerModules + 1, queue)), pLSHitsIdxs_buf(allocBufWrapper(devAccIn, maxPixelSegments, queue)), - pixelType_buf(allocBufWrapper<::lst::PixelType>(devAccIn, maxPixelSegments, queue)), + pixelType_buf(allocBufWrapper(devAccIn, maxPixelSegments, queue)), isQuad_buf(allocBufWrapper(devAccIn, maxPixelSegments, queue)), isDup_buf(allocBufWrapper(devAccIn, maxPixelSegments, queue)), partOfPT5_buf(allocBufWrapper(devAccIn, maxPixelSegments, queue)), @@ -186,20 +186,18 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { short side = modulesInGPU.sides[moduleIndex]; short rod = modulesInGPU.rods[moduleIndex]; - return (subdet == ::lst::Barrel) && - (((side != ::lst::Center) && (layer == 3)) || - ((side == ::lst::NegZ) && (((layer == 2) && (rod > 5)) || ((layer == 1) && (rod > 9)))) || - ((side == ::lst::PosZ) && (((layer == 2) && (rod < 8)) || ((layer == 1) && (rod < 4))))); + return (subdet == Barrel) && (((side != Center) && (layer == 3)) || + ((side == NegZ) && (((layer == 2) && (rod > 5)) || ((layer == 1) && (rod > 9)))) || + ((side == PosZ) && (((layer == 2) && (rod < 8)) || ((layer == 1) && (rod < 4))))); } ALPAKA_FN_ACC ALPAKA_FN_INLINE float isTighterTiltedModules_seg(short subdet, short layer, short side, short rod) { // The "tighter" tilted modules are the subset of tilted modules that have smaller spacing // This is the same as what was previously considered as"isNormalTiltedModules" // See Figure 9.1 of https://cds.cern.ch/record/2272264/files/CMS-TDR-014.pdf - return (subdet == ::lst::Barrel) && - (((side != ::lst::Center) && (layer == 3)) || - ((side == ::lst::NegZ) && (((layer == 2) && (rod > 5)) || ((layer == 1) && (rod > 9)))) || - ((side == ::lst::PosZ) && (((layer == 2) && (rod < 8)) || ((layer == 1) && (rod < 4))))); + return (subdet == Barrel) && (((side != Center) && (layer == 3)) || + ((side == NegZ) && (((layer == 2) && (rod > 5)) || ((layer == 1) && (rod > 9)))) || + ((side == PosZ) && (((layer == 2) && (rod < 8)) || ((layer == 1) && (rod < 4))))); } ALPAKA_FN_ACC ALPAKA_FN_INLINE float moduleGapSize_seg(short layer, short ring, short subdet, short side, short rod) { @@ -218,11 +216,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float moduleSeparation = 0; - if (subdet == ::lst::Barrel and side == ::lst::Center) { + if (subdet == Barrel and side == Center) { moduleSeparation = miniDeltaFlat[iL]; } else if (isTighterTiltedModules_seg(subdet, layer, side, rod)) { moduleSeparation = miniDeltaTilted[iL]; - } else if (subdet == ::lst::Endcap) { + } else if (subdet == Endcap) { moduleSeparation = miniDeltaEndcap[iL][iR]; } else //Loose tilted modules { @@ -250,11 +248,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float moduleSeparation = 0; - if (subdet == ::lst::Barrel and side == ::lst::Center) { + if (subdet == Barrel and side == Center) { moduleSeparation = miniDeltaFlat[iL]; } else if (isTighterTiltedModules_seg(modulesInGPU, moduleIndex)) { moduleSeparation = miniDeltaTilted[iL]; - } else if (subdet == ::lst::Endcap) { + } else if (subdet == Endcap) { moduleSeparation = miniDeltaEndcap[iL][iR]; } else //Loose tilted modules { @@ -281,7 +279,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { uint16_t outerLowerModuleIndex, unsigned int innerMDIndex, unsigned int outerMDIndex) { - float sdMuls = (modulesInGPU.subdets[innerLowerModuleIndex] == ::lst::Barrel) + float sdMuls = (modulesInGPU.subdets[innerLowerModuleIndex] == Barrel) ? kMiniMulsPtScaleBarrel[modulesInGPU.layers[innerLowerModuleIndex] - 1] * 3.f / ptCut : kMiniMulsPtScaleEndcap[modulesInGPU.layers[innerLowerModuleIndex] - 1] * 3.f / ptCut; @@ -291,10 +289,10 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { const float dAlpha_Bfield = alpaka::math::asin(acc, alpaka::math::min(acc, segmentDr * k2Rinv1GeVf / ptCut, kSinAlphaMax)); - bool isInnerTilted = modulesInGPU.subdets[innerLowerModuleIndex] == ::lst::Barrel and - modulesInGPU.sides[innerLowerModuleIndex] != ::lst::Center; - bool isOuterTilted = modulesInGPU.subdets[outerLowerModuleIndex] == ::lst::Barrel and - modulesInGPU.sides[outerLowerModuleIndex] != ::lst::Center; + bool isInnerTilted = + modulesInGPU.subdets[innerLowerModuleIndex] == Barrel and modulesInGPU.sides[innerLowerModuleIndex] != Center; + bool isOuterTilted = + modulesInGPU.subdets[outerLowerModuleIndex] == Barrel and modulesInGPU.sides[outerLowerModuleIndex] != Center; float drdzInner = modulesInGPU.drdzs[innerLowerModuleIndex]; float drdzOuter = modulesInGPU.drdzs[outerLowerModuleIndex]; @@ -315,14 +313,14 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float sdLumForInnerMini2; float sdLumForOuterMini2; - if (modulesInGPU.subdets[innerLowerModuleIndex] == ::lst::Barrel) { + if (modulesInGPU.subdets[innerLowerModuleIndex] == Barrel) { sdLumForInnerMini2 = innerminiTilt2 * (dAlpha_Bfield * dAlpha_Bfield); } else { sdLumForInnerMini2 = (mdsInGPU.dphis[innerMDIndex] * mdsInGPU.dphis[innerMDIndex]) * (kDeltaZLum * kDeltaZLum) / (mdsInGPU.dzs[innerMDIndex] * mdsInGPU.dzs[innerMDIndex]); } - if (modulesInGPU.subdets[outerLowerModuleIndex] == ::lst::Barrel) { + if (modulesInGPU.subdets[outerLowerModuleIndex] == Barrel) { sdLumForOuterMini2 = outerminiTilt2 * (dAlpha_Bfield * dAlpha_Bfield); } else { sdLumForOuterMini2 = (mdsInGPU.dphis[outerMDIndex] * mdsInGPU.dphis[outerMDIndex]) * (kDeltaZLum * kDeltaZLum) / @@ -332,23 +330,21 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { // Unique stuff for the segment dudes alone float dAlpha_res_inner = 0.02f / miniDelta * - (modulesInGPU.subdets[innerLowerModuleIndex] == ::lst::Barrel ? 1.0f : alpaka::math::abs(acc, zIn) / rtIn); + (modulesInGPU.subdets[innerLowerModuleIndex] == Barrel ? 1.0f : alpaka::math::abs(acc, zIn) / rtIn); float dAlpha_res_outer = 0.02f / miniDelta * - (modulesInGPU.subdets[outerLowerModuleIndex] == ::lst::Barrel ? 1.0f : alpaka::math::abs(acc, zOut) / rtOut); + (modulesInGPU.subdets[outerLowerModuleIndex] == Barrel ? 1.0f : alpaka::math::abs(acc, zOut) / rtOut); float dAlpha_res = dAlpha_res_inner + dAlpha_res_outer; - if (modulesInGPU.subdets[innerLowerModuleIndex] == ::lst::Barrel and - modulesInGPU.sides[innerLowerModuleIndex] == ::lst::Center) { + if (modulesInGPU.subdets[innerLowerModuleIndex] == Barrel and modulesInGPU.sides[innerLowerModuleIndex] == Center) { dAlphaThresholdValues[0] = dAlpha_Bfield + alpaka::math::sqrt(acc, dAlpha_res * dAlpha_res + sdMuls * sdMuls); } else { dAlphaThresholdValues[0] = dAlpha_Bfield + alpaka::math::sqrt(acc, dAlpha_res * dAlpha_res + sdMuls * sdMuls + sdLumForInnerMini2); } - if (modulesInGPU.subdets[outerLowerModuleIndex] == ::lst::Barrel and - modulesInGPU.sides[outerLowerModuleIndex] == ::lst::Center) { + if (modulesInGPU.subdets[outerLowerModuleIndex] == Barrel and modulesInGPU.sides[outerLowerModuleIndex] == Center) { dAlphaThresholdValues[1] = dAlpha_Bfield + alpaka::math::sqrt(acc, dAlpha_res * dAlpha_res + sdMuls * sdMuls); } else { dAlphaThresholdValues[1] = @@ -465,7 +461,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float& dPhiChange, float& dPhiChangeMin, float& dPhiChangeMax) { - float sdMuls = (modulesInGPU.subdets[innerLowerModuleIndex] == ::lst::Barrel) + float sdMuls = (modulesInGPU.subdets[innerLowerModuleIndex] == Barrel) ? kMiniMulsPtScaleBarrel[modulesInGPU.layers[innerLowerModuleIndex] - 1] * 3.f / ptCut : kMiniMulsPtScaleEndcap[modulesInGPU.layers[innerLowerModuleIndex] - 1] * 3.f / ptCut; @@ -567,8 +563,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { zOut = mdsInGPU.anchorZ[outerMDIndex]; rtOut = mdsInGPU.anchorRt[outerMDIndex]; - bool outerLayerEndcapTwoS = (modulesInGPU.subdets[outerLowerModuleIndex] == ::lst::Endcap) && - (modulesInGPU.moduleType[outerLowerModuleIndex] == ::lst::TwoS); + bool outerLayerEndcapTwoS = (modulesInGPU.subdets[outerLowerModuleIndex] == Endcap) && + (modulesInGPU.moduleType[outerLowerModuleIndex] == TwoS); float sdSlope = alpaka::math::asin(acc, alpaka::math::min(acc, rtOut * k2Rinv1GeVf / ptCut, kSinAlphaMax)); float disks2SMinRadius = 60.f; @@ -669,8 +665,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { float& dPhiChange, float& dPhiChangeMin, float& dPhiChangeMax) { - if (modulesInGPU.subdets[innerLowerModuleIndex] == ::lst::Barrel and - modulesInGPU.subdets[outerLowerModuleIndex] == ::lst::Barrel) { + if (modulesInGPU.subdets[innerLowerModuleIndex] == Barrel and + modulesInGPU.subdets[outerLowerModuleIndex] == Barrel) { return runSegmentDefaultAlgoBarrel(acc, modulesInGPU, mdsInGPU, diff --git a/RecoTracker/LSTCore/src/alpaka/Triplet.h b/RecoTracker/LSTCore/src/alpaka/Triplet.h index c5ac8bda543d8..5e1b352748573 100644 --- a/RecoTracker/LSTCore/src/alpaka/Triplet.h +++ b/RecoTracker/LSTCore/src/alpaka/Triplet.h @@ -280,8 +280,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int innerSegmentIndex, float& betaIn, float& betaInCut) { - bool isPSIn = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == ::lst::PS); - bool isPSOut = (modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::PS); + bool isPSIn = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == PS); + bool isPSOut = (modulesInGPU.moduleType[outerOuterLowerModuleIndex] == PS); float rtIn = mdsInGPU.anchorRt[firstMDIndex]; float rtMid = mdsInGPU.anchorRt[secondMDIndex]; @@ -378,8 +378,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { unsigned int outerSegmentIndex, float& betaIn, float& betaInCut) { - bool isPSIn = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == ::lst::PS); - bool isPSOut = (modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::PS); + bool isPSIn = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == PS); + bool isPSOut = (modulesInGPU.moduleType[outerOuterLowerModuleIndex] == PS); float rtIn = mdsInGPU.anchorRt[firstMDIndex]; float rtMid = mdsInGPU.anchorRt[secondMDIndex]; @@ -402,7 +402,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { return false; float dLum = alpaka::math::copysign(acc, kDeltaZLum, zIn); - bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::PS; + bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerOuterLowerModuleIndex] == PS; float rtGeom1 = isOutSgInnerMDPS ? kPixelPSZpitch : kStrip2SZpitch; float zGeom1 = alpaka::math::copysign(acc, zGeom, zIn); float rtLo = rtIn * (1.f + (zOut - zIn - zGeom1) / (zIn + zGeom1 + dLum) / dzDrtScale) - @@ -514,8 +514,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { return false; float dLum = alpaka::math::copysign(acc, kDeltaZLum, zIn); - bool isOutSgOuterMDPS = modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::PS; - bool isInSgInnerMDPS = modulesInGPU.moduleType[innerInnerLowerModuleIndex] == ::lst::PS; + bool isOutSgOuterMDPS = modulesInGPU.moduleType[outerOuterLowerModuleIndex] == PS; + bool isInSgInnerMDPS = modulesInGPU.moduleType[innerInnerLowerModuleIndex] == PS; float rtGeom = (isInSgInnerMDPS and isOutSgOuterMDPS) ? 2.f * kPixelPSZpitch : (isInSgInnerMDPS or isOutSgOuterMDPS) ? kPixelPSZpitch + kStrip2SZpitch @@ -529,7 +529,7 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { if ((rtOut < rtLo) || (rtOut > rtHi)) return false; - bool isInSgOuterMDPS = modulesInGPU.moduleType[outerOuterLowerModuleIndex] == ::lst::PS; + bool isInSgOuterMDPS = modulesInGPU.moduleType[outerOuterLowerModuleIndex] == PS; float drtSDIn = rtMid - rtIn; float dzSDIn = zMid - zIn; @@ -622,8 +622,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { short middleLowerModuleSubdet = modulesInGPU.subdets[middleLowerModuleIndex]; short outerOuterLowerModuleSubdet = modulesInGPU.subdets[outerOuterLowerModuleIndex]; - if (innerInnerLowerModuleSubdet == ::lst::Barrel and middleLowerModuleSubdet == ::lst::Barrel and - outerOuterLowerModuleSubdet == ::lst::Barrel) { + if (innerInnerLowerModuleSubdet == Barrel and middleLowerModuleSubdet == Barrel and + outerOuterLowerModuleSubdet == Barrel) { return passPointingConstraintBBB(acc, modulesInGPU, mdsInGPU, @@ -639,8 +639,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { innerSegmentIndex, betaIn, betaInCut); - } else if (innerInnerLowerModuleSubdet == ::lst::Barrel and middleLowerModuleSubdet == ::lst::Barrel and - outerOuterLowerModuleSubdet == ::lst::Endcap) { + } else if (innerInnerLowerModuleSubdet == Barrel and middleLowerModuleSubdet == Barrel and + outerOuterLowerModuleSubdet == Endcap) { return passPointingConstraintBBE(acc, modulesInGPU, mdsInGPU, @@ -658,8 +658,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { outerSegmentIndex, betaIn, betaInCut); - } else if (innerInnerLowerModuleSubdet == ::lst::Barrel and middleLowerModuleSubdet == ::lst::Endcap and - outerOuterLowerModuleSubdet == ::lst::Endcap) { + } else if (innerInnerLowerModuleSubdet == Barrel and middleLowerModuleSubdet == Endcap and + outerOuterLowerModuleSubdet == Endcap) { return passPointingConstraintBBE(acc, modulesInGPU, mdsInGPU, @@ -680,8 +680,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { } - else if (innerInnerLowerModuleSubdet == ::lst::Endcap and middleLowerModuleSubdet == ::lst::Endcap and - outerOuterLowerModuleSubdet == ::lst::Endcap) { + else if (innerInnerLowerModuleSubdet == Endcap and middleLowerModuleSubdet == Endcap and + outerOuterLowerModuleSubdet == Endcap) { return passPointingConstraintEEE(acc, modulesInGPU, mdsInGPU, diff --git a/RecoTracker/LSTCore/standalone/bin/lst.cc b/RecoTracker/LSTCore/standalone/bin/lst.cc index 1670adab72b81..344af542048f2 100644 --- a/RecoTracker/LSTCore/standalone/bin/lst.cc +++ b/RecoTracker/LSTCore/standalone/bin/lst.cc @@ -2,6 +2,8 @@ #include +using LSTEvent = ALPAKA_ACCELERATOR_NAMESPACE::lst::Event; + //___________________________________________________________________________________________________________________________________________________________________________________________ int main(int argc, char **argv) { //******************************************************************************** @@ -383,10 +385,9 @@ void run_lst() { full_timer.Reset(); full_timer.Start(); - std::vector events; + std::vector events; for (int s = 0; s < ana.streams; s++) { - ALPAKA_ACCELERATOR_NAMESPACE::lst::Event *event = - new ALPAKA_ACCELERATOR_NAMESPACE::lst::Event(ana.verbose >= 2, queues[s], &deviceESData); + LSTEvent *event = new LSTEvent(ana.verbose >= 2, queues[s], &deviceESData); events.push_back(event); } float timeForEventCreation = full_timer.RealTime() * 1000; diff --git a/RecoTracker/LSTCore/standalone/bin/lst.h b/RecoTracker/LSTCore/standalone/bin/lst.h index 4a5699a1dd59f..5a951552ba647 100644 --- a/RecoTracker/LSTCore/standalone/bin/lst.h +++ b/RecoTracker/LSTCore/standalone/bin/lst.h @@ -1,6 +1,9 @@ #ifndef lst_h #define lst_h +#include "Event.h" +#include "LST.h" + #include #include #include @@ -14,9 +17,6 @@ #include "rooutil.h" #include "cxxopts.h" -#include "Event.h" -#include "LST.h" - // Efficiency study modules #include "AnalysisConfig.h" #include "trkCore.h" diff --git a/RecoTracker/LSTCore/standalone/code/core/AccessHelper.cc b/RecoTracker/LSTCore/standalone/code/core/AccessHelper.cc index d27675221ed67..426a74babc4d1 100644 --- a/RecoTracker/LSTCore/standalone/code/core/AccessHelper.cc +++ b/RecoTracker/LSTCore/standalone/code/core/AccessHelper.cc @@ -31,7 +31,7 @@ std::vector getPixelHitsFrompLS(Event* event, unsigned int pLS) { Segments const* segments = event->getSegments().data(); MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); ObjectRanges const* rangesEvt = event->getRanges().data(); - ::lst::Modules const* modulesEvt = event->getModules().data(); + Modules const* modulesEvt = event->getModules().data(); const unsigned int pLS_offset = rangesEvt->segmentModuleIndices[*(modulesEvt->nLowerModules)]; unsigned int MD_1 = segments->mdIndices[2 * (pLS + pLS_offset)]; unsigned int MD_2 = segments->mdIndices[2 * (pLS + pLS_offset) + 1]; @@ -227,7 +227,7 @@ std::tuple, std::vector> getHitIdxsAndHi unsigned int getPixelLSFrompT3(Event* event, unsigned int pT3) { PixelTriplets const* pixelTriplets = event->getPixelTriplets().data(); ObjectRanges const* rangesEvt = event->getRanges().data(); - ::lst::Modules const* modulesEvt = event->getModules().data(); + Modules const* modulesEvt = event->getModules().data(); const unsigned int pLS_offset = rangesEvt->segmentModuleIndices[*(modulesEvt->nLowerModules)]; return pixelTriplets->pixelSegmentIndices[pT3] - pLS_offset; } @@ -316,7 +316,7 @@ std::tuple, std::vector> getHitIdxsAndHi unsigned int getPixelLSFrompT5(Event* event, unsigned int pT5) { PixelQuintuplets const* pixelQuintuplets = event->getPixelQuintuplets().data(); ObjectRanges const* rangesEvt = event->getRanges().data(); - ::lst::Modules const* modulesEvt = event->getModules().data(); + Modules const* modulesEvt = event->getModules().data(); const unsigned int pLS_offset = rangesEvt->segmentModuleIndices[*(modulesEvt->nLowerModules)]; return pixelQuintuplets->pixelIndices[pT5] - pLS_offset; } diff --git a/RecoTracker/LSTCore/standalone/code/core/AnalysisConfig.h b/RecoTracker/LSTCore/standalone/code/core/AnalysisConfig.h index ce7ce3824849e..8608bc95ed2fa 100644 --- a/RecoTracker/LSTCore/standalone/code/core/AnalysisConfig.h +++ b/RecoTracker/LSTCore/standalone/code/core/AnalysisConfig.h @@ -100,7 +100,7 @@ class AnalysisConfig { std::map>> moduleSimHits; std::map modulePopulation; - ::lst::ModuleConnectionMap moduleConnectiongMapLoose; + lst::ModuleConnectionMap moduleConnectiongMapLoose; // Boolean to trigger whether to run cut_value_ntupling bool do_cut_value_ntuple; diff --git a/RecoTracker/LSTCore/standalone/code/core/trkCore.h b/RecoTracker/LSTCore/standalone/code/core/trkCore.h index 139b8c5796e12..4a84a1c5f916a 100644 --- a/RecoTracker/LSTCore/standalone/code/core/trkCore.h +++ b/RecoTracker/LSTCore/standalone/code/core/trkCore.h @@ -1,13 +1,14 @@ #ifndef trkCore_h #define trkCore_h +#include "Event.h" + #include "Trktree.h" #include "TCanvas.h" #include "TSystem.h" #include "AnalysisConfig.h" #include "ModuleConnectionMap.h" #include "lst_math.h" -#include "Event.h" #include #include diff --git a/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc b/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc index 96a48481d1477..e12512f5c5c7d 100644 --- a/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc +++ b/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc @@ -294,7 +294,7 @@ void setPixelQuintupletOutputBranches(Event* event) { PixelQuintuplets const* pixelQuintuplets = event->getPixelQuintuplets().data(); Quintuplets const* quintuplets = event->getQuintuplets().data(); Segments const* segments = event->getSegments().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); int n_accepted_simtrk = ana.tx->getBranch>("sim_TC_matched").size(); unsigned int nPixelQuintuplets = @@ -368,7 +368,7 @@ void setPixelQuintupletOutputBranches(Event* event) { void setQuintupletOutputBranches(Event* event) { Quintuplets const* quintuplets = event->getQuintuplets().data(); ObjectRanges const* ranges = event->getRanges().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); int n_accepted_simtrk = ana.tx->getBranch>("sim_TC_matched").size(); std::vector sim_t5_matched(n_accepted_simtrk); @@ -438,7 +438,7 @@ void setQuintupletOutputBranches(Event* event) { //________________________________________________________________________________________________________________________________ void setPixelTripletOutputBranches(Event* event) { PixelTriplets const* pixelTriplets = event->getPixelTriplets().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); Segments const* segments = event->getSegments().data(); int n_accepted_simtrk = ana.tx->getBranch>("sim_TC_matched").size(); @@ -504,7 +504,7 @@ void setGnnNtupleBranches(Event* event) { Segments const* segments = event->getSegments().data(); MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); Hits const* hitsEvt = event->getHits().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); ObjectRanges const* ranges = event->getRanges().data(); TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); @@ -944,7 +944,7 @@ std::tuple, std::vectorgetModules().data(); + Modules const* modules = event->getModules().data(); ObjectRanges const* ranges = event->getRanges().data(); int nHits = 0; @@ -960,7 +960,7 @@ void printHitMultiplicities(Event* event) { //________________________________________________________________________________________________________________________________ void printMiniDoubletMultiplicities(Event* event) { MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); int nMiniDoublets = 0; int totOccupancyMiniDoublets = 0; @@ -988,7 +988,7 @@ void printAllObjects(Event* event) { void printMDs(Event* event) { MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); Hits const* hitsEvt = event->getHits().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); ObjectRanges const* ranges = event->getRanges().data(); // Then obtain the lower module index @@ -1011,7 +1011,7 @@ void printLSs(Event* event) { Segments const* segments = event->getSegments().data(); MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); Hits const* hitsEvt = event->getHits().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); ObjectRanges const* ranges = event->getRanges().data(); int nSegments = 0; @@ -1043,7 +1043,7 @@ void printpLSs(Event* event) { Segments const* segments = event->getSegments().data(); MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); Hits const* hitsEvt = event->getHits().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); ObjectRanges const* ranges = event->getRanges().data(); unsigned int i = *(modules->nLowerModules); @@ -1074,7 +1074,7 @@ void printT3s(Event* event) { Segments const* segments = event->getSegments().data(); MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); Hits const* hitsEvt = event->getHits().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); int nTriplets = 0; for (unsigned int i = 0; i < *(modules->nLowerModules); ++i) { // unsigned int idx = modules->lowerModuleIndices[i]; @@ -1116,7 +1116,7 @@ void debugPrintOutlierMultiplicities(Event* event) { Triplets const* triplets = event->getTriplets().data(); Segments const* segments = event->getSegments().data(); MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); - ::lst::Modules const* modules = event->getModules().data(); + Modules const* modules = event->getModules().data(); ObjectRanges const* ranges = event->getRanges().data(); //int nTrackCandidates = 0; for (unsigned int idx = 0; idx <= *(modules->nLowerModules); ++idx) {