Patatrack Shared Hit Cleaner Fix & Update

RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc

@@ -170,12 +170,17 @@ void CAHitNtupletGeneratorKernelsCPU::classifyTuples(HitsOnCPU const &hh, TkSoA
   kernel_fastDuplicateRemover(device_theCells_.get(), device_nCells_, tuples_d, tracks_d);
 
   // fill hit->track "map"
-  kernel_countHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
-  cms::cuda::launchFinalize(device_hitToTuple_.get(), cudaStream);
-  kernel_fillHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
+  if (params_.doSharedHitCut_ || params_.doStats_) {
+    kernel_countHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
+    cms::cuda::launchFinalize(device_hitToTuple_.get(), cudaStream);
+    kernel_fillHitInTracks(tuples_d, quality_d, device_hitToTuple_.get());
+  }
 
   // remove duplicates (tracks that share a hit)
-  kernel_tripletCleaner(hh.view(), tuples_d, tracks_d, quality_d, device_hitToTuple_.get());
+  if (params_.doSharedHitCut_) {
+    kernel_sharedHitCleaner(
+        hh.view(), tuples_d, tracks_d, quality_d, params_.minHitsForSharingCut_, device_hitToTuple_.get());
+  }
 
   if (params_.doStats_) {
     // counters (add flag???)

RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu

@@ -257,11 +257,12 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA
     kernel_fillHitInTracks<<<numberOfBlocks, blockSize, 0, cudaStream>>>(tuples_d, quality_d, device_hitToTuple_.get());
     cudaCheck(cudaGetLastError());
   }
-  if (params_.minHitsPerNtuplet_ < 4) {
+
+  if (params_.doSharedHitCut_) {
     // remove duplicates (tracks that share a hit)
     numberOfBlocks = (HitToTuple::capacity() + blockSize - 1) / blockSize;
-    kernel_tripletCleaner<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
-        hh.view(), tuples_d, tracks_d, quality_d, device_hitToTuple_.get());
+    kernel_sharedHitCleaner<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
+        hh.view(), tuples_d, tracks_d, quality_d, params_.minHitsForSharingCut_, device_hitToTuple_.get());
     cudaCheck(cudaGetLastError());
   }
 

RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.h

@@ -54,6 +54,7 @@ namespace cAHitNtupletGenerator {
     Params(bool onGPU,
            uint32_t minHitsPerNtuplet,
            uint32_t maxNumberOfDoublets,
+           uint16_t minHitsForSharingCuts,
            bool useRiemannFit,
            bool fit5as4,
            bool includeJumpingForwardDoublets,
@@ -64,16 +65,19 @@ namespace cAHitNtupletGenerator {
            bool doClusterCut,
            bool doZ0Cut,
            bool doPtCut,
+           bool doSharedHitCut,
            float ptmin,
            float CAThetaCutBarrel,
            float CAThetaCutForward,
            float hardCurvCut,
            float dcaCutInnerTriplet,
            float dcaCutOuterTriplet,
+
            QualityCuts const& cuts)
         : onGPU_(onGPU),
           minHitsPerNtuplet_(minHitsPerNtuplet),
           maxNumberOfDoublets_(maxNumberOfDoublets),
+          minHitsForSharingCut_(minHitsForSharingCuts),
           useRiemannFit_(useRiemannFit),
           fit5as4_(fit5as4),
           includeJumpingForwardDoublets_(includeJumpingForwardDoublets),
@@ -84,6 +88,7 @@ namespace cAHitNtupletGenerator {
           doClusterCut_(doClusterCut),
           doZ0Cut_(doZ0Cut),
           doPtCut_(doPtCut),
+          doSharedHitCut_(doSharedHitCut),
           ptmin_(ptmin),
           CAThetaCutBarrel_(CAThetaCutBarrel),
           CAThetaCutForward_(CAThetaCutForward),
@@ -95,6 +100,7 @@ namespace cAHitNtupletGenerator {
     const bool onGPU_;
     const uint32_t minHitsPerNtuplet_;
     const uint32_t maxNumberOfDoublets_;
+    const uint16_t minHitsForSharingCut_;
     const bool useRiemannFit_;
     const bool fit5as4_;
     const bool includeJumpingForwardDoublets_;
@@ -105,6 +111,7 @@ namespace cAHitNtupletGenerator {
     const bool doClusterCut_;
     const bool doZ0Cut_;
     const bool doPtCut_;
+    const bool doSharedHitCut_;
     const float ptmin_;
     const float CAThetaCutBarrel_;
     const float CAThetaCutForward_;

RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h

@@ -473,11 +473,12 @@ __global__ void kernel_doStatsForHitInTracks(CAHitNtupletGeneratorKernelsGPU::Hi
   }
 }
 
-__global__ void kernel_tripletCleaner(TrackingRecHit2DSOAView const *__restrict__ hhp,
-                                      HitContainer const *__restrict__ ptuples,
-                                      TkSoA const *__restrict__ ptracks,
-                                      Quality *__restrict__ quality,
-                                      CAHitNtupletGeneratorKernelsGPU::HitToTuple const *__restrict__ phitToTuple) {
+__global__ void kernel_sharedHitCleaner(TrackingRecHit2DSOAView const *__restrict__ hhp,
+                                        HitContainer const *__restrict__ ptuples,
+                                        TkSoA const *__restrict__ ptracks,
+                                        Quality *__restrict__ quality,
+                                        uint16_t nmin,
+                                        CAHitNtupletGeneratorKernelsGPU::HitToTuple const *__restrict__ phitToTuple) {
   constexpr auto bad = pixelTrack::Quality::bad;
   constexpr auto dup = pixelTrack::Quality::dup;
   // constexpr auto loose = trackQuality::loose;
@@ -486,6 +487,9 @@ __global__ void kernel_tripletCleaner(TrackingRecHit2DSOAView const *__restrict_
   auto const &foundNtuplets = *ptuples;
   auto const &tracks = *ptracks;
 
+  auto const &hh = *hhp;
+  int l1end = hh.hitsLayerStart()[1];
+
   int first = blockDim.x * blockIdx.x + threadIdx.x;
   for (int idx = first, ntot = hitToTuple.nbins(); idx < ntot; idx += gridDim.x * blockDim.x) {
     if (hitToTuple.size(idx) < 2)
@@ -503,6 +507,11 @@ __global__ void kernel_tripletCleaner(TrackingRecHit2DSOAView const *__restrict_
     // kill all tracks shorter than maxHn (only triplets???)
     for (auto it = hitToTuple.begin(idx); it != hitToTuple.end(idx); ++it) {
       uint32_t nh = foundNtuplets.size(*it);
+
+      //checking if shared hit is on bpix1 and if the tuple is short enough
+      if (idx < l1end and nh > nmin)
+        continue;
+
       if (maxNh != nh)
         quality[*it] = dup;
     }

RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc

@@ -57,6 +57,7 @@ CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU(const edm::ParameterSet&
     : m_params(cfg.getParameter<bool>("onGPU"),
                cfg.getParameter<unsigned int>("minHitsPerNtuplet"),
                cfg.getParameter<unsigned int>("maxNumberOfDoublets"),
+               cfg.getParameter<unsigned int>("minHitsForSharingCut"),
                cfg.getParameter<bool>("useRiemannFit"),
                cfg.getParameter<bool>("fit5as4"),
                cfg.getParameter<bool>("includeJumpingForwardDoublets"),
@@ -67,6 +68,7 @@ CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU(const edm::ParameterSet&
                cfg.getParameter<bool>("doClusterCut"),
                cfg.getParameter<bool>("doZ0Cut"),
                cfg.getParameter<bool>("doPtCut"),
+               cfg.getParameter<bool>("doSharedHitCut"),
                cfg.getParameter<double>("ptmin"),
                cfg.getParameter<double>("CAThetaCutBarrel"),
                cfg.getParameter<double>("CAThetaCutForward"),
@@ -141,12 +143,15 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::ParameterSetDescription&
   desc.add<bool>("fillStatistics", false);
   desc.add<unsigned int>("minHitsPerNtuplet", 4);
   desc.add<unsigned int>("maxNumberOfDoublets", caConstants::maxNumberOfDoublets);
+  desc.add<unsigned int>("minHitsForSharingCut", 5)
+      ->setComment("Maximum number of hits in a tuple to clean also if the shared hit is on bpx1");
   desc.add<bool>("includeJumpingForwardDoublets", false);
   desc.add<bool>("fit5as4", true);
   desc.add<bool>("doClusterCut", true);
   desc.add<bool>("doZ0Cut", true);
   desc.add<bool>("doPtCut", true);
   desc.add<bool>("useRiemannFit", false)->setComment("true for Riemann, false for BrokenLine");
+  desc.add<bool>("doSharedHitCut", true)->setComment("Sharing hit nTuples cleaning");
 
   edm::ParameterSetDescription trackQualityCuts;
   trackQualityCuts.add<double>("chi2MaxPt", 10.)->setComment("max pT used to determine the pT-dependent chi2 cut");