From 1ab05a65170248b157aad08512148d0c3aec8c51 Mon Sep 17 00:00:00 2001 From: Vincenzo Innocente Date: Sat, 6 Jun 2020 10:10:37 +0200 Subject: [PATCH 1/3] increase maxhit, count forlon doublets --- .../Track/interface/PixelTrackHeterogeneous.h | 2 +- .../interface/TrackingRecHit2DSOAView.h | 2 +- .../PixelTriplets/plugins/CAConstants.h | 2 +- .../plugins/CAHitNtupletGeneratorKernels.cc | 1 + .../plugins/CAHitNtupletGeneratorKernels.cu | 34 +++++++++++-------- .../CAHitNtupletGeneratorKernelsImpl.h | 9 +++-- 6 files changed, 30 insertions(+), 20 deletions(-) diff --git a/CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h b/CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h index d462be2c5dd7b..e79a32c21daa0 100644 --- a/CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h +++ b/CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h @@ -16,7 +16,7 @@ class TrackSoAT { static constexpr int32_t stride() { return S; } using Quality = trackQuality::Quality; - using hindex_type = uint16_t; + using hindex_type = uint32_t; using HitContainer = cms::cuda::OneToManyAssoc; // Always check quality is at least loose! diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h index 808feb2a4218f..c78fd9de09117 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h @@ -15,7 +15,7 @@ namespace pixelCPEforGPU { class TrackingRecHit2DSOAView { public: static constexpr uint32_t maxHits() { return gpuClustering::MaxNumClusters; } - using hindex_type = uint16_t; // if above is <=2^16 + using hindex_type = uint32_t; // if above is <=2^16 using Hist = cms::cuda::HistoContainer; diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h index fce0c23596137..66a5fa59e0248 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h @@ -44,7 +44,7 @@ namespace CAConstants { constexpr uint32_t maxTuples() { return maxNumberOfTuples(); } // types - using hindex_type = uint16_t; // FIXME from siPixelRecHitsHeterogeneousProduct + using hindex_type = uint32_t; // FIXME from siPixelRecHitsHeterogeneousProduct using tindex_type = uint16_t; // for tuples #ifndef ONLY_PHICUT diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc index 4eafb6dccd31c..02c61301aa57f 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc @@ -126,6 +126,7 @@ void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA * if (m_params.doStats_) { kernel_checkOverflows(tuples_d, device_tupleMultiplicity_.get(), + device_hitToTuple_.get(), device_hitTuple_apc_, device_theCells_.get(), device_nCells_, diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu index 541ab5ed905f5..befc918dddb71 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu @@ -125,21 +125,6 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA * cudaCheck(cudaGetLastError()); } - if (m_params.doStats_) { - numberOfBlocks = (std::max(nhits, m_params.maxNumberOfDoublets_) + blockSize - 1) / blockSize; - kernel_checkOverflows<<>>(tuples_d, - device_tupleMultiplicity_.get(), - device_hitTuple_apc_, - device_theCells_.get(), - device_nCells_, - device_theCellNeighbors_, - device_theCellTracks_, - device_isOuterHitOfCell_.get(), - nhits, - m_params.maxNumberOfDoublets_, - counters_); - cudaCheck(cudaGetLastError()); - } #ifdef GPU_DEBUG cudaDeviceSynchronize(); cudaCheck(cudaGetLastError()); @@ -265,6 +250,25 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA cudaCheck(cudaGetLastError()); } + + if (m_params.doStats_) { + auto nhits = hh.nHits(); + numberOfBlocks = (std::max(nhits, m_params.maxNumberOfDoublets_) + blockSize - 1) / blockSize; + kernel_checkOverflows<<>>(tuples_d, + device_tupleMultiplicity_.get(), + device_hitToTuple_.get(), + device_hitTuple_apc_, + device_theCells_.get(), + device_nCells_, + device_theCellNeighbors_, + device_theCellTracks_, + device_isOuterHitOfCell_.get(), + nhits, + m_params.maxNumberOfDoublets_, + counters_); + cudaCheck(cudaGetLastError()); + } + if (m_params.doStats_) { // counters (add flag???) numberOfBlocks = (HitToTuple::capacity() + blockSize - 1) / blockSize; diff --git a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h index 654b37c076f99..16d02a27fa303 100644 --- a/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h +++ b/RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h @@ -30,7 +30,8 @@ using TkSoA = pixelTrack::TrackSoA; using HitContainer = pixelTrack::HitContainer; __global__ void kernel_checkOverflows(HitContainer const *foundNtuplets, - CAConstants::TupleMultiplicity *tupleMultiplicity, + CAConstants::TupleMultiplicity const * tupleMultiplicity, + CAHitNtupletGeneratorKernelsGPU::HitToTuple const * hitToTuple, cms::cuda::AtomicPairCounter *apc, GPUCACell const *__restrict__ cells, uint32_t const *__restrict__ nCells, @@ -91,7 +92,11 @@ __global__ void kernel_checkOverflows(HitContainer const *foundNtuplets, atomicAdd(&c.nKilledCells, 1); if (0 == thisCell.theUsed) atomicAdd(&c.nEmptyCells, 1); - if (thisCell.tracks().empty()) + // if (thisCell.tracks().empty()) + // atomicAdd(&c.nZeroTrackCells, 1); + if (0==hitToTuple->size(thisCell.get_inner_hit_id()) && + 0==hitToTuple->size(thisCell.get_outer_hit_id()) + ) atomicAdd(&c.nZeroTrackCells, 1); } From 21d4bb4d257e21932fad49d0c50168e854cf5d21 Mon Sep 17 00:00:00 2001 From: Vincenzo Innocente Date: Sun, 7 Jun 2020 10:20:43 +0200 Subject: [PATCH 2/3] fix hit index in phi binner --- .../TrackingRecHit/interface/TrackingRecHit2DSOAView.h | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h index c78fd9de09117..bcffa0e1e6d2c 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h @@ -15,10 +15,12 @@ namespace pixelCPEforGPU { class TrackingRecHit2DSOAView { public: static constexpr uint32_t maxHits() { return gpuClustering::MaxNumClusters; } - using hindex_type = uint32_t; // if above is <=2^16 + using hindex_type = uint32_t; // if above is <=2^32 - using Hist = - cms::cuda::HistoContainer; + using PhiBinner = + cms::cuda::HistoContainer; + + using Hist = PhiBinner; // FIXME using AverageGeometry = phase1PixelTopology::AverageGeometry; @@ -93,7 +95,7 @@ class TrackingRecHit2DSOAView { uint32_t* m_hitsLayerStart; - Hist* m_hist; + PhiBinner * m_hist; // FIXME use a more descriptive name consistently uint32_t m_nHits; }; From c3b3d8cde560e723eb6cc931e8759e6ca8484952 Mon Sep 17 00:00:00 2001 From: Andrea Bocci Date: Tue, 1 Sep 2020 12:34:03 +0200 Subject: [PATCH 3/3] Remove misleading calls to std::move (#546) --- CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc | 2 +- .../SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc index ef229be4b9910..075d408a6f6fc 100644 --- a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc +++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc @@ -38,5 +38,5 @@ SiPixelDigiErrorsCUDA::HostDataError SiPixelDigiErrorsCUDA::dataErrorToHostAsync } auto err = *error_h; err.set_data(data.get()); - return HostDataError(std::move(err), std::move(data)); + return HostDataError(err, std::move(data)); } diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc index 8817606043a60..c5b568750ad7d 100644 --- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc +++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc @@ -52,7 +52,7 @@ void SiPixelDigiErrorsSoAFromCUDA::acquire(const edm::Event& iEvent, const auto& gpuDigiErrors = ctx.get(iEvent, digiErrorGetToken_); auto tmp = gpuDigiErrors.dataErrorToHostAsync(ctx.stream()); - error_ = std::move(tmp.first); + error_ = tmp.first; data_ = std::move(tmp.second); formatterErrors_ = &(gpuDigiErrors.formatterErrors()); }