Skip to content
This repository has been archived by the owner on Dec 9, 2024. It is now read-only.

Commit

Permalink
Merge pull request #296 from SegmentLinking/fix-dPhi
Browse files Browse the repository at this point in the history 
avoid the dPhi calculation in createTriplet
  • Loading branch information
@VourMa
VourMa authored Jul 8, 2023
2 parents 37d9d3c + 0acd6fc commit 6e6e2d0
Show file tree
Hide file tree
Showing 6 changed files with 48 additions and 53 deletions.
8 changes: 1 addition & 7 deletions SDL/Hit.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 );
Expand All @@ -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);
Expand Down
11 changes: 8 additions & 3 deletions SDL/MiniDoublet.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand All @@ -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));

Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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];
Expand Down Expand Up @@ -802,6 +805,8 @@ SDL::miniDoublets::miniDoublets()
anchorHighEdgeY = nullptr;
anchorLowEdgeX = nullptr;
anchorLowEdgeY = nullptr;
anchorHighEdgePhi = nullptr;
anchorLowEdgePhi = nullptr;
outerX = nullptr;
outerY = nullptr;
outerZ = nullptr;
Expand Down
2 changes: 2 additions & 0 deletions SDL/MiniDoublet.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,8 @@ namespace SDL
float* anchorHighEdgeY;
float* anchorLowEdgeX;
float* anchorLowEdgeY;
float* anchorLowEdgePhi;
float* anchorHighEdgePhi;

float* outerX;
float* outerY;
Expand Down
37 changes: 16 additions & 21 deletions SDL/Quintuplet.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand All @@ -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
Expand Down Expand Up @@ -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);
Expand All @@ -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]));
Expand All @@ -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;
Expand Down Expand Up @@ -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;
Expand All @@ -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]);
Expand All @@ -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;
Expand Down
10 changes: 5 additions & 5 deletions SDL/Segment.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand Down
Loading

0 comments on commit 6e6e2d0

Please sign in to comment.