diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc index 041254e78aa36..379aa5802f7dd 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc +++ b/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???) diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu index 179e73f3f23f8..160be9142ea4c 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu @@ -257,11 +257,12 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA kernel_fillHitInTracks<<>>(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<<>>( - hh.view(), tuples_d, tracks_d, quality_d, device_hitToTuple_.get()); + kernel_sharedHitCleaner<<>>( + hh.view(), tuples_d, tracks_d, quality_d, params_.minHitsForSharingCut_, device_hitToTuple_.get()); cudaCheck(cudaGetLastError()); } diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.h index d1a9f3d13a67f..f79119552e0bc 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.h +++ b/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_; diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h index 7c0cec51b8057..e55d7f124a9c5 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h +++ b/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; } diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc index c2c7c2b869752..5a86ece3590f6 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorOnGPU.cc @@ -57,6 +57,7 @@ CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU(const edm::ParameterSet& : m_params(cfg.getParameter("onGPU"), cfg.getParameter("minHitsPerNtuplet"), cfg.getParameter("maxNumberOfDoublets"), + cfg.getParameter("minHitsForSharingCut"), cfg.getParameter("useRiemannFit"), cfg.getParameter("fit5as4"), cfg.getParameter("includeJumpingForwardDoublets"), @@ -67,6 +68,7 @@ CAHitNtupletGeneratorOnGPU::CAHitNtupletGeneratorOnGPU(const edm::ParameterSet& cfg.getParameter("doClusterCut"), cfg.getParameter("doZ0Cut"), cfg.getParameter("doPtCut"), + cfg.getParameter("doSharedHitCut"), cfg.getParameter("ptmin"), cfg.getParameter("CAThetaCutBarrel"), cfg.getParameter("CAThetaCutForward"), @@ -141,12 +143,15 @@ void CAHitNtupletGeneratorOnGPU::fillDescriptions(edm::ParameterSetDescription& desc.add("fillStatistics", false); desc.add("minHitsPerNtuplet", 4); desc.add("maxNumberOfDoublets", caConstants::maxNumberOfDoublets); + desc.add("minHitsForSharingCut", 5) + ->setComment("Maximum number of hits in a tuple to clean also if the shared hit is on bpx1"); desc.add("includeJumpingForwardDoublets", false); desc.add("fit5as4", true); desc.add("doClusterCut", true); desc.add("doZ0Cut", true); desc.add("doPtCut", true); desc.add("useRiemannFit", false)->setComment("true for Riemann, false for BrokenLine"); + desc.add("doSharedHitCut", true)->setComment("Sharing hit nTuples cleaning"); edm::ParameterSetDescription trackQualityCuts; trackQualityCuts.add("chi2MaxPt", 10.)->setComment("max pT used to determine the pT-dependent chi2 cut");