Merge pull request #273 from tresreid/addObjects

VourMa · web-flow · commit fe2bae3f6986 · 2023-04-15T16:47:03.000+03:00
factorize addXXXObectsToEventExplicit into GPU and optional CPU functions
diff --git a/SDL/Event.cu b/SDL/Event.cu
diff --git a/SDL/Event.cuh b/SDL/Event.cuh
@@ -29,6 +29,7 @@ namespace SDL
     {
     private:
         cudaStream_t stream;
+        bool addObjects;
         std::array<unsigned int, 6> n_hits_by_layer_barrel_;
         std::array<unsigned int, 5> n_hits_by_layer_endcap_;
         std::array<unsigned int, 6> n_minidoublets_by_layer_barrel_;
@@ -72,21 +73,17 @@ namespace SDL
         int* superbinCPU;
         int8_t* pixelTypeCPU;
     public:
-        Event(cudaStream_t estream);
+        Event(cudaStream_t estream,bool verbose);
         ~Event();
         void resetEvent();
 
         void addHitToEvent(std::vector<float> x, std::vector<float> y, std::vector<float> z, std::vector<unsigned int> detId, std::vector<unsigned int> idxInNtuple); //call the appropriate hit function, then increment the counter here
         void addPixelSegmentToEvent(std::vector<unsigned int> hitIndices0,std::vector<unsigned int> hitIndices1,std::vector<unsigned int> hitIndices2,std::vector<unsigned int> hitIndices3, std::vector<float> dPhiChange, std::vector<float> ptIn, std::vector<float> ptErr, std::vector<float> px, std::vector<float> py, std::vector<float> pz, std::vector<float> eta, std::vector<float> etaErr, std::vector<float> phi, std::vector<int> charge, std::vector<unsigned int> seedIdx, std::vector<int> superbin, std::vector<int8_t> pixelType, std::vector<short> isQuad);
 
         /*functions that map the objects to the appropriate modules*/
-        void addMiniDoubletsToEvent();
-        void addSegmentsToEvent();
-        void addTripletsToEvent();
         void addMiniDoubletsToEventExplicit();
         void addSegmentsToEventExplicit();
         void addTripletsToEventExplicit();
-        void addQuintupletsToEvent();
         void addQuintupletsToEventExplicit();
         void resetObjectsInModule();
 
diff --git a/SDL/LST.cc b/SDL/LST.cc
@@ -10,7 +10,7 @@ void SDL::LST::eventSetup() {
   SDL::initModules(path);
 }
 
-void SDL::LST::run(cudaStream_t stream,
+void SDL::LST::run(cudaStream_t stream,bool verbose,
                    const std::vector<float> see_px,
                    const std::vector<float> see_py,
                    const std::vector<float> see_pz,
@@ -31,7 +31,7 @@ void SDL::LST::run(cudaStream_t stream,
                    const std::vector<float> ph2_x,
                    const std::vector<float> ph2_y,
                    const std::vector<float> ph2_z) {
-  auto event = SDL::Event(stream);
+  auto event = SDL::Event(stream,verbose);
   prepareInput(see_px, see_py, see_pz, see_dxy, see_dz, see_ptErr, see_etaErr, see_stateTrajGlbX, see_stateTrajGlbY, see_stateTrajGlbZ, see_stateTrajGlbPx, see_stateTrajGlbPy, see_stateTrajGlbPz, see_q, see_algo, see_hitIdx, ph2_detId, ph2_x, ph2_y, ph2_z);
 
   event.addHitToEvent(in_trkX_, in_trkY_, in_trkZ_, in_hitId_, in_hitIdxs_); // TODO : Need to fix the hitIdxs
diff --git a/SDL/LST.h b/SDL/LST.h
@@ -17,7 +17,7 @@ namespace SDL {
     LST();
 
     void eventSetup();
-    void run(cudaStream_t stream,
+    void run(cudaStream_t stream,bool verbose,
              const std::vector<float> see_px,
              const std::vector<float> see_py,
              const std::vector<float> see_pz,
diff --git a/SDL/Makefile b/SDL/Makefile
@@ -22,7 +22,7 @@ CXXFLAGS             =  -g --compiler-options -Wall --compiler-options -Wshadow
 ROOTCFLAGS           = --compiler-options -pthread --compiler-options -std=c++17 -m64 -I/cvmfs/cms.cern.ch/slc7_amd64_gcc900/cms/cmssw/CMSSW_11_2_0_pre5/external/slc7_amd64_gcc900/bin/../../../../../../../slc7_amd64_gcc900/lcg/root/6.20.06-ghbfee3/include
 LD                   = nvcc 
 SOFLAGS              = -g -shared --compiler-options -fPIC --cudart shared -arch=compute_70 -code=sm_72
-PRINTFLAG            = -DAddObjects -DT4FromT3 #-DWarnings
+PRINTFLAG            = -DT4FromT3 #-DWarnings
 DUPLICATES           = -DDUP_pLS -DDUP_T5 -DDUP_pT5 -DDUP_pT3 -DCrossclean_T5 -DCrossclean_pT3 -DFP16_Base -DFP16_dPhi
 MEMFLAG              =
 CACHEFLAG            =
diff --git a/SDL/MiniDoublet.cu b/SDL/MiniDoublet.cu
@@ -950,3 +950,21 @@ if(success)
         }
     }
 }
+__global__ void SDL::addMiniDoubletRangesToEventExplicit(struct modules& modulesInGPU, struct miniDoublets& mdsInGPU, struct objectRanges& rangesInGPU,struct hits& hitsInGPU)
+{
+    int gid = blockIdx.x * blockDim.x + threadIdx.x;
+    int np = gridDim.x * blockDim.x;
+    for(uint16_t i = gid; i < *modulesInGPU.nLowerModules; i+= np)
+    {
+        if(mdsInGPU.nMDs[i] == 0 or hitsInGPU.hitRanges[i * 2] == -1)
+        {
+            rangesInGPU.mdRanges[i * 2] = -1;
+            rangesInGPU.mdRanges[i * 2 + 1] = -1;
+        }
+        else
+        {
+            rangesInGPU.mdRanges[i * 2] = rangesInGPU.miniDoubletModuleIndices[i] ;
+            rangesInGPU.mdRanges[i * 2 + 1] = rangesInGPU.miniDoubletModuleIndices[i] + mdsInGPU.nMDs[i] - 1;
+        }
+    }
+}
diff --git a/SDL/MiniDoublet.cuh b/SDL/MiniDoublet.cuh
@@ -94,6 +94,7 @@ namespace SDL
 
     void createMDArrayRanges(struct modules& modulesInGPU, struct objectRanges& rangesInGPU, uint16_t& nLowerModules, unsigned int& nTotalMDs, cudaStream_t stream, const unsigned int& maxPixelMDs);
 
+    __global__ void addMiniDoubletRangesToEventExplicit(struct modules& modulesInGPU, struct miniDoublets& mdsInGPU, struct objectRanges& rangesInGPU, struct hits& hitsInGPU);
 
 //#ifdef CUT_VALUE_DEBUG
 //    CUDA_HOSTDEV void addMDToMemory(struct miniDoublets& mdsInGPU, struct hits& hitsInGPU, struct modules& modulesInGPU, unsigned int lowerHitIdx, unsigned int upperHitIdx, uint16_t& lowerModuleIdx, float dz, float drt, float dPhi, float dPhiChange, float shiftedX, float shiftedY, float shiftedZ, float noShiftedDz, float noShiftedDphi, float noShiftedDPhiChange, float dzCut, float drtCut, float miniCut, unsigned int idx);
diff --git a/SDL/Quintuplet.cu b/SDL/Quintuplet.cu
@@ -2068,3 +2068,21 @@ __device__ void SDL::runDeltaBetaIterationsT5(float& betaIn, float& betaOut, flo
 
     }
 }
+__global__ void SDL::addQuintupletRangesToEventExplicit(struct modules& modulesInGPU, struct quintuplets& quintupletsInGPU, struct objectRanges& rangesInGPU)
+{
+    int gid = blockIdx.x * blockDim.x + threadIdx.x;
+    int np = gridDim.x * blockDim.x;
+    for(uint16_t i = gid; i < *modulesInGPU.nLowerModules; i+= np)
+    {
+        if(quintupletsInGPU.nQuintuplets[i] == 0 or rangesInGPU.quintupletModuleIndices[i] == -1)
+        {
+            rangesInGPU.quintupletRanges[i * 2] = -1;
+            rangesInGPU.quintupletRanges[i * 2 + 1] = -1;
+        }
+       else
+        {
+            rangesInGPU.quintupletRanges[i * 2] = rangesInGPU.quintupletModuleIndices[i];
+            rangesInGPU.quintupletRanges[i * 2 + 1] = rangesInGPU.quintupletModuleIndices[i] + quintupletsInGPU.nQuintuplets[i] - 1;
+        }
+    }
+}
diff --git a/SDL/Quintuplet.cuh b/SDL/Quintuplet.cuh
@@ -61,6 +61,7 @@ namespace SDL
     void createQuintupletsInExplicitMemory(struct SDL::quintuplets& quintupletsInGPU, const unsigned int& maxQuintuplets, const uint16_t& nLowerModules, const uint16_t& nEligibleModules,cudaStream_t stream);
 
     __global__ void createEligibleModulesListForQuintupletsGPU(struct modules& modulesInGPU, struct triplets& tripletsInGPU, unsigned int* nTotalQuintuplets, struct objectRanges& rangesInGPU);
+    __global__ void addQuintupletRangesToEventExplicit(struct modules& modulesInGPU, struct quintuplets& quintupletsInGPU, struct objectRanges& rangesInGPU);
 
 //  CUDA_DEV void rmQuintupletToMemory(struct SDL::quintuplets& quintupletsInGPU, unsigned int quintupletIndex);
 
diff --git a/SDL/Segment.cu b/SDL/Segment.cu
@@ -869,3 +869,21 @@ __global__ void SDL::createSegmentsInGPUv2(struct SDL::modules& modulesInGPU, st
     }
     }
 }
+__global__ void SDL::addSegmentRangesToEventExplicit(struct modules& modulesInGPU, struct segments& segmentsInGPU, struct objectRanges& rangesInGPU)
+{
+    int gid = blockIdx.x * blockDim.x + threadIdx.x;
+    int np = gridDim.x * blockDim.x;
+    for(uint16_t i = gid; i < *modulesInGPU.nLowerModules; i+= np)
+    {
+        if(segmentsInGPU.nSegments[i] == 0)
+        {
+            rangesInGPU.segmentRanges[i * 2] = -1;
+            rangesInGPU.segmentRanges[i * 2 + 1] = -1;
+        }
+        else
+        {
+            rangesInGPU.segmentRanges[i * 2] = rangesInGPU.segmentModuleIndices[i];
+            rangesInGPU.segmentRanges[i * 2 + 1] = rangesInGPU.segmentModuleIndices[i] + segmentsInGPU.nSegments[i] - 1;
+        }
+    }
+}
diff --git a/SDL/Segment.cuh b/SDL/Segment.cuh
@@ -77,6 +77,8 @@ namespace SDL
     __global__ void createSegmentArrayRanges(struct modules& modulesInGPU, struct objectRanges& rangesInGPU, struct miniDoublets& mdsinGPU, unsigned int* nSegments);
 
 
+    __global__ void addSegmentRangesToEventExplicit(struct modules& modulesInGPU, struct segments& segmentsInGPU, struct objectRanges& rangesInGPU);
+    
     CUDA_DEV void dAlphaThreshold(float* dAlphaThresholdValues, struct modules& modulesInGPU, struct miniDoublets& 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, unsigned int& outerMDIndex);
     CUDA_DEV void addSegmentToMemory(struct 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, unsigned int idx);
 
diff --git a/SDL/Triplet.cu b/SDL/Triplet.cu
@@ -1442,3 +1442,21 @@ __device__ void SDL::runDeltaBetaIterationsT3(float& betaIn, float& betaOut, flo
 
     }
 }
+__global__ void SDL::addTripletRangesToEventExplicit(struct modules& modulesInGPU, struct triplets& tripletsInGPU, struct objectRanges& rangesInGPU)
+{
+    int gid = blockIdx.x * blockDim.x + threadIdx.x;
+    int np = gridDim.x * blockDim.x;
+    for(uint16_t i = gid; i < *modulesInGPU.nLowerModules; i+= np)
+    {
+        if(tripletsInGPU.nTriplets[i] == 0)
+        {
+            rangesInGPU.tripletRanges[i * 2] = -1;
+            rangesInGPU.tripletRanges[i * 2 + 1] = -1;
+        }
+        else
+        {
+            rangesInGPU.tripletRanges[i * 2] = rangesInGPU.tripletModuleIndices[i];
+            rangesInGPU.tripletRanges[i * 2 + 1] = rangesInGPU.tripletModuleIndices[i] +  tripletsInGPU.nTriplets[i] - 1;
+        }
+    }
+}
diff --git a/SDL/Triplet.cuh b/SDL/Triplet.cuh
@@ -74,6 +74,7 @@ namespace SDL
     };
 
     __global__ void createTripletArrayRanges(struct modules& modulesInGPU, struct objectRanges& rangesInGPU, struct segments& segmentsInGPU, unsigned int* nTotalTriplets);
+    __global__ void addTripletRangesToEventExplicit(struct modules& modulesInGPU, struct triplets& tripletsInGPU, struct objectRanges& rangesInGPU);
 
     void createTripletsInExplicitMemory(struct triplets& tripletsInGPU, unsigned int maxTriplets, uint16_t nLowerModules,cudaStream_t stream);
 #ifdef CUT_VALUE_DEBUG
diff --git a/bin/sdl.cc b/bin/sdl.cc
@@ -385,7 +385,7 @@ void run_sdl()
     {
 
         cudaStreamCreateWithFlags(&streams[s], cudaStreamNonBlocking);
-        SDL::Event *event = new SDL::Event(streams[s]);
+        SDL::Event *event = new SDL::Event(streams[s],ana.verbose>=2);
         ; //(streams[omp_get_thread_num()]);
         events.push_back(event);
     }
