Skip to content
This repository was archived by the owner on Dec 9, 2024. It is now read-only.
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 1 addition & 7 deletions SDL/Hit.cuh
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));
Comment thread
VourMa marked this conversation as resolved.
}
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
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]);
Comment on lines +201 to +202
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should we be applying the phi function here now?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we can, using phi(x,y)

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
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
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
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