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 #281 from YonsiG/rebase_T5_rzcut
Browse files Browse the repository at this point in the history
Rebase t5 rz chi2 cut
  • Loading branch information
VourMa authored May 10, 2023
2 parents ef224c0 + 9d5b64a commit 4d0158d
Show file tree
Hide file tree
Showing 7 changed files with 383 additions and 132 deletions.
1 change: 1 addition & 0 deletions SDL/Constants.cu
Original file line number Diff line number Diff line change
Expand Up @@ -17,3 +17,4 @@ CUDA_CONST_VAR const float SDL::deltaZLum = 15.0;
CUDA_CONST_VAR const float SDL::pixelPSZpitch = 0.15;
CUDA_CONST_VAR const float SDL::strip2SZpitch = 5.0;
CUDA_CONST_VAR const float SDL::pt_betaMax = 7.0f;
CUDA_CONST_VAR const float SDL::magnetic_field = 3.8112;
1 change: 1 addition & 0 deletions SDL/Constants.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -73,5 +73,6 @@ namespace SDL
extern CUDA_CONST_VAR const float pixelPSZpitch;
extern CUDA_CONST_VAR const float strip2SZpitch;
extern CUDA_CONST_VAR const float pt_betaMax;
extern CUDA_CONST_VAR const float magnetic_field;
}
#endif
3 changes: 1 addition & 2 deletions SDL/Event.cu
Original file line number Diff line number Diff line change
Expand Up @@ -1959,10 +1959,9 @@ SDL::quintuplets* SDL::Event::getQuintuplets()
quintupletsInCPU->eta = new FPX[nMemoryLocations];
quintupletsInCPU->phi = new FPX[nMemoryLocations];

quintupletsInCPU->rzChiSquared = new float[nMemoryLocations];
quintupletsInCPU->chiSquared = new float[nMemoryLocations];
quintupletsInCPU->nonAnchorChiSquared = new float[nMemoryLocations];

quintupletsInCPU->rzChiSquared = new float[nMemoryLocations];

cudaMemcpyAsync(quintupletsInCPU->nQuintuplets, quintupletsInGPU->nQuintuplets, nLowerModules * sizeof(unsigned int), cudaMemcpyDeviceToHost,stream);
cudaMemcpyAsync(quintupletsInCPU->totOccupancyQuintuplets, quintupletsInGPU->totOccupancyQuintuplets, nLowerModules * sizeof(unsigned int), cudaMemcpyDeviceToHost,stream);
Expand Down
2 changes: 0 additions & 2 deletions SDL/PixelTriplet.cu
Original file line number Diff line number Diff line change
Expand Up @@ -474,10 +474,8 @@ __device__ bool SDL::passPT3RZChiSquaredCuts(struct modules& modulesInGPU, uint1

__device__ float SDL::computePT3RZChiSquared(struct modules& modulesInGPU, uint16_t* lowerModuleIndices, float* rtPix, float* xPix, float* yPix, float* zPix, float* rts, float* xs, float* ys, float* zs, float pixelSegmentPt, float pixelSegmentPx, float pixelSegmentPy, float pixelSegmentPz, int pixelSegmentCharge)
{
float slope = (zPix[1] - zPix[0])/(rtPix[1] - rtPix[0]);
float residual = 0;
float error = 0;
//hardcoded array indices!!!
float RMSE = 0;

float Px=pixelSegmentPx, Py=pixelSegmentPy, Pz=pixelSegmentPz;
Expand Down
495 changes: 372 additions & 123 deletions SDL/Quintuplet.cu

Large diffs are not rendered by default.

9 changes: 5 additions & 4 deletions SDL/Quintuplet.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -40,6 +40,7 @@ namespace SDL
FPX* score_rphisum;
uint8_t* layer;
bool* isDup;
bool* TightCutFlag;
bool* partOfPT5;

float* regressionRadius;
Expand All @@ -65,12 +66,12 @@ namespace SDL

// CUDA_DEV void rmQuintupletToMemory(struct SDL::quintuplets& quintupletsInGPU, unsigned int quintupletIndex);

CUDA_DEV void addQuintupletToMemory(struct SDL::triplets& tripletsInGPU, struct SDL::quintuplets& quintupletsInGPU, unsigned int innerTripletIndex, unsigned int outerTripletIndex, uint16_t& lowerModule1, uint16_t& lowerModule2, 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 pt, float eta, float phi, float scores, uint8_t layer, unsigned int quintupletIndex);
CUDA_DEV void addQuintupletToMemory(struct SDL::triplets& tripletsInGPU, struct SDL::quintuplets& quintupletsInGPU, unsigned int innerTripletIndex, unsigned int outerTripletIndex, uint16_t& lowerModule1, uint16_t& lowerModule2, 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 pt, float eta, float phi, float scores, uint8_t layer, unsigned int quintupletIndex, bool TightCutFlag);

CUDA_DEV bool runQuintupletDefaultAlgo(struct SDL::modules& modulesInGPU, struct SDL::miniDoublets& mdsInGPU, struct SDL::segments& segmentsInGPU, struct SDL::triplets& tripletsInGPU, uint16_t& lowerModuleIndex1, uint16_t& lowerModuleIndex2, uint16_t& lowerModuleIndex3, uint16_t& lowerModuleIndex4, uint16_t& lowerModuleIndex5, unsigned int& innerTripletIndex, unsigned int& outerTripletIndex, float& innerRadius, float& outerRadius, float& bridgeRadius, float& regressionG, float& regressionF, float& regressionRadius, float& rzChiSquared, float& chiSquared, float& nonAnchorChiSquared);
CUDA_DEV bool runQuintupletDefaultAlgo(struct SDL::modules& modulesInGPU, struct SDL::miniDoublets& mdsInGPU, struct SDL::segments& segmentsInGPU, struct SDL::triplets& tripletsInGPU, uint16_t& lowerModuleIndex1, uint16_t& lowerModuleIndex2, uint16_t& lowerModuleIndex3, uint16_t& lowerModuleIndex4, uint16_t& lowerModuleIndex5, unsigned int& innerTripletIndex, unsigned int& outerTripletIndex, float& innerRadius, float& outerRadius, float& bridgeRadius, float& regressionG, float& regressionF, float& regressionRadius, float& rzChiSquared, float& chiSquared, float& nonAnchorChiSquared, bool& TightCutFlag);


CUDA_DEV bool passT5RZConstraint(struct SDL::modules& modulesInGPU, struct SDL::miniDoublets& mdsInGPU, unsigned int firstMDIndex, unsigned int secondMDIndex, unsigned int thirdMDIndex, unsigned int fourthMDIndex, unsigned int fifthMDIndex, uint16_t& lowerModuleIndex1, uint16_t& lowerModuleIndex2, uint16_t& lowerModuleIndex3, uint16_t& lowerModuleIndex4, uint16_t& lowerModuleIndex5);
CUDA_DEV bool passT5RZConstraint(struct SDL::modules& modulesInGPU, struct SDL::miniDoublets& mdsInGPU, unsigned int firstMDIndex, unsigned int secondMDIndex, unsigned int thirdMDIndex, unsigned int fourthMDIndex, unsigned int fifthMDIndex, uint16_t& lowerModuleIndex1, uint16_t& lowerModuleIndex2, uint16_t& lowerModuleIndex3, uint16_t& lowerModuleIndex4, uint16_t& lowerModuleIndex5, float& rzChiSquared, float inner_pt, float innerRadius, float g, float f, bool& TightCutFlag);

CUDA_DEV bool T5HasCommonMiniDoublet(struct triplets& tripletsInGPU, struct segments& segmentsInGPU, unsigned int innerTripletIndex, unsigned int outerTripletIndex);

Expand All @@ -88,7 +89,7 @@ namespace SDL

CUDA_DEV bool matchRadiiBBBEE23478(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& innerRadiusMin, float& innerRadiusMax, float& bridgeRadiusMin, float& bridgeRadiusMax, float& outerRadiusMin, float& outerRadiusMax);

CUDA_DEV bool matchRadiiBBBEE34578(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& innerRadiusMin, float& innerRadiusMax, float& bridgeRadiusMin, float& bridgeRadiusMax, float& outerRadiusMin, float& outerRadiusMax);
CUDA_DEV bool matchRadiiBBBEE34578(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& innerRadiusMin, float& innerRadiusMax, float& bridgeRadiusMin, float& bridgeRadiusMax, float& outerRadiusMin, float& outerRadiusMax);


CUDA_DEV bool matchRadiiBBEEE(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& innerRadiusMin, float& innerRadiusMax, float& bridgeRadiusMin, float& bridgeRadiusMax, float& outerRadiusMin, float& outerRadiusMax);
Expand Down
4 changes: 3 additions & 1 deletion SDL/TrackCandidate.cu
Original file line number Diff line number Diff line change
Expand Up @@ -355,7 +355,9 @@ __global__ void SDL::addT5asTrackCandidateInGPU(struct SDL::modules& modulesInGP
{
int quintupletIndex = rangesInGPU.quintupletModuleIndices[idx] + jdx;

if(quintupletsInGPU.isDup[quintupletIndex] or quintupletsInGPU.partOfPT5[quintupletIndex]) continue;
if (quintupletsInGPU.isDup[quintupletIndex] or quintupletsInGPU.partOfPT5[quintupletIndex]) continue;

if (!(quintupletsInGPU.TightCutFlag[quintupletIndex])) continue;

unsigned int trackCandidateIdx = atomicAdd(trackCandidatesInGPU.nTrackCandidates,1);
atomicAdd(trackCandidatesInGPU.nTrackCandidatesT5,1);
Expand Down

0 comments on commit 4d0158d

Please sign in to comment.