diff --git a/SDL/Hit.cuh b/SDL/Hit.cuh index 71f024d2..03a0ef3f 100644 --- a/SDL/Hit.cuh +++ b/SDL/Hit.cuh @@ -65,12 +65,6 @@ namespace SDL void addHitToMemory(struct hits& hitsInCPU,struct modules& modulesInGPU,float x, float y, float z, unsigned int detId, unsigned int idxInNtuple,cudaStream_t stream,struct objectRanges& rangesInGPU); CUDA_G void addHitToMemoryGPU(struct hits& hitsInCPU,struct modules& modulesInGPU,float x, float y, float z, unsigned int detId, unsigned int idxInNtuple,unsigned int moduleIndex, float phis,struct objectRanges& rangesInGPU); - CUDA_HOSTDEV inline float ATan2(float y, float x) { - if (x != 0) return atan2f(y, x); - if (y == 0) return 0; - if (y > 0) return float(M_PI) / 2.f; - else return -float(M_PI) / 2.f; - } CUDA_HOSTDEV inline float eta(float x, float y, float z) { float r3 = std::sqrt( x*x + y*y + z*z ); float rt = std::sqrt( x*x + y*y ); @@ -91,7 +85,7 @@ namespace SDL return x - n * float(2.f * float(M_PI)); } CUDA_HOSTDEV inline float phi(float x, float y) { - return phi_mpi_pi(float(M_PI) + ATan2(-y, -x)); + return phi_mpi_pi(float(M_PI) + atan2f(y,x)); } CUDA_HOSTDEV inline float deltaPhi(float x1, float y1, float x2, float y2) { float phi1 = phi(x1,y1); diff --git a/SDL/MiniDoublet.cu b/SDL/MiniDoublet.cu index 38a46fd2..45ad4abc 100644 --- a/SDL/MiniDoublet.cu +++ b/SDL/MiniDoublet.cu @@ -86,7 +86,7 @@ void SDL::createMDsInExplicitMemory(struct miniDoublets& mdsInGPU, unsigned int mdsInGPU.nMDs = (unsigned int*)cms::cuda::allocate_device(dev, (nLowerModules + 1) *sizeof(unsigned int),stream); mdsInGPU.totOccupancyMDs = (unsigned int*)cms::cuda::allocate_device(dev, (nLowerModules + 1) *sizeof(unsigned int),stream); mdsInGPU.anchorX = (float*)cms::cuda::allocate_device(dev, nMemoryLocations * 6 * sizeof(float), stream); - mdsInGPU.anchorHighEdgeX = (float*)cms::cuda::allocate_device(dev, nMemoryLocations * 4 * sizeof(float), stream); + mdsInGPU.anchorHighEdgeX = (float*)cms::cuda::allocate_device(dev, nMemoryLocations * 6 * sizeof(float), stream); mdsInGPU.outerX = (float*)cms::cuda::allocate_device(dev, nMemoryLocations * 6 * sizeof(float), stream); mdsInGPU.outerHighEdgeX = (float*)cms::cuda::allocate_device(dev, nMemoryLocations * 4 * sizeof(float), stream); @@ -99,7 +99,7 @@ void SDL::createMDsInExplicitMemory(struct miniDoublets& mdsInGPU, unsigned int cudaMalloc(&mdsInGPU.nMDs, (nLowerModules + 1) * sizeof(unsigned int)); cudaMalloc(&mdsInGPU.totOccupancyMDs, (nLowerModules + 1) * sizeof(unsigned int)); cudaMalloc(&mdsInGPU.anchorX, nMemoryLocations * 6 * sizeof(float)); - cudaMalloc(&mdsInGPU.anchorHighEdgeX, nMemoryLocations * 4 * sizeof(float)); + cudaMalloc(&mdsInGPU.anchorHighEdgeX, nMemoryLocations * 6 * sizeof(float)); cudaMalloc(&mdsInGPU.outerX, nMemoryLocations * 6 * sizeof(float)); cudaMalloc(&mdsInGPU.outerHighEdgeX, nMemoryLocations * 4 * sizeof(float)); @@ -129,6 +129,8 @@ void SDL::createMDsInExplicitMemory(struct miniDoublets& mdsInGPU, unsigned int mdsInGPU.anchorHighEdgeY = mdsInGPU.anchorHighEdgeX + nMemoryLocations; mdsInGPU.anchorLowEdgeX = mdsInGPU.anchorHighEdgeX + 2 * nMemoryLocations; mdsInGPU.anchorLowEdgeY = mdsInGPU.anchorHighEdgeX + 3 * nMemoryLocations; + mdsInGPU.anchorHighEdgePhi = mdsInGPU.anchorHighEdgeX + 4 * nMemoryLocations; + mdsInGPU.anchorLowEdgePhi = mdsInGPU.anchorHighEdgeX + 5 * nMemoryLocations; mdsInGPU.outerY = mdsInGPU.outerX + nMemoryLocations; mdsInGPU.outerZ = mdsInGPU.outerX + 2 * nMemoryLocations; @@ -196,7 +198,8 @@ __device__ void SDL::addMDToMemory(struct miniDoublets& mdsInGPU, struct hits& h mdsInGPU.anchorHighEdgeY[idx] = hitsInGPU.highEdgeYs[anchorHitIndex]; mdsInGPU.anchorLowEdgeX[idx] = hitsInGPU.lowEdgeXs[anchorHitIndex]; mdsInGPU.anchorLowEdgeY[idx] = hitsInGPU.lowEdgeYs[anchorHitIndex]; - + mdsInGPU.anchorHighEdgePhi[idx] = atan2f(mdsInGPU.anchorHighEdgeY[idx], mdsInGPU.anchorHighEdgeX[idx]); + mdsInGPU.anchorLowEdgePhi[idx] = atan2f(mdsInGPU.anchorLowEdgeY[idx], mdsInGPU.anchorLowEdgeX[idx]); mdsInGPU.outerX[idx] = hitsInGPU.xs[outerHitIndex]; mdsInGPU.outerY[idx] = hitsInGPU.ys[outerHitIndex]; mdsInGPU.outerZ[idx] = hitsInGPU.zs[outerHitIndex]; @@ -802,6 +805,8 @@ SDL::miniDoublets::miniDoublets() anchorHighEdgeY = nullptr; anchorLowEdgeX = nullptr; anchorLowEdgeY = nullptr; + anchorHighEdgePhi = nullptr; + anchorLowEdgePhi = nullptr; outerX = nullptr; outerY = nullptr; outerZ = nullptr; diff --git a/SDL/MiniDoublet.cuh b/SDL/MiniDoublet.cuh index 97e8b75c..903faf3f 100644 --- a/SDL/MiniDoublet.cuh +++ b/SDL/MiniDoublet.cuh @@ -59,6 +59,8 @@ namespace SDL float* anchorHighEdgeY; float* anchorLowEdgeX; float* anchorLowEdgeY; + float* anchorLowEdgePhi; + float* anchorHighEdgePhi; float* outerX; float* outerY; diff --git a/SDL/Quintuplet.cu b/SDL/Quintuplet.cu index 24e28295..277d85ee 100644 --- a/SDL/Quintuplet.cu +++ b/SDL/Quintuplet.cu @@ -1758,7 +1758,7 @@ __device__ bool SDL::runQuintupletDefaultAlgoBBBB(struct SDL::modules& modulesIn float sdlPVoff = 0.1f/rt_OutLo; sdlCut = alpha1GeV_OutLo + sqrtf(sdlMuls * sdlMuls + sdlPVoff * sdlPVoff); - deltaPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[secondMDIndex], mdsInGPU.anchorY[secondMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); + deltaPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[secondMDIndex]); // Cut #3: FIXME:deltaPhiPos can be tighter pass = pass and (fabsf(deltaPhiPos) <= sdlCut); if(not pass) return pass; @@ -1786,7 +1786,7 @@ __device__ bool SDL::runQuintupletDefaultAlgoBBBB(struct SDL::modules& modulesIn float alpha_OutUp,alpha_OutUp_highEdge,alpha_OutUp_lowEdge; - alpha_OutUp = SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); + alpha_OutUp = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorPhi[fourthMDIndex]); alpha_OutUp_highEdge = alpha_OutUp; alpha_OutUp_lowEdge = alpha_OutUp; @@ -1799,28 +1799,27 @@ __device__ bool SDL::runQuintupletDefaultAlgoBBBB(struct SDL::modules& modulesIn float tl_axis_lowEdge_x = tl_axis_x; float tl_axis_lowEdge_y = tl_axis_y; - betaIn = alpha_InLo - SDL::deltaPhi(mdsInGPU.anchorX[firstMDIndex], mdsInGPU.anchorY[firstMDIndex], tl_axis_x, tl_axis_y); + betaIn = alpha_InLo - SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[firstMDIndex]); float betaInRHmin = betaIn; float betaInRHmax = betaIn; - betaOut = -alpha_OutUp + SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], tl_axis_x, tl_axis_y); + betaOut = -alpha_OutUp + SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut; float betaOutRHmax = betaOut; if(isEC_lastLayer) { - alpha_OutUp_highEdge = SDL::deltaPhi(mdsInGPU.anchorHighEdgeX[fourthMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex], mdsInGPU.anchorHighEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); - alpha_OutUp_lowEdge = SDL::deltaPhi(mdsInGPU.anchorLowEdgeX[fourthMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex], mdsInGPU.anchorLowEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); + alpha_OutUp_highEdge = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorHighEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorHighEdgePhi[fourthMDIndex]); + alpha_OutUp_lowEdge = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorLowEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorLowEdgePhi[fourthMDIndex]); tl_axis_highEdge_x = mdsInGPU.anchorHighEdgeX[fourthMDIndex] - mdsInGPU.anchorX[firstMDIndex]; tl_axis_highEdge_y = mdsInGPU.anchorHighEdgeY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; tl_axis_lowEdge_x = mdsInGPU.anchorLowEdgeX[fourthMDIndex] - mdsInGPU.anchorX[firstMDIndex]; tl_axis_lowEdge_y = mdsInGPU.anchorLowEdgeY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; - - betaOutRHmin = -alpha_OutUp_highEdge + SDL::deltaPhi(mdsInGPU.anchorHighEdgeX[fourthMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex], tl_axis_highEdge_x, tl_axis_highEdge_y); - betaOutRHmax = -alpha_OutUp_lowEdge + SDL::deltaPhi(mdsInGPU.anchorLowEdgeX[fourthMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex], tl_axis_lowEdge_x, tl_axis_lowEdge_y); + betaOutRHmin = -alpha_OutUp_highEdge + SDL::phi_mpi_pi(SDL::phi(tl_axis_highEdge_x, tl_axis_highEdge_y)-mdsInGPU.anchorHighEdgePhi[fourthMDIndex]); + betaOutRHmax = -alpha_OutUp_lowEdge + SDL::phi_mpi_pi(SDL::phi(tl_axis_lowEdge_x, tl_axis_lowEdge_y)-mdsInGPU.anchorLowEdgePhi[fourthMDIndex]); } //beta computation @@ -1975,9 +1974,7 @@ __device__ bool SDL::runQuintupletDefaultAlgoBBEE(struct SDL::modules& modulesIn const float sdlPVoff = 0.1f / rt_OutLo; sdlCut = alpha1GeV_OutLo + sqrtf(sdlMuls * sdlMuls + sdlPVoff*sdlPVoff); - - deltaPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[secondMDIndex], mdsInGPU.anchorY[secondMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); - + deltaPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[secondMDIndex]); //Cut #4: deltaPhiPos can be tighter pass = pass and (fabsf(deltaPhiPos) <= sdlCut); @@ -2000,7 +1997,7 @@ __device__ bool SDL::runQuintupletDefaultAlgoBBEE(struct SDL::modules& modulesIn float sdIn_alpha_max = __H2F(segmentsInGPU.dPhiChangeMaxs[innerSegmentIndex]); float sdOut_alpha = sdIn_alpha; //weird - float sdOut_alphaOut = SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); + float sdOut_alphaOut = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorPhi[fourthMDIndex]); float sdOut_alphaOut_min = SDL::phi_mpi_pi(__H2F(segmentsInGPU.dPhiChangeMins[outerSegmentIndex]) - __H2F(segmentsInGPU.dPhiMins[outerSegmentIndex])); float sdOut_alphaOut_max = SDL::phi_mpi_pi(__H2F(segmentsInGPU.dPhiChangeMaxs[outerSegmentIndex]) - __H2F(segmentsInGPU.dPhiMaxs[outerSegmentIndex])); @@ -2009,11 +2006,11 @@ __device__ bool SDL::runQuintupletDefaultAlgoBBEE(struct SDL::modules& modulesIn float tl_axis_y = mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; float tl_axis_z = mdsInGPU.anchorZ[fourthMDIndex] - mdsInGPU.anchorZ[firstMDIndex]; - betaIn = sdIn_alpha - SDL::deltaPhi(mdsInGPU.anchorX[firstMDIndex], mdsInGPU.anchorY[firstMDIndex], tl_axis_x, tl_axis_y); + betaIn = sdIn_alpha - SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[firstMDIndex]); float betaInRHmin = betaIn; float betaInRHmax = betaIn; - betaOut = -sdOut_alphaOut + SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], tl_axis_x, tl_axis_y); + betaOut = -sdOut_alphaOut + SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut; float betaOutRHmax = betaOut; @@ -2196,7 +2193,7 @@ __device__ bool SDL::runQuintupletDefaultAlgoEEEE(struct SDL::modules& modulesIn float sdlPVoff = 0.1f/rtOut; sdlCut = alpha1GeV_OutLo + sqrtf(sdlMuls * sdlMuls + sdlPVoff * sdlPVoff); - deltaPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[secondMDIndex], mdsInGPU.anchorY[secondMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); + deltaPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[secondMDIndex]); pass = pass and (fabsf(deltaPhiPos) <= sdlCut); if(not pass) return pass; @@ -2216,8 +2213,7 @@ __device__ bool SDL::runQuintupletDefaultAlgoEEEE(struct SDL::modules& modulesIn float sdIn_alpha = __H2F(segmentsInGPU.dPhiChanges[innerSegmentIndex]); float sdOut_alpha = sdIn_alpha; //weird - float sdOut_dPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[thirdMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); - + float sdOut_dPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[thirdMDIndex]); float sdOut_dPhiChange = __H2F(segmentsInGPU.dPhiChanges[outerSegmentIndex]); float sdOut_dPhiChange_min = __H2F(segmentsInGPU.dPhiChangeMins[outerSegmentIndex]); float sdOut_dPhiChange_max = __H2F(segmentsInGPU.dPhiChangeMaxs[outerSegmentIndex]); @@ -2230,14 +2226,13 @@ __device__ bool SDL::runQuintupletDefaultAlgoEEEE(struct SDL::modules& modulesIn float tl_axis_y = mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; float tl_axis_z = mdsInGPU.anchorZ[fourthMDIndex] - mdsInGPU.anchorZ[firstMDIndex]; - betaIn = sdIn_alpha - SDL::deltaPhi(mdsInGPU.anchorX[firstMDIndex], mdsInGPU.anchorY[firstMDIndex], tl_axis_x, tl_axis_y); - + betaIn = sdIn_alpha - SDL::phi_mpi_pi(SDL::phi(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 = -sdOut_alphaOut + SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], tl_axis_x, tl_axis_y); + betaOut = -sdOut_alphaOut + SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut - sdOut_alphaOutRHmin + sdOut_alphaOut; float betaOutRHmax = betaOut - sdOut_alphaOutRHmax + sdOut_alphaOut; diff --git a/SDL/Segment.cu b/SDL/Segment.cu index c9ccfe7e..d8f53cc3 100644 --- a/SDL/Segment.cu +++ b/SDL/Segment.cu @@ -459,13 +459,13 @@ __device__ bool SDL::runSegmentDefaultAlgoEndcap(struct modules& modulesInGPU, s pass = pass and ((rtOut >= rtLo) & (rtOut <= rtHi)); if(not pass) return pass; - dPhi = deltaPhi(xIn, yIn, xOut, yOut); + dPhi = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[outerMDIndex]-mdsInGPU.anchorPhi[innerMDIndex]); sdCut = sdSlope; if(outerLayerEndcapTwoS) { - float dPhiPos_high = deltaPhi(xIn, yIn, xOutHigh, yOutHigh); - float dPhiPos_low = deltaPhi(xIn, yIn, xOutLow, yOutLow); + float dPhiPos_high = SDL::phi_mpi_pi(mdsInGPU.anchorHighEdgePhi[outerMDIndex]-mdsInGPU.anchorPhi[innerMDIndex]); + float dPhiPos_low = SDL::phi_mpi_pi(mdsInGPU.anchorLowEdgePhi[outerMDIndex]-mdsInGPU.anchorPhi[innerMDIndex]); dPhiMax = fabsf(dPhiPos_high) > fabsf(dPhiPos_low) ? dPhiPos_high : dPhiPos_low; dPhiMin = fabsf(dPhiPos_high) > fabsf(dPhiPos_low) ? dPhiPos_low : dPhiPos_high; @@ -543,12 +543,12 @@ __device__ bool SDL::runSegmentDefaultAlgoBarrel(struct modules& modulesInGPU, s sdCut = sdSlope + sqrtf(sdMuls * sdMuls + sdPVoff * sdPVoff); - dPhi = deltaPhi(xIn, yIn, xOut, yOut); + dPhi = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[outerMDIndex]-mdsInGPU.anchorPhi[innerMDIndex]); pass = pass and (fabsf(dPhi) <= sdCut); if(not pass) return pass; - dPhiChange = deltaPhiChange(xIn, yIn, xOut, yOut); + dPhiChange = SDL::phi_mpi_pi(SDL::phi(xOut-xIn, yOut-yIn)-mdsInGPU.anchorPhi[innerMDIndex]); pass = pass and (fabsf(dPhiChange) <= sdCut); if(not pass) return pass; diff --git a/SDL/Triplet.cu b/SDL/Triplet.cu index 7ab4ecf8..b69f9443 100644 --- a/SDL/Triplet.cu +++ b/SDL/Triplet.cu @@ -768,7 +768,7 @@ __device__ bool SDL::runTripletDefaultAlgoBBBB(struct SDL::modules& modulesInGPU float sdlPVoff = 0.1f/rt_OutLo; sdlCut = alpha1GeV_OutLo + sqrtf(sdlMuls * sdlMuls + sdlPVoff * sdlPVoff); - deltaPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[secondMDIndex], mdsInGPU.anchorY[secondMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); + deltaPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[secondMDIndex]); // Cut #3: FIXME:deltaPhiPos can be tighter pass = pass and (fabsf(deltaPhiPos) <= sdlCut); if(not pass) return pass; @@ -796,7 +796,7 @@ __device__ bool SDL::runTripletDefaultAlgoBBBB(struct SDL::modules& modulesInGPU float alpha_OutUp,alpha_OutUp_highEdge,alpha_OutUp_lowEdge; - alpha_OutUp = SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); + alpha_OutUp = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorPhi[fourthMDIndex]); alpha_OutUp_highEdge = alpha_OutUp; alpha_OutUp_lowEdge = alpha_OutUp; @@ -809,19 +809,19 @@ __device__ bool SDL::runTripletDefaultAlgoBBBB(struct SDL::modules& modulesInGPU float tl_axis_lowEdge_x = tl_axis_x; float tl_axis_lowEdge_y = tl_axis_y; - betaIn = alpha_InLo - SDL::deltaPhi(mdsInGPU.anchorX[firstMDIndex], mdsInGPU.anchorY[firstMDIndex], tl_axis_x, tl_axis_y); + betaIn = alpha_InLo - SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[firstMDIndex]); float betaInRHmin = betaIn; float betaInRHmax = betaIn; - betaOut = -alpha_OutUp + SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], tl_axis_x, tl_axis_y); + betaOut = -alpha_OutUp + SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut; float betaOutRHmax = betaOut; if(isEC_lastLayer) { - alpha_OutUp_highEdge = SDL::deltaPhi(mdsInGPU.anchorHighEdgeX[fourthMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex], mdsInGPU.anchorHighEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); - alpha_OutUp_lowEdge = SDL::deltaPhi(mdsInGPU.anchorLowEdgeX[fourthMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex], mdsInGPU.anchorLowEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); + alpha_OutUp_highEdge = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorHighEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorHighEdgePhi[fourthMDIndex]); + alpha_OutUp_lowEdge = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorLowEdgeX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorLowEdgePhi[fourthMDIndex]); tl_axis_highEdge_x = mdsInGPU.anchorHighEdgeX[fourthMDIndex] - mdsInGPU.anchorX[firstMDIndex]; tl_axis_highEdge_y = mdsInGPU.anchorHighEdgeY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; @@ -829,8 +829,8 @@ __device__ bool SDL::runTripletDefaultAlgoBBBB(struct SDL::modules& modulesInGPU tl_axis_lowEdge_y = mdsInGPU.anchorLowEdgeY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; - betaOutRHmin = -alpha_OutUp_highEdge + SDL::deltaPhi(mdsInGPU.anchorHighEdgeX[fourthMDIndex], mdsInGPU.anchorHighEdgeY[fourthMDIndex], tl_axis_highEdge_x, tl_axis_highEdge_y); - betaOutRHmax = -alpha_OutUp_lowEdge + SDL::deltaPhi(mdsInGPU.anchorLowEdgeX[fourthMDIndex], mdsInGPU.anchorLowEdgeY[fourthMDIndex], tl_axis_lowEdge_x, tl_axis_lowEdge_y); + betaOutRHmin = -alpha_OutUp_highEdge + SDL::phi_mpi_pi(SDL::phi(tl_axis_highEdge_x, tl_axis_highEdge_y)-mdsInGPU.anchorHighEdgePhi[fourthMDIndex]); + betaOutRHmax = -alpha_OutUp_lowEdge + SDL::phi_mpi_pi(SDL::phi(tl_axis_lowEdge_x, tl_axis_lowEdge_y)-mdsInGPU.anchorLowEdgePhi[fourthMDIndex]); } //beta computation @@ -986,8 +986,7 @@ __device__ bool SDL::runTripletDefaultAlgoBBEE(struct SDL::modules& modulesInGPU sdlCut = alpha1GeV_OutLo + sqrtf(sdlMuls * sdlMuls + sdlPVoff*sdlPVoff); - deltaPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[secondMDIndex], mdsInGPU.anchorY[secondMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); - + deltaPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[secondMDIndex]); //Cut #4: deltaPhiPos can be tighter pass = pass and (fabsf(deltaPhiPos) <= sdlCut); @@ -1010,7 +1009,7 @@ __device__ bool SDL::runTripletDefaultAlgoBBEE(struct SDL::modules& modulesInGPU float sdIn_alpha_max = __H2F(segmentsInGPU.dPhiChangeMaxs[innerSegmentIndex]); float sdOut_alpha = sdIn_alpha; //weird - float sdOut_alphaOut = SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex]); + float sdOut_alphaOut = SDL::phi_mpi_pi(SDL::phi(mdsInGPU.anchorX[fourthMDIndex] - mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[thirdMDIndex])-mdsInGPU.anchorPhi[fourthMDIndex]); float sdOut_alphaOut_min = SDL::phi_mpi_pi(__H2F(segmentsInGPU.dPhiChangeMins[outerSegmentIndex]) - __H2F(segmentsInGPU.dPhiMins[outerSegmentIndex])); float sdOut_alphaOut_max = SDL::phi_mpi_pi(__H2F(segmentsInGPU.dPhiChangeMaxs[outerSegmentIndex]) - __H2F(segmentsInGPU.dPhiMaxs[outerSegmentIndex])); @@ -1019,11 +1018,11 @@ __device__ bool SDL::runTripletDefaultAlgoBBEE(struct SDL::modules& modulesInGPU float tl_axis_y = mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; float tl_axis_z = mdsInGPU.anchorZ[fourthMDIndex] - mdsInGPU.anchorZ[firstMDIndex]; - betaIn = sdIn_alpha - SDL::deltaPhi(mdsInGPU.anchorX[firstMDIndex], mdsInGPU.anchorY[firstMDIndex], tl_axis_x, tl_axis_y); + betaIn = sdIn_alpha - SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[firstMDIndex]); float betaInRHmin = betaIn; float betaInRHmax = betaIn; - betaOut = -sdOut_alphaOut + SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], tl_axis_x, tl_axis_y); + betaOut = -sdOut_alphaOut + SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut; float betaOutRHmax = betaOut; @@ -1206,7 +1205,7 @@ __device__ bool SDL::runTripletDefaultAlgoEEEE(struct SDL::modules& modulesInGPU float sdlPVoff = 0.1f/rtOut; sdlCut = alpha1GeV_OutLo + sqrtf(sdlMuls * sdlMuls + sdlPVoff * sdlPVoff); - deltaPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[secondMDIndex], mdsInGPU.anchorY[secondMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); + deltaPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[secondMDIndex]); pass = pass and (fabsf(deltaPhiPos) <= sdlCut); if(not pass) return pass; @@ -1226,7 +1225,7 @@ __device__ bool SDL::runTripletDefaultAlgoEEEE(struct SDL::modules& modulesInGPU float sdIn_alpha = __H2F(segmentsInGPU.dPhiChanges[innerSegmentIndex]); float sdOut_alpha = sdIn_alpha; //weird - float sdOut_dPhiPos = SDL::deltaPhi(mdsInGPU.anchorX[thirdMDIndex], mdsInGPU.anchorY[thirdMDIndex], mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex]); + float sdOut_dPhiPos = SDL::phi_mpi_pi(mdsInGPU.anchorPhi[fourthMDIndex]-mdsInGPU.anchorPhi[thirdMDIndex]); float sdOut_dPhiChange = __H2F(segmentsInGPU.dPhiChanges[outerSegmentIndex]); float sdOut_dPhiChange_min = __H2F(segmentsInGPU.dPhiChangeMins[outerSegmentIndex]); @@ -1240,14 +1239,14 @@ __device__ bool SDL::runTripletDefaultAlgoEEEE(struct SDL::modules& modulesInGPU float tl_axis_y = mdsInGPU.anchorY[fourthMDIndex] - mdsInGPU.anchorY[firstMDIndex]; float tl_axis_z = mdsInGPU.anchorZ[fourthMDIndex] - mdsInGPU.anchorZ[firstMDIndex]; - betaIn = sdIn_alpha - SDL::deltaPhi(mdsInGPU.anchorX[firstMDIndex], mdsInGPU.anchorY[firstMDIndex], tl_axis_x, tl_axis_y); + betaIn = sdIn_alpha - SDL::phi_mpi_pi(SDL::phi(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 = -sdOut_alphaOut + SDL::deltaPhi(mdsInGPU.anchorX[fourthMDIndex], mdsInGPU.anchorY[fourthMDIndex], tl_axis_x, tl_axis_y); + betaOut = -sdOut_alphaOut + SDL::phi_mpi_pi(SDL::phi(tl_axis_x, tl_axis_y)-mdsInGPU.anchorPhi[fourthMDIndex]); float betaOutRHmin = betaOut - sdOut_alphaOutRHmin + sdOut_alphaOut; float betaOutRHmax = betaOut - sdOut_alphaOutRHmax + sdOut_alphaOut;