From 954a22ab377638ddbaf3f2922d7d95c1e8fa0dba Mon Sep 17 00:00:00 2001 From: Manos Vourliotis Date: Tue, 30 Jul 2024 03:31:08 -0700 Subject: [PATCH] Third batch of fixing function arguments --- RecoTracker/LSTCore/src/alpaka/Quintuplet.h | 789 ++++-------------- RecoTracker/LSTCore/src/alpaka/Segment.h | 257 ++---- .../LSTCore/src/alpaka/TrackCandidate.h | 74 +- 3 files changed, 287 insertions(+), 833 deletions(-) diff --git a/RecoTracker/LSTCore/src/alpaka/Quintuplet.h b/RecoTracker/LSTCore/src/alpaka/Quintuplet.h index 1d796a6613824..66061ae427849 100644 --- a/RecoTracker/LSTCore/src/alpaka/Quintuplet.h +++ b/RecoTracker/LSTCore/src/alpaka/Quintuplet.h @@ -104,10 +104,7 @@ namespace lst { Quintuplets data_; template - QuintupletsBuffer(unsigned int nTotalQuintuplets, - unsigned int nLowerModules, - TDevAcc const& devAccIn, - TQueue& queue) + QuintupletsBuffer(unsigned int nTotalQuintuplets, unsigned int nLowerModules, TDevAcc const& devAccIn, TQueue& queue) : tripletIndices_buf(allocBufWrapper(devAccIn, 2 * nTotalQuintuplets, queue)), lowerModuleIndices_buf(allocBufWrapper(devAccIn, Params_T5::kLayers * nTotalQuintuplets, queue)), nQuintuplets_buf(allocBufWrapper(devAccIn, nLowerModules, queue)), @@ -151,8 +148,8 @@ namespace lst { return ((firstMin <= secondMin) && (secondMin < firstMax)) || ((secondMin < firstMin) && (firstMin < secondMax)); }; - ALPAKA_FN_ACC ALPAKA_FN_INLINE void addQuintupletToMemory(struct lst::Triplets& tripletsInGPU, - struct lst::Quintuplets& quintupletsInGPU, + ALPAKA_FN_ACC ALPAKA_FN_INLINE void addQuintupletToMemory(lst::Triplets const& tripletsInGPU, + lst::Quintuplets& quintupletsInGPU, unsigned int innerTripletIndex, unsigned int outerTripletIndex, uint16_t lowerModule1, @@ -160,15 +157,15 @@ namespace lst { uint16_t lowerModule3, uint16_t lowerModule4, uint16_t lowerModule5, - float& innerRadius, - float& bridgeRadius, - float& outerRadius, - float& regressionG, - float& regressionF, - float& regressionRadius, - float& rzChiSquared, - float& rPhiChiSquared, - float& nonAnchorChiSquared, + float innerRadius, + float bridgeRadius, + float outerRadius, + float regressionG, + float regressionF, + float regressionRadius, + float rzChiSquared, + float rPhiChiSquared, + float nonAnchorChiSquared, float pt, float eta, float phi, @@ -234,7 +231,7 @@ namespace lst { }; //90% constraint - ALPAKA_FN_ACC ALPAKA_FN_INLINE bool passChiSquaredConstraint(struct lst::Modules& modulesInGPU, + ALPAKA_FN_ACC ALPAKA_FN_INLINE bool passChiSquaredConstraint(lst::Modules const& modulesInGPU, uint16_t lowerModuleIndex1, uint16_t lowerModuleIndex2, uint16_t lowerModuleIndex3, @@ -319,8 +316,8 @@ namespace lst { //bounds can be found at http://uaf-10.t2.ucsd.edu/~bsathian/SDL/T5_RZFix/t5_rz_thresholds.txt template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool passT5RZConstraint(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, unsigned int firstMDIndex, unsigned int secondMDIndex, unsigned int thirdMDIndex, @@ -753,8 +750,8 @@ namespace lst { }; template - ALPAKA_FN_ACC ALPAKA_FN_INLINE bool T5HasCommonMiniDoublet(struct lst::Triplets& tripletsInGPU, - struct lst::Segments& segmentsInGPU, + ALPAKA_FN_ACC ALPAKA_FN_INLINE bool T5HasCommonMiniDoublet(lst::Triplets const& tripletsInGPU, + lst::Segments const& segmentsInGPU, unsigned int innerTripletIndex, unsigned int outerTripletIndex) { unsigned int innerOuterSegmentIndex = tripletsInGPU.segmentIndices[2 * innerTripletIndex + 1]; @@ -801,24 +798,15 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBBBEE12378(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius, + float bridgeRadiusMin2S, + float bridgeRadiusMax2S) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 0.178f; float bridgeInvRadiusErrorBound = 0.507f; - float outerInvRadiusErrorBound = 7.655f; innerInvRadiusMax = (1.f + innerInvRadiusErrorBound) / innerRadius; innerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - innerInvRadiusErrorBound) / innerRadius); @@ -826,9 +814,6 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(innerInvRadiusMin, innerInvRadiusMax, alpaka::math::min(acc, bridgeInvRadiusMin, 1.0f / bridgeRadiusMax2S), @@ -838,23 +823,17 @@ namespace lst { /*bounds for high Pt taken from : http://uaf-10.t2.ucsd.edu/~bsathian/SDL/T5_efficiency/efficiencies/new_efficiencies/efficiencies_20210513_T5_recovering_high_Pt_efficiencies/highE_radius_matching/highE_bounds.txt */ template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBBBBB(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 0.1512f; float bridgeInvRadiusErrorBound = 0.1781f; - float outerInvRadiusErrorBound = 0.1840f; if (innerRadius > 2.0f / (2.f * k2Rinv1GeVf)) { innerInvRadiusErrorBound = 0.4449f; bridgeInvRadiusErrorBound = 0.4033f; - outerInvRadiusErrorBound = 0.8016f; } innerInvRadiusMax = (1.f + innerInvRadiusErrorBound) / innerRadius; @@ -863,37 +842,22 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax); }; template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBBBBE(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 0.1781f; float bridgeInvRadiusErrorBound = 0.2167f; - float outerInvRadiusErrorBound = 1.1116f; if (innerRadius > 2.0f / (2.f * k2Rinv1GeVf)) { innerInvRadiusErrorBound = 0.4750f; bridgeInvRadiusErrorBound = 0.3903f; - outerInvRadiusErrorBound = 15.2120f; } innerInvRadiusMax = (1.f + innerInvRadiusErrorBound) / innerRadius; @@ -902,75 +866,20 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax); }; - template - ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBBBEE(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { - float innerInvRadiusErrorBound = 0.1840f; - float bridgeInvRadiusErrorBound = 0.5971f; - float outerInvRadiusErrorBound = 11.7102f; - - if (innerRadius > 2.0f / (2.f * k2Rinv1GeVf)) //as good as no selections - { - innerInvRadiusErrorBound = 1.0412f; - outerInvRadiusErrorBound = 32.2737f; - bridgeInvRadiusErrorBound = 10.9688f; - } - - innerInvRadiusMax = (1.f + innerInvRadiusErrorBound) / innerRadius; - innerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - innerInvRadiusErrorBound) / innerRadius); - - bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; - bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - - return checkIntervalOverlap(innerInvRadiusMin, - innerInvRadiusMax, - alpaka::math::min(acc, bridgeInvRadiusMin, 1.0f / bridgeRadiusMax2S), - alpaka::math::max(acc, bridgeInvRadiusMax, 1.0f / bridgeRadiusMin2S)); - }; - template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBBBEE23478(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius, + float bridgeRadiusMin2S, + float bridgeRadiusMax2S) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 0.2097f; float bridgeInvRadiusErrorBound = 0.8557f; - float outerInvRadiusErrorBound = 24.0450f; innerInvRadiusMax = (1.f + innerInvRadiusErrorBound) / innerRadius; innerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - innerInvRadiusErrorBound) / innerRadius); @@ -978,9 +887,6 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(innerInvRadiusMin, innerInvRadiusMax, alpaka::math::min(acc, bridgeInvRadiusMin, 1.0f / bridgeRadiusMax2S), @@ -989,24 +895,15 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBBBEE34578(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius, + float bridgeRadiusMin2S, + float bridgeRadiusMax2S) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 0.066f; float bridgeInvRadiusErrorBound = 0.617f; - float outerInvRadiusErrorBound = 2.688f; innerInvRadiusMax = (1.f + innerInvRadiusErrorBound) / innerRadius; innerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - innerInvRadiusErrorBound) / innerRadius); @@ -1014,9 +911,6 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(innerInvRadiusMin, innerInvRadiusMax, alpaka::math::min(acc, bridgeInvRadiusMin, 1.0f / bridgeRadiusMax2S), @@ -1025,29 +919,19 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBBEEE(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius, + float bridgeRadiusMin2S, + float bridgeRadiusMax2S) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 0.6376f; float bridgeInvRadiusErrorBound = 2.1381f; - float outerInvRadiusErrorBound = 20.4179f; if (innerRadius > 2.0f / (2.f * k2Rinv1GeVf)) //as good as no selections! { innerInvRadiusErrorBound = 12.9173f; - outerInvRadiusErrorBound = 25.6702f; bridgeInvRadiusErrorBound = 5.1700f; } @@ -1057,9 +941,6 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(innerInvRadiusMin, innerInvRadiusMax, alpaka::math::min(acc, bridgeInvRadiusMin, 1.0f / bridgeRadiusMax2S), @@ -1068,28 +949,20 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiBEEEE(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius, + float innerRadiusMin2S, + float innerRadiusMax2S, + float bridgeRadiusMin2S, + float bridgeRadiusMax2S) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 1.9382f; float bridgeInvRadiusErrorBound = 3.7280f; - float outerInvRadiusErrorBound = 5.7030f; if (innerRadius > 2.0f / (2.f * k2Rinv1GeVf)) { innerInvRadiusErrorBound = 23.2713f; - outerInvRadiusErrorBound = 24.0450f; bridgeInvRadiusErrorBound = 21.7980f; } @@ -1099,9 +972,6 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(alpaka::math::min(acc, innerInvRadiusMin, 1.0 / innerRadiusMax2S), alpaka::math::max(acc, innerInvRadiusMax, 1.0 / innerRadiusMin2S), alpaka::math::min(acc, bridgeInvRadiusMin, 1.0 / bridgeRadiusMax2S), @@ -1110,29 +980,21 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool matchRadiiEEEEE(TAcc const& acc, - const float& innerRadius, - const float& bridgeRadius, - const float& outerRadius, - const float& innerRadiusMin2S, - const float& innerRadiusMax2S, - const float& bridgeRadiusMin2S, - const float& bridgeRadiusMax2S, - const float& outerRadiusMin2S, - const float& outerRadiusMax2S, - float& innerInvRadiusMin, - float& innerInvRadiusMax, - float& bridgeInvRadiusMin, - float& bridgeInvRadiusMax, - float& outerInvRadiusMin, - float& outerInvRadiusMax) { + float innerRadius, + float bridgeRadius, + float outerRadius, + float innerRadiusMin2S, + float innerRadiusMax2S, + float bridgeRadiusMin2S, + float bridgeRadiusMax2S) { + float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax; + float innerInvRadiusErrorBound = 1.9382f; float bridgeInvRadiusErrorBound = 2.2091f; - float outerInvRadiusErrorBound = 7.4084f; if (innerRadius > 2.0f / (2.f * k2Rinv1GeVf)) { innerInvRadiusErrorBound = 22.5226f; bridgeInvRadiusErrorBound = 21.0966f; - outerInvRadiusErrorBound = 19.1252f; } innerInvRadiusMax = (1.f + innerInvRadiusErrorBound) / innerRadius; @@ -1141,9 +1003,6 @@ namespace lst { bridgeInvRadiusMax = (1.f + bridgeInvRadiusErrorBound) / bridgeRadius; bridgeInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - bridgeInvRadiusErrorBound) / bridgeRadius); - outerInvRadiusMax = (1.f + outerInvRadiusErrorBound) / outerRadius; - outerInvRadiusMin = alpaka::math::max(acc, 0.f, (1.f - outerInvRadiusErrorBound) / outerRadius); - return checkIntervalOverlap(alpaka::math::min(acc, innerInvRadiusMin, 1.0 / innerRadiusMax2S), alpaka::math::max(acc, innerInvRadiusMax, 1.0 / innerRadiusMin2S), alpaka::math::min(acc, bridgeInvRadiusMin, 1.0 / bridgeRadiusMax2S), @@ -1152,7 +1011,7 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE void computeSigmasForRegression(TAcc const& acc, - lst::Modules& modulesInGPU, + lst::Modules const& modulesInGPU, const uint16_t* lowerModuleIndices, float* delta1, float* delta2, @@ -1385,7 +1244,7 @@ namespace lst { ALPAKA_FN_ACC ALPAKA_FN_INLINE void runDeltaBetaIterationsT5(TAcc const& acc, float& betaIn, float& betaOut, - float& betaAv, + float betaAv, float& pt_beta, float sdIn_dr, float sdOut_dr, @@ -1489,9 +1348,9 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool runQuintupletDefaultAlgoBBBB(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, - struct lst::Segments& segmentsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, + lst::Segments const& segmentsInGPU, uint16_t innerInnerLowerModuleIndex, uint16_t innerOuterLowerModuleIndex, uint16_t outerInnerLowerModuleIndex, @@ -1501,22 +1360,7 @@ namespace lst { unsigned int firstMDIndex, unsigned int secondMDIndex, unsigned int thirdMDIndex, - unsigned int fourthMDIndex, - float& zOut, - float& rtOut, - float& deltaPhiPos, - float& dPhi, - float& betaIn, - float& betaOut, - float& pt_beta, - float& zLo, - float& zHi, - float& zLoPointed, - float& zHiPointed, - float& sdlCut, - float& betaInCut, - float& betaOutCut, - float& deltaBetaCut) { + unsigned int fourthMDIndex) { bool isPS_InLo = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == lst::PS); bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == lst::PS); @@ -1537,14 +1381,12 @@ namespace lst { float zpitch_InLo = (isPS_InLo ? lst::kPixelPSZpitch : lst::kStrip2SZpitch); float zpitch_OutLo = (isPS_OutLo ? lst::kPixelPSZpitch : lst::kStrip2SZpitch); - zHi = z_InLo + (z_InLo + lst::kDeltaZLum) * (rtRatio_OutLoInLo - 1.f) * (z_InLo < 0.f ? 1.f : dzDrtScale) + - (zpitch_InLo + zpitch_OutLo); - zLo = z_InLo + (z_InLo - lst::kDeltaZLum) * (rtRatio_OutLoInLo - 1.f) * (z_InLo > 0.f ? 1.f : dzDrtScale) - - (zpitch_InLo + zpitch_OutLo); + float zHi = z_InLo + (z_InLo + lst::kDeltaZLum) * (rtRatio_OutLoInLo - 1.f) * (z_InLo < 0.f ? 1.f : dzDrtScale) + + (zpitch_InLo + zpitch_OutLo); + float zLo = z_InLo + (z_InLo - lst::kDeltaZLum) * (rtRatio_OutLoInLo - 1.f) * (z_InLo > 0.f ? 1.f : dzDrtScale) - + (zpitch_InLo + zpitch_OutLo); //Cut 1 - z compatibility - zOut = z_OutLo; - rtOut = rt_OutLo; if ((z_OutLo < zLo) || (z_OutLo > zHi)) return false; @@ -1569,17 +1411,17 @@ namespace lst { const float zWindow = dzErr / drt_InSeg * drt_OutLo_InLo + (zpitch_InLo + zpitch_OutLo); //FIXME for lst::ptCut lower than ~0.8 need to add curv path correction - zLoPointed = z_InLo + dzMean * (z_InLo > 0.f ? 1.f : dzDrtScale) - zWindow; - zHiPointed = z_InLo + dzMean * (z_InLo < 0.f ? 1.f : dzDrtScale) + zWindow; + float zLoPointed = z_InLo + dzMean * (z_InLo > 0.f ? 1.f : dzDrtScale) - zWindow; + float zHiPointed = z_InLo + dzMean * (z_InLo < 0.f ? 1.f : dzDrtScale) + zWindow; // Cut #2: Pointed Z (Inner segment two MD points to outer segment inner MD) if ((z_OutLo < zLoPointed) || (z_OutLo > zHiPointed)) return false; float sdlPVoff = 0.1f / rt_OutLo; - sdlCut = alpha1GeV_OutLo + alpaka::math::sqrt(acc, sdlMuls2 + sdlPVoff * sdlPVoff); + float sdlCut = alpha1GeV_OutLo + alpaka::math::sqrt(acc, sdlMuls2 + sdlPVoff * sdlPVoff); - deltaPhiPos = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[fourthMDIndex] - mdsInGPU.anchorPhi[secondMDIndex]); + float deltaPhiPos = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[fourthMDIndex] - mdsInGPU.anchorPhi[secondMDIndex]); // Cut #3: FIXME:deltaPhiPos can be tighter if (alpaka::math::abs(acc, deltaPhiPos) > sdlCut) return false; @@ -1589,7 +1431,7 @@ namespace lst { float diffX = mdsInGPU.anchorX[thirdMDIndex] - mdsInGPU.anchorX[firstMDIndex]; float diffY = mdsInGPU.anchorY[thirdMDIndex] - mdsInGPU.anchorY[firstMDIndex]; - dPhi = lst::deltaPhi(acc, midPointX, midPointY, diffX, diffY); + float dPhi = lst::deltaPhi(acc, midPointX, midPointY, diffX, diffY); // Cut #4: deltaPhiChange if (alpaka::math::abs(acc, dPhi) > sdlCut) @@ -1621,11 +1463,12 @@ namespace lst { float tl_axis_lowEdge_x = tl_axis_x; float tl_axis_lowEdge_y = tl_axis_y; - betaIn = alpha_InLo - lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[firstMDIndex]); + float betaIn = + alpha_InLo - lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[firstMDIndex]); float betaInRHmin = betaIn; float betaInRHmax = betaIn; - betaOut = + float betaOut = -alpha_OutUp + lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut; @@ -1669,18 +1512,19 @@ namespace lst { (mdsInGPU.anchorX[secondMDIndex] - mdsInGPU.anchorX[firstMDIndex]) + (mdsInGPU.anchorY[secondMDIndex] - mdsInGPU.anchorY[firstMDIndex]) * (mdsInGPU.anchorY[secondMDIndex] - mdsInGPU.anchorY[firstMDIndex])); - betaInCut = alpaka::math::asin( - acc, - alpaka::math::min( - acc, (-rt_InSeg * corrF + drt_tl_axis) * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + - (0.02f / drt_InSeg); + float betaInCut = + alpaka::math::asin( + acc, + alpaka::math::min( + acc, (-rt_InSeg * corrF + drt_tl_axis) * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + + (0.02f / drt_InSeg); //Cut #5: first beta cut if (alpaka::math::abs(acc, betaInRHmin) >= betaInCut) return false; float betaAv = 0.5f * (betaIn + betaOut); - pt_beta = drt_tl_axis * lst::k2Rinv1GeVf / alpaka::math::sin(acc, betaAv); + float pt_beta = drt_tl_axis * lst::k2Rinv1GeVf / alpaka::math::sin(acc, betaAv); int lIn = 5; int lOut = isEC_lastLayer ? 11 : 5; float sdOut_dr = alpaka::math::sqrt(acc, @@ -1735,10 +1579,10 @@ namespace lst { const float dBetaROut2 = dBetaROut * dBetaROut; - //FIXME: need faster version - betaOutCut = alpaka::math::asin( - acc, alpaka::math::min(acc, drt_tl_axis * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + - (0.02f / sdOut_d) + alpaka::math::sqrt(acc, dBetaLum2 + dBetaMuls2); + float betaOutCut = + alpaka::math::asin(acc, + alpaka::math::min(acc, drt_tl_axis * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + + (0.02f / sdOut_d) + alpaka::math::sqrt(acc, dBetaLum2 + dBetaMuls2); //Cut #6: The real beta cut if (alpaka::math::abs(acc, betaOut) >= betaOutCut) @@ -1752,15 +1596,14 @@ namespace lst { (alpaka::math::abs(acc, betaInRHmin - betaInRHmax) + alpaka::math::abs(acc, betaOutRHmin - betaOutRHmax))); float dBeta = betaIn - betaOut; - deltaBetaCut = alpaka::math::sqrt(acc, dBetaCut2); return dBeta * dBeta <= dBetaCut2; }; template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool runQuintupletDefaultAlgoBBEE(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, - struct lst::Segments& segmentsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, + lst::Segments const& segmentsInGPU, uint16_t innerInnerLowerModuleIndex, uint16_t innerOuterLowerModuleIndex, uint16_t outerInnerLowerModuleIndex, @@ -1770,22 +1613,7 @@ namespace lst { unsigned int firstMDIndex, unsigned int secondMDIndex, unsigned int thirdMDIndex, - unsigned int fourthMDIndex, - float& zOut, - float& rtOut, - float& deltaPhiPos, - float& dPhi, - float& betaIn, - float& betaOut, - float& pt_beta, - float& zLo, - float& rtLo, - float& rtHi, - float& sdlCut, - float& betaInCut, - float& betaOutCut, - float& deltaBetaCut, - float& kZ) { + unsigned int fourthMDIndex) { bool isPS_InLo = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == lst::PS); bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == lst::PS); @@ -1800,15 +1628,12 @@ namespace lst { float alpha1GeV_OutLo = alpaka::math::asin(acc, alpaka::math::min(acc, rt_OutLo * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)); - float rtRatio_OutLoInLo = rt_OutLo / rt_InLo; // Outer segment beginning rt divided by inner segment beginning rt; float dzDrtScale = alpaka::math::tan(acc, alpha1GeV_OutLo) / alpha1GeV_OutLo; // The track can bend in r-z plane slightly float zpitch_InLo = (isPS_InLo ? lst::kPixelPSZpitch : lst::kStrip2SZpitch); float zpitch_OutLo = (isPS_OutLo ? lst::kPixelPSZpitch : lst::kStrip2SZpitch); float zGeom = zpitch_InLo + zpitch_OutLo; - zLo = z_InLo + (z_InLo - lst::kDeltaZLum) * (rtRatio_OutLoInLo - 1.f) * (z_InLo > 0.f ? 1.f : dzDrtScale) - zGeom; - // Cut #0: Preliminary (Only here in endcap case) if (z_InLo * z_OutLo <= 0) return false; @@ -1817,10 +1642,9 @@ namespace lst { bool isOutSgInnerMDPS = modulesInGPU.moduleType[outerInnerLowerModuleIndex] == lst::PS; float rtGeom1 = isOutSgInnerMDPS ? lst::kPixelPSZpitch : lst::kStrip2SZpitch; float zGeom1 = alpaka::math::copysign(acc, zGeom, z_InLo); - rtLo = rt_InLo * (1.f + (z_OutLo - z_InLo - zGeom1) / (z_InLo + zGeom1 + dLum) / dzDrtScale) - - rtGeom1; //slope correction only on the lower end - zOut = z_OutLo; - rtOut = rt_OutLo; + float rtLo = rt_InLo * (1.f + (z_OutLo - z_InLo - zGeom1) / (z_InLo + zGeom1 + dLum) / dzDrtScale) - + rtGeom1; //slope correction only on the lower end + float rtOut = rt_OutLo; //Cut #1: rt condition if (rtOut < rtLo) @@ -1830,7 +1654,7 @@ namespace lst { if (zInForHi * z_InLo < 0) { zInForHi = alpaka::math::copysign(acc, 0.1f, z_InLo); } - rtHi = rt_InLo * (1.f + (z_OutLo - z_InLo + zGeom1) / zInForHi) + rtGeom1; + float rtHi = rt_InLo * (1.f + (z_OutLo - z_InLo + zGeom1) / zInForHi) + rtGeom1; //Cut #2: rt condition if ((rt_OutLo < rtLo) || (rt_OutLo > rtHi)) @@ -1846,7 +1670,7 @@ namespace lst { const float dzOutInAbs = alpaka::math::abs(acc, z_OutLo - z_InLo); const float multDzDr = dzOutInAbs * coshEta / (coshEta * coshEta - 1.f); const float zGeom1_another = lst::kPixelPSZpitch; - kZ = (z_OutLo - z_InLo) / dzSDIn; + float kZ = (z_OutLo - z_InLo) / dzSDIn; float drtErr = zGeom1_another * zGeom1_another * drtSDIn * drtSDIn / dzSDIn / dzSDIn * (1.f - 2.f * kZ + 2.f * kZ * kZ); const float sdlThetaMulsF2 = @@ -1860,9 +1684,9 @@ namespace lst { return false; const float sdlPVoff = 0.1f / rt_OutLo; - sdlCut = alpha1GeV_OutLo + alpaka::math::sqrt(acc, sdlMuls2 + sdlPVoff * sdlPVoff); + float sdlCut = alpha1GeV_OutLo + alpaka::math::sqrt(acc, sdlMuls2 + sdlPVoff * sdlPVoff); - deltaPhiPos = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[fourthMDIndex] - mdsInGPU.anchorPhi[secondMDIndex]); + float deltaPhiPos = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[fourthMDIndex] - mdsInGPU.anchorPhi[secondMDIndex]); //Cut #4: deltaPhiPos can be tighter if (alpaka::math::abs(acc, deltaPhiPos) > sdlCut) @@ -1873,7 +1697,7 @@ namespace lst { float diffX = mdsInGPU.anchorX[thirdMDIndex] - mdsInGPU.anchorX[firstMDIndex]; float diffY = mdsInGPU.anchorY[thirdMDIndex] - mdsInGPU.anchorY[firstMDIndex]; - dPhi = lst::deltaPhi(acc, midPointX, midPointY, diffX, diffY); + float dPhi = lst::deltaPhi(acc, midPointX, midPointY, diffX, diffY); // Cut #5: deltaPhiChange if (alpaka::math::abs(acc, dPhi) > sdlCut) return false; @@ -1897,11 +1721,12 @@ namespace lst { float tl_axis_x = mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[firstMDIndex]; float tl_axis_y = mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; - betaIn = sdIn_alpha - lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[firstMDIndex]); + float betaIn = + sdIn_alpha - lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[firstMDIndex]); float betaInRHmin = betaIn; float betaInRHmax = betaIn; - betaOut = + float betaOut = -sdOut_alphaOut + lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut; @@ -1940,7 +1765,7 @@ namespace lst { float dr = alpaka::math::sqrt(acc, tl_axis_x * tl_axis_x + tl_axis_y * tl_axis_y); const float corrF = 1.f; - betaInCut = + float betaInCut = alpaka::math::asin( acc, alpaka::math::min(acc, (-sdIn_dr * corrF + dr) * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + (0.02f / sdIn_d); @@ -1950,7 +1775,7 @@ namespace lst { return false; float betaAv = 0.5f * (betaIn + betaOut); - pt_beta = dr * lst::k2Rinv1GeVf / alpaka::math::sin(acc, betaAv); + float pt_beta = dr * lst::k2Rinv1GeVf / alpaka::math::sin(acc, betaAv); float lIn = 5; float lOut = 11; @@ -2006,8 +1831,7 @@ namespace lst { } const float dBetaROut2 = dBetaROut * dBetaROut; - //FIXME: need faster version - betaOutCut = + float betaOutCut = alpaka::math::asin(acc, alpaka::math::min(acc, dr * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + (0.02f / sdOut_d) + alpaka::math::sqrt(acc, dBetaLum2 + dBetaMuls2); @@ -2022,16 +1846,15 @@ namespace lst { (alpaka::math::abs(acc, betaInRHmin - betaInRHmax) + alpaka::math::abs(acc, betaOutRHmin - betaOutRHmax)) * (alpaka::math::abs(acc, betaInRHmin - betaInRHmax) + alpaka::math::abs(acc, betaOutRHmin - betaOutRHmax))); float dBeta = betaIn - betaOut; - deltaBetaCut = alpaka::math::sqrt(acc, dBetaCut2); //Cut #7: Cut on dBet return dBeta * dBeta <= dBetaCut2; }; template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool runQuintupletDefaultAlgoEEEE(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, - struct lst::Segments& segmentsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, + lst::Segments const& segmentsInGPU, uint16_t innerInnerLowerModuleIndex, uint16_t innerOuterLowerModuleIndex, uint16_t outerInnerLowerModuleIndex, @@ -2041,25 +1864,7 @@ namespace lst { unsigned int firstMDIndex, unsigned int secondMDIndex, unsigned int thirdMDIndex, - unsigned int fourthMDIndex, - float& zOut, - float& rtOut, - float& deltaPhiPos, - float& dPhi, - float& betaIn, - float& betaOut, - float& pt_beta, - float& zLo, - float& rtLo, - float& rtHi, - float& sdlCut, - float& betaInCut, - float& betaOutCut, - float& deltaBetaCut, - float& kZ) { - bool isPS_InLo = (modulesInGPU.moduleType[innerInnerLowerModuleIndex] == lst::PS); - bool isPS_OutLo = (modulesInGPU.moduleType[outerInnerLowerModuleIndex] == lst::PS); - + unsigned int fourthMDIndex) { float rt_InLo = mdsInGPU.anchorRt[firstMDIndex]; float rt_InOut = mdsInGPU.anchorRt[secondMDIndex]; float rt_OutLo = mdsInGPU.anchorRt[thirdMDIndex]; @@ -2071,15 +1876,8 @@ namespace lst { float alpha1GeV_OutLo = alpaka::math::asin(acc, alpaka::math::min(acc, rt_OutLo * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)); - float rtRatio_OutLoInLo = rt_OutLo / rt_InLo; // Outer segment beginning rt divided by inner segment beginning rt; float dzDrtScale = alpaka::math::tan(acc, alpha1GeV_OutLo) / alpha1GeV_OutLo; // The track can bend in r-z plane slightly - float zpitch_InLo = (isPS_InLo ? lst::kPixelPSZpitch : lst::kStrip2SZpitch); - float zpitch_OutLo = (isPS_OutLo ? lst::kPixelPSZpitch : lst::kStrip2SZpitch); - float zGeom = zpitch_InLo + zpitch_OutLo; - - zLo = z_InLo + (z_InLo - lst::kDeltaZLum) * (rtRatio_OutLoInLo - 1.f) * (z_InLo > 0.f ? 1.f : dzDrtScale) - - zGeom; //slope-correction only on outer end // Cut #0: Preliminary (Only here in endcap case) if ((z_InLo * z_OutLo) <= 0) @@ -2094,14 +1892,13 @@ namespace lst { : 2.f * lst::kStrip2SZpitch; float dz = z_OutLo - z_InLo; - rtLo = rt_InLo * (1.f + dz / (z_InLo + dLum) / dzDrtScale) - rtGeom; //slope correction only on the lower end + float rtLo = rt_InLo * (1.f + dz / (z_InLo + dLum) / dzDrtScale) - rtGeom; //slope correction only on the lower end - zOut = z_OutLo; - rtOut = rt_OutLo; + float rtOut = rt_OutLo; //Cut #1: rt condition - rtHi = rt_InLo * (1.f + dz / (z_InLo - dLum)) + rtGeom; + float rtHi = rt_InLo * (1.f + dz / (z_InLo - dLum)) + rtGeom; if ((rtOut < rtLo) || (rtOut > rtHi)) return false; @@ -2116,7 +1913,7 @@ namespace lst { float dzOutInAbs = alpaka::math::abs(acc, z_OutLo - z_InLo); float multDzDr = dzOutInAbs * coshEta / (coshEta * coshEta - 1.f); - kZ = (z_OutLo - z_InLo) / dzSDIn; + float kZ = (z_OutLo - z_InLo) / dzSDIn; float sdlThetaMulsF2 = (kMulsInGeV * kMulsInGeV) * (0.1f + 0.2f * (rt_OutLo - rt_InLo) / 50.f); float sdlMuls2 = sdlThetaMulsF2 * 9.f / (lst::ptCut * lst::ptCut) * 16.f; @@ -2141,9 +1938,9 @@ namespace lst { } float sdlPVoff = 0.1f / rtOut; - sdlCut = alpha1GeV_OutLo + alpaka::math::sqrt(acc, sdlMuls2 + sdlPVoff * sdlPVoff); + float sdlCut = alpha1GeV_OutLo + alpaka::math::sqrt(acc, sdlMuls2 + sdlPVoff * sdlPVoff); - deltaPhiPos = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[fourthMDIndex] - mdsInGPU.anchorPhi[secondMDIndex]); + float deltaPhiPos = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[fourthMDIndex] - mdsInGPU.anchorPhi[secondMDIndex]); if (alpaka::math::abs(acc, deltaPhiPos) > sdlCut) return false; @@ -2153,7 +1950,7 @@ namespace lst { float diffX = mdsInGPU.anchorX[thirdMDIndex] - mdsInGPU.anchorX[firstMDIndex]; float diffY = mdsInGPU.anchorY[thirdMDIndex] - mdsInGPU.anchorY[firstMDIndex]; - dPhi = lst::deltaPhi(acc, midPointX, midPointY, diffX, diffY); + float dPhi = lst::deltaPhi(acc, midPointX, midPointY, diffX, diffY); // Cut #5: deltaPhiChange if (alpaka::math::abs(acc, dPhi) > sdlCut) @@ -2174,14 +1971,15 @@ namespace lst { float tl_axis_x = mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[firstMDIndex]; float tl_axis_y = mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; - betaIn = sdIn_alpha - lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[firstMDIndex]); + float betaIn = + sdIn_alpha - lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[firstMDIndex]); float sdIn_alphaRHmin = __H2F(segmentsInGPU.dPhiChangeMins[innerSegmentIndex]); float sdIn_alphaRHmax = __H2F(segmentsInGPU.dPhiChangeMaxs[innerSegmentIndex]); float betaInRHmin = betaIn + sdIn_alphaRHmin - sdIn_alpha; float betaInRHmax = betaIn + sdIn_alphaRHmax - sdIn_alpha; - betaOut = + float betaOut = -sdOut_alphaOut + lst::phi_mpi_pi(acc, lst::phi(acc, tl_axis_x, tl_axis_y) - mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut - sdOut_alphaOutRHmin + sdOut_alphaOut; @@ -2208,7 +2006,7 @@ namespace lst { float dr = alpaka::math::sqrt(acc, tl_axis_x * tl_axis_x + tl_axis_y * tl_axis_y); const float corrF = 1.f; - betaInCut = + float betaInCut = alpaka::math::asin( acc, alpaka::math::min(acc, (-sdIn_dr * corrF + dr) * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + (0.02f / sdIn_d); @@ -2218,7 +2016,7 @@ namespace lst { return false; float betaAv = 0.5f * (betaIn + betaOut); - pt_beta = dr * lst::k2Rinv1GeVf / alpaka::math::sin(acc, betaAv); + float pt_beta = dr * lst::k2Rinv1GeVf / alpaka::math::sin(acc, betaAv); int lIn = 11; //endcap int lOut = 13; //endcap @@ -2262,8 +2060,7 @@ namespace lst { const float dBetaRIn2 = 0; // TODO-RH float dBetaROut2 = 0; //TODO-RH - //FIXME: need faster version - betaOutCut = + float betaOutCut = alpaka::math::asin(acc, alpaka::math::min(acc, dr * lst::k2Rinv1GeVf / lst::ptCut, lst::kSinAlphaMax)) + (0.02f / sdOut_d) + alpaka::math::sqrt(acc, dBetaLum2 + dBetaMuls2); @@ -2279,16 +2076,14 @@ namespace lst { (alpaka::math::abs(acc, betaInRHmin - betaInRHmax) + alpaka::math::abs(acc, betaOutRHmin - betaOutRHmax))); float dBeta = betaIn - betaOut; //Cut #7: Cut on dBeta - deltaBetaCut = alpaka::math::sqrt(acc, dBetaCut2); - return dBeta * dBeta <= dBetaCut2; }; template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool runQuintupletAlgoSelector(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, - struct lst::Segments& segmentsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, + lst::Segments const& segmentsInGPU, uint16_t innerInnerLowerModuleIndex, uint16_t innerOuterLowerModuleIndex, uint16_t outerInnerLowerModuleIndex, @@ -2298,34 +2093,7 @@ namespace lst { unsigned int firstMDIndex, unsigned int secondMDIndex, unsigned int thirdMDIndex, - unsigned int fourthMDIndex, - float& zOut, - float& rtOut, - float& deltaPhiPos, - float& deltaPhi, - float& betaIn, - float& betaOut, - float& pt_beta, - float& zLo, - float& zHi, - float& rtLo, - float& rtHi, - float& zLoPointed, - float& zHiPointed, - float& sdlCut, - float& betaInCut, - float& betaOutCut, - float& deltaBetaCut, - float& kZ) { - zLo = -999; - zHi = -999; - rtLo = -999; - rtHi = -999; - zLoPointed = -999; - zHiPointed = -999; - kZ = -999; - betaInCut = -999; - + unsigned int fourthMDIndex) { short innerInnerLowerModuleSubdet = modulesInGPU.subdets[innerInnerLowerModuleIndex]; short innerOuterLowerModuleSubdet = modulesInGPU.subdets[innerOuterLowerModuleIndex]; short outerInnerLowerModuleSubdet = modulesInGPU.subdets[outerInnerLowerModuleIndex]; @@ -2346,22 +2114,7 @@ namespace lst { firstMDIndex, secondMDIndex, thirdMDIndex, - fourthMDIndex, - zOut, - rtOut, - deltaPhiPos, - deltaPhi, - betaIn, - betaOut, - pt_beta, - zLo, - zHi, - zLoPointed, - zHiPointed, - sdlCut, - betaInCut, - betaOutCut, - deltaBetaCut); + fourthMDIndex); } else if (innerInnerLowerModuleSubdet == lst::Barrel and innerOuterLowerModuleSubdet == lst::Barrel and outerInnerLowerModuleSubdet == lst::Endcap and outerOuterLowerModuleSubdet == lst::Endcap) { return runQuintupletDefaultAlgoBBEE(acc, @@ -2377,22 +2130,7 @@ namespace lst { firstMDIndex, secondMDIndex, thirdMDIndex, - fourthMDIndex, - zOut, - rtOut, - deltaPhiPos, - deltaPhi, - betaIn, - betaOut, - pt_beta, - zLo, - rtLo, - rtHi, - sdlCut, - betaInCut, - betaOutCut, - deltaBetaCut, - kZ); + fourthMDIndex); } else if (innerInnerLowerModuleSubdet == lst::Barrel and innerOuterLowerModuleSubdet == lst::Barrel and outerInnerLowerModuleSubdet == lst::Barrel and outerOuterLowerModuleSubdet == lst::Endcap) { return runQuintupletDefaultAlgoBBBB(acc, @@ -2408,22 +2146,7 @@ namespace lst { firstMDIndex, secondMDIndex, thirdMDIndex, - fourthMDIndex, - zOut, - rtOut, - deltaPhiPos, - deltaPhi, - betaIn, - betaOut, - pt_beta, - zLo, - zHi, - zLoPointed, - zHiPointed, - sdlCut, - betaInCut, - betaOutCut, - deltaBetaCut); + fourthMDIndex); } else if (innerInnerLowerModuleSubdet == lst::Barrel and innerOuterLowerModuleSubdet == lst::Endcap and outerInnerLowerModuleSubdet == lst::Endcap and outerOuterLowerModuleSubdet == lst::Endcap) { return runQuintupletDefaultAlgoBBEE(acc, @@ -2439,22 +2162,7 @@ namespace lst { firstMDIndex, secondMDIndex, thirdMDIndex, - fourthMDIndex, - zOut, - rtOut, - deltaPhiPos, - deltaPhi, - betaIn, - betaOut, - pt_beta, - zLo, - rtLo, - rtHi, - sdlCut, - betaInCut, - betaOutCut, - deltaBetaCut, - kZ); + fourthMDIndex); } else if (innerInnerLowerModuleSubdet == lst::Endcap and innerOuterLowerModuleSubdet == lst::Endcap and outerInnerLowerModuleSubdet == lst::Endcap and outerOuterLowerModuleSubdet == lst::Endcap) { return runQuintupletDefaultAlgoEEEE(acc, @@ -2470,22 +2178,7 @@ namespace lst { firstMDIndex, secondMDIndex, thirdMDIndex, - fourthMDIndex, - zOut, - rtOut, - deltaPhiPos, - deltaPhi, - betaIn, - betaOut, - pt_beta, - zLo, - rtLo, - rtHi, - sdlCut, - betaInCut, - betaOutCut, - deltaBetaCut, - kZ); + fourthMDIndex); } return false; @@ -2528,10 +2221,6 @@ namespace lst { if (innerOuterOuterMiniDoubletIndex != outerInnerInnerMiniDoubletIndex) return false; - //apply T4 criteria between segments 1 and 3 - float zOut, rtOut, deltaPhiPos, deltaPhi, betaIn, betaOut, pt_beta; //temp stuff - float zLo, zHi, rtLo, rtHi, zLoPointed, zHiPointed, sdlCut, betaInCut, betaOutCut, deltaBetaCut, kZ; - unsigned int firstMDIndex = segmentsInGPU.mdIndices[2 * firstSegmentIndex]; unsigned int secondMDIndex = segmentsInGPU.mdIndices[2 * secondSegmentIndex]; unsigned int thirdMDIndex = segmentsInGPU.mdIndices[2 * secondSegmentIndex + 1]; @@ -2551,25 +2240,7 @@ namespace lst { firstMDIndex, secondMDIndex, thirdMDIndex, - fourthMDIndex, - zOut, - rtOut, - deltaPhiPos, - deltaPhi, - betaIn, - betaOut, - pt_beta, - zLo, - zHi, - rtLo, - rtHi, - zLoPointed, - zHiPointed, - sdlCut, - betaInCut, - betaOutCut, - deltaBetaCut, - kZ)) + fourthMDIndex)) return false; if (not runQuintupletAlgoSelector(acc, @@ -2585,25 +2256,7 @@ namespace lst { firstMDIndex, secondMDIndex, fourthMDIndex, - fifthMDIndex, - zOut, - rtOut, - deltaPhiPos, - deltaPhi, - betaIn, - betaOut, - pt_beta, - zLo, - zHi, - rtLo, - rtHi, - zLoPointed, - zHiPointed, - sdlCut, - betaInCut, - betaOutCut, - deltaBetaCut, - kZ)) + fifthMDIndex)) return false; float x1 = mdsInGPU.anchorX[firstMDIndex]; @@ -2722,9 +2375,6 @@ namespace lst { if (innerRadius < 0.95f * ptCut / (2.f * k2Rinv1GeVf)) return false; - float innerInvRadiusMin, innerInvRadiusMax, bridgeInvRadiusMin, bridgeInvRadiusMax, outerInvRadiusMin, - outerInvRadiusMax; - //split by category bool matchedRadii; if (modulesInGPU.subdets[lowerModuleIndex1] == lst::Barrel and @@ -2732,93 +2382,27 @@ namespace lst { modulesInGPU.subdets[lowerModuleIndex3] == lst::Barrel and modulesInGPU.subdets[lowerModuleIndex4] == lst::Barrel and modulesInGPU.subdets[lowerModuleIndex5] == lst::Barrel) { - matchedRadii = matchRadiiBBBBB(acc, - innerRadius, - bridgeRadius, - outerRadius, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + 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) { - matchedRadii = matchRadiiBBBBE(acc, - innerRadius, - bridgeRadius, - outerRadius, - innerRadiusMin2S, - innerRadiusMax2S, - bridgeRadiusMin2S, - bridgeRadiusMax2S, - outerRadiusMin2S, - outerRadiusMax2S, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + 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) { if (modulesInGPU.layers[lowerModuleIndex1] == 1) { - matchedRadii = matchRadiiBBBEE12378(acc, - innerRadius, - bridgeRadius, - outerRadius, - innerRadiusMin2S, - innerRadiusMax2S, - bridgeRadiusMin2S, - bridgeRadiusMax2S, - outerRadiusMin2S, - outerRadiusMax2S, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + matchedRadii = + matchRadiiBBBEE12378(acc, innerRadius, bridgeRadius, outerRadius, bridgeRadiusMin2S, bridgeRadiusMax2S); } else if (modulesInGPU.layers[lowerModuleIndex1] == 2) { - matchedRadii = matchRadiiBBBEE23478(acc, - innerRadius, - bridgeRadius, - outerRadius, - innerRadiusMin2S, - innerRadiusMax2S, - bridgeRadiusMin2S, - bridgeRadiusMax2S, - outerRadiusMin2S, - outerRadiusMax2S, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + matchedRadii = + matchRadiiBBBEE23478(acc, innerRadius, bridgeRadius, outerRadius, bridgeRadiusMin2S, bridgeRadiusMax2S); } else { - matchedRadii = matchRadiiBBBEE34578(acc, - innerRadius, - bridgeRadius, - outerRadius, - innerRadiusMin2S, - innerRadiusMax2S, - bridgeRadiusMin2S, - bridgeRadiusMax2S, - outerRadiusMin2S, - outerRadiusMax2S, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + matchedRadii = + matchRadiiBBBEE34578(acc, innerRadius, bridgeRadius, outerRadius, bridgeRadiusMin2S, bridgeRadiusMax2S); } } @@ -2827,22 +2411,7 @@ namespace lst { modulesInGPU.subdets[lowerModuleIndex3] == lst::Endcap and modulesInGPU.subdets[lowerModuleIndex4] == lst::Endcap and modulesInGPU.subdets[lowerModuleIndex5] == lst::Endcap) { - matchedRadii = matchRadiiBBEEE(acc, - innerRadius, - bridgeRadius, - outerRadius, - innerRadiusMin2S, - innerRadiusMax2S, - bridgeRadiusMin2S, - bridgeRadiusMax2S, - outerRadiusMin2S, - outerRadiusMax2S, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + 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 @@ -2855,15 +2424,7 @@ namespace lst { innerRadiusMin2S, innerRadiusMax2S, bridgeRadiusMin2S, - bridgeRadiusMax2S, - outerRadiusMin2S, - outerRadiusMax2S, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + bridgeRadiusMax2S); } else { matchedRadii = matchRadiiEEEEE(acc, innerRadius, @@ -2872,15 +2433,7 @@ namespace lst { innerRadiusMin2S, innerRadiusMax2S, bridgeRadiusMin2S, - bridgeRadiusMax2S, - outerRadiusMin2S, - outerRadiusMax2S, - innerInvRadiusMin, - innerInvRadiusMax, - bridgeInvRadiusMin, - bridgeInvRadiusMax, - outerInvRadiusMin, - outerInvRadiusMax); + bridgeRadiusMax2S); } //compute regression radius right here - this computation is expensive!!! @@ -2985,12 +2538,12 @@ namespace lst { struct createQuintupletsInGPUv2 { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::MiniDoublets mdsInGPU, - struct lst::Segments segmentsInGPU, - struct lst::Triplets tripletsInGPU, - struct lst::Quintuplets quintupletsInGPU, - struct lst::ObjectRanges rangesInGPU, + lst::Modules modulesInGPU, + lst::MiniDoublets mdsInGPU, + lst::Segments segmentsInGPU, + lst::Triplets tripletsInGPU, + lst::Quintuplets quintupletsInGPU, + lst::ObjectRanges rangesInGPU, uint16_t nEligibleT5Modules) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -3114,9 +2667,9 @@ namespace lst { struct createEligibleModulesListForQuintupletsGPU { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::Triplets tripletsInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::Modules modulesInGPU, + lst::Triplets tripletsInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -3212,9 +2765,9 @@ namespace lst { struct addQuintupletRangesToEventExplicit { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::Quintuplets quintupletsInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::Modules modulesInGPU, + lst::Quintuplets quintupletsInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); diff --git a/RecoTracker/LSTCore/src/alpaka/Segment.h b/RecoTracker/LSTCore/src/alpaka/Segment.h index 93a17e3d72e46..643ee9b0dddd3 100644 --- a/RecoTracker/LSTCore/src/alpaka/Segment.h +++ b/RecoTracker/LSTCore/src/alpaka/Segment.h @@ -175,7 +175,7 @@ namespace lst { inline void setData(SegmentsBuffer& buf) { data_.setData(buf); } }; - ALPAKA_FN_ACC ALPAKA_FN_INLINE float isTighterTiltedModules_seg(struct lst::Modules& modulesInGPU, + ALPAKA_FN_ACC ALPAKA_FN_INLINE float isTighterTiltedModules_seg(lst::Modules const& modulesInGPU, unsigned int moduleIndex) { // 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" @@ -199,8 +199,7 @@ namespace lst { ((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) { + ALPAKA_FN_ACC ALPAKA_FN_INLINE float moduleGapSize_seg(short layer, short ring, short subdet, short side, short rod) { static constexpr float miniDeltaTilted[3] = {0.26f, 0.26f, 0.26f}; static constexpr float miniDeltaFlat[6] = {0.26f, 0.16f, 0.16f, 0.18f, 0.18f, 0.18f}; static constexpr float miniDeltaLooseTilted[3] = {0.4f, 0.4f, 0.4f}; @@ -230,8 +229,7 @@ namespace lst { return moduleSeparation; }; - ALPAKA_FN_ACC ALPAKA_FN_INLINE float moduleGapSize_seg(struct lst::Modules& modulesInGPU, - unsigned int moduleIndex) { + ALPAKA_FN_ACC ALPAKA_FN_INLINE float moduleGapSize_seg(lst::Modules const& modulesInGPU, unsigned int moduleIndex) { static constexpr float miniDeltaTilted[3] = {0.26f, 0.26f, 0.26f}; static constexpr float miniDeltaFlat[6] = {0.26f, 0.16f, 0.16f, 0.18f, 0.18f, 0.18f}; static constexpr float miniDeltaLooseTilted[3] = {0.4f, 0.4f, 0.4f}; @@ -266,16 +264,16 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE void dAlphaThreshold(TAcc const& acc, float* dAlphaThresholdValues, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, - float& xIn, - float& yIn, - float& zIn, - float& rtIn, - float& xOut, - float& yOut, - float& zOut, - float& rtOut, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, + float xIn, + float yIn, + float zIn, + float rtIn, + float xOut, + float yOut, + float zOut, + float rtOut, uint16_t innerLowerModuleIndex, uint16_t outerLowerModuleIndex, unsigned int innerMDIndex, @@ -295,8 +293,8 @@ namespace lst { bool isOuterTilted = modulesInGPU.subdets[outerLowerModuleIndex] == lst::Barrel and modulesInGPU.sides[outerLowerModuleIndex] != lst::Center; - const float& drdzInner = modulesInGPU.drdzs[innerLowerModuleIndex]; - const float& drdzOuter = modulesInGPU.drdzs[outerLowerModuleIndex]; + float drdzInner = modulesInGPU.drdzs[innerLowerModuleIndex]; + float drdzOuter = modulesInGPU.drdzs[outerLowerModuleIndex]; float innerModuleGapSize = lst::moduleGapSize_seg(modulesInGPU, innerLowerModuleIndex); float outerModuleGapSize = lst::moduleGapSize_seg(modulesInGPU, outerLowerModuleIndex); const float innerminiTilt2 = isInnerTilted @@ -358,24 +356,20 @@ namespace lst { dAlphaThresholdValues[2] = dAlpha_Bfield + alpaka::math::sqrt(acc, dAlpha_res * dAlpha_res + sdMuls * sdMuls); }; - ALPAKA_FN_ACC ALPAKA_FN_INLINE void addSegmentToMemory(struct lst::Segments& segmentsInGPU, + ALPAKA_FN_ACC ALPAKA_FN_INLINE void addSegmentToMemory(lst::Segments& segmentsInGPU, unsigned int lowerMDIndex, unsigned int upperMDIndex, uint16_t innerLowerModuleIndex, uint16_t outerLowerModuleIndex, unsigned int innerMDAnchorHitIndex, unsigned int outerMDAnchorHitIndex, - float& dPhi, - float& dPhiMin, - float& dPhiMax, - float& dPhiChange, - float& dPhiChangeMin, - float& dPhiChangeMax, + float dPhi, + float dPhiMin, + float dPhiMax, + float dPhiChange, + float dPhiChangeMin, + float dPhiChangeMax, unsigned int idx) { - //idx will be computed in the kernel, which is the index into which the - //segment will be written - //nSegments will be incremented in the kernel - //printf("seg: %u %u %u %u\n",lowerMDIndex, upperMDIndex,innerLowerModuleIndex,outerLowerModuleIndex); segmentsInGPU.mdIndices[idx * 2] = lowerMDIndex; segmentsInGPU.mdIndices[idx * 2 + 1] = upperMDIndex; segmentsInGPU.innerLowerModuleIndices[idx] = innerLowerModuleIndex; @@ -393,8 +387,8 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE void addPixelSegmentToMemory(TAcc const& acc, - struct lst::Segments& segmentsInGPU, - struct lst::MiniDoublets& mdsInGPU, + lst::Segments& segmentsInGPU, + lst::MiniDoublets const& mdsInGPU, unsigned int innerMDIndex, unsigned int outerMDIndex, uint16_t pixelModuleIndex, @@ -456,36 +450,23 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool runSegmentDefaultAlgoBarrel(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, uint16_t innerLowerModuleIndex, uint16_t outerLowerModuleIndex, - unsigned int& innerMDIndex, - unsigned int& outerMDIndex, - float& zIn, - float& zOut, - float& rtIn, - float& rtOut, + unsigned int innerMDIndex, + unsigned int outerMDIndex, float& dPhi, float& dPhiMin, float& dPhiMax, float& dPhiChange, float& dPhiChangeMin, - float& dPhiChangeMax, - float& dAlphaInnerMDSegment, - float& dAlphaOuterMDSegment, - float& dAlphaInnerMDOuterMD, - float& zLo, - float& zHi, - float& sdCut, - float& dAlphaInnerMDSegmentThreshold, - float& dAlphaOuterMDSegmentThreshold, - float& dAlphaInnerMDOuterMDThreshold) { + float& dPhiChangeMax) { float sdMuls = (modulesInGPU.subdets[innerLowerModuleIndex] == lst::Barrel) ? kMiniMulsPtScaleBarrel[modulesInGPU.layers[innerLowerModuleIndex] - 1] * 3.f / ptCut : kMiniMulsPtScaleEndcap[modulesInGPU.layers[innerLowerModuleIndex] - 1] * 3.f / ptCut; - float xIn, yIn, xOut, yOut; + float xIn, yIn, zIn, rtIn, xOut, yOut, zOut, rtOut; xIn = mdsInGPU.anchorX[innerMDIndex]; yIn = mdsInGPU.anchorY[innerMDIndex]; @@ -503,14 +484,14 @@ namespace lst { const float zGeom = modulesInGPU.layers[innerLowerModuleIndex] <= 2 ? 2.f * kPixelPSZpitch : 2.f * kStrip2SZpitch; - zLo = zIn + (zIn - kDeltaZLum) * (rtOut / rtIn - 1.f) * (zIn > 0.f ? 1.f : dzDrtScale) - - zGeom; //slope-correction only on outer end - zHi = zIn + (zIn + kDeltaZLum) * (rtOut / rtIn - 1.f) * (zIn < 0.f ? 1.f : dzDrtScale) + zGeom; + float zLo = zIn + (zIn - kDeltaZLum) * (rtOut / rtIn - 1.f) * (zIn > 0.f ? 1.f : dzDrtScale) - + zGeom; //slope-correction only on outer end + float zHi = zIn + (zIn + kDeltaZLum) * (rtOut / rtIn - 1.f) * (zIn < 0.f ? 1.f : dzDrtScale) + zGeom; if ((zOut < zLo) || (zOut > zHi)) return false; - sdCut = sdSlope + alpaka::math::sqrt(acc, sdMuls * sdMuls + sdPVoff * sdPVoff); + float sdCut = sdSlope + alpaka::math::sqrt(acc, sdMuls * sdMuls + sdPVoff * sdPVoff); dPhi = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[outerMDIndex] - mdsInGPU.anchorPhi[innerMDIndex]); @@ -542,13 +523,13 @@ namespace lst { float innerMDAlpha = mdsInGPU.dphichanges[innerMDIndex]; float outerMDAlpha = mdsInGPU.dphichanges[outerMDIndex]; - dAlphaInnerMDSegment = innerMDAlpha - dPhiChange; - dAlphaOuterMDSegment = outerMDAlpha - dPhiChange; - dAlphaInnerMDOuterMD = innerMDAlpha - outerMDAlpha; + float dAlphaInnerMDSegment = innerMDAlpha - dPhiChange; + float dAlphaOuterMDSegment = outerMDAlpha - dPhiChange; + float dAlphaInnerMDOuterMD = innerMDAlpha - outerMDAlpha; - dAlphaInnerMDSegmentThreshold = dAlphaThresholdValues[0]; - dAlphaOuterMDSegmentThreshold = dAlphaThresholdValues[1]; - dAlphaInnerMDOuterMDThreshold = dAlphaThresholdValues[2]; + float dAlphaInnerMDSegmentThreshold = dAlphaThresholdValues[0]; + float dAlphaOuterMDSegmentThreshold = dAlphaThresholdValues[1]; + float dAlphaInnerMDOuterMDThreshold = dAlphaThresholdValues[2]; if (alpaka::math::abs(acc, dAlphaInnerMDSegment) >= dAlphaInnerMDSegmentThreshold) return false; @@ -559,33 +540,19 @@ namespace lst { template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool runSegmentDefaultAlgoEndcap(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, uint16_t innerLowerModuleIndex, uint16_t outerLowerModuleIndex, unsigned int innerMDIndex, unsigned int outerMDIndex, - float& zIn, - float& zOut, - float& rtIn, - float& rtOut, float& dPhi, float& dPhiMin, float& dPhiMax, float& dPhiChange, float& dPhiChangeMin, - float& dPhiChangeMax, - float& dAlphaInnerMDSegment, - float& dAlphaOuterMDSegment, - float& rtLo, - float& rtHi, - float& sdCut, - float& dAlphaInnerMDSegmentThreshold, - float& dAlphaOuterMDSegmentThreshold, - float& dAlphaInnerMDOuterMDThreshold, - float& dAlphaInnerMDOuterMD) { - float xIn, yIn; - float xOut, yOut; + float& dPhiChangeMax) { + float xIn, yIn, zIn, rtIn, xOut, yOut, zOut, rtOut; xIn = mdsInGPU.anchorX[innerMDIndex]; yIn = mdsInGPU.anchorY[innerMDIndex]; @@ -617,10 +584,10 @@ namespace lst { float dLum = alpaka::math::copysign(acc, kDeltaZLum, zIn); float drtDzScale = sdSlope / alpaka::math::tan(acc, sdSlope); - rtLo = alpaka::math::max( + float rtLo = alpaka::math::max( acc, rtIn * (1.f + dz / (zIn + dLum) * drtDzScale) - rtGeom, rtIn - 0.5f * rtGeom); //rt should increase - rtHi = rtIn * (zOut - dLum) / (zIn - dLum) + - rtGeom; //dLum for luminous; rGeom for measurement size; no tanTheta_loc(pt) correction + float rtHi = rtIn * (zOut - dLum) / (zIn - dLum) + + rtGeom; //dLum for luminous; rGeom for measurement size; no tanTheta_loc(pt) correction // Completeness if ((rtOut < rtLo) || (rtOut > rtHi)) @@ -628,7 +595,7 @@ namespace lst { dPhi = lst::phi_mpi_pi(acc, mdsInGPU.anchorPhi[outerMDIndex] - mdsInGPU.anchorPhi[innerMDIndex]); - sdCut = sdSlope; + float sdCut = sdSlope; if (outerLayerEndcapTwoS) { float dPhiPos_high = lst::phi_mpi_pi(acc, mdsInGPU.anchorHighEdgePhi[outerMDIndex] - mdsInGPU.anchorPhi[innerMDIndex]); @@ -670,57 +637,37 @@ namespace lst { innerMDIndex, outerMDIndex); - dAlphaInnerMDSegmentThreshold = dAlphaThresholdValues[0]; - dAlphaOuterMDSegmentThreshold = dAlphaThresholdValues[1]; - dAlphaInnerMDOuterMDThreshold = dAlphaThresholdValues[2]; - float innerMDAlpha = mdsInGPU.dphichanges[innerMDIndex]; float outerMDAlpha = mdsInGPU.dphichanges[outerMDIndex]; - dAlphaInnerMDSegment = innerMDAlpha - dPhiChange; - dAlphaOuterMDSegment = outerMDAlpha - dPhiChange; - dAlphaInnerMDOuterMD = innerMDAlpha - outerMDAlpha; + float dAlphaInnerMDSegment = innerMDAlpha - dPhiChange; + float dAlphaOuterMDSegment = outerMDAlpha - dPhiChange; + float dAlphaInnerMDOuterMD = innerMDAlpha - outerMDAlpha; + + float dAlphaInnerMDSegmentThreshold = dAlphaThresholdValues[0]; + float dAlphaOuterMDSegmentThreshold = dAlphaThresholdValues[1]; + float dAlphaInnerMDOuterMDThreshold = dAlphaThresholdValues[2]; - if (alpaka::math::abs(acc, dAlphaInnerMDSegment) >= dAlphaThresholdValues[0]) + if (alpaka::math::abs(acc, dAlphaInnerMDSegment) >= dAlphaInnerMDSegmentThreshold) return false; - if (alpaka::math::abs(acc, dAlphaOuterMDSegment) >= dAlphaThresholdValues[1]) + if (alpaka::math::abs(acc, dAlphaOuterMDSegment) >= dAlphaOuterMDSegmentThreshold) return false; - return alpaka::math::abs(acc, dAlphaInnerMDOuterMD) < dAlphaThresholdValues[2]; + return alpaka::math::abs(acc, dAlphaInnerMDOuterMD) < dAlphaInnerMDOuterMDThreshold; }; template ALPAKA_FN_ACC ALPAKA_FN_INLINE bool runSegmentDefaultAlgo(TAcc const& acc, - struct lst::Modules& modulesInGPU, - struct lst::MiniDoublets& mdsInGPU, + lst::Modules const& modulesInGPU, + lst::MiniDoublets const& mdsInGPU, uint16_t innerLowerModuleIndex, uint16_t outerLowerModuleIndex, unsigned int innerMDIndex, unsigned int outerMDIndex, - float& zIn, - float& zOut, - float& rtIn, - float& rtOut, float& dPhi, float& dPhiMin, float& dPhiMax, float& dPhiChange, float& dPhiChangeMin, - float& dPhiChangeMax, - float& dAlphaInnerMDSegment, - float& dAlphaOuterMDSegment, - float& dAlphaInnerMDOuterMD, - float& zLo, - float& zHi, - float& rtLo, - float& rtHi, - float& sdCut, - float& dAlphaInnerMDSegmentThreshold, - float& dAlphaOuterMDSegmentThreshold, - float& dAlphaInnerMDOuterMDThreshold) { - zLo = -999.f; - zHi = -999.f; - rtLo = -999.f; - rtHi = -999.f; - + float& dPhiChangeMax) { if (modulesInGPU.subdets[innerLowerModuleIndex] == lst::Barrel and modulesInGPU.subdets[outerLowerModuleIndex] == lst::Barrel) { return runSegmentDefaultAlgoBarrel(acc, @@ -730,25 +677,12 @@ namespace lst { outerLowerModuleIndex, innerMDIndex, outerMDIndex, - zIn, - zOut, - rtIn, - rtOut, dPhi, dPhiMin, dPhiMax, dPhiChange, dPhiChangeMin, - dPhiChangeMax, - dAlphaInnerMDSegment, - dAlphaOuterMDSegment, - dAlphaInnerMDOuterMD, - zLo, - zHi, - sdCut, - dAlphaInnerMDSegmentThreshold, - dAlphaOuterMDSegmentThreshold, - dAlphaInnerMDOuterMDThreshold); + dPhiChangeMax); } else { return runSegmentDefaultAlgoEndcap(acc, modulesInGPU, @@ -757,35 +691,22 @@ namespace lst { outerLowerModuleIndex, innerMDIndex, outerMDIndex, - zIn, - zOut, - rtIn, - rtOut, dPhi, dPhiMin, dPhiMax, dPhiChange, dPhiChangeMin, - dPhiChangeMax, - dAlphaInnerMDSegment, - dAlphaOuterMDSegment, - dAlphaInnerMDOuterMD, - rtLo, - rtHi, - sdCut, - dAlphaInnerMDSegmentThreshold, - dAlphaOuterMDSegmentThreshold, - dAlphaInnerMDOuterMDThreshold); + dPhiChangeMax); } }; struct createSegmentsInGPUv2 { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::MiniDoublets mdsInGPU, - struct lst::Segments segmentsInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::Modules modulesInGPU, + lst::MiniDoublets mdsInGPU, + lst::Segments segmentsInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalBlockIdx = alpaka::getIdx(acc); auto const blockThreadIdx = alpaka::getIdx(acc); auto const gridBlockExtent = alpaka::getWorkDiv(acc); @@ -819,8 +740,7 @@ namespace lst { unsigned int innerMDIndex = rangesInGPU.mdRanges[innerLowerModuleIndex * 2] + innerMDArrayIdx; unsigned int outerMDIndex = rangesInGPU.mdRanges[outerLowerModuleIndex * 2] + outerMDArrayIdx; - float zIn, zOut, rtIn, rtOut, dPhi, dPhiMin, dPhiMax, dPhiChange, dPhiChangeMin, dPhiChangeMax, - dAlphaInnerMDSegment, dAlphaOuterMDSegment, dAlphaInnerMDOuterMD; + float dPhi, dPhiMin, dPhiMax, dPhiChange, dPhiChangeMin, dPhiChangeMax; unsigned int innerMiniDoubletAnchorHitIndex = mdsInGPU.anchorHitIndices[innerMDIndex]; unsigned int outerMiniDoubletAnchorHitIndex = mdsInGPU.anchorHitIndices[outerMDIndex]; @@ -828,8 +748,6 @@ namespace lst { dPhiMax = 0; dPhiChangeMin = 0; dPhiChangeMax = 0; - float zLo, zHi, rtLo, rtHi, sdCut, dAlphaInnerMDSegmentThreshold, dAlphaOuterMDSegmentThreshold, - dAlphaInnerMDOuterMDThreshold; if (runSegmentDefaultAlgo(acc, modulesInGPU, mdsInGPU, @@ -837,27 +755,12 @@ namespace lst { outerLowerModuleIndex, innerMDIndex, outerMDIndex, - zIn, - zOut, - rtIn, - rtOut, dPhi, dPhiMin, dPhiMax, dPhiChange, dPhiChangeMin, - dPhiChangeMax, - dAlphaInnerMDSegment, - dAlphaOuterMDSegment, - dAlphaInnerMDOuterMD, - zLo, - zHi, - rtLo, - rtHi, - sdCut, - dAlphaInnerMDSegmentThreshold, - dAlphaOuterMDSegmentThreshold, - dAlphaInnerMDOuterMDThreshold)) { + dPhiChangeMax)) { unsigned int totOccupancySegments = alpaka::atomicOp( acc, &segmentsInGPU.totOccupancySegments[innerLowerModuleIndex], 1u); if (static_cast(totOccupancySegments) >= rangesInGPU.segmentModuleOccupancy[innerLowerModuleIndex]) { @@ -894,9 +797,9 @@ namespace lst { struct createSegmentArrayRanges { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::ObjectRanges rangesInGPU, - struct lst::MiniDoublets mdsInGPU) const { + lst::Modules modulesInGPU, + lst::ObjectRanges rangesInGPU, + lst::MiniDoublets mdsInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -992,9 +895,9 @@ namespace lst { struct addSegmentRangesToEventExplicit { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::Segments segmentsInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::Modules modulesInGPU, + lst::Segments segmentsInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -1013,11 +916,11 @@ namespace lst { struct addPixelSegmentToEventKernel { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::ObjectRanges rangesInGPU, - struct lst::Hits hitsInGPU, - struct lst::MiniDoublets mdsInGPU, - struct lst::Segments segmentsInGPU, + lst::Modules modulesInGPU, + lst::ObjectRanges rangesInGPU, + lst::Hits hitsInGPU, + lst::MiniDoublets mdsInGPU, + lst::Segments segmentsInGPU, unsigned int* hitIndices0, unsigned int* hitIndices1, unsigned int* hitIndices2, diff --git a/RecoTracker/LSTCore/src/alpaka/TrackCandidate.h b/RecoTracker/LSTCore/src/alpaka/TrackCandidate.h index 7451f5e47dc80..ede4dd9471e8e 100644 --- a/RecoTracker/LSTCore/src/alpaka/TrackCandidate.h +++ b/RecoTracker/LSTCore/src/alpaka/TrackCandidate.h @@ -109,7 +109,7 @@ namespace lst { inline void setData(TrackCandidatesBuffer& buf) { data_.setData(buf); } }; - ALPAKA_FN_ACC ALPAKA_FN_INLINE void addpLSTrackCandidateToMemory(struct lst::TrackCandidates& trackCandidatesInGPU, + ALPAKA_FN_ACC ALPAKA_FN_INLINE void addpLSTrackCandidateToMemory(lst::TrackCandidates& trackCandidatesInGPU, unsigned int trackletIndex, unsigned int trackCandidateIndex, uint4 hitIndices, @@ -128,7 +128,7 @@ namespace lst { trackCandidatesInGPU.hitIndices[Params_pT5::kHits * trackCandidateIndex + 3] = hitIndices.w; }; - ALPAKA_FN_ACC ALPAKA_FN_INLINE void addTrackCandidateToMemory(struct lst::TrackCandidates& trackCandidatesInGPU, + ALPAKA_FN_ACC ALPAKA_FN_INLINE void addTrackCandidateToMemory(lst::TrackCandidates& trackCandidatesInGPU, short trackCandidateType, unsigned int innerTrackletIndex, unsigned int outerTrackletIndex, @@ -167,9 +167,9 @@ namespace lst { ALPAKA_FN_ACC ALPAKA_FN_INLINE int checkPixelHits(unsigned int ix, unsigned int jx, - struct lst::MiniDoublets& mdsInGPU, - struct lst::Segments& segmentsInGPU, - struct lst::Hits& hitsInGPU) { + lst::MiniDoublets const& mdsInGPU, + lst::Segments const& segmentsInGPU, + lst::Hits const& hitsInGPU) { int phits1[Params_pLS::kHits]; int phits2[Params_pLS::kHits]; @@ -208,11 +208,11 @@ namespace lst { struct crossCleanpT3 { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::ObjectRanges rangesInGPU, - struct lst::PixelTriplets pixelTripletsInGPU, - struct lst::Segments segmentsInGPU, - struct lst::PixelQuintuplets pixelQuintupletsInGPU) const { + lst::Modules modulesInGPU, + lst::ObjectRanges rangesInGPU, + lst::PixelTriplets pixelTripletsInGPU, + lst::Segments segmentsInGPU, + lst::PixelQuintuplets pixelQuintupletsInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -249,11 +249,11 @@ namespace lst { struct crossCleanT5 { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::Quintuplets quintupletsInGPU, - struct lst::PixelQuintuplets pixelQuintupletsInGPU, - struct lst::PixelTriplets pixelTripletsInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::Modules modulesInGPU, + lst::Quintuplets quintupletsInGPU, + lst::PixelQuintuplets pixelQuintupletsInGPU, + lst::PixelTriplets pixelTripletsInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -301,19 +301,17 @@ namespace lst { } }; - // Using Matt's block for the outer loop and thread for inner loop trick here! - // This will eliminate the need for another kernel just for adding the pLS, because we can __syncthreads() struct crossCleanpLS { template ALPAKA_FN_ACC void operator()(TAcc const& acc, - struct lst::Modules modulesInGPU, - struct lst::ObjectRanges rangesInGPU, - struct lst::PixelTriplets pixelTripletsInGPU, - struct lst::TrackCandidates trackCandidatesInGPU, - struct lst::Segments segmentsInGPU, - struct lst::MiniDoublets mdsInGPU, - struct lst::Hits hitsInGPU, - struct lst::Quintuplets quintupletsInGPU) const { + lst::Modules modulesInGPU, + lst::ObjectRanges rangesInGPU, + lst::PixelTriplets pixelTripletsInGPU, + lst::TrackCandidates trackCandidatesInGPU, + lst::Segments segmentsInGPU, + lst::MiniDoublets mdsInGPU, + lst::Hits hitsInGPU, + lst::Quintuplets quintupletsInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -388,10 +386,10 @@ namespace lst { template ALPAKA_FN_ACC void operator()(TAcc const& acc, uint16_t nLowerModules, - struct lst::PixelTriplets pixelTripletsInGPU, - struct lst::TrackCandidates trackCandidatesInGPU, - struct lst::Segments segmentsInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::PixelTriplets pixelTripletsInGPU, + lst::TrackCandidates trackCandidatesInGPU, + lst::Segments segmentsInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -440,9 +438,9 @@ namespace lst { template ALPAKA_FN_ACC void operator()(TAcc const& acc, uint16_t nLowerModules, - struct lst::Quintuplets quintupletsInGPU, - struct lst::TrackCandidates trackCandidatesInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::Quintuplets quintupletsInGPU, + lst::TrackCandidates trackCandidatesInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -494,8 +492,8 @@ namespace lst { template ALPAKA_FN_ACC void operator()(TAcc const& acc, uint16_t nLowerModules, - struct lst::TrackCandidates trackCandidatesInGPU, - struct lst::Segments segmentsInGPU, + lst::TrackCandidates trackCandidatesInGPU, + lst::Segments segmentsInGPU, bool tc_pls_triplets) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc); @@ -533,10 +531,10 @@ namespace lst { template ALPAKA_FN_ACC void operator()(TAcc const& acc, uint16_t nLowerModules, - struct lst::PixelQuintuplets pixelQuintupletsInGPU, - struct lst::TrackCandidates trackCandidatesInGPU, - struct lst::Segments segmentsInGPU, - struct lst::ObjectRanges rangesInGPU) const { + lst::PixelQuintuplets pixelQuintupletsInGPU, + lst::TrackCandidates trackCandidatesInGPU, + lst::Segments segmentsInGPU, + lst::ObjectRanges rangesInGPU) const { auto const globalThreadIdx = alpaka::getIdx(acc); auto const gridThreadExtent = alpaka::getWorkDiv(acc);