diff --git a/RecoTracker/LST/plugins/LSTOutputConverter.cc b/RecoTracker/LST/plugins/LSTOutputConverter.cc index e21b456962b84..c40603c9c1e5d 100644 --- a/RecoTracker/LST/plugins/LSTOutputConverter.cc +++ b/RecoTracker/LST/plugins/LSTOutputConverter.cc @@ -57,16 +57,14 @@ class LSTOutputConverter : public edm::global::EDProducer<> { }; LSTOutputConverter::LSTOutputConverter(edm::ParameterSet const& iConfig) - : lstOutputToken_(consumes(iConfig.getParameter("lstOutput"))), - lstPhase2OTHitsInputToken_{consumes(iConfig.getParameter("phase2OTHits"))}, - lstPixelSeedToken_{consumes(iConfig.getParameter("lstPixelSeeds"))}, + : lstOutputToken_(consumes(iConfig.getParameter("lstOutput"))), + lstPhase2OTHitsInputToken_{consumes(iConfig.getParameter("phase2OTHits"))}, + lstPixelSeedToken_{consumes(iConfig.getParameter("lstPixelSeeds"))}, includeT5s_(iConfig.getParameter("includeT5s")), includeNonpLSTSs_(iConfig.getParameter("includeNonpLSTSs")), mfToken_(esConsumes()), - propagatorAlongToken_{ - esConsumes(iConfig.getParameter("propagatorAlong"))}, - propagatorOppositeToken_{esConsumes( - iConfig.getParameter("propagatorOpposite"))}, + propagatorAlongToken_{esConsumes(iConfig.getParameter("propagatorAlong"))}, + propagatorOppositeToken_{esConsumes(iConfig.getParameter("propagatorOpposite"))}, tGeomToken_(esConsumes()), seedCreator_(SeedCreatorFactory::get()->create("SeedFromConsecutiveHitsCreator", iConfig.getParameter("SeedCreatorPSet"), @@ -77,15 +75,15 @@ LSTOutputConverter::LSTOutputConverter(edm::ParameterSet const& iConfig) // - The minimal set for TCs is t5TCsLST, pTTCsLST and pLSTCsLST. // That would complicate the handling of collections though, // so it is deferred to when we have a clearer picture of what's needed. - trajectorySeedPutToken_(produces("")), - trajectorySeedpLSPutToken_(produces("pLSTSsLST")), - trackCandidatePutToken_(produces("")), - trackCandidatepTCPutToken_(produces("pTCsLST")), - trackCandidateT5TCPutToken_(produces("t5TCsLST")), - trackCandidateNopLSTCPutToken_(produces("nopLSTCsLST")), - trackCandidatepTTCPutToken_(produces("pTTCsLST")), - trackCandidatepLSTCPutToken_(produces("pLSTCsLST")), - seedStopInfoPutToken_(produces>()) {} + trajectorySeedPutToken_(produces("")), + trajectorySeedpLSPutToken_(produces("pLSTSsLST")), + trackCandidatePutToken_(produces("")), + trackCandidatepTCPutToken_(produces("pTCsLST")), + trackCandidateT5TCPutToken_(produces("t5TCsLST")), + trackCandidateNopLSTCPutToken_(produces("nopLSTCsLST")), + trackCandidatepTTCPutToken_(produces("pTTCsLST")), + trackCandidatepLSTCPutToken_(produces("pLSTCsLST")), + seedStopInfoPutToken_(produces()) {} void LSTOutputConverter::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; diff --git a/RecoTracker/LST/plugins/LSTPhase2OTHitsInputProducer.cc b/RecoTracker/LST/plugins/LSTPhase2OTHitsInputProducer.cc index 3fd0a76770e56..a0fcc72f598b6 100644 --- a/RecoTracker/LST/plugins/LSTPhase2OTHitsInputProducer.cc +++ b/RecoTracker/LST/plugins/LSTPhase2OTHitsInputProducer.cc @@ -21,9 +21,8 @@ class LSTPhase2OTHitsInputProducer : public edm::global::EDProducer<> { }; LSTPhase2OTHitsInputProducer::LSTPhase2OTHitsInputProducer(edm::ParameterSet const& iConfig) - : phase2OTRecHitToken_( - consumes(iConfig.getParameter("phase2OTRecHits"))), - lstPhase2OTHitsInputPutToken_(produces()) {} + : phase2OTRecHitToken_(consumes(iConfig.getParameter("phase2OTRecHits"))), + lstPhase2OTHitsInputPutToken_(produces()) {} void LSTPhase2OTHitsInputProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; diff --git a/RecoTracker/LST/plugins/LSTPixelSeedInputProducer.cc b/RecoTracker/LST/plugins/LSTPixelSeedInputProducer.cc index 098d8731d62c8..819baf78c6aa4 100644 --- a/RecoTracker/LST/plugins/LSTPixelSeedInputProducer.cc +++ b/RecoTracker/LST/plugins/LSTPixelSeedInputProducer.cc @@ -39,9 +39,9 @@ class LSTPixelSeedInputProducer : public edm::global::EDProducer<> { LSTPixelSeedInputProducer::LSTPixelSeedInputProducer(edm::ParameterSet const& iConfig) : mfToken_(esConsumes()), - beamSpotToken_(consumes(iConfig.getParameter("beamSpot"))), - lstPixelSeedInputPutToken_(produces()), - lstPixelSeedsPutToken_(produces()) { + beamSpotToken_(consumes(iConfig.getParameter("beamSpot"))), + lstPixelSeedInputPutToken_(produces()), + lstPixelSeedsPutToken_(produces()) { seedTokens_ = edm::vector_transform(iConfig.getParameter>("seedTracks"), [&](const edm::InputTag& tag) { return consumes>(tag); }); } 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 e92ff549dffd1..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,16 +21,13 @@ #include "RecoTracker/Record/interface/TrackerRecoGeometryRecord.h" -#include "RecoTracker/LSTCore/interface/alpaka/LST.h" - namespace ALPAKA_ACCELERATOR_NAMESPACE { class LSTProducer : public stream::SynchronizingEDProducer<> { public: LSTProducer(edm::ParameterSet const& config) - : lstPixelSeedInputToken_{consumes(config.getParameter("pixelSeedInput"))}, - lstPhase2OTHitsInputToken_{ - consumes(config.getParameter("phase2OTHitsInput"))}, + : lstPixelSeedInputToken_{consumes(config.getParameter("pixelSeedInput"))}, + lstPhase2OTHitsInputToken_{consumes(config.getParameter("phase2OTHitsInput"))}, lstESToken_{esConsumes()}, verbose_(config.getParameter("verbose")), nopLSDupClean_(config.getParameter("nopLSDupClean")), @@ -87,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/Constants.h b/RecoTracker/LSTCore/interface/Constants.h index c0c342b6ad8a0..350857ac0b2e5 100644 --- a/RecoTracker/LSTCore/interface/Constants.h +++ b/RecoTracker/LSTCore/interface/Constants.h @@ -32,6 +32,9 @@ namespace lst { alpaka_common::Vec1D(static_cast(nElements))); } + // Named constants for pixelTypes + enum PixelType : int8_t { kInvalid = -1, kHighPt = 0, kLowPtPosCurv = 1, kLowPtNegCurv = 2 }; + // If a compile time flag does not define PT_CUT, default to 0.8 (GeV) #ifndef PT_CUT constexpr float PT_CUT = 0.8f; @@ -52,7 +55,7 @@ namespace lst { constexpr unsigned int size_superbins = 45000; - //defining the constant host device variables right up here + // Defining the constant host device variables right up here // Currently pixel tracks treated as LSs with 2 double layers (IT layers 1+2 and 3+4) and 4 hits. To be potentially handled better in the future. struct Params_pLS { static constexpr int kLayers = 2, kHits = 4; diff --git a/RecoTracker/LSTCore/interface/EndcapGeometry.h b/RecoTracker/LSTCore/interface/EndcapGeometry.h index 1a84d89abf90d..b8c44c14fb143 100644 --- a/RecoTracker/LSTCore/interface/EndcapGeometry.h +++ b/RecoTracker/LSTCore/interface/EndcapGeometry.h @@ -2,12 +2,8 @@ #define RecoTracker_LSTCore_interface_EndcapGeometry_h #include -#include -#include -#include #include #include -#include namespace lst { class EndcapGeometry { diff --git a/RecoTracker/LSTCore/interface/EndcapGeometryBuffer.h b/RecoTracker/LSTCore/interface/EndcapGeometryBuffer.h index 2c6df9ab2773c..ce037b026fc22 100644 --- a/RecoTracker/LSTCore/interface/EndcapGeometryBuffer.h +++ b/RecoTracker/LSTCore/interface/EndcapGeometryBuffer.h @@ -28,7 +28,6 @@ namespace lst { struct EndcapGeometryBuffer { Buf geoMapDetId_buf; Buf geoMapPhi_buf; - EndcapGeometryDev data_; EndcapGeometryBuffer(TDev const& dev, unsigned int nEndCapMap) : geoMapDetId_buf(allocBufWrapper(dev, nEndCapMap)), @@ -49,6 +48,9 @@ namespace lst { } inline EndcapGeometryDev const* data() const { return &data_; } + + private: + EndcapGeometryDev data_; }; } // namespace lst diff --git a/RecoTracker/LSTCore/interface/ModuleConnectionMap.h b/RecoTracker/LSTCore/interface/ModuleConnectionMap.h index 1d4445d3b423e..63c3496523c0d 100644 --- a/RecoTracker/LSTCore/interface/ModuleConnectionMap.h +++ b/RecoTracker/LSTCore/interface/ModuleConnectionMap.h @@ -1,12 +1,10 @@ #ifndef RecoTracker_LSTCore_interface_ModuleConnectionMap_h #define RecoTracker_LSTCore_interface_ModuleConnectionMap_h -#include -#include -#include +#include #include -#include -#include +#include +#include namespace lst { class ModuleConnectionMap { diff --git a/RecoTracker/LSTCore/interface/PixelMap.h b/RecoTracker/LSTCore/interface/PixelMap.h index 73d88bc64f7ec..a0fd89387e7e4 100644 --- a/RecoTracker/LSTCore/interface/PixelMap.h +++ b/RecoTracker/LSTCore/interface/PixelMap.h @@ -17,7 +17,7 @@ namespace lst { std::vector connectedPixelsIndexNeg; std::vector connectedPixelsSizesNeg; - int* pixelType; + const int* pixelType; PixelMap(unsigned int sizef = size_superbins) : pixelModuleIndex(0), diff --git a/RecoTracker/LSTCore/interface/TiltedGeometry.h b/RecoTracker/LSTCore/interface/TiltedGeometry.h index 420000dd38aa0..7a17106195522 100644 --- a/RecoTracker/LSTCore/interface/TiltedGeometry.h +++ b/RecoTracker/LSTCore/interface/TiltedGeometry.h @@ -1,13 +1,9 @@ #ifndef RecoTracker_LSTCore_interface_TiltedGeometry_h #define RecoTracker_LSTCore_interface_TiltedGeometry_h -#include #include -#include -#include -#include #include -#include +#include namespace lst { class TiltedGeometry { 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 0e4c64d2535df..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 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/EndcapGeometry.cc b/RecoTracker/LSTCore/src/EndcapGeometry.cc index d35b2520022ce..17e72379bb2ec 100644 --- a/RecoTracker/LSTCore/src/EndcapGeometry.cc +++ b/RecoTracker/LSTCore/src/EndcapGeometry.cc @@ -1,5 +1,10 @@ #include "RecoTracker/LSTCore/interface/EndcapGeometry.h" +#include +#include +#include +#include + lst::EndcapGeometry::EndcapGeometry(std::string const& filename) { load(filename); } void lst::EndcapGeometry::load(std::string const& filename) { diff --git a/RecoTracker/LSTCore/src/ModuleConnectionMap.cc b/RecoTracker/LSTCore/src/ModuleConnectionMap.cc index d1b68b7f485bb..881b2a66f6216 100644 --- a/RecoTracker/LSTCore/src/ModuleConnectionMap.cc +++ b/RecoTracker/LSTCore/src/ModuleConnectionMap.cc @@ -1,5 +1,10 @@ #include "RecoTracker/LSTCore/interface/ModuleConnectionMap.h" +#include +#include +#include +#include + lst::ModuleConnectionMap::ModuleConnectionMap() {} lst::ModuleConnectionMap::ModuleConnectionMap(std::string const& filename) { load(filename); } diff --git a/RecoTracker/LSTCore/src/TiltedGeometry.cc b/RecoTracker/LSTCore/src/TiltedGeometry.cc index a3442147939c3..d65a9a4a5f7b9 100644 --- a/RecoTracker/LSTCore/src/TiltedGeometry.cc +++ b/RecoTracker/LSTCore/src/TiltedGeometry.cc @@ -1,5 +1,10 @@ #include "RecoTracker/LSTCore/interface/TiltedGeometry.h" +#include +#include +#include +#include + lst::TiltedGeometry::TiltedGeometry(std::string const& filename) { load(filename); } void lst::TiltedGeometry::load(std::string const& filename) { diff --git a/RecoTracker/LSTCore/src/alpaka/Event.dev.cc b/RecoTracker/LSTCore/src/alpaka/Event.dev.cc index 62629bb08fc52..659591b836ec9 100644 --- a/RecoTracker/LSTCore/src/alpaka/Event.dev.cc +++ b/RecoTracker/LSTCore/src/alpaka/Event.dev.cc @@ -10,28 +10,8 @@ using Acc3D = ALPAKA_ACCELERATOR_NAMESPACE::Acc3D; using namespace ALPAKA_ACCELERATOR_NAMESPACE::lst; void Event::initSync(bool verbose) { - alpaka::wait(queue); // other calls can be asynchronous - addObjects = verbose; - hitsInGPU = nullptr; - mdsInGPU = nullptr; - segmentsInGPU = nullptr; - tripletsInGPU = nullptr; - quintupletsInGPU = nullptr; - trackCandidatesInGPU = nullptr; - pixelTripletsInGPU = nullptr; - pixelQuintupletsInGPU = nullptr; - rangesInGPU = nullptr; - - hitsInCPU = nullptr; - rangesInCPU = nullptr; - mdsInCPU = nullptr; - segmentsInCPU = nullptr; - tripletsInCPU = nullptr; - trackCandidatesInCPU = nullptr; - modulesInCPU = nullptr; - quintupletsInCPU = nullptr; - pixelTripletsInCPU = nullptr; - pixelQuintupletsInCPU = nullptr; + alpaka::wait(queue_); // other calls can be asynchronous + addObjects_ = verbose; //reset the arrays for (int i = 0; i < 6; i++) { @@ -53,7 +33,7 @@ void Event::initSync(bool verbose) { } void Event::resetEventSync() { - alpaka::wait(queue); // synchronize to reset consistently + alpaka::wait(queue_); // synchronize to reset consistently //reset the arrays for (int i = 0; i < 6; i++) { n_hits_by_layer_barrel_[i] = 0; @@ -71,92 +51,35 @@ void Event::resetEventSync() { n_quintuplets_by_layer_endcap_[i] = 0; } } - if (hitsInGPU) { - delete hitsInGPU; - delete hitsBuffers; - hitsInGPU = nullptr; - } - if (mdsInGPU) { - delete mdsInGPU; - delete miniDoubletsBuffers; - mdsInGPU = nullptr; - } - if (rangesInGPU) { - delete rangesInGPU; - delete rangesBuffers; - rangesInGPU = nullptr; - } - if (segmentsInGPU) { - delete segmentsInGPU; - delete segmentsBuffers; - segmentsInGPU = nullptr; - } - if (tripletsInGPU) { - delete tripletsInGPU; - delete tripletsBuffers; - tripletsInGPU = nullptr; - } - if (quintupletsInGPU) { - delete quintupletsInGPU; - delete quintupletsBuffers; - quintupletsInGPU = nullptr; - } - if (trackCandidatesInGPU) { - delete trackCandidatesInGPU; - delete trackCandidatesBuffers; - trackCandidatesInGPU = nullptr; - } - if (pixelTripletsInGPU) { - delete pixelTripletsInGPU; - delete pixelTripletsBuffers; - pixelTripletsInGPU = nullptr; - } - if (pixelQuintupletsInGPU) { - delete pixelQuintupletsInGPU; - delete pixelQuintupletsBuffers; - pixelQuintupletsInGPU = nullptr; - } - - if (hitsInCPU != nullptr) { - delete hitsInCPU; - hitsInCPU = nullptr; - } - if (rangesInCPU != nullptr) { - delete rangesInCPU; - rangesInCPU = nullptr; - } - if (mdsInCPU != nullptr) { - delete mdsInCPU; - mdsInCPU = nullptr; - } - if (segmentsInCPU != nullptr) { - delete segmentsInCPU; - segmentsInCPU = nullptr; - } - if (tripletsInCPU != nullptr) { - delete tripletsInCPU; - tripletsInCPU = nullptr; - } - if (quintupletsInCPU != nullptr) { - delete quintupletsInCPU; - quintupletsInCPU = nullptr; - } - if (pixelTripletsInCPU != nullptr) { - delete pixelTripletsInCPU; - pixelTripletsInCPU = nullptr; - } - if (pixelQuintupletsInCPU != nullptr) { - delete pixelQuintupletsInCPU; - pixelQuintupletsInCPU = nullptr; - } - if (trackCandidatesInCPU != nullptr) { - delete trackCandidatesInCPU; - trackCandidatesInCPU = nullptr; - } - if (modulesInCPU != nullptr) { - delete modulesInCPU; - modulesInCPU = nullptr; - } + hitsInGPU_.reset(); + hitsBuffers_.reset(); + mdsInGPU_.reset(); + miniDoubletsBuffers_.reset(); + rangesInGPU_.reset(); + rangesBuffers_.reset(); + segmentsInGPU_.reset(); + segmentsBuffers_.reset(); + tripletsInGPU_.reset(); + tripletsBuffers_.reset(); + quintupletsInGPU_.reset(); + quintupletsBuffers_.reset(); + trackCandidatesInGPU_.reset(); + trackCandidatesBuffers_.reset(); + pixelTripletsInGPU_.reset(); + pixelTripletsBuffers_.reset(); + pixelQuintupletsInGPU_.reset(); + pixelQuintupletsBuffers_.reset(); + + hitsInCPU_.reset(); + rangesInCPU_.reset(); + mdsInCPU_.reset(); + segmentsInCPU_.reset(); + tripletsInCPU_.reset(); + quintupletsInCPU_.reset(); + pixelTripletsInCPU_.reset(); + pixelQuintupletsInCPU_.reset(); + trackCandidatesInCPU_.reset(); + modulesInCPU_.reset(); } void Event::addHitToEvent(std::vector const& x, @@ -168,45 +91,45 @@ void Event::addHitToEvent(std::vector const& x, unsigned int nHits = x.size(); // Initialize space on device/host for next event. - if (hitsInGPU == nullptr) { - hitsInGPU = new Hits(); - hitsBuffers = new HitsBuffer(nModules_, nHits, devAcc, queue); - hitsInGPU->setData(*hitsBuffers); + if (!hitsInGPU_) { + hitsInGPU_.emplace(); + hitsBuffers_.emplace(nModules_, nHits, devAcc_, queue_); + hitsInGPU_->setData(*hitsBuffers_); } - if (rangesInGPU == nullptr) { - rangesInGPU = new ObjectRanges(); - rangesBuffers = new ObjectRangesBuffer(nModules_, nLowerModules_, devAcc, queue); - rangesInGPU->setData(*rangesBuffers); + if (!rangesInGPU_) { + rangesInGPU_.emplace(); + rangesBuffers_.emplace(nModules_, nLowerModules_, devAcc_, queue_); + rangesInGPU_->setData(*rangesBuffers_); } // Need a view here before transferring to the device. - auto nHits_view = alpaka::createView(devHost, &nHits, (Idx)1u); + auto nHits_view = alpaka::createView(cms::alpakatools::host(), &nHits, (Idx)1u); // Copy the host arrays to the GPU. - alpaka::memcpy(queue, hitsBuffers->xs_buf, x, nHits); - alpaka::memcpy(queue, hitsBuffers->ys_buf, y, nHits); - alpaka::memcpy(queue, hitsBuffers->zs_buf, z, nHits); - alpaka::memcpy(queue, hitsBuffers->detid_buf, detId, nHits); - alpaka::memcpy(queue, hitsBuffers->idxs_buf, idxInNtuple, nHits); - alpaka::memcpy(queue, hitsBuffers->nHits_buf, nHits_view); - alpaka::wait(queue); // FIXME: remove synch after inputs refactored to be in pinned memory + alpaka::memcpy(queue_, hitsBuffers_->xs_buf, x, nHits); + alpaka::memcpy(queue_, hitsBuffers_->ys_buf, y, nHits); + alpaka::memcpy(queue_, hitsBuffers_->zs_buf, z, nHits); + alpaka::memcpy(queue_, hitsBuffers_->detid_buf, detId, nHits); + alpaka::memcpy(queue_, hitsBuffers_->idxs_buf, idxInNtuple, nHits); + alpaka::memcpy(queue_, hitsBuffers_->nHits_buf, nHits_view); + alpaka::wait(queue_); // FIXME: remove synch after inputs refactored to be in pinned memory Vec3D const threadsPerBlock1{1, 1, 256}; Vec3D const blocksPerGrid1{1, 1, max_blocks}; WorkDiv3D const hit_loop_workdiv = createWorkDiv(blocksPerGrid1, threadsPerBlock1, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, hit_loop_workdiv, HitLoopKernel{}, - ::lst::Endcap, - ::lst::TwoS, + Endcap, + TwoS, nModules_, nEndCapMap_, endcapGeometryBuffers_.geoMapDetId_buf.data(), endcapGeometryBuffers_.geoMapPhi_buf.data(), *modulesBuffers_.data(), - *hitsInGPU, + *hitsInGPU_, nHits); Vec3D const threadsPerBlock2{1, 1, 256}; @@ -214,7 +137,7 @@ void Event::addHitToEvent(std::vector const& x, WorkDiv3D const module_ranges_workdiv = createWorkDiv(blocksPerGrid2, threadsPerBlock2, elementsPerThread); alpaka::exec( - queue, module_ranges_workdiv, ModuleRangesKernel{}, *modulesBuffers_.data(), *hitsInGPU, nLowerModules_); + queue_, module_ranges_workdiv, ModuleRangesKernel{}, *modulesBuffers_.data(), *hitsInGPU_, nLowerModules_); } void Event::addPixelSegmentToEvent(std::vector const& hitIndices0, @@ -233,7 +156,7 @@ void Event::addPixelSegmentToEvent(std::vector const& hitIndices0, std::vector const& charge, std::vector const& seedIdx, std::vector const& superbin, - std::vector const& pixelType, + std::vector const& pixelType, std::vector const& isQuad) { unsigned int size = ptIn.size(); @@ -249,121 +172,120 @@ void Event::addPixelSegmentToEvent(std::vector const& hitIndices0, unsigned int mdSize = 2 * size; uint16_t pixelModuleIndex = pixelMapping_.pixelModuleIndex; - if (mdsInGPU == nullptr) { - // Create a view for the element nLowerModules_ inside rangesBuffers->miniDoubletModuleOccupancy + if (!mdsInGPU_) { + // Create a view for the element nLowerModules_ inside rangesBuffers_->miniDoubletModuleOccupancy auto dst_view_miniDoubletModuleOccupancy = - alpaka::createSubView(rangesBuffers->miniDoubletModuleOccupancy_buf, (Idx)1u, (Idx)nLowerModules_); + alpaka::createSubView(rangesBuffers_->miniDoubletModuleOccupancy_buf, (Idx)1u, (Idx)nLowerModules_); // Create a host buffer for a value to be passed to the device - auto pixelMaxMDs_buf_h = cms::alpakatools::make_host_buffer(queue, (Idx)1u); + auto pixelMaxMDs_buf_h = cms::alpakatools::make_host_buffer(queue_, (Idx)1u); *pixelMaxMDs_buf_h.data() = n_max_pixel_md_per_modules; - alpaka::memcpy(queue, dst_view_miniDoubletModuleOccupancy, pixelMaxMDs_buf_h); + alpaka::memcpy(queue_, dst_view_miniDoubletModuleOccupancy, pixelMaxMDs_buf_h); WorkDiv1D const createMDArrayRangesGPU_workDiv = createWorkDiv({1}, {1024}, {1}); alpaka::exec( - queue, createMDArrayRangesGPU_workDiv, CreateMDArrayRangesGPU{}, *modulesBuffers_.data(), *rangesInGPU); + queue_, createMDArrayRangesGPU_workDiv, CreateMDArrayRangesGPU{}, *modulesBuffers_.data(), *rangesInGPU_); - auto nTotalMDs_buf_h = cms::alpakatools::make_host_buffer(queue, (Idx)1u); - alpaka::memcpy(queue, nTotalMDs_buf_h, rangesBuffers->device_nTotalMDs_buf); - alpaka::wait(queue); // wait to get the data before manipulation + auto nTotalMDs_buf_h = cms::alpakatools::make_host_buffer(queue_, (Idx)1u); + alpaka::memcpy(queue_, nTotalMDs_buf_h, rangesBuffers_->device_nTotalMDs_buf); + alpaka::wait(queue_); // wait to get the data before manipulation *nTotalMDs_buf_h.data() += n_max_pixel_md_per_modules; unsigned int nTotalMDs = *nTotalMDs_buf_h.data(); - mdsInGPU = new MiniDoublets(); - miniDoubletsBuffers = new MiniDoubletsBuffer(nTotalMDs, nLowerModules_, devAcc, queue); - mdsInGPU->setData(*miniDoubletsBuffers); + mdsInGPU_.emplace(); + miniDoubletsBuffers_.emplace(nTotalMDs, nLowerModules_, devAcc_, queue_); + mdsInGPU_->setData(*miniDoubletsBuffers_); - alpaka::memcpy(queue, miniDoubletsBuffers->nMemoryLocations_buf, nTotalMDs_buf_h); + alpaka::memcpy(queue_, miniDoubletsBuffers_->nMemoryLocations_buf, nTotalMDs_buf_h); } - if (segmentsInGPU == nullptr) { + if (!segmentsInGPU_) { // can be optimized here: because we didn't distinguish pixel segments and outer-tracker segments and call them both "segments", so they use the index continuously. // If we want to further study the memory footprint in detail, we can separate the two and allocate different memories to them WorkDiv1D const createSegmentArrayRanges_workDiv = createWorkDiv({1}, {1024}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, createSegmentArrayRanges_workDiv, CreateSegmentArrayRanges{}, *modulesBuffers_.data(), - *rangesInGPU, - *mdsInGPU); + *rangesInGPU_, + *mdsInGPU_); - auto nTotalSegments_view = alpaka::createView(devHost, &nTotalSegments_, (Idx)1u); + auto nTotalSegments_view = alpaka::createView(cms::alpakatools::host(), &nTotalSegments_, (Idx)1u); - alpaka::memcpy(queue, nTotalSegments_view, rangesBuffers->device_nTotalSegs_buf); - alpaka::wait(queue); // wait to get the value before manipulation + alpaka::memcpy(queue_, nTotalSegments_view, rangesBuffers_->device_nTotalSegs_buf); + alpaka::wait(queue_); // wait to get the value before manipulation nTotalSegments_ += n_max_pixel_segments_per_module; - segmentsInGPU = new Segments(); - segmentsBuffers = - new SegmentsBuffer(nTotalSegments_, nLowerModules_, n_max_pixel_segments_per_module, devAcc, queue); - segmentsInGPU->setData(*segmentsBuffers); - - alpaka::memcpy(queue, segmentsBuffers->nMemoryLocations_buf, nTotalSegments_view); - } - - auto hitIndices0_dev = allocBufWrapper(devAcc, size, queue); - auto hitIndices1_dev = allocBufWrapper(devAcc, size, queue); - auto hitIndices2_dev = allocBufWrapper(devAcc, size, queue); - auto hitIndices3_dev = allocBufWrapper(devAcc, size, queue); - auto dPhiChange_dev = allocBufWrapper(devAcc, size, queue); - - alpaka::memcpy(queue, hitIndices0_dev, hitIndices0, size); - alpaka::memcpy(queue, hitIndices1_dev, hitIndices1, size); - alpaka::memcpy(queue, hitIndices2_dev, hitIndices2, size); - alpaka::memcpy(queue, hitIndices3_dev, hitIndices3, size); - alpaka::memcpy(queue, dPhiChange_dev, dPhiChange, size); - - alpaka::memcpy(queue, segmentsBuffers->ptIn_buf, ptIn, size); - alpaka::memcpy(queue, segmentsBuffers->ptErr_buf, ptErr, size); - alpaka::memcpy(queue, segmentsBuffers->px_buf, px, size); - alpaka::memcpy(queue, segmentsBuffers->py_buf, py, size); - alpaka::memcpy(queue, segmentsBuffers->pz_buf, pz, size); - alpaka::memcpy(queue, segmentsBuffers->etaErr_buf, etaErr, size); - alpaka::memcpy(queue, segmentsBuffers->isQuad_buf, isQuad, size); - alpaka::memcpy(queue, segmentsBuffers->eta_buf, eta, size); - alpaka::memcpy(queue, segmentsBuffers->phi_buf, phi, size); - alpaka::memcpy(queue, segmentsBuffers->charge_buf, charge, size); - alpaka::memcpy(queue, segmentsBuffers->seedIdx_buf, seedIdx, size); - alpaka::memcpy(queue, segmentsBuffers->superbin_buf, superbin, size); - alpaka::memcpy(queue, segmentsBuffers->pixelType_buf, pixelType, size); + segmentsInGPU_.emplace(); + segmentsBuffers_.emplace(nTotalSegments_, nLowerModules_, n_max_pixel_segments_per_module, devAcc_, queue_); + segmentsInGPU_->setData(*segmentsBuffers_); + + alpaka::memcpy(queue_, segmentsBuffers_->nMemoryLocations_buf, nTotalSegments_view); + } + + auto hitIndices0_dev = allocBufWrapper(devAcc_, size, queue_); + auto hitIndices1_dev = allocBufWrapper(devAcc_, size, queue_); + auto hitIndices2_dev = allocBufWrapper(devAcc_, size, queue_); + auto hitIndices3_dev = allocBufWrapper(devAcc_, size, queue_); + auto dPhiChange_dev = allocBufWrapper(devAcc_, size, queue_); + + alpaka::memcpy(queue_, hitIndices0_dev, hitIndices0, size); + alpaka::memcpy(queue_, hitIndices1_dev, hitIndices1, size); + alpaka::memcpy(queue_, hitIndices2_dev, hitIndices2, size); + alpaka::memcpy(queue_, hitIndices3_dev, hitIndices3, size); + alpaka::memcpy(queue_, dPhiChange_dev, dPhiChange, size); + + alpaka::memcpy(queue_, segmentsBuffers_->ptIn_buf, ptIn, size); + alpaka::memcpy(queue_, segmentsBuffers_->ptErr_buf, ptErr, size); + alpaka::memcpy(queue_, segmentsBuffers_->px_buf, px, size); + alpaka::memcpy(queue_, segmentsBuffers_->py_buf, py, size); + alpaka::memcpy(queue_, segmentsBuffers_->pz_buf, pz, size); + alpaka::memcpy(queue_, segmentsBuffers_->etaErr_buf, etaErr, size); + alpaka::memcpy(queue_, segmentsBuffers_->isQuad_buf, isQuad, size); + alpaka::memcpy(queue_, segmentsBuffers_->eta_buf, eta, size); + alpaka::memcpy(queue_, segmentsBuffers_->phi_buf, phi, size); + alpaka::memcpy(queue_, segmentsBuffers_->charge_buf, charge, size); + alpaka::memcpy(queue_, segmentsBuffers_->seedIdx_buf, seedIdx, size); + alpaka::memcpy(queue_, segmentsBuffers_->superbin_buf, superbin, size); + alpaka::memcpy(queue_, segmentsBuffers_->pixelType_buf, pixelType, size); // Create source views for size and mdSize - auto src_view_size = alpaka::createView(devHost, &size, (Idx)1u); - auto src_view_mdSize = alpaka::createView(devHost, &mdSize, (Idx)1u); + auto src_view_size = alpaka::createView(cms::alpakatools::host(), &size, (Idx)1u); + auto src_view_mdSize = alpaka::createView(cms::alpakatools::host(), &mdSize, (Idx)1u); - auto dst_view_segments = alpaka::createSubView(segmentsBuffers->nSegments_buf, (Idx)1u, (Idx)pixelModuleIndex); - alpaka::memcpy(queue, dst_view_segments, src_view_size); + auto dst_view_segments = alpaka::createSubView(segmentsBuffers_->nSegments_buf, (Idx)1u, (Idx)pixelModuleIndex); + alpaka::memcpy(queue_, dst_view_segments, src_view_size); auto dst_view_totOccupancySegments = - alpaka::createSubView(segmentsBuffers->totOccupancySegments_buf, (Idx)1u, (Idx)pixelModuleIndex); - alpaka::memcpy(queue, dst_view_totOccupancySegments, src_view_size); + alpaka::createSubView(segmentsBuffers_->totOccupancySegments_buf, (Idx)1u, (Idx)pixelModuleIndex); + alpaka::memcpy(queue_, dst_view_totOccupancySegments, src_view_size); - auto dst_view_nMDs = alpaka::createSubView(miniDoubletsBuffers->nMDs_buf, (Idx)1u, (Idx)pixelModuleIndex); - alpaka::memcpy(queue, dst_view_nMDs, src_view_mdSize); + auto dst_view_nMDs = alpaka::createSubView(miniDoubletsBuffers_->nMDs_buf, (Idx)1u, (Idx)pixelModuleIndex); + alpaka::memcpy(queue_, dst_view_nMDs, src_view_mdSize); auto dst_view_totOccupancyMDs = - alpaka::createSubView(miniDoubletsBuffers->totOccupancyMDs_buf, (Idx)1u, (Idx)pixelModuleIndex); - alpaka::memcpy(queue, dst_view_totOccupancyMDs, src_view_mdSize); + alpaka::createSubView(miniDoubletsBuffers_->totOccupancyMDs_buf, (Idx)1u, (Idx)pixelModuleIndex); + alpaka::memcpy(queue_, dst_view_totOccupancyMDs, src_view_mdSize); - alpaka::wait(queue); // FIXME: remove synch after inputs refactored to be in pinned memory + alpaka::wait(queue_); // FIXME: remove synch after inputs refactored to be in pinned memory Vec3D const threadsPerBlock{1, 1, 256}; Vec3D const blocksPerGrid{1, 1, max_blocks}; WorkDiv3D const addPixelSegmentToEvent_workdiv = createWorkDiv(blocksPerGrid, threadsPerBlock, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, addPixelSegmentToEvent_workdiv, AddPixelSegmentToEventKernel{}, *modulesBuffers_.data(), - *rangesInGPU, - *hitsInGPU, - *mdsInGPU, - *segmentsInGPU, + *rangesInGPU_, + *hitsInGPU_, + *mdsInGPU_, + *segmentsInGPU_, hitIndices0_dev.data(), hitIndices1_dev.data(), hitIndices2_dev.data(), @@ -374,32 +296,32 @@ void Event::addPixelSegmentToEvent(std::vector const& hitIndices0, } void Event::createMiniDoublets() { - // Create a view for the element nLowerModules_ inside rangesBuffers->miniDoubletModuleOccupancy + // Create a view for the element nLowerModules_ inside rangesBuffers_->miniDoubletModuleOccupancy auto dst_view_miniDoubletModuleOccupancy = - alpaka::createSubView(rangesBuffers->miniDoubletModuleOccupancy_buf, (Idx)1u, (Idx)nLowerModules_); + alpaka::createSubView(rangesBuffers_->miniDoubletModuleOccupancy_buf, (Idx)1u, (Idx)nLowerModules_); // Create a host buffer for a value to be passed to the device - auto pixelMaxMDs_buf_h = cms::alpakatools::make_host_buffer(queue, (Idx)1u); + auto pixelMaxMDs_buf_h = cms::alpakatools::make_host_buffer(queue_, (Idx)1u); *pixelMaxMDs_buf_h.data() = n_max_pixel_md_per_modules; - alpaka::memcpy(queue, dst_view_miniDoubletModuleOccupancy, pixelMaxMDs_buf_h); + alpaka::memcpy(queue_, dst_view_miniDoubletModuleOccupancy, pixelMaxMDs_buf_h); WorkDiv1D const createMDArrayRangesGPU_workDiv = createWorkDiv({1}, {1024}, {1}); alpaka::exec( - queue, createMDArrayRangesGPU_workDiv, CreateMDArrayRangesGPU{}, *modulesBuffers_.data(), *rangesInGPU); + queue_, createMDArrayRangesGPU_workDiv, CreateMDArrayRangesGPU{}, *modulesBuffers_.data(), *rangesInGPU_); - auto nTotalMDs_buf_h = cms::alpakatools::make_host_buffer(queue, (Idx)1u); - alpaka::memcpy(queue, nTotalMDs_buf_h, rangesBuffers->device_nTotalMDs_buf); - alpaka::wait(queue); // wait to get the data before manipulation + auto nTotalMDs_buf_h = cms::alpakatools::make_host_buffer(queue_, (Idx)1u); + alpaka::memcpy(queue_, nTotalMDs_buf_h, rangesBuffers_->device_nTotalMDs_buf); + alpaka::wait(queue_); // wait to get the data before manipulation *nTotalMDs_buf_h.data() += n_max_pixel_md_per_modules; unsigned int nTotalMDs = *nTotalMDs_buf_h.data(); - if (mdsInGPU == nullptr) { - mdsInGPU = new MiniDoublets(); - miniDoubletsBuffers = new MiniDoubletsBuffer(nTotalMDs, nLowerModules_, devAcc, queue); - mdsInGPU->setData(*miniDoubletsBuffers); + if (!mdsInGPU_) { + mdsInGPU_.emplace(); + miniDoubletsBuffers_.emplace(nTotalMDs, nLowerModules_, devAcc_, queue_); + mdsInGPU_->setData(*miniDoubletsBuffers_); } Vec3D const threadsPerBlockCreateMDInGPU{1, 16, 32}; @@ -407,35 +329,34 @@ void Event::createMiniDoublets() { WorkDiv3D const createMiniDoubletsInGPUv2_workDiv = createWorkDiv(blocksPerGridCreateMDInGPU, threadsPerBlockCreateMDInGPU, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, createMiniDoubletsInGPUv2_workDiv, CreateMiniDoubletsInGPUv2{}, *modulesBuffers_.data(), - *hitsInGPU, - *mdsInGPU, - *rangesInGPU); + *hitsInGPU_, + *mdsInGPU_, + *rangesInGPU_); WorkDiv1D const addMiniDoubletRangesToEventExplicit_workDiv = createWorkDiv({1}, {1024}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, addMiniDoubletRangesToEventExplicit_workDiv, AddMiniDoubletRangesToEventExplicit{}, *modulesBuffers_.data(), - *mdsInGPU, - *rangesInGPU, - *hitsInGPU); + *mdsInGPU_, + *rangesInGPU_, + *hitsInGPU_); - if (addObjects) { + if (addObjects_) { addMiniDoubletsToEventExplicit(); } } void Event::createSegmentsWithModuleMap() { - if (segmentsInGPU == nullptr) { - segmentsInGPU = new Segments(); - segmentsBuffers = - new SegmentsBuffer(nTotalSegments_, nLowerModules_, n_max_pixel_segments_per_module, devAcc, queue); - segmentsInGPU->setData(*segmentsBuffers); + if (!segmentsInGPU_) { + segmentsInGPU_.emplace(); + segmentsBuffers_.emplace(nTotalSegments_, nLowerModules_, n_max_pixel_segments_per_module, devAcc_, queue_); + segmentsInGPU_->setData(*segmentsBuffers_); } Vec3D const threadsPerBlockCreateSeg{1, 1, 64}; @@ -443,71 +364,71 @@ void Event::createSegmentsWithModuleMap() { WorkDiv3D const createSegmentsInGPUv2_workDiv = createWorkDiv(blocksPerGridCreateSeg, threadsPerBlockCreateSeg, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, createSegmentsInGPUv2_workDiv, CreateSegmentsInGPUv2{}, *modulesBuffers_.data(), - *mdsInGPU, - *segmentsInGPU, - *rangesInGPU); + *mdsInGPU_, + *segmentsInGPU_, + *rangesInGPU_); WorkDiv1D const addSegmentRangesToEventExplicit_workDiv = createWorkDiv({1}, {1024}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, addSegmentRangesToEventExplicit_workDiv, AddSegmentRangesToEventExplicit{}, *modulesBuffers_.data(), - *segmentsInGPU, - *rangesInGPU); + *segmentsInGPU_, + *rangesInGPU_); - if (addObjects) { + if (addObjects_) { addSegmentsToEventExplicit(); } } void Event::createTriplets() { - if (tripletsInGPU == nullptr) { + if (!tripletsInGPU_) { WorkDiv1D const createTripletArrayRanges_workDiv = createWorkDiv({1}, {1024}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, createTripletArrayRanges_workDiv, CreateTripletArrayRanges{}, *modulesBuffers_.data(), - *rangesInGPU, - *segmentsInGPU); + *rangesInGPU_, + *segmentsInGPU_); // TODO: Why are we pulling this back down only to put it back on the device in a new struct? - auto maxTriplets_buf_h = cms::alpakatools::make_host_buffer(queue, (Idx)1u); + auto maxTriplets_buf_h = cms::alpakatools::make_host_buffer(queue_, (Idx)1u); - alpaka::memcpy(queue, maxTriplets_buf_h, rangesBuffers->device_nTotalTrips_buf); - alpaka::wait(queue); // wait to get the value before using it + alpaka::memcpy(queue_, maxTriplets_buf_h, rangesBuffers_->device_nTotalTrips_buf); + alpaka::wait(queue_); // wait to get the value before using it - tripletsInGPU = new Triplets(); - tripletsBuffers = new TripletsBuffer(*maxTriplets_buf_h.data(), nLowerModules_, devAcc, queue); - tripletsInGPU->setData(*tripletsBuffers); + tripletsInGPU_.emplace(); + tripletsBuffers_.emplace(*maxTriplets_buf_h.data(), nLowerModules_, devAcc_, queue_); + tripletsInGPU_->setData(*tripletsBuffers_); - alpaka::memcpy(queue, tripletsBuffers->nMemoryLocations_buf, maxTriplets_buf_h); + alpaka::memcpy(queue_, tripletsBuffers_->nMemoryLocations_buf, maxTriplets_buf_h); } uint16_t nonZeroModules = 0; unsigned int max_InnerSeg = 0; // Allocate and copy nSegments from device to host (only nLowerModules in OT, not the +1 with pLSs) - auto nSegments_buf_h = cms::alpakatools::make_host_buffer(queue, nLowerModules_); - alpaka::memcpy(queue, nSegments_buf_h, segmentsBuffers->nSegments_buf, nLowerModules_); + auto nSegments_buf_h = cms::alpakatools::make_host_buffer(queue_, nLowerModules_); + alpaka::memcpy(queue_, nSegments_buf_h, segmentsBuffers_->nSegments_buf, nLowerModules_); // ... same for module_nConnectedModules // FIXME: replace by ES host data - auto module_nConnectedModules_buf_h = cms::alpakatools::make_host_buffer(queue, nLowerModules_); - alpaka::memcpy(queue, module_nConnectedModules_buf_h, modulesBuffers_.nConnectedModules_buf, nLowerModules_); + auto module_nConnectedModules_buf_h = cms::alpakatools::make_host_buffer(queue_, nLowerModules_); + alpaka::memcpy(queue_, module_nConnectedModules_buf_h, modulesBuffers_.nConnectedModules_buf, nLowerModules_); - alpaka::wait(queue); // wait for nSegments and module_nConnectedModules before using + alpaka::wait(queue_); // wait for nSegments and module_nConnectedModules before using auto const* nSegments = nSegments_buf_h.data(); auto const* module_nConnectedModules = module_nConnectedModules_buf_h.data(); // Allocate host index and fill it directly - auto index_buf_h = cms::alpakatools::make_host_buffer(queue, nLowerModules_); + auto index_buf_h = cms::alpakatools::make_host_buffer(queue_, nLowerModules_); auto* index = index_buf_h.data(); for (uint16_t innerLowerModuleIndex = 0; innerLowerModuleIndex < nLowerModules_; innerLowerModuleIndex++) { @@ -521,45 +442,44 @@ void Event::createTriplets() { } // Allocate and copy to device index - auto index_gpu_buf = allocBufWrapper(devAcc, nLowerModules_, queue); - alpaka::memcpy(queue, index_gpu_buf, index_buf_h, nonZeroModules); + auto index_gpu_buf = allocBufWrapper(devAcc_, nLowerModules_, queue_); + alpaka::memcpy(queue_, index_gpu_buf, index_buf_h, nonZeroModules); Vec3D const threadsPerBlockCreateTrip{1, 16, 16}; Vec3D const blocksPerGridCreateTrip{max_blocks, 1, 1}; WorkDiv3D const createTripletsInGPUv2_workDiv = createWorkDiv(blocksPerGridCreateTrip, threadsPerBlockCreateTrip, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, createTripletsInGPUv2_workDiv, CreateTripletsInGPUv2{}, *modulesBuffers_.data(), - *mdsInGPU, - *segmentsInGPU, - *tripletsInGPU, - *rangesInGPU, + *mdsInGPU_, + *segmentsInGPU_, + *tripletsInGPU_, + *rangesInGPU_, index_gpu_buf.data(), nonZeroModules); WorkDiv1D const addTripletRangesToEventExplicit_workDiv = createWorkDiv({1}, {1024}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, addTripletRangesToEventExplicit_workDiv, AddTripletRangesToEventExplicit{}, *modulesBuffers_.data(), - *tripletsInGPU, - *rangesInGPU); + *tripletsInGPU_, + *rangesInGPU_); - if (addObjects) { + if (addObjects_) { addTripletsToEventExplicit(); } } void Event::createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets) { - if (trackCandidatesInGPU == nullptr) { - trackCandidatesInGPU = new TrackCandidates(); - trackCandidatesBuffers = new TrackCandidatesBuffer( - n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, devAcc, queue); - trackCandidatesInGPU->setData(*trackCandidatesBuffers); + if (!trackCandidatesInGPU_) { + trackCandidatesInGPU_.emplace(); + trackCandidatesBuffers_.emplace(n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, devAcc_, queue_); + trackCandidatesInGPU_->setData(*trackCandidatesBuffers_); } Vec3D const threadsPerBlock_crossCleanpT3{1, 16, 64}; @@ -567,30 +487,30 @@ void Event::createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets) { WorkDiv3D const crossCleanpT3_workDiv = createWorkDiv(blocksPerGrid_crossCleanpT3, threadsPerBlock_crossCleanpT3, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, crossCleanpT3_workDiv, CrossCleanpT3{}, *modulesBuffers_.data(), - *rangesInGPU, - *pixelTripletsInGPU, - *segmentsInGPU, - *pixelQuintupletsInGPU); + *rangesInGPU_, + *pixelTripletsInGPU_, + *segmentsInGPU_, + *pixelQuintupletsInGPU_); WorkDiv1D const addpT3asTrackCandidatesInGPU_workDiv = createWorkDiv({1}, {512}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, addpT3asTrackCandidatesInGPU_workDiv, AddpT3asTrackCandidatesInGPU{}, nLowerModules_, - *pixelTripletsInGPU, - *trackCandidatesInGPU, - *segmentsInGPU, - *rangesInGPU); + *pixelTripletsInGPU_, + *trackCandidatesInGPU_, + *segmentsInGPU_, + *rangesInGPU_); // Pull nEligibleT5Modules from the device. - auto nEligibleModules_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nEligibleModules_buf_h, rangesBuffers->nEligibleT5Modules_buf); - alpaka::wait(queue); // wait to get the value before using + auto nEligibleModules_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nEligibleModules_buf_h, rangesBuffers_->nEligibleT5Modules_buf); + alpaka::wait(queue_); // wait to get the value before using auto const nEligibleModules = *nEligibleModules_buf_h.data(); Vec3D const threadsPerBlockRemoveDupQuints{1, 16, 32}; @@ -598,38 +518,38 @@ void Event::createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets) { WorkDiv3D const removeDupQuintupletsInGPUBeforeTC_workDiv = createWorkDiv(blocksPerGridRemoveDupQuints, threadsPerBlockRemoveDupQuints, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, removeDupQuintupletsInGPUBeforeTC_workDiv, RemoveDupQuintupletsInGPUBeforeTC{}, - *quintupletsInGPU, - *rangesInGPU); + *quintupletsInGPU_, + *rangesInGPU_); Vec3D const threadsPerBlock_crossCleanT5{32, 1, 32}; Vec3D const blocksPerGrid_crossCleanT5{(13296 / 32) + 1, 1, max_blocks}; WorkDiv3D const crossCleanT5_workDiv = createWorkDiv(blocksPerGrid_crossCleanT5, threadsPerBlock_crossCleanT5, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, crossCleanT5_workDiv, CrossCleanT5{}, *modulesBuffers_.data(), - *quintupletsInGPU, - *pixelQuintupletsInGPU, - *pixelTripletsInGPU, - *rangesInGPU); + *quintupletsInGPU_, + *pixelQuintupletsInGPU_, + *pixelTripletsInGPU_, + *rangesInGPU_); Vec3D const threadsPerBlock_addT5asTrackCandidateInGPU{1, 8, 128}; Vec3D const blocksPerGrid_addT5asTrackCandidateInGPU{1, 8, 10}; WorkDiv3D const addT5asTrackCandidateInGPU_workDiv = createWorkDiv( blocksPerGrid_addT5asTrackCandidateInGPU, threadsPerBlock_addT5asTrackCandidateInGPU, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, addT5asTrackCandidateInGPU_workDiv, AddT5asTrackCandidateInGPU{}, nLowerModules_, - *quintupletsInGPU, - *trackCandidatesInGPU, - *rangesInGPU); + *quintupletsInGPU_, + *trackCandidatesInGPU_, + *rangesInGPU_); if (!no_pls_dupclean) { Vec3D const threadsPerBlockCheckHitspLS{1, 16, 16}; @@ -637,7 +557,7 @@ void Event::createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets) { WorkDiv3D const checkHitspLS_workDiv = createWorkDiv(blocksPerGridCheckHitspLS, threadsPerBlockCheckHitspLS, elementsPerThread); - alpaka::exec(queue, checkHitspLS_workDiv, CheckHitspLS{}, *modulesBuffers_.data(), *segmentsInGPU, true); + alpaka::exec(queue_, checkHitspLS_workDiv, CheckHitspLS{}, *modulesBuffers_.data(), *segmentsInGPU_, true); } Vec3D const threadsPerBlock_crossCleanpLS{1, 16, 32}; @@ -645,41 +565,41 @@ void Event::createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets) { WorkDiv3D const crossCleanpLS_workDiv = createWorkDiv(blocksPerGrid_crossCleanpLS, threadsPerBlock_crossCleanpLS, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, crossCleanpLS_workDiv, CrossCleanpLS{}, *modulesBuffers_.data(), - *rangesInGPU, - *pixelTripletsInGPU, - *trackCandidatesInGPU, - *segmentsInGPU, - *mdsInGPU, - *hitsInGPU, - *quintupletsInGPU); + *rangesInGPU_, + *pixelTripletsInGPU_, + *trackCandidatesInGPU_, + *segmentsInGPU_, + *mdsInGPU_, + *hitsInGPU_, + *quintupletsInGPU_); Vec3D const threadsPerBlock_addpLSasTrackCandidateInGPU{1, 1, 384}; Vec3D const blocksPerGrid_addpLSasTrackCandidateInGPU{1, 1, max_blocks}; WorkDiv3D const addpLSasTrackCandidateInGPU_workDiv = createWorkDiv( blocksPerGrid_addpLSasTrackCandidateInGPU, threadsPerBlock_addpLSasTrackCandidateInGPU, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, addpLSasTrackCandidateInGPU_workDiv, AddpLSasTrackCandidateInGPU{}, nLowerModules_, - *trackCandidatesInGPU, - *segmentsInGPU, + *trackCandidatesInGPU_, + *segmentsInGPU_, tc_pls_triplets); // Check if either n_max_pixel_track_candidates or n_max_nonpixel_track_candidates was reached - auto nTrackCanpT5Host_buf = allocBufWrapper(devHost, 1, queue); - auto nTrackCanpT3Host_buf = allocBufWrapper(devHost, 1, queue); - auto nTrackCanpLSHost_buf = allocBufWrapper(devHost, 1, queue); - auto nTrackCanT5Host_buf = allocBufWrapper(devHost, 1, queue); - alpaka::memcpy(queue, nTrackCanpT5Host_buf, trackCandidatesBuffers->nTrackCandidatespT5_buf); - alpaka::memcpy(queue, nTrackCanpT3Host_buf, trackCandidatesBuffers->nTrackCandidatespT3_buf); - alpaka::memcpy(queue, nTrackCanpLSHost_buf, trackCandidatesBuffers->nTrackCandidatespLS_buf); - alpaka::memcpy(queue, nTrackCanT5Host_buf, trackCandidatesBuffers->nTrackCandidatesT5_buf); - alpaka::wait(queue); // wait to get the values before using them + auto nTrackCanpT5Host_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); + auto nTrackCanpT3Host_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); + auto nTrackCanpLSHost_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); + auto nTrackCanT5Host_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); + alpaka::memcpy(queue_, nTrackCanpT5Host_buf, trackCandidatesBuffers_->nTrackCandidatespT5_buf); + alpaka::memcpy(queue_, nTrackCanpT3Host_buf, trackCandidatesBuffers_->nTrackCandidatespT3_buf); + alpaka::memcpy(queue_, nTrackCanpLSHost_buf, trackCandidatesBuffers_->nTrackCandidatespLS_buf); + alpaka::memcpy(queue_, nTrackCanT5Host_buf, trackCandidatesBuffers_->nTrackCandidatesT5_buf); + alpaka::wait(queue_); // wait to get the values before using them auto nTrackCandidatespT5 = *nTrackCanpT5Host_buf.data(); auto nTrackCandidatespT3 = *nTrackCanpT3Host_buf.data(); @@ -697,33 +617,33 @@ void Event::createTrackCandidates(bool no_pls_dupclean, bool tc_pls_triplets) { } void Event::createPixelTriplets() { - if (pixelTripletsInGPU == nullptr) { - pixelTripletsInGPU = new PixelTriplets(); - pixelTripletsBuffers = new PixelTripletsBuffer(n_max_pixel_triplets, devAcc, queue); - pixelTripletsInGPU->setData(*pixelTripletsBuffers); + if (!pixelTripletsInGPU_) { + pixelTripletsInGPU_.emplace(); + pixelTripletsBuffers_.emplace(n_max_pixel_triplets, devAcc_, queue_); + pixelTripletsInGPU_->setData(*pixelTripletsBuffers_); } - auto superbins_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); - auto pixelTypes_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); + auto superbins_buf = allocBufWrapper(cms::alpakatools::host(), n_max_pixel_segments_per_module, queue_); + auto pixelTypes_buf = allocBufWrapper(cms::alpakatools::host(), n_max_pixel_segments_per_module, queue_); - alpaka::memcpy(queue, superbins_buf, segmentsBuffers->superbin_buf); - alpaka::memcpy(queue, pixelTypes_buf, segmentsBuffers->pixelType_buf); + alpaka::memcpy(queue_, superbins_buf, segmentsBuffers_->superbin_buf); + alpaka::memcpy(queue_, pixelTypes_buf, segmentsBuffers_->pixelType_buf); auto const* superbins = superbins_buf.data(); auto const* pixelTypes = pixelTypes_buf.data(); unsigned int nInnerSegments; - auto nInnerSegments_src_view = alpaka::createView(devHost, &nInnerSegments, (size_t)1u); + auto nInnerSegments_src_view = alpaka::createView(cms::alpakatools::host(), &nInnerSegments, (size_t)1u); // Create a sub-view for the device buffer - auto dev_view_nSegments = alpaka::createSubView(segmentsBuffers->nSegments_buf, (Idx)1u, (Idx)nLowerModules_); + auto dev_view_nSegments = alpaka::createSubView(segmentsBuffers_->nSegments_buf, (Idx)1u, (Idx)nLowerModules_); - alpaka::memcpy(queue, nInnerSegments_src_view, dev_view_nSegments); - alpaka::wait(queue); // wait to get nInnerSegments (also superbins and pixelTypes) before using + alpaka::memcpy(queue_, nInnerSegments_src_view, dev_view_nSegments); + alpaka::wait(queue_); // wait to get nInnerSegments (also superbins and pixelTypes) before using - auto connectedPixelSize_host_buf = allocBufWrapper(devHost, nInnerSegments, queue); - auto connectedPixelIndex_host_buf = allocBufWrapper(devHost, nInnerSegments, queue); - auto connectedPixelSize_dev_buf = allocBufWrapper(devAcc, nInnerSegments, queue); - auto connectedPixelIndex_dev_buf = allocBufWrapper(devAcc, nInnerSegments, queue); + auto connectedPixelSize_host_buf = allocBufWrapper(cms::alpakatools::host(), nInnerSegments, queue_); + auto connectedPixelIndex_host_buf = allocBufWrapper(cms::alpakatools::host(), nInnerSegments, queue_); + auto connectedPixelSize_dev_buf = allocBufWrapper(devAcc_, nInnerSegments, queue_); + auto connectedPixelIndex_dev_buf = allocBufWrapper(devAcc_, nInnerSegments, queue_); unsigned int* connectedPixelSize_host = connectedPixelSize_host_buf.data(); unsigned int* connectedPixelIndex_host = connectedPixelIndex_host_buf.data(); @@ -736,60 +656,67 @@ 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 - int8_t 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 > 2) or (pixelType < 0)) { + if ((superbin < 0) or (superbin >= (int)size_superbins) or + ((pixelType != PixelType::kHighPt) and (pixelType != PixelType::kLowPtPosCurv) and + (pixelType != PixelType::kLowPtNegCurv))) { connectedPixelSize_host[i] = 0; connectedPixelIndex_host[i] = 0; continue; } // Used pixel type to select correct size-index arrays - if (pixelType == 0) { - connectedPixelSize_host[i] = - pixelMapping_.connectedPixelsSizes[superbin]; // number of connected modules to this pixel - auto connectedIdxBase = pixelMapping_.connectedPixelsIndex[superbin]; - connectedPixelIndex_host[i] = - connectedIdxBase; // index to get start of connected modules for this superbin in map - } else if (pixelType == 1) { - connectedPixelSize_host[i] = - pixelMapping_.connectedPixelsSizesPos[superbin]; // number of pixel connected modules - auto connectedIdxBase = pixelMapping_.connectedPixelsIndexPos[superbin] + pixelIndexOffsetPos; - connectedPixelIndex_host[i] = connectedIdxBase; // index to get start of connected pixel modules - } else if (pixelType == 2) { - connectedPixelSize_host[i] = - pixelMapping_.connectedPixelsSizesNeg[superbin]; // number of pixel connected modules - auto connectedIdxBase = pixelMapping_.connectedPixelsIndexNeg[superbin] + pixelIndexOffsetNeg; - connectedPixelIndex_host[i] = connectedIdxBase; // index to get start of connected pixel modules + switch (pixelType) { + case PixelType::kInvalid: + break; + 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 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 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 + connectedPixelIndex_host[i] = pixelMapping_.connectedPixelsIndexNeg[superbin] + pixelIndexOffsetNeg; + break; } } - alpaka::memcpy(queue, connectedPixelSize_dev_buf, connectedPixelSize_host_buf, nInnerSegments); - alpaka::memcpy(queue, connectedPixelIndex_dev_buf, connectedPixelIndex_host_buf, nInnerSegments); + alpaka::memcpy(queue_, connectedPixelSize_dev_buf, connectedPixelSize_host_buf, nInnerSegments); + alpaka::memcpy(queue_, connectedPixelIndex_dev_buf, connectedPixelIndex_host_buf, nInnerSegments); Vec3D const threadsPerBlock{1, 4, 32}; Vec3D const blocksPerGrid{16 /* above median of connected modules*/, 4096, 1}; WorkDiv3D const createPixelTripletsInGPUFromMapv2_workDiv = createWorkDiv(blocksPerGrid, threadsPerBlock, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, createPixelTripletsInGPUFromMapv2_workDiv, CreatePixelTripletsInGPUFromMapv2{}, *modulesBuffers_.data(), - *rangesInGPU, - *mdsInGPU, - *segmentsInGPU, - *tripletsInGPU, - *pixelTripletsInGPU, + *rangesInGPU_, + *mdsInGPU_, + *segmentsInGPU_, + *tripletsInGPU_, + *pixelTripletsInGPU_, connectedPixelSize_dev_buf.data(), connectedPixelIndex_dev_buf.data(), nInnerSegments); #ifdef WARNINGS - auto nPixelTriplets_buf = allocBufWrapper(devHost, 1, queue); + auto nPixelTriplets_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); - alpaka::memcpy(queue, nPixelTriplets_buf, pixelTripletsBuffers->nPixelTriplets_buf); - alpaka::wait(queue); // wait to get the value before using it + alpaka::memcpy(queue_, nPixelTriplets_buf, pixelTripletsBuffers_->nPixelTriplets_buf); + alpaka::wait(queue_); // wait to get the value before using it std::cout << "number of pixel triplets = " << *nPixelTriplets_buf.data() << std::endl; #endif @@ -802,35 +729,35 @@ void Event::createPixelTriplets() { createWorkDiv(blocksPerGridDupPixTrip, threadsPerBlockDupPixTrip, elementsPerThread); alpaka::exec( - queue, removeDupPixelTripletsInGPUFromMap_workDiv, RemoveDupPixelTripletsInGPUFromMap{}, *pixelTripletsInGPU); + queue_, removeDupPixelTripletsInGPUFromMap_workDiv, RemoveDupPixelTripletsInGPUFromMap{}, *pixelTripletsInGPU_); } void Event::createQuintuplets() { WorkDiv1D const createEligibleModulesListForQuintupletsGPU_workDiv = createWorkDiv({1}, {1024}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, createEligibleModulesListForQuintupletsGPU_workDiv, CreateEligibleModulesListForQuintupletsGPU{}, *modulesBuffers_.data(), - *tripletsInGPU, - *rangesInGPU); + *tripletsInGPU_, + *rangesInGPU_); - auto nEligibleT5Modules_buf = allocBufWrapper(devHost, 1, queue); - auto nTotalQuintuplets_buf = allocBufWrapper(devHost, 1, queue); + auto nEligibleT5Modules_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); + auto nTotalQuintuplets_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); - alpaka::memcpy(queue, nEligibleT5Modules_buf, rangesBuffers->nEligibleT5Modules_buf); - alpaka::memcpy(queue, nTotalQuintuplets_buf, rangesBuffers->device_nTotalQuints_buf); - alpaka::wait(queue); // wait for the values before using them + alpaka::memcpy(queue_, nEligibleT5Modules_buf, rangesBuffers_->nEligibleT5Modules_buf); + alpaka::memcpy(queue_, nTotalQuintuplets_buf, rangesBuffers_->device_nTotalQuints_buf); + alpaka::wait(queue_); // wait for the values before using them auto nEligibleT5Modules = *nEligibleT5Modules_buf.data(); auto nTotalQuintuplets = *nTotalQuintuplets_buf.data(); - if (quintupletsInGPU == nullptr) { - quintupletsInGPU = new Quintuplets(); - quintupletsBuffers = new QuintupletsBuffer(nTotalQuintuplets, nLowerModules_, devAcc, queue); - quintupletsInGPU->setData(*quintupletsBuffers); + if (!quintupletsInGPU_) { + quintupletsInGPU_.emplace(); + quintupletsBuffers_.emplace(nTotalQuintuplets, nLowerModules_, devAcc_, queue_); + quintupletsInGPU_->setData(*quintupletsBuffers_); - alpaka::memcpy(queue, quintupletsBuffers->nMemoryLocations_buf, nTotalQuintuplets_buf); + alpaka::memcpy(queue_, quintupletsBuffers_->nMemoryLocations_buf, nTotalQuintuplets_buf); } Vec3D const threadsPerBlockQuints{1, 8, 32}; @@ -838,15 +765,15 @@ void Event::createQuintuplets() { WorkDiv3D const createQuintupletsInGPUv2_workDiv = createWorkDiv(blocksPerGridQuints, threadsPerBlockQuints, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, createQuintupletsInGPUv2_workDiv, CreateQuintupletsInGPUv2{}, *modulesBuffers_.data(), - *mdsInGPU, - *segmentsInGPU, - *tripletsInGPU, - *quintupletsInGPU, - *rangesInGPU, + *mdsInGPU_, + *segmentsInGPU_, + *tripletsInGPU_, + *quintupletsInGPU_, + *rangesInGPU_, nEligibleT5Modules); Vec3D const threadsPerBlockDupQuint{1, 16, 16}; @@ -854,23 +781,23 @@ void Event::createQuintuplets() { WorkDiv3D const removeDupQuintupletsInGPUAfterBuild_workDiv = createWorkDiv(blocksPerGridDupQuint, threadsPerBlockDupQuint, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, removeDupQuintupletsInGPUAfterBuild_workDiv, RemoveDupQuintupletsInGPUAfterBuild{}, *modulesBuffers_.data(), - *quintupletsInGPU, - *rangesInGPU); + *quintupletsInGPU_, + *rangesInGPU_); WorkDiv1D const addQuintupletRangesToEventExplicit_workDiv = createWorkDiv({1}, {1024}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, addQuintupletRangesToEventExplicit_workDiv, AddQuintupletRangesToEventExplicit{}, *modulesBuffers_.data(), - *quintupletsInGPU, - *rangesInGPU); + *quintupletsInGPU_, + *rangesInGPU_); - if (addObjects) { + if (addObjects_) { addQuintupletsToEventExplicit(); } } @@ -882,44 +809,43 @@ void Event::pixelLineSegmentCleaning(bool no_pls_dupclean) { WorkDiv3D const checkHitspLS_workDiv = createWorkDiv(blocksPerGridCheckHitspLS, threadsPerBlockCheckHitspLS, elementsPerThread); - alpaka::exec(queue, checkHitspLS_workDiv, CheckHitspLS{}, *modulesBuffers_.data(), *segmentsInGPU, false); + alpaka::exec(queue_, checkHitspLS_workDiv, CheckHitspLS{}, *modulesBuffers_.data(), *segmentsInGPU_, false); } } void Event::createPixelQuintuplets() { - if (pixelQuintupletsInGPU == nullptr) { - pixelQuintupletsInGPU = new PixelQuintuplets(); - pixelQuintupletsBuffers = new PixelQuintupletsBuffer(n_max_pixel_quintuplets, devAcc, queue); - pixelQuintupletsInGPU->setData(*pixelQuintupletsBuffers); + if (!pixelQuintupletsInGPU_) { + pixelQuintupletsInGPU_.emplace(); + pixelQuintupletsBuffers_.emplace(n_max_pixel_quintuplets, devAcc_, queue_); + pixelQuintupletsInGPU_->setData(*pixelQuintupletsBuffers_); } - if (trackCandidatesInGPU == nullptr) { - trackCandidatesInGPU = new TrackCandidates(); - trackCandidatesBuffers = new TrackCandidatesBuffer( - n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, devAcc, queue); - trackCandidatesInGPU->setData(*trackCandidatesBuffers); + if (!trackCandidatesInGPU_) { + trackCandidatesInGPU_.emplace(); + trackCandidatesBuffers_.emplace(n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, devAcc_, queue_); + trackCandidatesInGPU_->setData(*trackCandidatesBuffers_); } - auto superbins_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); - auto pixelTypes_buf = allocBufWrapper(devHost, n_max_pixel_segments_per_module, queue); + auto superbins_buf = allocBufWrapper(cms::alpakatools::host(), n_max_pixel_segments_per_module, queue_); + auto pixelTypes_buf = allocBufWrapper(cms::alpakatools::host(), n_max_pixel_segments_per_module, queue_); - alpaka::memcpy(queue, superbins_buf, segmentsBuffers->superbin_buf); - alpaka::memcpy(queue, pixelTypes_buf, segmentsBuffers->pixelType_buf); + alpaka::memcpy(queue_, superbins_buf, segmentsBuffers_->superbin_buf); + alpaka::memcpy(queue_, pixelTypes_buf, segmentsBuffers_->pixelType_buf); auto const* superbins = superbins_buf.data(); auto const* pixelTypes = pixelTypes_buf.data(); unsigned int nInnerSegments; - auto nInnerSegments_src_view = alpaka::createView(devHost, &nInnerSegments, (size_t)1u); + auto nInnerSegments_src_view = alpaka::createView(cms::alpakatools::host(), &nInnerSegments, (size_t)1u); // Create a sub-view for the device buffer - auto dev_view_nSegments = alpaka::createSubView(segmentsBuffers->nSegments_buf, (Idx)1u, (Idx)nLowerModules_); + auto dev_view_nSegments = alpaka::createSubView(segmentsBuffers_->nSegments_buf, (Idx)1u, (Idx)nLowerModules_); - alpaka::memcpy(queue, nInnerSegments_src_view, dev_view_nSegments); - alpaka::wait(queue); // wait to get nInnerSegments (also superbins and pixelTypes) before using + alpaka::memcpy(queue_, nInnerSegments_src_view, dev_view_nSegments); + alpaka::wait(queue_); // wait to get nInnerSegments (also superbins and pixelTypes) before using - auto connectedPixelSize_host_buf = allocBufWrapper(devHost, nInnerSegments, queue); - auto connectedPixelIndex_host_buf = allocBufWrapper(devHost, nInnerSegments, queue); - auto connectedPixelSize_dev_buf = allocBufWrapper(devAcc, nInnerSegments, queue); - auto connectedPixelIndex_dev_buf = allocBufWrapper(devAcc, nInnerSegments, queue); + auto connectedPixelSize_host_buf = allocBufWrapper(cms::alpakatools::host(), nInnerSegments, queue_); + auto connectedPixelIndex_host_buf = allocBufWrapper(cms::alpakatools::host(), nInnerSegments, queue_); + auto connectedPixelSize_dev_buf = allocBufWrapper(devAcc_, nInnerSegments, queue_); + auto connectedPixelIndex_dev_buf = allocBufWrapper(devAcc_, nInnerSegments, queue_); auto* connectedPixelSize_host = connectedPixelSize_host_buf.data(); auto* connectedPixelIndex_host = connectedPixelIndex_host_buf.data(); @@ -931,98 +857,109 @@ void Event::createPixelQuintuplets() { // Loop over # pLS for (unsigned int i = 0; i < nInnerSegments; i++) { - int8_t 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 > 2) or (pixelType < 0)) { - connectedPixelIndex_host[i] = 0; + 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 != PixelType::kHighPt) and (pixelType != PixelType::kLowPtPosCurv) and + (pixelType != PixelType::kLowPtNegCurv))) { connectedPixelSize_host[i] = 0; + connectedPixelIndex_host[i] = 0; continue; } + // Used pixel type to select correct size-index arrays - if (pixelType == 0) { - connectedPixelSize_host[i] = - pixelMapping_.connectedPixelsSizes[superbin]; //number of connected modules to this pixel - unsigned int connectedIdxBase = pixelMapping_.connectedPixelsIndex[superbin]; - connectedPixelIndex_host[i] = connectedIdxBase; - } else if (pixelType == 1) { - connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizesPos[superbin]; //number of pixel connected modules - unsigned int connectedIdxBase = pixelMapping_.connectedPixelsIndexPos[superbin] + pixelIndexOffsetPos; - connectedPixelIndex_host[i] = connectedIdxBase; - } else if (pixelType == 2) { - connectedPixelSize_host[i] = pixelMapping_.connectedPixelsSizesNeg[superbin]; //number of pixel connected modules - unsigned int connectedIdxBase = pixelMapping_.connectedPixelsIndexNeg[superbin] + pixelIndexOffsetNeg; - connectedPixelIndex_host[i] = connectedIdxBase; + switch (pixelType) { + case PixelType::kInvalid: + break; + 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 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 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 + connectedPixelIndex_host[i] = pixelMapping_.connectedPixelsIndexNeg[superbin] + pixelIndexOffsetNeg; + break; } } - alpaka::memcpy(queue, connectedPixelSize_dev_buf, connectedPixelSize_host_buf, nInnerSegments); - alpaka::memcpy(queue, connectedPixelIndex_dev_buf, connectedPixelIndex_host_buf, nInnerSegments); + alpaka::memcpy(queue_, connectedPixelSize_dev_buf, connectedPixelSize_host_buf, nInnerSegments); + alpaka::memcpy(queue_, connectedPixelIndex_dev_buf, connectedPixelIndex_host_buf, nInnerSegments); Vec3D const threadsPerBlockCreatePixQuints{1, 16, 16}; Vec3D const blocksPerGridCreatePixQuints{16, max_blocks, 1}; WorkDiv3D const createPixelQuintupletsInGPUFromMapv2_workDiv = createWorkDiv(blocksPerGridCreatePixQuints, threadsPerBlockCreatePixQuints, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, createPixelQuintupletsInGPUFromMapv2_workDiv, CreatePixelQuintupletsInGPUFromMapv2{}, *modulesBuffers_.data(), - *mdsInGPU, - *segmentsInGPU, - *tripletsInGPU, - *quintupletsInGPU, - *pixelQuintupletsInGPU, + *mdsInGPU_, + *segmentsInGPU_, + *tripletsInGPU_, + *quintupletsInGPU_, + *pixelQuintupletsInGPU_, connectedPixelSize_dev_buf.data(), connectedPixelIndex_dev_buf.data(), nInnerSegments, - *rangesInGPU); + *rangesInGPU_); Vec3D const threadsPerBlockDupPix{1, 16, 16}; Vec3D const blocksPerGridDupPix{1, max_blocks, 1}; WorkDiv3D const removeDupPixelQuintupletsInGPUFromMap_workDiv = createWorkDiv(blocksPerGridDupPix, threadsPerBlockDupPix, elementsPerThread); - alpaka::exec(queue, + alpaka::exec(queue_, removeDupPixelQuintupletsInGPUFromMap_workDiv, RemoveDupPixelQuintupletsInGPUFromMap{}, - *pixelQuintupletsInGPU); + *pixelQuintupletsInGPU_); WorkDiv1D const addpT5asTrackCandidateInGPU_workDiv = createWorkDiv({1}, {256}, {1}); - alpaka::exec(queue, + alpaka::exec(queue_, addpT5asTrackCandidateInGPU_workDiv, AddpT5asTrackCandidateInGPU{}, nLowerModules_, - *pixelQuintupletsInGPU, - *trackCandidatesInGPU, - *segmentsInGPU, - *rangesInGPU); + *pixelQuintupletsInGPU_, + *trackCandidatesInGPU_, + *segmentsInGPU_, + *rangesInGPU_); #ifdef WARNINGS - auto nPixelQuintuplets_buf = allocBufWrapper(devHost, 1, queue); + auto nPixelQuintuplets_buf = allocBufWrapper(cms::alpakatools::host(), 1, queue_); - alpaka::memcpy(queue, nPixelQuintuplets_buf, pixelQuintupletsBuffers->nPixelQuintuplets_buf); - alpaka::wait(queue); // wait to get the value before using it + alpaka::memcpy(queue_, nPixelQuintuplets_buf, pixelQuintupletsBuffers_->nPixelQuintuplets_buf); + alpaka::wait(queue_); // wait to get the value before using it std::cout << "number of pixel quintuplets = " << *nPixelQuintuplets_buf.data() << std::endl; #endif } void Event::addMiniDoubletsToEventExplicit() { - auto nMDsCPU_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, nMDsCPU_buf, miniDoubletsBuffers->nMDs_buf, nLowerModules_); + auto nMDsCPU_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, nMDsCPU_buf, miniDoubletsBuffers_->nMDs_buf, nLowerModules_); // FIXME: replace by ES host data - auto module_subdets_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_subdets_buf, modulesBuffers_.subdets_buf, nLowerModules_); + auto module_subdets_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_subdets_buf, modulesBuffers_.subdets_buf, nLowerModules_); - auto module_layers_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); + auto module_layers_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); - auto module_hitRanges_buf = allocBufWrapper(devHost, nLowerModules_ * 2, queue); - alpaka::memcpy(queue, module_hitRanges_buf, hitsBuffers->hitRanges_buf, nLowerModules_ * 2u); + auto module_hitRanges_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_ * 2, queue_); + alpaka::memcpy(queue_, module_hitRanges_buf, hitsBuffers_->hitRanges_buf, nLowerModules_ * 2u); - alpaka::wait(queue); // wait for inputs before using them + alpaka::wait(queue_); // wait for inputs before using them auto const* nMDsCPU = nMDsCPU_buf.data(); auto const* module_subdets = module_subdets_buf.data(); @@ -1031,7 +968,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]; @@ -1041,17 +978,17 @@ void Event::addMiniDoubletsToEventExplicit() { } void Event::addSegmentsToEventExplicit() { - auto nSegmentsCPU_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, nSegmentsCPU_buf, segmentsBuffers->nSegments_buf, nLowerModules_); + auto nSegmentsCPU_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, nSegmentsCPU_buf, segmentsBuffers_->nSegments_buf, nLowerModules_); // FIXME: replace by ES host data - auto module_subdets_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_subdets_buf, modulesBuffers_.subdets_buf, nLowerModules_); + auto module_subdets_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_subdets_buf, modulesBuffers_.subdets_buf, nLowerModules_); - auto module_layers_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); + auto module_layers_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); - alpaka::wait(queue); // wait for inputs before using them + alpaka::wait(queue_); // wait for inputs before using them auto const* nSegmentsCPU = nSegmentsCPU_buf.data(); auto const* module_subdets = module_subdets_buf.data(); @@ -1059,7 +996,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]; @@ -1069,20 +1006,20 @@ void Event::addSegmentsToEventExplicit() { } void Event::addQuintupletsToEventExplicit() { - auto nQuintupletsCPU_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, nQuintupletsCPU_buf, quintupletsBuffers->nQuintuplets_buf); + auto nQuintupletsCPU_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, nQuintupletsCPU_buf, quintupletsBuffers_->nQuintuplets_buf); // FIXME: replace by ES host data - auto module_subdets_buf = allocBufWrapper(devHost, nModules_, queue); - alpaka::memcpy(queue, module_subdets_buf, modulesBuffers_.subdets_buf, nModules_); + auto module_subdets_buf = allocBufWrapper(cms::alpakatools::host(), nModules_, queue_); + alpaka::memcpy(queue_, module_subdets_buf, modulesBuffers_.subdets_buf, nModules_); - auto module_layers_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); + auto module_layers_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); - auto module_quintupletModuleIndices_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_quintupletModuleIndices_buf, rangesBuffers->quintupletModuleIndices_buf); + auto module_quintupletModuleIndices_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_quintupletModuleIndices_buf, rangesBuffers_->quintupletModuleIndices_buf); - alpaka::wait(queue); // wait for inputs before using them + alpaka::wait(queue_); // wait for inputs before using them auto const* nQuintupletsCPU = nQuintupletsCPU_buf.data(); auto const* module_subdets = module_subdets_buf.data(); @@ -1091,7 +1028,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]; @@ -1101,17 +1038,17 @@ void Event::addQuintupletsToEventExplicit() { } void Event::addTripletsToEventExplicit() { - auto nTripletsCPU_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, nTripletsCPU_buf, tripletsBuffers->nTriplets_buf); + auto nTripletsCPU_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, nTripletsCPU_buf, tripletsBuffers_->nTriplets_buf); // FIXME: replace by ES host data - auto module_subdets_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_subdets_buf, modulesBuffers_.subdets_buf, nLowerModules_); + auto module_subdets_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_subdets_buf, modulesBuffers_.subdets_buf, nLowerModules_); - auto module_layers_buf = allocBufWrapper(devHost, nLowerModules_, queue); - alpaka::memcpy(queue, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); + auto module_layers_buf = allocBufWrapper(cms::alpakatools::host(), nLowerModules_, queue_); + alpaka::memcpy(queue_, module_layers_buf, modulesBuffers_.layers_buf, nLowerModules_); - alpaka::wait(queue); // wait for inputs before using them + alpaka::wait(queue_); // wait for inputs before using them auto const* nTripletsCPU = nTripletsCPU_buf.data(); auto const* module_subdets = module_subdets_buf.data(); @@ -1119,7 +1056,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]; @@ -1225,17 +1162,19 @@ unsigned int Event::getNumberOfTripletsByLayerBarrel(unsigned int layer) { retur unsigned int Event::getNumberOfTripletsByLayerEndcap(unsigned int layer) { return n_triplets_by_layer_endcap_[layer]; } int Event::getNumberOfPixelTriplets() { - auto nPixelTriplets_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nPixelTriplets_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nPixelTriplets_buf_h, pixelTripletsBuffers->nPixelTriplets_buf); + alpaka::memcpy(queue_, nPixelTriplets_buf_h, pixelTripletsBuffers_->nPixelTriplets_buf); + alpaka::wait(queue_); return *nPixelTriplets_buf_h.data(); } int Event::getNumberOfPixelQuintuplets() { - auto nPixelQuintuplets_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nPixelQuintuplets_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nPixelQuintuplets_buf_h, pixelQuintupletsBuffers->nPixelQuintuplets_buf); + alpaka::memcpy(queue_, nPixelQuintuplets_buf_h, pixelQuintupletsBuffers_->nPixelQuintuplets_buf); + alpaka::wait(queue_); return *nPixelQuintuplets_buf_h.data(); } @@ -1268,401 +1207,417 @@ unsigned int Event::getNumberOfQuintupletsByLayerEndcap(unsigned int layer) { } int Event::getNumberOfTrackCandidates() { - auto nTrackCandidates_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nTrackCandidates_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nTrackCandidates_buf_h, trackCandidatesBuffers->nTrackCandidates_buf); + alpaka::memcpy(queue_, nTrackCandidates_buf_h, trackCandidatesBuffers_->nTrackCandidates_buf); + alpaka::wait(queue_); return *nTrackCandidates_buf_h.data(); } int Event::getNumberOfPT5TrackCandidates() { - auto nTrackCandidatesPT5_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nTrackCandidatesPT5_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nTrackCandidatesPT5_buf_h, trackCandidatesBuffers->nTrackCandidatespT5_buf); - alpaka::wait(queue); + alpaka::memcpy(queue_, nTrackCandidatesPT5_buf_h, trackCandidatesBuffers_->nTrackCandidatespT5_buf); + alpaka::wait(queue_); return *nTrackCandidatesPT5_buf_h.data(); } int Event::getNumberOfPT3TrackCandidates() { - auto nTrackCandidatesPT3_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nTrackCandidatesPT3_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nTrackCandidatesPT3_buf_h, trackCandidatesBuffers->nTrackCandidatespT3_buf); + alpaka::memcpy(queue_, nTrackCandidatesPT3_buf_h, trackCandidatesBuffers_->nTrackCandidatespT3_buf); + alpaka::wait(queue_); return *nTrackCandidatesPT3_buf_h.data(); } int Event::getNumberOfPLSTrackCandidates() { - auto nTrackCandidatesPLS_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nTrackCandidatesPLS_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nTrackCandidatesPLS_buf_h, trackCandidatesBuffers->nTrackCandidatespLS_buf); + alpaka::memcpy(queue_, nTrackCandidatesPLS_buf_h, trackCandidatesBuffers_->nTrackCandidatespLS_buf); + alpaka::wait(queue_); return *nTrackCandidatesPLS_buf_h.data(); } int Event::getNumberOfPixelTrackCandidates() { - auto nTrackCandidates_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - auto nTrackCandidatesT5_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nTrackCandidates_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + auto nTrackCandidatesT5_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nTrackCandidates_buf_h, trackCandidatesBuffers->nTrackCandidates_buf); - alpaka::memcpy(queue, nTrackCandidatesT5_buf_h, trackCandidatesBuffers->nTrackCandidatesT5_buf); + alpaka::memcpy(queue_, nTrackCandidates_buf_h, trackCandidatesBuffers_->nTrackCandidates_buf); + alpaka::memcpy(queue_, nTrackCandidatesT5_buf_h, trackCandidatesBuffers_->nTrackCandidatesT5_buf); + alpaka::wait(queue_); return (*nTrackCandidates_buf_h.data()) - (*nTrackCandidatesT5_buf_h.data()); } int Event::getNumberOfT5TrackCandidates() { - auto nTrackCandidatesT5_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); + auto nTrackCandidatesT5_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); - alpaka::memcpy(queue, nTrackCandidatesT5_buf_h, trackCandidatesBuffers->nTrackCandidatesT5_buf); + alpaka::memcpy(queue_, nTrackCandidatesT5_buf_h, trackCandidatesBuffers_->nTrackCandidatesT5_buf); + alpaka::wait(queue_); return *nTrackCandidatesT5_buf_h.data(); } -HitsBuffer* Event::getHits(bool sync) //std::shared_ptr should take care of garbage collection +HitsBuffer& Event::getHits(bool sync) //std::shared_ptr should take care of garbage collection { - if (hitsInCPU == nullptr) { - auto nHits_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nHits_buf_h, hitsBuffers->nHits_buf); - alpaka::wait(queue); // wait for the value before using + if (!hitsInCPU_) { + auto nHits_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nHits_buf_h, hitsBuffers_->nHits_buf); + alpaka::wait(queue_); // wait for the value before using auto const nHits = *nHits_buf_h.data(); - hitsInCPU = new HitsBuffer(nModules_, nHits, devHost, queue); - hitsInCPU->setData(*hitsInCPU); - - *hitsInCPU->nHits_buf.data() = nHits; - alpaka::memcpy(queue, hitsInCPU->idxs_buf, hitsBuffers->idxs_buf, nHits); - alpaka::memcpy(queue, hitsInCPU->detid_buf, hitsBuffers->detid_buf, nHits); - alpaka::memcpy(queue, hitsInCPU->xs_buf, hitsBuffers->xs_buf, nHits); - alpaka::memcpy(queue, hitsInCPU->ys_buf, hitsBuffers->ys_buf, nHits); - alpaka::memcpy(queue, hitsInCPU->zs_buf, hitsBuffers->zs_buf, nHits); - alpaka::memcpy(queue, hitsInCPU->moduleIndices_buf, hitsBuffers->moduleIndices_buf, nHits); + hitsInCPU_.emplace(nModules_, nHits, cms::alpakatools::host(), queue_); + hitsInCPU_->setData(*hitsInCPU_); + + alpaka::memcpy(queue_, hitsInCPU_->nHits_buf, hitsBuffers_->nHits_buf); + alpaka::memcpy(queue_, hitsInCPU_->idxs_buf, hitsBuffers_->idxs_buf, nHits); + alpaka::memcpy(queue_, hitsInCPU_->detid_buf, hitsBuffers_->detid_buf, nHits); + alpaka::memcpy(queue_, hitsInCPU_->xs_buf, hitsBuffers_->xs_buf, nHits); + alpaka::memcpy(queue_, hitsInCPU_->ys_buf, hitsBuffers_->ys_buf, nHits); + alpaka::memcpy(queue_, hitsInCPU_->zs_buf, hitsBuffers_->zs_buf, nHits); + alpaka::memcpy(queue_, hitsInCPU_->moduleIndices_buf, hitsBuffers_->moduleIndices_buf, nHits); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return hitsInCPU; + return hitsInCPU_.value(); } -HitsBuffer* Event::getHitsInCMSSW(bool sync) { - if (hitsInCPU == nullptr) { - auto nHits_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nHits_buf_h, hitsBuffers->nHits_buf); - alpaka::wait(queue); // wait for the value before using +HitsBuffer& Event::getHitsInCMSSW(bool sync) { + if (!hitsInCPU_) { + auto nHits_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nHits_buf_h, hitsBuffers_->nHits_buf); + alpaka::wait(queue_); // wait for the value before using auto const nHits = *nHits_buf_h.data(); - hitsInCPU = new HitsBuffer(nModules_, nHits, devHost, queue); - hitsInCPU->setData(*hitsInCPU); + hitsInCPU_.emplace(nModules_, nHits, cms::alpakatools::host(), queue_); + hitsInCPU_->setData(*hitsInCPU_); - *hitsInCPU->nHits_buf.data() = nHits; - alpaka::memcpy(queue, hitsInCPU->idxs_buf, hitsBuffers->idxs_buf, nHits); + alpaka::memcpy(queue_, hitsInCPU_->nHits_buf, hitsBuffers_->nHits_buf); + alpaka::memcpy(queue_, hitsInCPU_->idxs_buf, hitsBuffers_->idxs_buf, nHits); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return hitsInCPU; + return hitsInCPU_.value(); } -ObjectRangesBuffer* Event::getRanges(bool sync) { - if (rangesInCPU == nullptr) { - rangesInCPU = new ObjectRangesBuffer(nModules_, nLowerModules_, devHost, queue); - rangesInCPU->setData(*rangesInCPU); +ObjectRangesBuffer& Event::getRanges(bool sync) { + if (!rangesInCPU_) { + rangesInCPU_.emplace(nModules_, nLowerModules_, cms::alpakatools::host(), queue_); + rangesInCPU_->setData(*rangesInCPU_); - alpaka::memcpy(queue, rangesInCPU->hitRanges_buf, rangesBuffers->hitRanges_buf); - alpaka::memcpy(queue, rangesInCPU->quintupletModuleIndices_buf, rangesBuffers->quintupletModuleIndices_buf); - alpaka::memcpy(queue, rangesInCPU->miniDoubletModuleIndices_buf, rangesBuffers->miniDoubletModuleIndices_buf); - alpaka::memcpy(queue, rangesInCPU->segmentModuleIndices_buf, rangesBuffers->segmentModuleIndices_buf); - alpaka::memcpy(queue, rangesInCPU->tripletModuleIndices_buf, rangesBuffers->tripletModuleIndices_buf); + alpaka::memcpy(queue_, rangesInCPU_->hitRanges_buf, rangesBuffers_->hitRanges_buf); + alpaka::memcpy(queue_, rangesInCPU_->quintupletModuleIndices_buf, rangesBuffers_->quintupletModuleIndices_buf); + alpaka::memcpy(queue_, rangesInCPU_->miniDoubletModuleIndices_buf, rangesBuffers_->miniDoubletModuleIndices_buf); + alpaka::memcpy(queue_, rangesInCPU_->segmentModuleIndices_buf, rangesBuffers_->segmentModuleIndices_buf); + alpaka::memcpy(queue_, rangesInCPU_->tripletModuleIndices_buf, rangesBuffers_->tripletModuleIndices_buf); if (sync) - alpaka::wait(queue); // wait to get completed host data + alpaka::wait(queue_); // wait to get completed host data } - return rangesInCPU; + return rangesInCPU_.value(); } -MiniDoubletsBuffer* Event::getMiniDoublets(bool sync) { - if (mdsInCPU == nullptr) { - // Get nMemoryLocations parameter to initialize host based mdsInCPU - auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nMemHost_buf_h, miniDoubletsBuffers->nMemoryLocations_buf); - alpaka::wait(queue); // wait for the value before using +MiniDoubletsBuffer& Event::getMiniDoublets(bool sync) { + if (!mdsInCPU_) { + // Get nMemoryLocations parameter to initialize host based mdsInCPU_ + auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nMemHost_buf_h, miniDoubletsBuffers_->nMemoryLocations_buf); + alpaka::wait(queue_); // wait for the value before using auto const nMemHost = *nMemHost_buf_h.data(); - mdsInCPU = new MiniDoubletsBuffer(nMemHost, nLowerModules_, devHost, queue); - mdsInCPU->setData(*mdsInCPU); - - *mdsInCPU->nMemoryLocations_buf.data() = nMemHost; - alpaka::memcpy(queue, mdsInCPU->anchorHitIndices_buf, miniDoubletsBuffers->anchorHitIndices_buf, nMemHost); - alpaka::memcpy(queue, mdsInCPU->outerHitIndices_buf, miniDoubletsBuffers->outerHitIndices_buf, nMemHost); - alpaka::memcpy(queue, mdsInCPU->dphichanges_buf, miniDoubletsBuffers->dphichanges_buf, nMemHost); - alpaka::memcpy(queue, mdsInCPU->nMDs_buf, miniDoubletsBuffers->nMDs_buf); - alpaka::memcpy(queue, mdsInCPU->totOccupancyMDs_buf, miniDoubletsBuffers->totOccupancyMDs_buf); + mdsInCPU_.emplace(nMemHost, nLowerModules_, cms::alpakatools::host(), queue_); + mdsInCPU_->setData(*mdsInCPU_); + + alpaka::memcpy(queue_, mdsInCPU_->nMemoryLocations_buf, miniDoubletsBuffers_->nMemoryLocations_buf); + alpaka::memcpy(queue_, mdsInCPU_->anchorHitIndices_buf, miniDoubletsBuffers_->anchorHitIndices_buf, nMemHost); + alpaka::memcpy(queue_, mdsInCPU_->outerHitIndices_buf, miniDoubletsBuffers_->outerHitIndices_buf, nMemHost); + alpaka::memcpy(queue_, mdsInCPU_->dphichanges_buf, miniDoubletsBuffers_->dphichanges_buf, nMemHost); + alpaka::memcpy(queue_, mdsInCPU_->nMDs_buf, miniDoubletsBuffers_->nMDs_buf); + alpaka::memcpy(queue_, mdsInCPU_->totOccupancyMDs_buf, miniDoubletsBuffers_->totOccupancyMDs_buf); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return mdsInCPU; + return mdsInCPU_.value(); } -SegmentsBuffer* Event::getSegments(bool sync) { - if (segmentsInCPU == nullptr) { - // Get nMemoryLocations parameter to initialize host based segmentsInCPU - auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nMemHost_buf_h, segmentsBuffers->nMemoryLocations_buf); - alpaka::wait(queue); // wait for the value before using +SegmentsBuffer& Event::getSegments(bool sync) { + if (!segmentsInCPU_) { + // Get nMemoryLocations parameter to initialize host based segmentsInCPU_ + auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nMemHost_buf_h, segmentsBuffers_->nMemoryLocations_buf); + alpaka::wait(queue_); // wait for the value before using auto const nMemHost = *nMemHost_buf_h.data(); - segmentsInCPU = - new SegmentsBuffer(nMemHost, nLowerModules_, n_max_pixel_segments_per_module, devHost, queue); - segmentsInCPU->setData(*segmentsInCPU); - - *segmentsInCPU->nMemoryLocations_buf.data() = nMemHost; - alpaka::memcpy(queue, segmentsInCPU->nSegments_buf, segmentsBuffers->nSegments_buf); - alpaka::memcpy(queue, segmentsInCPU->mdIndices_buf, segmentsBuffers->mdIndices_buf, 2u * nMemHost); - alpaka::memcpy(queue, - segmentsInCPU->innerMiniDoubletAnchorHitIndices_buf, - segmentsBuffers->innerMiniDoubletAnchorHitIndices_buf, + segmentsInCPU_.emplace(nMemHost, nLowerModules_, n_max_pixel_segments_per_module, cms::alpakatools::host(), queue_); + segmentsInCPU_->setData(*segmentsInCPU_); + + alpaka::memcpy(queue_, segmentsInCPU_->nMemoryLocations_buf, segmentsBuffers_->nMemoryLocations_buf); + alpaka::memcpy(queue_, segmentsInCPU_->nSegments_buf, segmentsBuffers_->nSegments_buf); + alpaka::memcpy(queue_, segmentsInCPU_->mdIndices_buf, segmentsBuffers_->mdIndices_buf, 2u * nMemHost); + alpaka::memcpy(queue_, + segmentsInCPU_->innerMiniDoubletAnchorHitIndices_buf, + segmentsBuffers_->innerMiniDoubletAnchorHitIndices_buf, nMemHost); - alpaka::memcpy(queue, - segmentsInCPU->outerMiniDoubletAnchorHitIndices_buf, - segmentsBuffers->outerMiniDoubletAnchorHitIndices_buf, + alpaka::memcpy(queue_, + segmentsInCPU_->outerMiniDoubletAnchorHitIndices_buf, + segmentsBuffers_->outerMiniDoubletAnchorHitIndices_buf, nMemHost); - alpaka::memcpy(queue, segmentsInCPU->totOccupancySegments_buf, segmentsBuffers->totOccupancySegments_buf); - alpaka::memcpy(queue, segmentsInCPU->ptIn_buf, segmentsBuffers->ptIn_buf); - alpaka::memcpy(queue, segmentsInCPU->eta_buf, segmentsBuffers->eta_buf); - alpaka::memcpy(queue, segmentsInCPU->phi_buf, segmentsBuffers->phi_buf); - alpaka::memcpy(queue, segmentsInCPU->seedIdx_buf, segmentsBuffers->seedIdx_buf); - alpaka::memcpy(queue, segmentsInCPU->isDup_buf, segmentsBuffers->isDup_buf); - alpaka::memcpy(queue, segmentsInCPU->isQuad_buf, segmentsBuffers->isQuad_buf); - alpaka::memcpy(queue, segmentsInCPU->score_buf, segmentsBuffers->score_buf); + alpaka::memcpy(queue_, segmentsInCPU_->totOccupancySegments_buf, segmentsBuffers_->totOccupancySegments_buf); + alpaka::memcpy(queue_, segmentsInCPU_->ptIn_buf, segmentsBuffers_->ptIn_buf); + alpaka::memcpy(queue_, segmentsInCPU_->eta_buf, segmentsBuffers_->eta_buf); + alpaka::memcpy(queue_, segmentsInCPU_->phi_buf, segmentsBuffers_->phi_buf); + alpaka::memcpy(queue_, segmentsInCPU_->seedIdx_buf, segmentsBuffers_->seedIdx_buf); + alpaka::memcpy(queue_, segmentsInCPU_->isDup_buf, segmentsBuffers_->isDup_buf); + alpaka::memcpy(queue_, segmentsInCPU_->isQuad_buf, segmentsBuffers_->isQuad_buf); + alpaka::memcpy(queue_, segmentsInCPU_->score_buf, segmentsBuffers_->score_buf); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return segmentsInCPU; + return segmentsInCPU_.value(); } -TripletsBuffer* Event::getTriplets(bool sync) { - if (tripletsInCPU == nullptr) { - // Get nMemoryLocations parameter to initialize host based tripletsInCPU - auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nMemHost_buf_h, tripletsBuffers->nMemoryLocations_buf); - alpaka::wait(queue); // wait for the value before using +TripletsBuffer& Event::getTriplets(bool sync) { + if (!tripletsInCPU_) { + // Get nMemoryLocations parameter to initialize host based tripletsInCPU_ + auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nMemHost_buf_h, tripletsBuffers_->nMemoryLocations_buf); + alpaka::wait(queue_); // wait for the value before using auto const nMemHost = *nMemHost_buf_h.data(); - tripletsInCPU = new TripletsBuffer(nMemHost, nLowerModules_, devHost, queue); - tripletsInCPU->setData(*tripletsInCPU); + tripletsInCPU_.emplace(nMemHost, nLowerModules_, cms::alpakatools::host(), queue_); + tripletsInCPU_->setData(*tripletsInCPU_); - *tripletsInCPU->nMemoryLocations_buf.data() = nMemHost; + alpaka::memcpy(queue_, tripletsInCPU_->nMemoryLocations_buf, tripletsBuffers_->nMemoryLocations_buf); #ifdef CUT_VALUE_DEBUG - alpaka::memcpy(queue, tripletsInCPU->zOut_buf, tripletsBuffers->zOut_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->zLo_buf, tripletsBuffers->zLo_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->zHi_buf, tripletsBuffers->zHi_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->zLoPointed_buf, tripletsBuffers->zLoPointed_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->zHiPointed_buf, tripletsBuffers->zHiPointed_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->dPhiCut_buf, tripletsBuffers->dPhiCut_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->betaInCut_buf, tripletsBuffers->betaInCut_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->rtLo_buf, tripletsBuffers->rtLo_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->rtHi_buf, tripletsBuffers->rtHi_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->zOut_buf, tripletsBuffers_->zOut_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->zLo_buf, tripletsBuffers_->zLo_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->zHi_buf, tripletsBuffers_->zHi_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->zLoPointed_buf, tripletsBuffers_->zLoPointed_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->zHiPointed_buf, tripletsBuffers_->zHiPointed_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->dPhiCut_buf, tripletsBuffers_->dPhiCut_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->betaInCut_buf, tripletsBuffers_->betaInCut_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->rtLo_buf, tripletsBuffers_->rtLo_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->rtHi_buf, tripletsBuffers_->rtHi_buf, nMemHost); #endif - alpaka::memcpy(queue, tripletsInCPU->hitIndices_buf, tripletsBuffers->hitIndices_buf, Params_T3::kHits * nMemHost); alpaka::memcpy( - queue, tripletsInCPU->logicalLayers_buf, tripletsBuffers->logicalLayers_buf, Params_T3::kLayers * nMemHost); - alpaka::memcpy(queue, tripletsInCPU->segmentIndices_buf, tripletsBuffers->segmentIndices_buf, 2 * nMemHost); - alpaka::memcpy(queue, tripletsInCPU->betaIn_buf, tripletsBuffers->betaIn_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->circleRadius_buf, tripletsBuffers->circleRadius_buf, nMemHost); - alpaka::memcpy(queue, tripletsInCPU->nTriplets_buf, tripletsBuffers->nTriplets_buf); - alpaka::memcpy(queue, tripletsInCPU->totOccupancyTriplets_buf, tripletsBuffers->totOccupancyTriplets_buf); + queue_, tripletsInCPU_->hitIndices_buf, tripletsBuffers_->hitIndices_buf, Params_T3::kHits * nMemHost); + alpaka::memcpy( + queue_, tripletsInCPU_->logicalLayers_buf, tripletsBuffers_->logicalLayers_buf, Params_T3::kLayers * nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->segmentIndices_buf, tripletsBuffers_->segmentIndices_buf, 2 * nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->betaIn_buf, tripletsBuffers_->betaIn_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->circleRadius_buf, tripletsBuffers_->circleRadius_buf, nMemHost); + alpaka::memcpy(queue_, tripletsInCPU_->nTriplets_buf, tripletsBuffers_->nTriplets_buf); + alpaka::memcpy(queue_, tripletsInCPU_->totOccupancyTriplets_buf, tripletsBuffers_->totOccupancyTriplets_buf); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return tripletsInCPU; + return tripletsInCPU_.value(); } -QuintupletsBuffer* Event::getQuintuplets(bool sync) { - if (quintupletsInCPU == nullptr) { - // Get nMemoryLocations parameter to initialize host based quintupletsInCPU - auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nMemHost_buf_h, quintupletsBuffers->nMemoryLocations_buf); - alpaka::wait(queue); // wait for the value before using +QuintupletsBuffer& Event::getQuintuplets(bool sync) { + if (!quintupletsInCPU_) { + // Get nMemoryLocations parameter to initialize host based quintupletsInCPU_ + auto nMemHost_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nMemHost_buf_h, quintupletsBuffers_->nMemoryLocations_buf); + alpaka::wait(queue_); // wait for the value before using auto const nMemHost = *nMemHost_buf_h.data(); - quintupletsInCPU = new QuintupletsBuffer(nMemHost, nLowerModules_, devHost, queue); - quintupletsInCPU->setData(*quintupletsInCPU); + quintupletsInCPU_.emplace(nMemHost, nLowerModules_, cms::alpakatools::host(), queue_); + quintupletsInCPU_->setData(*quintupletsInCPU_); - *quintupletsInCPU->nMemoryLocations_buf.data() = nMemHost; - alpaka::memcpy(queue, quintupletsInCPU->nQuintuplets_buf, quintupletsBuffers->nQuintuplets_buf); + alpaka::memcpy(queue_, quintupletsInCPU_->nMemoryLocations_buf, quintupletsBuffers_->nMemoryLocations_buf); + alpaka::memcpy(queue_, quintupletsInCPU_->nQuintuplets_buf, quintupletsBuffers_->nQuintuplets_buf); alpaka::memcpy( - queue, quintupletsInCPU->totOccupancyQuintuplets_buf, quintupletsBuffers->totOccupancyQuintuplets_buf); - alpaka::memcpy(queue, quintupletsInCPU->tripletIndices_buf, quintupletsBuffers->tripletIndices_buf, 2 * nMemHost); - alpaka::memcpy(queue, - quintupletsInCPU->lowerModuleIndices_buf, - quintupletsBuffers->lowerModuleIndices_buf, + queue_, quintupletsInCPU_->totOccupancyQuintuplets_buf, quintupletsBuffers_->totOccupancyQuintuplets_buf); + alpaka::memcpy( + queue_, quintupletsInCPU_->tripletIndices_buf, quintupletsBuffers_->tripletIndices_buf, 2 * nMemHost); + alpaka::memcpy(queue_, + quintupletsInCPU_->lowerModuleIndices_buf, + quintupletsBuffers_->lowerModuleIndices_buf, Params_T5::kLayers * nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->innerRadius_buf, quintupletsBuffers->innerRadius_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->bridgeRadius_buf, quintupletsBuffers->bridgeRadius_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->outerRadius_buf, quintupletsBuffers->outerRadius_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->isDup_buf, quintupletsBuffers->isDup_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->score_rphisum_buf, quintupletsBuffers->score_rphisum_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->eta_buf, quintupletsBuffers->eta_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->phi_buf, quintupletsBuffers->phi_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->chiSquared_buf, quintupletsBuffers->chiSquared_buf, nMemHost); - alpaka::memcpy(queue, quintupletsInCPU->rzChiSquared_buf, quintupletsBuffers->rzChiSquared_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->innerRadius_buf, quintupletsBuffers_->innerRadius_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->bridgeRadius_buf, quintupletsBuffers_->bridgeRadius_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->outerRadius_buf, quintupletsBuffers_->outerRadius_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->isDup_buf, quintupletsBuffers_->isDup_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->score_rphisum_buf, quintupletsBuffers_->score_rphisum_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->eta_buf, quintupletsBuffers_->eta_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->phi_buf, quintupletsBuffers_->phi_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->chiSquared_buf, quintupletsBuffers_->chiSquared_buf, nMemHost); + alpaka::memcpy(queue_, quintupletsInCPU_->rzChiSquared_buf, quintupletsBuffers_->rzChiSquared_buf, nMemHost); alpaka::memcpy( - queue, quintupletsInCPU->nonAnchorChiSquared_buf, quintupletsBuffers->nonAnchorChiSquared_buf, nMemHost); + queue_, quintupletsInCPU_->nonAnchorChiSquared_buf, quintupletsBuffers_->nonAnchorChiSquared_buf, nMemHost); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return quintupletsInCPU; + return quintupletsInCPU_.value(); } -PixelTripletsBuffer* Event::getPixelTriplets(bool sync) { - if (pixelTripletsInCPU == nullptr) { - // Get nPixelTriplets parameter to initialize host based quintupletsInCPU - auto nPixelTriplets_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nPixelTriplets_buf_h, pixelTripletsBuffers->nPixelTriplets_buf); - alpaka::wait(queue); // wait for the value before using +PixelTripletsBuffer& Event::getPixelTriplets(bool sync) { + if (!pixelTripletsInCPU_) { + // Get nPixelTriplets parameter to initialize host based quintupletsInCPU_ + auto nPixelTriplets_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nPixelTriplets_buf_h, pixelTripletsBuffers_->nPixelTriplets_buf); + alpaka::wait(queue_); // wait for the value before using auto const nPixelTriplets = *nPixelTriplets_buf_h.data(); - pixelTripletsInCPU = new PixelTripletsBuffer(nPixelTriplets, devHost, queue); - pixelTripletsInCPU->setData(*pixelTripletsInCPU); + pixelTripletsInCPU_.emplace(nPixelTriplets, cms::alpakatools::host(), queue_); + pixelTripletsInCPU_->setData(*pixelTripletsInCPU_); - *pixelTripletsInCPU->nPixelTriplets_buf.data() = nPixelTriplets; + alpaka::memcpy(queue_, pixelTripletsInCPU_->nPixelTriplets_buf, pixelTripletsBuffers_->nPixelTriplets_buf); + alpaka::memcpy(queue_, + pixelTripletsInCPU_->totOccupancyPixelTriplets_buf, + pixelTripletsBuffers_->totOccupancyPixelTriplets_buf); alpaka::memcpy( - queue, pixelTripletsInCPU->totOccupancyPixelTriplets_buf, pixelTripletsBuffers->totOccupancyPixelTriplets_buf); - alpaka::memcpy(queue, pixelTripletsInCPU->rzChiSquared_buf, pixelTripletsBuffers->rzChiSquared_buf, nPixelTriplets); + queue_, pixelTripletsInCPU_->rzChiSquared_buf, pixelTripletsBuffers_->rzChiSquared_buf, nPixelTriplets); alpaka::memcpy( - queue, pixelTripletsInCPU->rPhiChiSquared_buf, pixelTripletsBuffers->rPhiChiSquared_buf, nPixelTriplets); - alpaka::memcpy(queue, - pixelTripletsInCPU->rPhiChiSquaredInwards_buf, - pixelTripletsBuffers->rPhiChiSquaredInwards_buf, + queue_, pixelTripletsInCPU_->rPhiChiSquared_buf, pixelTripletsBuffers_->rPhiChiSquared_buf, nPixelTriplets); + alpaka::memcpy(queue_, + pixelTripletsInCPU_->rPhiChiSquaredInwards_buf, + pixelTripletsBuffers_->rPhiChiSquaredInwards_buf, nPixelTriplets); alpaka::memcpy( - queue, pixelTripletsInCPU->tripletIndices_buf, pixelTripletsBuffers->tripletIndices_buf, nPixelTriplets); - alpaka::memcpy(queue, - pixelTripletsInCPU->pixelSegmentIndices_buf, - pixelTripletsBuffers->pixelSegmentIndices_buf, + queue_, pixelTripletsInCPU_->tripletIndices_buf, pixelTripletsBuffers_->tripletIndices_buf, nPixelTriplets); + alpaka::memcpy(queue_, + pixelTripletsInCPU_->pixelSegmentIndices_buf, + pixelTripletsBuffers_->pixelSegmentIndices_buf, nPixelTriplets); - alpaka::memcpy(queue, pixelTripletsInCPU->pixelRadius_buf, pixelTripletsBuffers->pixelRadius_buf, nPixelTriplets); alpaka::memcpy( - queue, pixelTripletsInCPU->tripletRadius_buf, pixelTripletsBuffers->tripletRadius_buf, nPixelTriplets); - alpaka::memcpy(queue, pixelTripletsInCPU->isDup_buf, pixelTripletsBuffers->isDup_buf, nPixelTriplets); - alpaka::memcpy(queue, pixelTripletsInCPU->eta_buf, pixelTripletsBuffers->eta_buf, nPixelTriplets); - alpaka::memcpy(queue, pixelTripletsInCPU->phi_buf, pixelTripletsBuffers->phi_buf, nPixelTriplets); - alpaka::memcpy(queue, pixelTripletsInCPU->score_buf, pixelTripletsBuffers->score_buf, nPixelTriplets); + queue_, pixelTripletsInCPU_->pixelRadius_buf, pixelTripletsBuffers_->pixelRadius_buf, nPixelTriplets); + alpaka::memcpy( + queue_, pixelTripletsInCPU_->tripletRadius_buf, pixelTripletsBuffers_->tripletRadius_buf, nPixelTriplets); + alpaka::memcpy(queue_, pixelTripletsInCPU_->isDup_buf, pixelTripletsBuffers_->isDup_buf, nPixelTriplets); + alpaka::memcpy(queue_, pixelTripletsInCPU_->eta_buf, pixelTripletsBuffers_->eta_buf, nPixelTriplets); + alpaka::memcpy(queue_, pixelTripletsInCPU_->phi_buf, pixelTripletsBuffers_->phi_buf, nPixelTriplets); + alpaka::memcpy(queue_, pixelTripletsInCPU_->score_buf, pixelTripletsBuffers_->score_buf, nPixelTriplets); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return pixelTripletsInCPU; + return pixelTripletsInCPU_.value(); } -PixelQuintupletsBuffer* Event::getPixelQuintuplets(bool sync) { - if (pixelQuintupletsInCPU == nullptr) { - // Get nPixelQuintuplets parameter to initialize host based quintupletsInCPU - auto nPixelQuintuplets_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nPixelQuintuplets_buf_h, pixelQuintupletsBuffers->nPixelQuintuplets_buf); - alpaka::wait(queue); // wait for the value before using +PixelQuintupletsBuffer& Event::getPixelQuintuplets(bool sync) { + if (!pixelQuintupletsInCPU_) { + // Get nPixelQuintuplets parameter to initialize host based quintupletsInCPU_ + auto nPixelQuintuplets_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nPixelQuintuplets_buf_h, pixelQuintupletsBuffers_->nPixelQuintuplets_buf); + alpaka::wait(queue_); // wait for the value before using auto const nPixelQuintuplets = *nPixelQuintuplets_buf_h.data(); - pixelQuintupletsInCPU = new PixelQuintupletsBuffer(nPixelQuintuplets, devHost, queue); - pixelQuintupletsInCPU->setData(*pixelQuintupletsInCPU); + pixelQuintupletsInCPU_.emplace(nPixelQuintuplets, cms::alpakatools::host(), queue_); + pixelQuintupletsInCPU_->setData(*pixelQuintupletsInCPU_); - *pixelQuintupletsInCPU->nPixelQuintuplets_buf.data() = nPixelQuintuplets; - alpaka::memcpy(queue, - pixelQuintupletsInCPU->totOccupancyPixelQuintuplets_buf, - pixelQuintupletsBuffers->totOccupancyPixelQuintuplets_buf); alpaka::memcpy( - queue, pixelQuintupletsInCPU->rzChiSquared_buf, pixelQuintupletsBuffers->rzChiSquared_buf, nPixelQuintuplets); - alpaka::memcpy(queue, - pixelQuintupletsInCPU->rPhiChiSquared_buf, - pixelQuintupletsBuffers->rPhiChiSquared_buf, + queue_, pixelQuintupletsInCPU_->nPixelQuintuplets_buf, pixelQuintupletsBuffers_->nPixelQuintuplets_buf); + alpaka::memcpy(queue_, + pixelQuintupletsInCPU_->totOccupancyPixelQuintuplets_buf, + pixelQuintupletsBuffers_->totOccupancyPixelQuintuplets_buf); + alpaka::memcpy(queue_, + pixelQuintupletsInCPU_->rzChiSquared_buf, + pixelQuintupletsBuffers_->rzChiSquared_buf, nPixelQuintuplets); - alpaka::memcpy(queue, - pixelQuintupletsInCPU->rPhiChiSquaredInwards_buf, - pixelQuintupletsBuffers->rPhiChiSquaredInwards_buf, + alpaka::memcpy(queue_, + pixelQuintupletsInCPU_->rPhiChiSquared_buf, + pixelQuintupletsBuffers_->rPhiChiSquared_buf, + nPixelQuintuplets); + alpaka::memcpy(queue_, + pixelQuintupletsInCPU_->rPhiChiSquaredInwards_buf, + pixelQuintupletsBuffers_->rPhiChiSquaredInwards_buf, + nPixelQuintuplets); + alpaka::memcpy(queue_, + pixelQuintupletsInCPU_->pixelIndices_buf, + pixelQuintupletsBuffers_->pixelIndices_buf, nPixelQuintuplets); alpaka::memcpy( - queue, pixelQuintupletsInCPU->pixelIndices_buf, pixelQuintupletsBuffers->pixelIndices_buf, nPixelQuintuplets); - alpaka::memcpy( - queue, pixelQuintupletsInCPU->T5Indices_buf, pixelQuintupletsBuffers->T5Indices_buf, nPixelQuintuplets); - alpaka::memcpy(queue, pixelQuintupletsInCPU->isDup_buf, pixelQuintupletsBuffers->isDup_buf, nPixelQuintuplets); - alpaka::memcpy(queue, pixelQuintupletsInCPU->score_buf, pixelQuintupletsBuffers->score_buf, nPixelQuintuplets); + queue_, pixelQuintupletsInCPU_->T5Indices_buf, pixelQuintupletsBuffers_->T5Indices_buf, nPixelQuintuplets); + alpaka::memcpy(queue_, pixelQuintupletsInCPU_->isDup_buf, pixelQuintupletsBuffers_->isDup_buf, nPixelQuintuplets); + alpaka::memcpy(queue_, pixelQuintupletsInCPU_->score_buf, pixelQuintupletsBuffers_->score_buf, nPixelQuintuplets); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return pixelQuintupletsInCPU; + return pixelQuintupletsInCPU_.value(); } -TrackCandidatesBuffer* Event::getTrackCandidates(bool sync) { - if (trackCandidatesInCPU == nullptr) { - // Get nTrackCanHost parameter to initialize host based trackCandidatesInCPU - auto nTrackCanHost_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nTrackCanHost_buf_h, trackCandidatesBuffers->nTrackCandidates_buf); - alpaka::wait(queue); +TrackCandidatesBuffer& Event::getTrackCandidates(bool sync) { + if (!trackCandidatesInCPU_) { + // Get nTrackCanHost parameter to initialize host based trackCandidatesInCPU_ + auto nTrackCanHost_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nTrackCanHost_buf_h, trackCandidatesBuffers_->nTrackCandidates_buf); + trackCandidatesInCPU_.emplace( + n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, cms::alpakatools::host(), queue_); + trackCandidatesInCPU_->setData(*trackCandidatesInCPU_); + alpaka::wait(queue_); // wait here before we get nTrackCanHost and trackCandidatesInCPU_ becomes usable auto const nTrackCanHost = *nTrackCanHost_buf_h.data(); - trackCandidatesInCPU = new TrackCandidatesBuffer( - n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, devHost, queue); - trackCandidatesInCPU->setData(*trackCandidatesInCPU); - - *trackCandidatesInCPU->nTrackCandidates_buf.data() = nTrackCanHost; - alpaka::memcpy(queue, - trackCandidatesInCPU->hitIndices_buf, - trackCandidatesBuffers->hitIndices_buf, + + *trackCandidatesInCPU_->nTrackCandidates_buf.data() = nTrackCanHost; + alpaka::memcpy(queue_, + trackCandidatesInCPU_->hitIndices_buf, + trackCandidatesBuffers_->hitIndices_buf, Params_pT5::kHits * nTrackCanHost); alpaka::memcpy( - queue, trackCandidatesInCPU->pixelSeedIndex_buf, trackCandidatesBuffers->pixelSeedIndex_buf, nTrackCanHost); - alpaka::memcpy(queue, - trackCandidatesInCPU->logicalLayers_buf, - trackCandidatesBuffers->logicalLayers_buf, + queue_, trackCandidatesInCPU_->pixelSeedIndex_buf, trackCandidatesBuffers_->pixelSeedIndex_buf, nTrackCanHost); + alpaka::memcpy(queue_, + trackCandidatesInCPU_->logicalLayers_buf, + trackCandidatesBuffers_->logicalLayers_buf, Params_pT5::kLayers * nTrackCanHost); - alpaka::memcpy(queue, - trackCandidatesInCPU->directObjectIndices_buf, - trackCandidatesBuffers->directObjectIndices_buf, + alpaka::memcpy(queue_, + trackCandidatesInCPU_->directObjectIndices_buf, + trackCandidatesBuffers_->directObjectIndices_buf, nTrackCanHost); - alpaka::memcpy( - queue, trackCandidatesInCPU->objectIndices_buf, trackCandidatesBuffers->objectIndices_buf, 2 * nTrackCanHost); - alpaka::memcpy(queue, - trackCandidatesInCPU->trackCandidateType_buf, - trackCandidatesBuffers->trackCandidateType_buf, + alpaka::memcpy(queue_, + trackCandidatesInCPU_->objectIndices_buf, + trackCandidatesBuffers_->objectIndices_buf, + 2 * nTrackCanHost); + alpaka::memcpy(queue_, + trackCandidatesInCPU_->trackCandidateType_buf, + trackCandidatesBuffers_->trackCandidateType_buf, nTrackCanHost); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return trackCandidatesInCPU; + return trackCandidatesInCPU_.value(); } -TrackCandidatesBuffer* Event::getTrackCandidatesInCMSSW(bool sync) { - if (trackCandidatesInCPU == nullptr) { - // Get nTrackCanHost parameter to initialize host based trackCandidatesInCPU - auto nTrackCanHost_buf_h = cms::alpakatools::make_host_buffer(queue, 1u); - alpaka::memcpy(queue, nTrackCanHost_buf_h, trackCandidatesBuffers->nTrackCandidates_buf); - alpaka::wait(queue); // wait for the value before using +TrackCandidatesBuffer& Event::getTrackCandidatesInCMSSW(bool sync) { + if (!trackCandidatesInCPU_) { + // Get nTrackCanHost parameter to initialize host based trackCandidatesInCPU_ + auto nTrackCanHost_buf_h = cms::alpakatools::make_host_buffer(queue_, 1u); + alpaka::memcpy(queue_, nTrackCanHost_buf_h, trackCandidatesBuffers_->nTrackCandidates_buf); + trackCandidatesInCPU_.emplace( + n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, cms::alpakatools::host(), queue_); + trackCandidatesInCPU_->setData(*trackCandidatesInCPU_); + alpaka::wait(queue_); // wait for the value before using and trackCandidatesInCPU_ becomes usable auto const nTrackCanHost = *nTrackCanHost_buf_h.data(); - trackCandidatesInCPU = new TrackCandidatesBuffer( - n_max_nonpixel_track_candidates + n_max_pixel_track_candidates, devHost, queue); - trackCandidatesInCPU->setData(*trackCandidatesInCPU); - - *trackCandidatesInCPU->nTrackCandidates_buf.data() = nTrackCanHost; - alpaka::memcpy(queue, - trackCandidatesInCPU->hitIndices_buf, - trackCandidatesBuffers->hitIndices_buf, + + *trackCandidatesInCPU_->nTrackCandidates_buf.data() = nTrackCanHost; + alpaka::memcpy(queue_, + trackCandidatesInCPU_->hitIndices_buf, + trackCandidatesBuffers_->hitIndices_buf, Params_pT5::kHits * nTrackCanHost); alpaka::memcpy( - queue, trackCandidatesInCPU->pixelSeedIndex_buf, trackCandidatesBuffers->pixelSeedIndex_buf, nTrackCanHost); - alpaka::memcpy(queue, - trackCandidatesInCPU->trackCandidateType_buf, - trackCandidatesBuffers->trackCandidateType_buf, + queue_, trackCandidatesInCPU_->pixelSeedIndex_buf, trackCandidatesBuffers_->pixelSeedIndex_buf, nTrackCanHost); + alpaka::memcpy(queue_, + trackCandidatesInCPU_->trackCandidateType_buf, + trackCandidatesBuffers_->trackCandidateType_buf, nTrackCanHost); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return trackCandidatesInCPU; + return trackCandidatesInCPU_.value(); } -ModulesBuffer* Event::getModules(bool isFull, bool sync) { - if (modulesInCPU == nullptr) { +ModulesBuffer& Event::getModules(bool isFull, bool sync) { + if (!modulesInCPU_) { // The last input here is just a small placeholder for the allocation. - modulesInCPU = new ModulesBuffer(devHost, nModules_, nPixels_); + modulesInCPU_.emplace(cms::alpakatools::host(), nModules_, nPixels_); - modulesInCPU->copyFromSrc(queue, modulesBuffers_, isFull); + modulesInCPU_->copyFromSrc(queue_, modulesBuffers_, isFull); if (sync) - alpaka::wait(queue); // host consumers expect filled data + alpaka::wait(queue_); // host consumers expect filled data } - return modulesInCPU; + return modulesInCPU_.value(); } diff --git a/RecoTracker/LSTCore/src/alpaka/Event.h b/RecoTracker/LSTCore/src/alpaka/Event.h index 2ad8e150ece88..2b09565cf4176 100644 --- a/RecoTracker/LSTCore/src/alpaka/Event.h +++ b/RecoTracker/LSTCore/src/alpaka/Event.h @@ -1,6 +1,8 @@ #ifndef RecoTracker_LSTCore_src_alpaka_Event_h #define RecoTracker_LSTCore_src_alpaka_Event_h +#include + #include "RecoTracker/LSTCore/interface/alpaka/Constants.h" #include "RecoTracker/LSTCore/interface/alpaka/LST.h" #include "RecoTracker/LSTCore/interface/Module.h" @@ -17,193 +19,180 @@ #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 - ObjectRanges* rangesInGPU; - ObjectRangesBuffer* rangesBuffers; - Hits* hitsInGPU; - HitsBuffer* hitsBuffers; - MiniDoublets* mdsInGPU; - MiniDoubletsBuffer* miniDoubletsBuffers; - Segments* segmentsInGPU; - SegmentsBuffer* segmentsBuffers; - Triplets* tripletsInGPU; - TripletsBuffer* tripletsBuffers; - Quintuplets* quintupletsInGPU; - QuintupletsBuffer* quintupletsBuffers; - TrackCandidates* trackCandidatesInGPU; - TrackCandidatesBuffer* trackCandidatesBuffers; - PixelTriplets* pixelTripletsInGPU; - PixelTripletsBuffer* pixelTripletsBuffers; - PixelQuintuplets* pixelQuintupletsInGPU; - PixelQuintupletsBuffer* pixelQuintupletsBuffers; - - //CPU interface stuff - ObjectRangesBuffer* rangesInCPU; - HitsBuffer* hitsInCPU; - MiniDoubletsBuffer* mdsInCPU; - SegmentsBuffer* segmentsInCPU; - TripletsBuffer* tripletsInCPU; - TrackCandidatesBuffer* trackCandidatesInCPU; - ModulesBuffer* modulesInCPU; - QuintupletsBuffer* quintupletsInCPU; - PixelTripletsBuffer* pixelTripletsInCPU; - PixelQuintupletsBuffer* pixelQuintupletsInCPU; - - void initSync(bool verbose); - - int* superbinCPU; - int8_t* pixelTypeCPU; - - 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 lst - -} // namespace ALPAKA_ACCELERATOR_NAMESPACE +namespace ALPAKA_ACCELERATOR_NAMESPACE::lst { + + class Event { + private: + Queue& queue_; + Device devAcc_; + 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& q, const LSTESData* deviceESData) + : queue_(q), + devAcc_(alpaka::getDev(q)), + 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 e847eb892af8c..65543720a1d34 100644 --- a/RecoTracker/LSTCore/src/alpaka/LST.dev.cc +++ b/RecoTracker/LSTCore/src/alpaka/LST.dev.cc @@ -2,9 +2,10 @@ #include "Event.h" -using namespace ALPAKA_ACCELERATOR_NAMESPACE; +using namespace ALPAKA_ACCELERATOR_NAMESPACE::lst; #include "Math/Vector3D.h" +#include "Math/VectorUtil.h" using XYZVector = ROOT::Math::XYZVector; namespace { @@ -19,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; @@ -77,14 +78,13 @@ void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::prepareInput(std::vector con std::vector hitIdxs(ph2_detId.size()); std::vector superbin_vec; - std::vector pixelType_vec; + std::vector pixelType_vec; std::vector isQuad_vec; std::iota(hitIdxs.begin(), hitIdxs.end(), 0); const int hit_size = trkX.size(); for (size_t iSeed = 0; iSeed < n_see; iSeed++) { XYZVector p3LH(see_stateTrajGlbPx[iSeed], see_stateTrajGlbPy[iSeed], see_stateTrajGlbPz[iSeed]); - XYZVector p3LH_helper(see_stateTrajGlbPx[iSeed], see_stateTrajGlbPy[iSeed], see_stateTrajGlbPz[iSeed]); float ptIn = p3LH.rho(); float eta = p3LH.eta(); float ptErr = see_ptErr[iSeed]; @@ -94,22 +94,23 @@ void ALPAKA_ACCELERATOR_NAMESPACE::lst::LST::prepareInput(std::vector con XYZVector p3PCA(see_px[iSeed], see_py[iSeed], see_pz[iSeed]); XYZVector r3PCA(calculateR3FromPCA(p3PCA, see_dxy[iSeed], see_dz[iSeed])); - float pixelSegmentDeltaPhiChange = (r3LH - p3LH_helper).phi(); //FIXME: this looks like a bug + // The charge could be used directly in the line below + float pixelSegmentDeltaPhiChange = ROOT::Math::VectorUtil::DeltaPhi(p3LH, r3LH); float etaErr = see_etaErr[iSeed]; float px = p3LH.x(); float py = p3LH.y(); float pz = p3LH.z(); int charge = see_q[iSeed]; - int pixtype = -1; + PixelType pixtype = PixelType::kInvalid; if (ptIn >= 2.0) - pixtype = 0; + pixtype = PixelType::kHighPt; else if (ptIn >= (0.8 - 2 * ptErr) and ptIn < 2.0) { if (pixelSegmentDeltaPhiChange >= 0) - pixtype = 1; + pixtype = PixelType::kLowPtPosCurv; else - pixtype = 2; + pixtype = PixelType::kLowPtNegCurv; } else continue; @@ -211,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; @@ -246,14 +247,14 @@ 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; std::vector tc_trackCandidateType; - HitsBuffer& hitsInGPU = (*event.getHitsInCMSSW(false)); // sync on next line - TrackCandidates const* trackCandidates = event.getTrackCandidatesInCMSSW()->data(); + HitsBuffer& hitsInGPU = event.getHitsInCMSSW(false); // sync on next line + TrackCandidates const* trackCandidates = event.getTrackCandidatesInCMSSW().data(); unsigned int nTrackCandidates = *trackCandidates->nTrackCandidates; @@ -274,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 b74de58f3c233..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; - int8_t* 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(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 060308e4dabab..c0e52d0a0d194 100644 --- a/RecoTracker/LSTCore/standalone/bin/lst.cc +++ b/RecoTracker/LSTCore/standalone/bin/lst.cc @@ -2,6 +2,9 @@ #include +using LSTEvent = ALPAKA_ACCELERATOR_NAMESPACE::lst::Event; +using namespace ::lst; + //___________________________________________________________________________________________________________________________________________________________________________________________ int main(int argc, char **argv) { //******************************************************************************** @@ -338,7 +341,7 @@ void run_lst() { std::vector> out_charge_vec; std::vector> out_seedIdx_vec; std::vector> out_superbin_vec; - std::vector> out_pixelType_vec; + std::vector> out_pixelType_vec; std::vector> out_isQuad_vec; std::vector evt_num; std::vector file_name; @@ -383,10 +386,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 eb48917952a38..426a74babc4d1 100644 --- a/RecoTracker/LSTCore/standalone/code/core/AccessHelper.cc +++ b/RecoTracker/LSTCore/standalone/code/core/AccessHelper.cc @@ -9,7 +9,7 @@ using namespace ALPAKA_ACCELERATOR_NAMESPACE::lst; //____________________________________________________________________________________________ std::tuple, std::vector> convertHitsToHitIdxsAndHitTypes( Event* event, std::vector hits) { - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); std::vector hitidxs; std::vector hittypes; for (auto& hit : hits) { @@ -28,10 +28,10 @@ std::tuple, std::vector> convertHitsToHi //____________________________________________________________________________________________ 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(); + Segments const* segments = event->getSegments().data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + ObjectRanges const* rangesEvt = event->getRanges().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]; @@ -47,7 +47,7 @@ std::vector getPixelHitsFrompLS(Event* event, unsigned int pLS) { //____________________________________________________________________________________________ std::vector getPixelHitIdxsFrompLS(Event* event, unsigned int pLS) { - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); std::vector hits = getPixelHitsFrompLS(event, pLS); std::vector hitidxs; for (auto& hit : hits) @@ -74,7 +74,7 @@ std::tuple, std::vector> getHitIdxsAndHi //____________________________________________________________________________________________ std::vector getHitsFromMD(Event* event, unsigned int MD) { - MiniDoublets const* miniDoublets = event->getMiniDoublets()->data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); unsigned int hit_1 = miniDoublets->anchorHitIndices[MD]; unsigned int hit_2 = miniDoublets->outerHitIndices[MD]; return {hit_1, hit_2}; @@ -92,7 +92,7 @@ std::tuple, std::vector> getHitIdxsAndHi //____________________________________________________________________________________________ std::vector getMDsFromLS(Event* event, unsigned int LS) { - Segments const* segments = event->getSegments()->data(); + Segments const* segments = event->getSegments().data(); unsigned int MD_1 = segments->mdIndices[2 * LS]; unsigned int MD_2 = segments->mdIndices[2 * LS + 1]; return {MD_1, MD_2}; @@ -118,7 +118,7 @@ std::tuple, std::vector> getHitIdxsAndHi //____________________________________________________________________________________________ std::vector getLSsFromT3(Event* event, unsigned int T3) { - Triplets const* triplets = event->getTriplets()->data(); + Triplets const* triplets = event->getTriplets().data(); unsigned int LS_1 = triplets->segmentIndices[2 * T3]; unsigned int LS_2 = triplets->segmentIndices[2 * T3 + 1]; return {LS_1, LS_2}; @@ -153,7 +153,7 @@ std::tuple, std::vector> getHitIdxsAndHi //____________________________________________________________________________________________ std::vector getT3sFromT5(Event* event, unsigned int T5) { - Quintuplets const* quintuplets = event->getQuintuplets()->data(); + Quintuplets const* quintuplets = event->getQuintuplets().data(); unsigned int T3_1 = quintuplets->tripletIndices[2 * T5]; unsigned int T3_2 = quintuplets->tripletIndices[2 * T5 + 1]; return {T3_1, T3_2}; @@ -190,7 +190,7 @@ std::vector getHitsFromT5(Event* event, unsigned int T5) { //____________________________________________________________________________________________ std::vector getHitIdxsFromT5(Event* event, unsigned int T5) { - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); std::vector hits = getHitsFromT5(event, T5); std::vector hitidxs; for (auto& hit : hits) @@ -201,7 +201,7 @@ std::vector getHitIdxsFromT5(Event* event, unsigned int T5) { std::vector getModuleIdxsFromT5(Event* event, unsigned int T5) { std::vector hits = getHitsFromT5(event, T5); std::vector module_idxs; - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); for (auto& hitIdx : hits) { module_idxs.push_back(hitsEvt->moduleIndices[hitIdx]); } @@ -225,16 +225,16 @@ 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(); + PixelTriplets const* pixelTriplets = event->getPixelTriplets().data(); + ObjectRanges const* rangesEvt = event->getRanges().data(); + Modules const* modulesEvt = event->getModules().data(); const unsigned int pLS_offset = rangesEvt->segmentModuleIndices[*(modulesEvt->nLowerModules)]; return pixelTriplets->pixelSegmentIndices[pT3] - pLS_offset; } //____________________________________________________________________________________________ unsigned int getT3FrompT3(Event* event, unsigned int pT3) { - PixelTriplets const* pixelTriplets = event->getPixelTriplets()->data(); + PixelTriplets const* pixelTriplets = event->getPixelTriplets().data(); return pixelTriplets->tripletIndices[pT3]; } @@ -274,7 +274,7 @@ std::vector getHitsFrompT3(Event* event, unsigned int pT3) { //____________________________________________________________________________________________ std::vector getHitIdxsFrompT3(Event* event, unsigned int pT3) { - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); std::vector hits = getHitsFrompT3(event, pT3); std::vector hitidxs; for (auto& hit : hits) @@ -285,7 +285,7 @@ std::vector getHitIdxsFrompT3(Event* event, unsigned int pT3) { std::vector getModuleIdxsFrompT3(Event* event, unsigned int pT3) { std::vector hits = getOuterTrackerHitsFrompT3(event, pT3); std::vector module_idxs; - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); for (auto& hitIdx : hits) { module_idxs.push_back(hitsEvt->moduleIndices[hitIdx]); } @@ -314,16 +314,16 @@ 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(); + PixelQuintuplets const* pixelQuintuplets = event->getPixelQuintuplets().data(); + ObjectRanges const* rangesEvt = event->getRanges().data(); + Modules const* modulesEvt = event->getModules().data(); const unsigned int pLS_offset = rangesEvt->segmentModuleIndices[*(modulesEvt->nLowerModules)]; return pixelQuintuplets->pixelIndices[pT5] - pLS_offset; } //____________________________________________________________________________________________ unsigned int getT5FrompT5(Event* event, unsigned int pT5) { - PixelQuintuplets const* pixelQuintuplets = event->getPixelQuintuplets()->data(); + PixelQuintuplets const* pixelQuintuplets = event->getPixelQuintuplets().data(); return pixelQuintuplets->T5Indices[pT5]; } @@ -369,7 +369,7 @@ std::vector getHitsFrompT5(Event* event, unsigned int pT5) { //____________________________________________________________________________________________ std::vector getHitIdxsFrompT5(Event* event, unsigned int pT5) { - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); std::vector hits = getHitsFrompT5(event, pT5); std::vector hitidxs; for (auto& hit : hits) @@ -381,7 +381,7 @@ std::vector getHitIdxsFrompT5(Event* event, unsigned int pT5) { std::vector getModuleIdxsFrompT5(Event* event, unsigned int pT5) { std::vector hits = getOuterTrackerHitsFrompT5(event, pT5); std::vector module_idxs; - Hits const* hitsEvt = event->getHits()->data(); + Hits const* hitsEvt = event->getHits().data(); for (auto& hitIdx : hits) { module_idxs.push_back(hitsEvt->moduleIndices[hitIdx]); } @@ -412,7 +412,7 @@ std::tuple, std::vector> getHitIdxsAndHi //____________________________________________________________________________________________ std::vector getLSsFromTC(Event* event, unsigned int TC) { // Get the type of the track candidate - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); short type = trackCandidates->trackCandidateType[TC]; unsigned int objidx = trackCandidates->directObjectIndices[TC]; switch (type) { @@ -435,7 +435,7 @@ std::vector getLSsFromTC(Event* event, unsigned int TC) { std::tuple, std::vector> getHitIdxsAndHitTypesFromTC(Event* event, unsigned TC) { // Get the type of the track candidate - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); short type = trackCandidates->trackCandidateType[TC]; unsigned int objidx = trackCandidates->directObjectIndices[TC]; switch (type) { 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.cc b/RecoTracker/LSTCore/standalone/code/core/trkCore.cc index 3841affaaf059..ffb2e7de205ac 100644 --- a/RecoTracker/LSTCore/standalone/code/core/trkCore.cc +++ b/RecoTracker/LSTCore/standalone/code/core/trkCore.cc @@ -610,7 +610,7 @@ void addInputsToLineSegmentTrackingPreLoad(std::vector> &out_ std::vector> &out_charge_vec, std::vector> &out_seedIdx_vec, std::vector> &out_superbin_vec, - std::vector> &out_pixelType_vec, + std::vector> &out_pixelType_vec, std::vector> &out_isQuad_vec) { unsigned int count = 0; auto n_see = trk.see_stateTrajGlbPx().size(); @@ -651,7 +651,7 @@ void addInputsToLineSegmentTrackingPreLoad(std::vector> &out_ std::vector hitIdxs(trk.ph2_detId().size()); std::vector superbin_vec; - std::vector pixelType_vec; + std::vector pixelType_vec; std::vector isQuad_vec; std::iota(hitIdxs.begin(), hitIdxs.end(), 0); const int hit_size = trkX.size(); @@ -710,6 +710,7 @@ void addInputsToLineSegmentTrackingPreLoad(std::vector> &out_ TVector3 seedSD_r3 = r3LH; TVector3 seedSD_p3 = p3LH; + // The charge could be used directly in the line below float pixelSegmentDeltaPhiChange = r3LH.DeltaPhi(p3LH); float etaErr = trk.see_etaErr()[iSeed]; float px = p3LH.X(); @@ -718,17 +719,14 @@ void addInputsToLineSegmentTrackingPreLoad(std::vector> &out_ int charge = trk.see_q()[iSeed]; unsigned int seedIdx = iSeed; - // get pixel superbin - // int ptbin = -1; - int pixtype = -1; - if (ptIn >= 2.0) { /*ptbin = 1;*/ - pixtype = 0; + PixelType pixtype = PixelType::kInvalid; + if (ptIn >= 2.0) { + pixtype = PixelType::kHighPt; } else if (ptIn >= (PT_CUT - 2 * ptErr) and ptIn < 2.0) { - // ptbin = 0; if (pixelSegmentDeltaPhiChange >= 0) { - pixtype = 1; + pixtype = PixelType::kLowPtPosCurv; } else { - pixtype = 2; + pixtype = PixelType::kLowPtNegCurv; } } else { continue; @@ -868,7 +866,7 @@ float addInputsToEventPreLoad(LSTEvent *event, std::vector charge_vec, std::vector seedIdx_vec, std::vector superbin_vec, - std::vector pixelType_vec, + std::vector pixelType_vec, std::vector isQuad_vec) { TStopwatch my_timer; @@ -1140,212 +1138,3 @@ void writeMetaData() { TNamed tracklooper_path("tracklooper_path", ana.track_looper_dir_path.Data()); tracklooper_path.Write(); } - -// ---------------------------------- =========================================== ---------------------------------------------- -// ---------------------------------- =========================================== ---------------------------------------------- -// ---------------------------------- =========================================== ---------------------------------------------- -// ---------------------------------- =========================================== ---------------------------------------------- -// ---------------------------------- =========================================== ---------------------------------------------- - -// DEPRECATED FUNCTIONS - -//__________________________________________________________________________________________ -[[deprecated]] float addInputsToLineSegmentTracking(LSTEvent &event, bool useOMP) { - TStopwatch my_timer; - if (ana.verbose >= 2) - std::cout << "Loading Inputs (i.e. outer tracker hits, and pixel line segements) to the Line Segment Tracking.... " - << std::endl; - my_timer.Start(); - - unsigned int count = 0; - std::vector px_vec; - std::vector py_vec; - std::vector pz_vec; - std::vector hitIndices_vec0; - std::vector hitIndices_vec1; - std::vector hitIndices_vec2; - std::vector hitIndices_vec3; - std::vector ptIn_vec; - std::vector ptErr_vec; - std::vector etaErr_vec; - std::vector eta_vec; - std::vector phi_vec; - std::vector charge_vec; - std::vector seedIdx_vec; - std::vector deltaPhi_vec; - std::vector trkX = trk.ph2_x(); - std::vector trkY = trk.ph2_y(); - std::vector trkZ = trk.ph2_z(); - std::vector hitId = trk.ph2_detId(); - std::vector hitIdxs(trk.ph2_detId().size()); - std::vector superbin_vec; - std::vector pixelType_vec; - std::vector isQuad_vec; - std::iota(hitIdxs.begin(), hitIdxs.end(), 0); - const int hit_size = trkX.size(); - - for (size_t iSeed = 0; iSeed < trk.see_stateTrajGlbPx().size(); ++iSeed) { - bool good_seed_type = false; - if (trk.see_algo()[iSeed] == 4) - good_seed_type = true; - // if (trk.see_algo()[iSeed] == 5) good_seed_type = true; - // if (trk.see_algo()[iSeed] == 7) good_seed_type = true; - if (trk.see_algo()[iSeed] == 22) - good_seed_type = true; - // if (trk.see_algo()[iSeed] == 23) good_seed_type = true; - // if (trk.see_algo()[iSeed] == 24) good_seed_type = true; - if (not good_seed_type) - continue; - - TVector3 p3LH(trk.see_stateTrajGlbPx()[iSeed], trk.see_stateTrajGlbPy()[iSeed], trk.see_stateTrajGlbPz()[iSeed]); - float ptIn = p3LH.Pt(); - float ptErr = trk.see_ptErr()[iSeed]; - float eta = p3LH.Eta(); - - if ((ptIn > 0.8 - 2 * ptErr)) { - TVector3 r3LH(trk.see_stateTrajGlbX()[iSeed], trk.see_stateTrajGlbY()[iSeed], trk.see_stateTrajGlbZ()[iSeed]); - TVector3 p3PCA(trk.see_px()[iSeed], trk.see_py()[iSeed], trk.see_pz()[iSeed]); - TVector3 r3PCA(calculateR3FromPCA(p3PCA, trk.see_dxy()[iSeed], trk.see_dz()[iSeed])); - - TVector3 seedSD_mdRef_r3 = r3PCA; - TVector3 seedSD_mdOut_r3 = r3LH; - TVector3 seedSD_r3 = r3LH; - TVector3 seedSD_p3 = p3LH; - - float pixelSegmentDeltaPhiChange = r3LH.DeltaPhi(p3LH); - float etaErr = trk.see_etaErr()[iSeed]; - float px = p3LH.X(); - float py = p3LH.Y(); - float pz = p3LH.Z(); - float phi = p3LH.Phi(); - int charge = trk.see_q()[iSeed]; - unsigned int seedIdx = iSeed; - // extra bit - - // get pixel superbin - // int ptbin = -1; - int pixtype = -1; - if (ptIn >= 2.0) { /*ptbin = 1;*/ - pixtype = 0; - } else if (ptIn >= (0.8 - 2 * ptErr) and ptIn < 2.0) { - // ptbin = 0; - if (pixelSegmentDeltaPhiChange >= 0) { - pixtype = 1; - } else { - pixtype = 2; - } - } else { - continue; - } - - unsigned int hitIdx0 = hit_size + count; - count++; - - unsigned int hitIdx1 = hit_size + count; - count++; - - unsigned int hitIdx2 = hit_size + count; - count++; - - unsigned int hitIdx3; - if (trk.see_hitIdx()[iSeed].size() <= 3) { - hitIdx3 = hitIdx2; - } else { - hitIdx3 = hit_size + count; - count++; - } - - trkX.push_back(r3PCA.X()); - trkY.push_back(r3PCA.Y()); - trkZ.push_back(r3PCA.Z()); - trkX.push_back(p3PCA.Pt()); - float p3PCA_Eta = p3PCA.Eta(); - trkY.push_back(p3PCA_Eta); - float p3PCA_Phi = p3PCA.Phi(); - trkZ.push_back(p3PCA_Phi); - trkX.push_back(r3LH.X()); - trkY.push_back(r3LH.Y()); - trkZ.push_back(r3LH.Z()); - hitId.push_back(1); - hitId.push_back(1); - hitId.push_back(1); - if (trk.see_hitIdx()[iSeed].size() > 3) { - trkX.push_back(r3LH.X()); - trkY.push_back(trk.see_dxy()[iSeed]); - trkZ.push_back(trk.see_dz()[iSeed]); - hitId.push_back(1); - } - px_vec.push_back(px); - py_vec.push_back(py); - pz_vec.push_back(pz); - - hitIndices_vec0.push_back(hitIdx0); - hitIndices_vec1.push_back(hitIdx1); - hitIndices_vec2.push_back(hitIdx2); - hitIndices_vec3.push_back(hitIdx3); - ptIn_vec.push_back(ptIn); - ptErr_vec.push_back(ptErr); - etaErr_vec.push_back(etaErr); - eta_vec.push_back(eta); - phi_vec.push_back(phi); - charge_vec.push_back(charge); - seedIdx_vec.push_back(seedIdx); - deltaPhi_vec.push_back(pixelSegmentDeltaPhiChange); - - // For matching with sim tracks - hitIdxs.push_back(trk.see_hitIdx()[iSeed][0]); - hitIdxs.push_back(trk.see_hitIdx()[iSeed][1]); - hitIdxs.push_back(trk.see_hitIdx()[iSeed][2]); - char isQuad = false; - if (trk.see_hitIdx()[iSeed].size() > 3) { - isQuad = true; - hitIdxs.push_back(trk.see_hitIdx()[iSeed].size() > 3 ? trk.see_hitIdx()[iSeed][3] : trk.see_hitIdx()[iSeed][2]); - } - // if (pt < 0){ ptbin = 0;} - float neta = 25.; - float nphi = 72.; - float nz = 25.; - int etabin = (p3PCA_Eta + 2.6) / ((2 * 2.6) / neta); - int phibin = (p3PCA_Phi + 3.14159265358979323846) / ((2. * 3.14159265358979323846) / nphi); - int dzbin = (trk.see_dz()[iSeed] + 30) / (2 * 30 / nz); - int isuperbin = - /*(nz * nphi * neta) * ptbin + (removed since pt bin is determined by pixelType)*/ (nz * nphi) * etabin + - (nz)*phibin + dzbin; - // if(isuperbin<0 || isuperbin>=44900){printf("isuperbin %d %d %d %d %f\n",isuperbin,etabin,phibin,dzbin,p3PCA.Eta());} - superbin_vec.push_back(isuperbin); - pixelType_vec.push_back(pixtype); - isQuad_vec.push_back(isQuad); - } - } - - event.addHitToEvent(trkX, trkY, trkZ, hitId, hitIdxs); - event.addPixelSegmentToEvent(hitIndices_vec0, - hitIndices_vec1, - hitIndices_vec2, - hitIndices_vec3, - deltaPhi_vec, - ptIn_vec, - ptErr_vec, - px_vec, - py_vec, - pz_vec, - eta_vec, - etaErr_vec, - phi_vec, - charge_vec, - seedIdx_vec, - superbin_vec, - pixelType_vec, - isQuad_vec); - - event.wait(); // device side event calls are asynchronous: wait to measure time or print - float hit_loading_elapsed = my_timer.RealTime(); - if (ana.verbose >= 2) - std::cout << "Loading inputs processing time: " << hit_loading_elapsed << " secs" << std::endl; - return hit_loading_elapsed; -} - -//__________________________________________________________________________________________ -[[deprecated]] float addInputsToLineSegmentTrackingUsingExplicitMemory(LSTEvent &event) { - return addInputsToLineSegmentTracking(event, true); -} diff --git a/RecoTracker/LSTCore/standalone/code/core/trkCore.h b/RecoTracker/LSTCore/standalone/code/core/trkCore.h index 66d5c10baf431..be18f02af2329 100644 --- a/RecoTracker/LSTCore/standalone/code/core/trkCore.h +++ b/RecoTracker/LSTCore/standalone/code/core/trkCore.h @@ -1,17 +1,19 @@ #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 using LSTEvent = ALPAKA_ACCELERATOR_NAMESPACE::lst::Event; +using ::lst::PixelType; // --------------------- ======================== --------------------- @@ -68,7 +70,7 @@ void addInputsToLineSegmentTrackingPreLoad(std::vector> &out_ std::vector> &out_charge_vec, std::vector> &out_seedIdx_vec, std::vector> &out_superbin_vec, - std::vector> &out_pixelType_vec, + std::vector> &out_pixelType_vec, std::vector> &out_isQuad_vec); float addInputsToEventPreLoad(LSTEvent *event, @@ -94,7 +96,7 @@ float addInputsToEventPreLoad(LSTEvent *event, std::vector charge_vec, std::vector seedIdx_vec, std::vector superbin_vec, - std::vector pixelType_vec, + std::vector pixelType_vec, std::vector isQuad_vec); void printTimingInformation(std::vector> &timing_information, float fullTime, float fullavg); diff --git a/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc b/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc index 7c330a768a175..e12512f5c5c7d 100644 --- a/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc +++ b/RecoTracker/LSTCore/standalone/code/core/write_lst_ntuple.cc @@ -226,7 +226,7 @@ void setOutputBranches(Event* event) { std::vector> tc_matched_simIdx; // ============ Track candidates ============= - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); unsigned int nTrackCandidates = *trackCandidates->nTrackCandidates; for (unsigned int idx = 0; idx < nTrackCandidates; idx++) { // Compute reco quantities of track candidate based on final object @@ -291,10 +291,10 @@ void setOptionalOutputBranches(Event* event) { //________________________________________________________________________________________________________________________________ void setPixelQuintupletOutputBranches(Event* event) { // ============ pT5 ============= - 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(); + PixelQuintuplets const* pixelQuintuplets = event->getPixelQuintuplets().data(); + Quintuplets const* quintuplets = event->getQuintuplets().data(); + Segments const* segments = event->getSegments().data(); + Modules const* modules = event->getModules().data(); int n_accepted_simtrk = ana.tx->getBranch>("sim_TC_matched").size(); unsigned int nPixelQuintuplets = @@ -366,9 +366,9 @@ 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(); + Quintuplets const* quintuplets = event->getQuintuplets().data(); + ObjectRanges const* ranges = event->getRanges().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); @@ -437,9 +437,9 @@ void setQuintupletOutputBranches(Event* event) { //________________________________________________________________________________________________________________________________ void setPixelTripletOutputBranches(Event* event) { - PixelTriplets const* pixelTriplets = event->getPixelTriplets()->data(); - ::lst::Modules const* modules = event->getModules()->data(); - Segments const* segments = event->getSegments()->data(); + PixelTriplets const* pixelTriplets = event->getPixelTriplets().data(); + Modules const* modules = event->getModules().data(); + Segments const* segments = event->getSegments().data(); int n_accepted_simtrk = ana.tx->getBranch>("sim_TC_matched").size(); unsigned int nPixelTriplets = *pixelTriplets->nPixelTriplets; @@ -501,12 +501,12 @@ void setPixelTripletOutputBranches(Event* event) { //________________________________________________________________________________________________________________________________ void setGnnNtupleBranches(Event* event) { // Get relevant information - 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(); - ObjectRanges const* ranges = event->getRanges()->data(); - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); + Segments const* segments = event->getSegments().data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + Hits const* hitsEvt = event->getHits().data(); + Modules const* modules = event->getModules().data(); + ObjectRanges const* ranges = event->getRanges().data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); std::set mds_used_in_sg; std::map md_index_map; @@ -642,8 +642,8 @@ void setGnnNtupleBranches(Event* event) { //________________________________________________________________________________________________________________________________ void setGnnNtupleMiniDoublet(Event* event, unsigned int MD) { // Get relevant information - MiniDoublets const* miniDoublets = event->getMiniDoublets()->data(); - Hits const* hitsEvt = event->getHits()->data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + Hits const* hitsEvt = event->getHits().data(); // Get the hit indices unsigned int hit0 = miniDoublets->anchorHitIndices[MD]; @@ -710,7 +710,7 @@ void setGnnNtupleMiniDoublet(Event* event, unsigned int MD) { //________________________________________________________________________________________________________________________________ std::tuple> parseTrackCandidate(Event* event, unsigned int idx) { // Get the type of the track candidate - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); short type = trackCandidates->trackCandidateType[idx]; enum { pT5 = 7, pT3 = 5, T5 = 4, pLS = 8 }; @@ -744,9 +744,9 @@ std::tuple> parseTrackCandidate( std::tuple, std::vector> parsepT5(Event* event, unsigned int idx) { // Get relevant information - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); - Quintuplets const* quintuplets = event->getQuintuplets()->data(); - Segments const* segments = event->getSegments()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); + Quintuplets const* quintuplets = event->getQuintuplets().data(); + Segments const* segments = event->getSegments().data(); // // pictorial representation of a pT5 @@ -856,9 +856,9 @@ std::tuple, std::vector, std::vector> parsepT3(Event* event, unsigned int idx) { // Get relevant information - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); - Triplets const* triplets = event->getTriplets()->data(); - Segments const* segments = event->getSegments()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); + Triplets const* triplets = event->getTriplets().data(); + Segments const* segments = event->getSegments().data(); // // pictorial representation of a pT3 @@ -890,8 +890,8 @@ std::tuple, std::vector, std::vector> parseT5(Event* event, unsigned int idx) { - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); - Quintuplets const* quintuplets = event->getQuintuplets()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); + Quintuplets const* quintuplets = event->getQuintuplets().data(); unsigned int T5 = trackCandidates->directObjectIndices[idx]; std::vector hits = getHitsFromT5(event, T5); @@ -924,8 +924,8 @@ std::tuple, std::vector, std::vector> parsepLS(Event* event, unsigned int idx) { - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); - Segments const* segments = event->getSegments()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); + Segments const* segments = event->getSegments().data(); // Getting pLS index unsigned int pLS = trackCandidates->directObjectIndices[idx]; @@ -944,8 +944,8 @@ std::tuple, std::vectorgetModules()->data(); - ObjectRanges const* ranges = event->getRanges()->data(); + Modules const* modules = event->getModules().data(); + ObjectRanges const* ranges = event->getRanges().data(); int nHits = 0; for (unsigned int idx = 0; idx <= *(modules->nLowerModules); @@ -959,8 +959,8 @@ void printHitMultiplicities(Event* event) { //________________________________________________________________________________________________________________________________ void printMiniDoubletMultiplicities(Event* event) { - MiniDoublets const* miniDoublets = event->getMiniDoublets()->data(); - ::lst::Modules const* modules = event->getModules()->data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + Modules const* modules = event->getModules().data(); int nMiniDoublets = 0; int totOccupancyMiniDoublets = 0; @@ -986,10 +986,10 @@ 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(); - ObjectRanges const* ranges = event->getRanges()->data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + Hits const* hitsEvt = event->getHits().data(); + Modules const* modules = event->getModules().data(); + ObjectRanges const* ranges = event->getRanges().data(); // Then obtain the lower module index for (unsigned int idx = 0; idx <= *(modules->nLowerModules); ++idx) { @@ -1008,11 +1008,11 @@ void printMDs(Event* event) { //________________________________________________________________________________________________________________________________ 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(); - ObjectRanges const* ranges = event->getRanges()->data(); + Segments const* segments = event->getSegments().data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + Hits const* hitsEvt = event->getHits().data(); + Modules const* modules = event->getModules().data(); + ObjectRanges const* ranges = event->getRanges().data(); int nSegments = 0; for (unsigned int i = 0; i < *(modules->nLowerModules); ++i) { @@ -1040,11 +1040,11 @@ void printLSs(Event* event) { //________________________________________________________________________________________________________________________________ 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(); - ObjectRanges const* ranges = event->getRanges()->data(); + Segments const* segments = event->getSegments().data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + Hits const* hitsEvt = event->getHits().data(); + Modules const* modules = event->getModules().data(); + ObjectRanges const* ranges = event->getRanges().data(); unsigned int i = *(modules->nLowerModules); unsigned int idx = i; //modules->lowerModuleIndices[i]; @@ -1070,11 +1070,11 @@ void printpLSs(Event* event) { //________________________________________________________________________________________________________________________________ void printT3s(Event* event) { - Triplets const* triplets = event->getTriplets()->data(); - 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(); + Triplets const* triplets = event->getTriplets().data(); + Segments const* segments = event->getSegments().data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().data(); + Hits const* hitsEvt = event->getHits().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]; @@ -1112,12 +1112,12 @@ void printT3s(Event* event) { //________________________________________________________________________________________________________________________________ void debugPrintOutlierMultiplicities(Event* event) { - TrackCandidates const* trackCandidates = event->getTrackCandidates()->data(); - 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(); - ObjectRanges const* ranges = event->getRanges()->data(); + TrackCandidates const* trackCandidates = event->getTrackCandidates().data(); + Triplets const* triplets = event->getTriplets().data(); + Segments const* segments = event->getSegments().data(); + MiniDoublets const* miniDoublets = event->getMiniDoublets().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) { if (trackCandidates->nTrackCandidates[idx] > 50000) {