From f45232edcf9a666762566e0fa90fe082bb9ec876 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 20 Nov 2020 15:24:27 +0100 Subject: [PATCH 01/12] Remove SiPixelDigiHeterogeneousConverter as obsolete Should have been removed as part of cms-patatrack#100. --- .../SiPixelDigiHeterogeneousConverter.cc | 95 ------------------- 1 file changed, 95 deletions(-) delete mode 100644 RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigiHeterogeneousConverter.cc diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigiHeterogeneousConverter.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigiHeterogeneousConverter.cc deleted file mode 100644 index 568e9272468a7..0000000000000 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigiHeterogeneousConverter.cc +++ /dev/null @@ -1,95 +0,0 @@ -#include "FWCore/Framework/interface/Event.h" -#include "FWCore/Framework/interface/MakerMacros.h" -#include "FWCore/Framework/interface/global/EDProducer.h" -#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h" -#include "FWCore/ParameterSet/interface/ParameterSetDescription.h" -#include "FWCore/ParameterSet/interface/ParameterSet.h" - -#include "DataFormats/Common/interface/DetSetVector.h" -#include "DataFormats/DetId/interface/DetIdCollection.h" -#include "DataFormats/SiPixelDetId/interface/PixelFEDChannel.h" -#include "DataFormats/SiPixelDigi/interface/PixelDigi.h" -#include "DataFormats/SiPixelRawData/interface/SiPixelRawDataError.h" - -/** - * This is very stupid but currently the easiest way to go forward as - * one can't replace and EDProducer with an EDAlias in the - * configuration... - */ - -class SiPixelDigiHeterogeneousConverter : public edm::global::EDProducer<> { -public: - explicit SiPixelDigiHeterogeneousConverter(edm::ParameterSet const& iConfig); - ~SiPixelDigiHeterogeneousConverter() override = default; - - static void fillDescriptions(edm::ConfigurationDescriptions& descriptions); - -private: - void produce(edm::StreamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const override; - - edm::EDGetTokenT> token_collection_; - edm::EDGetTokenT> token_errorcollection_; - edm::EDGetTokenT token_tkerror_detidcollection_; - edm::EDGetTokenT token_usererror_detidcollection_; - edm::EDGetTokenT> token_disabled_channelcollection_; - bool includeErrors_; -}; - -SiPixelDigiHeterogeneousConverter::SiPixelDigiHeterogeneousConverter(edm::ParameterSet const& iConfig) - : includeErrors_(iConfig.getParameter("includeErrors")) { - auto src = iConfig.getParameter("src"); - - token_collection_ = consumes>(src); - produces>(); - if (includeErrors_) { - token_errorcollection_ = consumes>(src); - produces>(); - - token_tkerror_detidcollection_ = consumes(src); - produces(); - - token_usererror_detidcollection_ = consumes(edm::InputTag(src.label(), "UserErrorModules")); - produces("UserErrorModules"); - - token_disabled_channelcollection_ = consumes>(src); - produces>(); - } -} - -void SiPixelDigiHeterogeneousConverter::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { - edm::ParameterSetDescription desc; - desc.add("src", edm::InputTag("siPixelClustersPreSplitting")); - desc.add("includeErrors", true); - - descriptions.addWithDefaultLabel(desc); -} - -namespace { - template - void copy(edm::Event& iEvent, const edm::EDGetTokenT& token) { - edm::Handle h; - iEvent.getByToken(token, h); - iEvent.put(std::make_unique(*h)); - } - - template - void copy(edm::Event& iEvent, const edm::EDGetTokenT& token, const std::string& instance) { - edm::Handle h; - iEvent.getByToken(token, h); - iEvent.put(std::make_unique(*h), instance); - } -} // namespace - -void SiPixelDigiHeterogeneousConverter::produce(edm::StreamID, - edm::Event& iEvent, - const edm::EventSetup& iSetup) const { - copy(iEvent, token_collection_); - if (includeErrors_) { - copy(iEvent, token_errorcollection_); - copy(iEvent, token_tkerror_detidcollection_); - copy(iEvent, token_usererror_detidcollection_, "UserErrorModules"); - copy(iEvent, token_disabled_channelcollection_); - } -} - -DEFINE_FWK_MODULE(SiPixelDigiHeterogeneousConverter); From 8fa6e1fe289551e0c334282754e08d267c70402b Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 20 Nov 2020 16:11:46 +0100 Subject: [PATCH 02/12] Address review comments for SiPixelClustersCUDA - Remove commented out default constructor and private: from DeviceConstView * This is perhaps the best compromise between non-default constructors not being preferred for device allocations, and the use case in SiPixelRecHitSoAFromLegacy (for the expected life time of this class) - Remove const getters without c_ prefix - Improve constructor parameter name - Use more initializer list - initialize nClusters_h --- .../SiPixelCluster/interface/SiPixelClustersCUDA.h | 14 ++------------ .../SiPixelCluster/src/SiPixelClustersCUDA.cc | 11 +++++------ .../SiPixelRecHits/plugins/PixelRecHits.cu | 4 ++-- 3 files changed, 9 insertions(+), 20 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h index dbfb5ff5e1761..47dd32e50ec32 100644 --- a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h +++ b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h @@ -10,7 +10,7 @@ class SiPixelClustersCUDA { public: SiPixelClustersCUDA() = default; - explicit SiPixelClustersCUDA(size_t maxClusters, cudaStream_t stream); + explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream); ~SiPixelClustersCUDA() = default; SiPixelClustersCUDA(const SiPixelClustersCUDA &) = delete; @@ -27,11 +27,6 @@ class SiPixelClustersCUDA { uint32_t *moduleId() { return moduleId_d.get(); } uint32_t *clusModuleStart() { return clusModuleStart_d.get(); } - uint32_t const *moduleStart() const { return moduleStart_d.get(); } - uint32_t const *clusInModule() const { return clusInModule_d.get(); } - uint32_t const *moduleId() const { return moduleId_d.get(); } - uint32_t const *clusModuleStart() const { return clusModuleStart_d.get(); } - uint32_t const *c_moduleStart() const { return moduleStart_d.get(); } uint32_t const *c_clusInModule() const { return clusInModule_d.get(); } uint32_t const *c_moduleId() const { return moduleId_d.get(); } @@ -39,16 +34,11 @@ class SiPixelClustersCUDA { class DeviceConstView { public: - // DeviceConstView() = default; - __device__ __forceinline__ uint32_t moduleStart(int i) const { return __ldg(moduleStart_ + i); } __device__ __forceinline__ uint32_t clusInModule(int i) const { return __ldg(clusInModule_ + i); } __device__ __forceinline__ uint32_t moduleId(int i) const { return __ldg(moduleId_ + i); } __device__ __forceinline__ uint32_t clusModuleStart(int i) const { return __ldg(clusModuleStart_ + i); } - friend SiPixelClustersCUDA; - - // private: uint32_t const *moduleStart_; uint32_t const *clusInModule_; uint32_t const *moduleId_; @@ -67,7 +57,7 @@ class SiPixelClustersCUDA { cms::cuda::device::unique_ptr view_d; // "me" pointer - uint32_t nClusters_h; + uint32_t nClusters_h = 0; }; #endif diff --git a/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc b/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc index 7bef9d0d8a52f..5e53f49570bb4 100644 --- a/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc +++ b/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc @@ -4,12 +4,11 @@ #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" -SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxClusters, cudaStream_t stream) { - moduleStart_d = cms::cuda::make_device_unique(maxClusters + 1, stream); - clusInModule_d = cms::cuda::make_device_unique(maxClusters, stream); - moduleId_d = cms::cuda::make_device_unique(maxClusters, stream); - clusModuleStart_d = cms::cuda::make_device_unique(maxClusters + 1, stream); - +SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream) + : moduleStart_d(cms::cuda::make_device_unique(maxModules + 1, stream)), + clusInModule_d(cms::cuda::make_device_unique(maxModules, stream)), + moduleId_d(cms::cuda::make_device_unique(maxModules, stream)), + clusModuleStart_d(cms::cuda::make_device_unique(maxModules + 1, stream)) { auto view = cms::cuda::make_host_unique(stream); view->moduleStart_ = moduleStart_d.get(); view->clusInModule_ = clusInModule_d.get(); diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu index aae20b54e150d..9abb10897961b 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu @@ -39,7 +39,7 @@ namespace pixelgpudetails { pixelCPEforGPU::ParamsOnGPU const* cpeParams, cudaStream_t stream) const { auto nHits = clusters_d.nClusters(); - TrackingRecHit2DCUDA hits_d(nHits, cpeParams, clusters_d.clusModuleStart(), stream); + TrackingRecHit2DCUDA hits_d(nHits, cpeParams, clusters_d.c_clusModuleStart(), stream); int threadsPerBlock = 128; int blocks = digis_d.nModules(); // active modules (with digis) @@ -58,7 +58,7 @@ namespace pixelgpudetails { // assuming full warp of threads is better than a smaller number... if (nHits) { - setHitsLayerStart<<<1, 32, 0, stream>>>(clusters_d.clusModuleStart(), cpeParams, hits_d.hitsLayerStart()); + setHitsLayerStart<<<1, 32, 0, stream>>>(clusters_d.c_clusModuleStart(), cpeParams, hits_d.hitsLayerStart()); cudaCheck(cudaGetLastError()); } From 7a5fc716cb2ad0052f64004244a62447e37e75df Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 20 Nov 2020 18:36:39 +0100 Subject: [PATCH 03/12] Address review comments for SiPixelDigiErrorsCUDA - Use type alias - Remove unnecessary method - Use more initializer list --- .../SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h | 14 +++++++------- .../SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc | 9 ++++----- 2 files changed, 11 insertions(+), 12 deletions(-) diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h index aa06e8dbbd57d..766a0bacd4b4d 100644 --- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h +++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h @@ -10,6 +10,8 @@ class SiPixelDigiErrorsCUDA { public: + using PixelErrorCompactVector = cms::cuda::SimpleVector; + SiPixelDigiErrorsCUDA() = default; explicit SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream); ~SiPixelDigiErrorsCUDA() = default; @@ -21,20 +23,18 @@ class SiPixelDigiErrorsCUDA { const PixelFormatterErrors& formatterErrors() const { return formatterErrors_h; } - cms::cuda::SimpleVector* error() { return error_d.get(); } - cms::cuda::SimpleVector const* error() const { return error_d.get(); } - cms::cuda::SimpleVector const* c_error() const { return error_d.get(); } + PixelErrorCompactVector* error() { return error_d.get(); } + PixelErrorCompactVector const* c_error() const { return error_d.get(); } - using HostDataError = - std::pair, cms::cuda::host::unique_ptr>; + using HostDataError = std::pair>; HostDataError dataErrorToHostAsync(cudaStream_t stream) const; void copyErrorToHostAsync(cudaStream_t stream); private: cms::cuda::device::unique_ptr data_d; - cms::cuda::device::unique_ptr> error_d; - cms::cuda::host::unique_ptr> error_h; + cms::cuda::device::unique_ptr error_d; + cms::cuda::host::unique_ptr error_h; PixelFormatterErrors formatterErrors_h; }; diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc index 075d408a6f6fc..b328d2239efea 100644 --- a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc +++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc @@ -8,13 +8,12 @@ #include SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream) - : formatterErrors_h(std::move(errors)) { - error_d = cms::cuda::make_device_unique>(stream); - data_d = cms::cuda::make_device_unique(maxFedWords, stream); - + : data_d(cms::cuda::make_device_unique(maxFedWords, stream)), + error_d(cms::cuda::make_device_unique(stream)), + error_h(cms::cuda::make_host_unique(stream)), + formatterErrors_h(std::move(errors)) { cms::cuda::memsetAsync(data_d, 0x00, maxFedWords, stream); - error_h = cms::cuda::make_host_unique>(stream); cms::cuda::make_SimpleVector(error_h.get(), maxFedWords, data_d.get()); assert(error_h->empty()); assert(error_h->capacity() == static_cast(maxFedWords)); From 86f33fcdd2d2a3c640353f97fe6fc2079747cd04 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 20 Nov 2020 18:41:49 +0100 Subject: [PATCH 04/12] Address review comments for SiPixelDigisCUDA - Remove unnecessary methods - Remove commented out default constructor and private: from DeviceConstView - Add comments for remaining SiPixelDigisCUDA member arrays --- .../SiPixelDigi/interface/SiPixelDigisCUDA.h | 17 ++--------------- 1 file changed, 2 insertions(+), 15 deletions(-) diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h index 04207f3e0b385..f07f000e4f387 100644 --- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h +++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h @@ -34,14 +34,6 @@ class SiPixelDigisCUDA { uint32_t *pdigi() { return pdigi_d.get(); } uint32_t *rawIdArr() { return rawIdArr_d.get(); } - uint16_t const *xx() const { return xx_d.get(); } - uint16_t const *yy() const { return yy_d.get(); } - uint16_t const *adc() const { return adc_d.get(); } - uint16_t const *moduleInd() const { return moduleInd_d.get(); } - int32_t const *clus() const { return clus_d.get(); } - uint32_t const *pdigi() const { return pdigi_d.get(); } - uint32_t const *rawIdArr() const { return rawIdArr_d.get(); } - uint16_t const *c_xx() const { return xx_d.get(); } uint16_t const *c_yy() const { return yy_d.get(); } uint16_t const *c_adc() const { return adc_d.get(); } @@ -57,17 +49,12 @@ class SiPixelDigisCUDA { class DeviceConstView { public: - // DeviceConstView() = default; - __device__ __forceinline__ uint16_t xx(int i) const { return __ldg(xx_ + i); } __device__ __forceinline__ uint16_t yy(int i) const { return __ldg(yy_ + i); } __device__ __forceinline__ uint16_t adc(int i) const { return __ldg(adc_ + i); } __device__ __forceinline__ uint16_t moduleInd(int i) const { return __ldg(moduleInd_ + i); } __device__ __forceinline__ int32_t clus(int i) const { return __ldg(clus_ + i); } - friend class SiPixelDigisCUDA; - - // private: uint16_t const *xx_; uint16_t const *yy_; uint16_t const *adc_; @@ -88,8 +75,8 @@ class SiPixelDigisCUDA { // These are for CPU output; should we (eventually) place them to a // separate product? - cms::cuda::device::unique_ptr pdigi_d; - cms::cuda::device::unique_ptr rawIdArr_d; + cms::cuda::device::unique_ptr pdigi_d; // packed digi (row, col, adc) of each pixel + cms::cuda::device::unique_ptr rawIdArr_d; // DetId of each pixel uint32_t nModules_h = 0; uint32_t nDigis_h = 0; From f044a97c61305e6d51b77f3a4f90c400f1ec3a23 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 20 Nov 2020 22:11:12 +0100 Subject: [PATCH 05/12] Move PixelErrorsCompact and SiPixelDigiErrorsSoa to DataFormats/SiPixelRawData, rename classes --- .../interface/SiPixelDigiErrorsCUDA.h | 23 ++++++++------- .../SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc | 10 +++---- .../SiPixelDigi/interface/PixelErrors.h | 21 -------------- .../interface/SiPixelDigiErrorsSoA.h | 28 ------------------ .../SiPixelDigi/src/SiPixelDigiErrorsSoA.cc | 10 ------- DataFormats/SiPixelDigi/src/classes.h | 1 - DataFormats/SiPixelDigi/src/classes_def.xml | 3 -- .../interface/SiPixelErrorCompact.h | 13 +++++++++ .../interface/SiPixelErrorsSoA.h | 29 +++++++++++++++++++ .../interface/SiPixelFormatterErrors.h | 12 ++++++++ .../SiPixelRawData/src/SiPixelErrorsSoA.cc | 8 +++++ DataFormats/SiPixelRawData/src/classes.h | 1 + .../SiPixelRawData/src/classes_def.xml | 3 ++ .../plugins/SiPixelDigiErrorsFromSoA.cc | 8 ++--- .../plugins/SiPixelDigiErrorsSoAFromCUDA.cc | 14 ++++----- .../plugins/SiPixelRawToClusterGPUKernel.cu | 10 +++---- .../plugins/SiPixelRawToClusterGPUKernel.h | 5 ++-- 17 files changed, 102 insertions(+), 97 deletions(-) delete mode 100644 DataFormats/SiPixelDigi/interface/PixelErrors.h delete mode 100644 DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h delete mode 100644 DataFormats/SiPixelDigi/src/SiPixelDigiErrorsSoA.cc create mode 100644 DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h create mode 100644 DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h create mode 100644 DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h create mode 100644 DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h index 766a0bacd4b4d..fe4ee7d228161 100644 --- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h +++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h @@ -1,7 +1,8 @@ #ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h #define CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h -#include "DataFormats/SiPixelDigi/interface/PixelErrors.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h" #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/SimpleVector.h" @@ -10,10 +11,10 @@ class SiPixelDigiErrorsCUDA { public: - using PixelErrorCompactVector = cms::cuda::SimpleVector; + using SiPixelErrorCompactVector = cms::cuda::SimpleVector; SiPixelDigiErrorsCUDA() = default; - explicit SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream); + explicit SiPixelDigiErrorsCUDA(size_t maxFedWords, SiPixelFormatterErrors errors, cudaStream_t stream); ~SiPixelDigiErrorsCUDA() = default; SiPixelDigiErrorsCUDA(const SiPixelDigiErrorsCUDA&) = delete; @@ -21,21 +22,21 @@ class SiPixelDigiErrorsCUDA { SiPixelDigiErrorsCUDA(SiPixelDigiErrorsCUDA&&) = default; SiPixelDigiErrorsCUDA& operator=(SiPixelDigiErrorsCUDA&&) = default; - const PixelFormatterErrors& formatterErrors() const { return formatterErrors_h; } + const SiPixelFormatterErrors& formatterErrors() const { return formatterErrors_h; } - PixelErrorCompactVector* error() { return error_d.get(); } - PixelErrorCompactVector const* c_error() const { return error_d.get(); } + SiPixelErrorCompactVector* error() { return error_d.get(); } + SiPixelErrorCompactVector const* c_error() const { return error_d.get(); } - using HostDataError = std::pair>; + using HostDataError = std::pair>; HostDataError dataErrorToHostAsync(cudaStream_t stream) const; void copyErrorToHostAsync(cudaStream_t stream); private: - cms::cuda::device::unique_ptr data_d; - cms::cuda::device::unique_ptr error_d; - cms::cuda::host::unique_ptr error_h; - PixelFormatterErrors formatterErrors_h; + cms::cuda::device::unique_ptr data_d; + cms::cuda::device::unique_ptr error_d; + cms::cuda::host::unique_ptr error_h; + SiPixelFormatterErrors formatterErrors_h; }; #endif diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc index b328d2239efea..70bf2e8aa19f5 100644 --- a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc +++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc @@ -7,10 +7,10 @@ #include -SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream) - : data_d(cms::cuda::make_device_unique(maxFedWords, stream)), - error_d(cms::cuda::make_device_unique(stream)), - error_h(cms::cuda::make_host_unique(stream)), +SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, SiPixelFormatterErrors errors, cudaStream_t stream) + : data_d(cms::cuda::make_device_unique(maxFedWords, stream)), + error_d(cms::cuda::make_device_unique(stream)), + error_h(cms::cuda::make_host_unique(stream)), formatterErrors_h(std::move(errors)) { cms::cuda::memsetAsync(data_d, 0x00, maxFedWords, stream); @@ -29,7 +29,7 @@ SiPixelDigiErrorsCUDA::HostDataError SiPixelDigiErrorsCUDA::dataErrorToHostAsync // On one hand size() could be sufficient. On the other hand, if // someone copies the SimpleVector<>, (s)he might expect the data // buffer to actually have space for capacity() elements. - auto data = cms::cuda::make_host_unique(error_h->capacity(), stream); + auto data = cms::cuda::make_host_unique(error_h->capacity(), stream); // but transfer only the required amount if (not error_h->empty()) { diff --git a/DataFormats/SiPixelDigi/interface/PixelErrors.h b/DataFormats/SiPixelDigi/interface/PixelErrors.h deleted file mode 100644 index 073b9962deaaa..0000000000000 --- a/DataFormats/SiPixelDigi/interface/PixelErrors.h +++ /dev/null @@ -1,21 +0,0 @@ -#ifndef DataFormats_SiPixelDigi_interface_PixelErrors_h -#define DataFormats_SiPixelDigi_interface_PixelErrors_h - -#include -#include - -#include "DataFormats/SiPixelRawData/interface/SiPixelRawDataError.h" -#include "FWCore/Utilities/interface/typedefs.h" - -// Better ideas for the placement of these? - -struct PixelErrorCompact { - uint32_t rawId; - uint32_t word; - uint8_t errorType; - uint8_t fedId; -}; - -using PixelFormatterErrors = std::map>; - -#endif // DataFormats_SiPixelDigi_interface_PixelErrors_h diff --git a/DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h b/DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h deleted file mode 100644 index ee1227ed4fae1..0000000000000 --- a/DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h +++ /dev/null @@ -1,28 +0,0 @@ -#ifndef DataFormats_SiPixelDigi_interface_SiPixelDigiErrorsSoA_h -#define DataFormats_SiPixelDigi_interface_SiPixelDigiErrorsSoA_h - -#include "DataFormats/SiPixelDigi/interface/PixelErrors.h" - -#include -#include - -class SiPixelDigiErrorsSoA { -public: - SiPixelDigiErrorsSoA() = default; - explicit SiPixelDigiErrorsSoA(size_t nErrors, const PixelErrorCompact *error, const PixelFormatterErrors *err); - ~SiPixelDigiErrorsSoA() = default; - - auto size() const { return error_.size(); } - - const PixelFormatterErrors *formatterErrors() const { return formatterErrors_; } - - const PixelErrorCompact &error(size_t i) const { return error_[i]; } - - const std::vector &errorVector() const { return error_; } - -private: - std::vector error_; - const PixelFormatterErrors *formatterErrors_ = nullptr; -}; - -#endif diff --git a/DataFormats/SiPixelDigi/src/SiPixelDigiErrorsSoA.cc b/DataFormats/SiPixelDigi/src/SiPixelDigiErrorsSoA.cc deleted file mode 100644 index a93bd7d3774f3..0000000000000 --- a/DataFormats/SiPixelDigi/src/SiPixelDigiErrorsSoA.cc +++ /dev/null @@ -1,10 +0,0 @@ -#include "DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h" - -#include - -SiPixelDigiErrorsSoA::SiPixelDigiErrorsSoA(size_t nErrors, - const PixelErrorCompact *error, - const PixelFormatterErrors *err) - : error_(error, error + nErrors), formatterErrors_(err) { - assert(error_.size() == nErrors); -} diff --git a/DataFormats/SiPixelDigi/src/classes.h b/DataFormats/SiPixelDigi/src/classes.h index ba68d3289e8cd..1360ee6e469d9 100644 --- a/DataFormats/SiPixelDigi/src/classes.h +++ b/DataFormats/SiPixelDigi/src/classes.h @@ -6,7 +6,6 @@ #include "DataFormats/SiPixelDigi/interface/SiPixelCalibDigi.h" #include "DataFormats/SiPixelDigi/interface/SiPixelCalibDigiError.h" #include "DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h" -#include "DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h" #include "DataFormats/Common/interface/Wrapper.h" #include "DataFormats/Common/interface/DetSetVector.h" #include "DataFormats/Common/interface/DetSetVectorNew.h" diff --git a/DataFormats/SiPixelDigi/src/classes_def.xml b/DataFormats/SiPixelDigi/src/classes_def.xml index 8cabbd3f3f06e..e6bc08de161fa 100755 --- a/DataFormats/SiPixelDigi/src/classes_def.xml +++ b/DataFormats/SiPixelDigi/src/classes_def.xml @@ -52,7 +52,4 @@ - - - diff --git a/DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h b/DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h new file mode 100644 index 0000000000000..0b1a80868594f --- /dev/null +++ b/DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h @@ -0,0 +1,13 @@ +#ifndef DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h +#define DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h + +#include + +struct SiPixelErrorCompact { + uint32_t rawId; + uint32_t word; + uint8_t errorType; + uint8_t fedId; +}; + +#endif // DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h diff --git a/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h b/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h new file mode 100644 index 0000000000000..837b8a0634e6d --- /dev/null +++ b/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h @@ -0,0 +1,29 @@ +#ifndef DataFormats_SiPixelDigi_interface_SiPixelErrorsSoA_h +#define DataFormats_SiPixelDigi_interface_SiPixelErrorsSoA_h + +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h" + +#include +#include + +class SiPixelErrorsSoA { +public: + SiPixelErrorsSoA() = default; + explicit SiPixelErrorsSoA(size_t nErrors, const SiPixelErrorCompact *error, const SiPixelFormatterErrors *err); + ~SiPixelErrorsSoA() = default; + + auto size() const { return error_.size(); } + + const SiPixelFormatterErrors *formatterErrors() const { return formatterErrors_; } + + const SiPixelErrorCompact &error(size_t i) const { return error_[i]; } + + const std::vector &errorVector() const { return error_; } + +private: + std::vector error_; + const SiPixelFormatterErrors *formatterErrors_ = nullptr; +}; + +#endif diff --git a/DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h b/DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h new file mode 100644 index 0000000000000..9d372737300d4 --- /dev/null +++ b/DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h @@ -0,0 +1,12 @@ +#ifndef DataFormats_SiPixelRawData_interface_SiPixelFormatterErrors_h +#define DataFormats_SiPixelRawData_interface_SiPixelFormatterErrors_h + +#include +#include + +#include "DataFormats/SiPixelRawData/interface/SiPixelRawDataError.h" +#include "FWCore/Utilities/interface/typedefs.h" + +using SiPixelFormatterErrors = std::map>; + +#endif // DataFormats_SiPixelRawData_interface_SiPixelFormatterErrors_h diff --git a/DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc b/DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc new file mode 100644 index 0000000000000..a63bb179c59cd --- /dev/null +++ b/DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc @@ -0,0 +1,8 @@ +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h" + +#include + +SiPixelErrorsSoA::SiPixelErrorsSoA(size_t nErrors, const SiPixelErrorCompact *error, const SiPixelFormatterErrors *err) + : error_(error, error + nErrors), formatterErrors_(err) { + assert(error_.size() == nErrors); +} diff --git a/DataFormats/SiPixelRawData/src/classes.h b/DataFormats/SiPixelRawData/src/classes.h index 73768cc373013..7a07e9f35f388 100644 --- a/DataFormats/SiPixelRawData/src/classes.h +++ b/DataFormats/SiPixelRawData/src/classes.h @@ -2,6 +2,7 @@ #define SIPIXELRAWDATA_CLASSES_H #include "DataFormats/SiPixelRawData/interface/SiPixelRawDataError.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h" #include "DataFormats/Common/interface/Wrapper.h" #include "DataFormats/Common/interface/DetSetVector.h" #include diff --git a/DataFormats/SiPixelRawData/src/classes_def.xml b/DataFormats/SiPixelRawData/src/classes_def.xml index 827d4b1191cf6..fd2b5dcf27965 100644 --- a/DataFormats/SiPixelRawData/src/classes_def.xml +++ b/DataFormats/SiPixelRawData/src/classes_def.xml @@ -14,4 +14,7 @@ + + + diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc index ea381948ec352..7a49646d7a9a1 100644 --- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc +++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsFromSoA.cc @@ -6,7 +6,7 @@ #include "DataFormats/DetId/interface/DetIdCollection.h" #include "DataFormats/SiPixelDetId/interface/PixelFEDChannel.h" #include "DataFormats/SiPixelDigi/interface/PixelDigi.h" -#include "DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h" #include "EventFilter/SiPixelRawToDigi/interface/PixelDataFormatter.h" #include "FWCore/Framework/interface/ESTransientHandle.h" #include "FWCore/Framework/interface/ESWatcher.h" @@ -31,7 +31,7 @@ class SiPixelDigiErrorsFromSoA : public edm::stream::EDProducer<> { void produce(edm::Event& iEvent, const edm::EventSetup& iSetup) override; const edm::ESGetToken cablingToken_; - const edm::EDGetTokenT digiErrorSoAGetToken_; + const edm::EDGetTokenT digiErrorSoAGetToken_; const edm::EDPutTokenT> errorPutToken_; const edm::EDPutTokenT tkErrorPutToken_; const edm::EDPutTokenT userErrorPutToken_; @@ -48,7 +48,7 @@ class SiPixelDigiErrorsFromSoA : public edm::stream::EDProducer<> { SiPixelDigiErrorsFromSoA::SiPixelDigiErrorsFromSoA(const edm::ParameterSet& iConfig) : cablingToken_(esConsumes(edm::ESInputTag("", iConfig.getParameter("CablingMapLabel")))), - digiErrorSoAGetToken_{consumes(iConfig.getParameter("digiErrorSoASrc"))}, + digiErrorSoAGetToken_{consumes(iConfig.getParameter("digiErrorSoASrc"))}, errorPutToken_{produces>()}, tkErrorPutToken_{produces()}, userErrorPutToken_{produces("UserErrorModules")}, @@ -95,7 +95,7 @@ void SiPixelDigiErrorsFromSoA::produce(edm::Event& iEvent, const edm::EventSetup auto size = digiErrors.size(); for (auto i = 0U; i < size; i++) { - PixelErrorCompact err = digiErrors.error(i); + SiPixelErrorCompact err = digiErrors.error(i); if (err.errorType != 0) { SiPixelRawDataError error(err.word, err.errorType, err.fedId + 1200); errors[err.rawId].push_back(error); diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc index c5b568750ad7d..f2c7d0de5fe24 100644 --- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc +++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc @@ -1,6 +1,6 @@ #include "CUDADataFormats/Common/interface/Product.h" #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h" -#include "DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h" #include "FWCore/Framework/interface/EventSetup.h" #include "FWCore/Framework/interface/Event.h" #include "FWCore/Framework/interface/MakerMacros.h" @@ -25,17 +25,17 @@ class SiPixelDigiErrorsSoAFromCUDA : public edm::stream::EDProducer> digiErrorGetToken_; - edm::EDPutTokenT digiErrorPutToken_; + edm::EDPutTokenT digiErrorPutToken_; - cms::cuda::host::unique_ptr data_; - cms::cuda::SimpleVector error_; - const PixelFormatterErrors* formatterErrors_ = nullptr; + cms::cuda::host::unique_ptr data_; + cms::cuda::SimpleVector error_; + const SiPixelFormatterErrors* formatterErrors_ = nullptr; }; SiPixelDigiErrorsSoAFromCUDA::SiPixelDigiErrorsSoAFromCUDA(const edm::ParameterSet& iConfig) : digiErrorGetToken_( consumes>(iConfig.getParameter("src"))), - digiErrorPutToken_(produces()) {} + digiErrorPutToken_(produces()) {} void SiPixelDigiErrorsSoAFromCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) { edm::ParameterSetDescription desc; @@ -70,7 +70,7 @@ void SiPixelDigiErrorsSoAFromCUDA::produce(edm::Event& iEvent, const edm::EventS // use cudaMallocHost without a GPU... iEvent.emplace(digiErrorPutToken_, error_.size(), error_.data(), formatterErrors_); - error_ = cms::cuda::make_SimpleVector(0, nullptr); + error_ = cms::cuda::make_SimpleVector(0, nullptr); data_.reset(); formatterErrors_ = nullptr; } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index f14808dda1e2b..a522fe0bd966a 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -362,7 +362,7 @@ namespace pixelgpudetails { uint32_t *pdigi, uint32_t *rawIdArr, uint16_t *moduleId, - cms::cuda::SimpleVector *err, + cms::cuda::SimpleVector *err, bool useQualityInfo, bool includeErrors, bool debug) { @@ -397,7 +397,7 @@ namespace pixelgpudetails { skipROC = (roc < pixelgpudetails::maxROCIndex) ? false : (errorType != 0); if (includeErrors and skipROC) { uint32_t rID = getErrRawID(fedId, ww, errorType, cablingMap, debug); - err->push_back(PixelErrorCompact{rID, ww, errorType, fedId}); + err->push_back(SiPixelErrorCompact{rID, ww, errorType, fedId}); continue; } @@ -441,7 +441,7 @@ namespace pixelgpudetails { if (includeErrors) { if (not rocRowColIsValid(row, col)) { uint8_t error = conversionError(fedId, 3, debug); //use the device function and fill the arrays - err->push_back(PixelErrorCompact{rawId, ww, error, fedId}); + err->push_back(SiPixelErrorCompact{rawId, ww, error, fedId}); if (debug) printf("BPIX1 Error status: %i\n", error); continue; @@ -457,7 +457,7 @@ namespace pixelgpudetails { localPix.col = col; if (includeErrors and not dcolIsValid(dcol, pxid)) { uint8_t error = conversionError(fedId, 3, debug); - err->push_back(PixelErrorCompact{rawId, ww, error, fedId}); + err->push_back(SiPixelErrorCompact{rawId, ww, error, fedId}); if (debug) printf("Error status: %i %d %d %d %d\n", error, dcol, pxid, fedId, roc); continue; @@ -528,7 +528,7 @@ namespace pixelgpudetails { const unsigned char *modToUnp, const SiPixelGainForHLTonGPU *gains, const WordFedAppender &wordFed, - PixelFormatterErrors &&errors, + SiPixelFormatterErrors &&errors, const uint32_t wordCounter, const uint32_t fedCounter, bool useQualityInfo, diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h index d214e7784af48..1f3983e4f4889 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h @@ -11,7 +11,8 @@ #include "HeterogeneousCore/CUDAUtilities/interface/SimpleVector.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" -#include "DataFormats/SiPixelDigi/interface/PixelErrors.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h" +#include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h" struct SiPixelFedCablingMapGPU; class SiPixelGainForHLTonGPU; @@ -173,7 +174,7 @@ namespace pixelgpudetails { const unsigned char* modToUnp, const SiPixelGainForHLTonGPU* gains, const WordFedAppender& wordFed, - PixelFormatterErrors&& errors, + SiPixelFormatterErrors&& errors, const uint32_t wordCounter, const uint32_t fedCounter, bool useQualityInfo, From 7eff58dc0e5615929ceac115acebc2e73f493e27 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 20 Nov 2020 23:18:44 +0100 Subject: [PATCH 06/12] Address review comments for SiPixelErrorsSoA - Remove redundant assert - Move constructor inline --- DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h | 3 ++- DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc | 8 -------- 2 files changed, 2 insertions(+), 9 deletions(-) delete mode 100644 DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc diff --git a/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h b/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h index 837b8a0634e6d..c72c19ce5fda4 100644 --- a/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h +++ b/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h @@ -10,7 +10,8 @@ class SiPixelErrorsSoA { public: SiPixelErrorsSoA() = default; - explicit SiPixelErrorsSoA(size_t nErrors, const SiPixelErrorCompact *error, const SiPixelFormatterErrors *err); + explicit SiPixelErrorsSoA(size_t nErrors, const SiPixelErrorCompact *error, const SiPixelFormatterErrors *err) + : error_(error, error + nErrors), formatterErrors_(err) {} ~SiPixelErrorsSoA() = default; auto size() const { return error_.size(); } diff --git a/DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc b/DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc deleted file mode 100644 index a63bb179c59cd..0000000000000 --- a/DataFormats/SiPixelRawData/src/SiPixelErrorsSoA.cc +++ /dev/null @@ -1,8 +0,0 @@ -#include "DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h" - -#include - -SiPixelErrorsSoA::SiPixelErrorsSoA(size_t nErrors, const SiPixelErrorCompact *error, const SiPixelFormatterErrors *err) - : error_(error, error + nErrors), formatterErrors_(err) { - assert(error_.size() == nErrors); -} From ccacaf241ef6ccdf068b657e1fa838a539094bd1 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 20 Nov 2020 23:20:27 +0100 Subject: [PATCH 07/12] Address review comments for SiPixelDigisSoA - Remove redundant assert - Add comments --- .../SiPixelDigi/interface/SiPixelDigisSoA.h | 14 ++++++++++---- DataFormats/SiPixelDigi/src/SiPixelDigisSoA.cc | 4 +--- 2 files changed, 11 insertions(+), 7 deletions(-) diff --git a/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h b/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h index 50e863f03ff02..6c016155b1cb0 100644 --- a/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h +++ b/DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h @@ -4,6 +4,12 @@ #include #include +// The main purpose of this class is to deliver digi and cluster data +// from an EDProducer that transfers the data from GPU to host to an +// EDProducer that converts the SoA to legacy data products. The class +// is independent of any GPU technology, and in prunciple could be +// produced by host code, and be used for other purposes than +// conversion-to-legacy as well. class SiPixelDigisSoA { public: SiPixelDigisSoA() = default; @@ -24,10 +30,10 @@ class SiPixelDigisSoA { const std::vector& clusVector() const { return clus_; } private: - std::vector pdigi_; - std::vector rawIdArr_; - std::vector adc_; - std::vector clus_; + std::vector pdigi_; // packed digi (row, col, adc) of each pixel + std::vector rawIdArr_; // DetId of each pixel + std::vector adc_; // ADC of each pixel + std::vector clus_; // cluster id of each pixel }; #endif diff --git a/DataFormats/SiPixelDigi/src/SiPixelDigisSoA.cc b/DataFormats/SiPixelDigi/src/SiPixelDigisSoA.cc index 992c98f450616..b95c004a50a25 100644 --- a/DataFormats/SiPixelDigi/src/SiPixelDigisSoA.cc +++ b/DataFormats/SiPixelDigi/src/SiPixelDigisSoA.cc @@ -7,6 +7,4 @@ SiPixelDigisSoA::SiPixelDigisSoA( : pdigi_(pdigi, pdigi + nDigis), rawIdArr_(rawIdArr, rawIdArr + nDigis), adc_(adc, adc + nDigis), - clus_(clus, clus + nDigis) { - assert(pdigi_.size() == nDigis); -} + clus_(clus, clus + nDigis) {} From 402adeadf79c7cb56184ccd6e5271490bae220ef Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Sat, 21 Nov 2020 01:32:37 +0100 Subject: [PATCH 08/12] Enable if constexpr also for CUDA in TrackingRecHit2DHeterogeneous --- .../interface/TrackingRecHit2DHeterogeneous.h | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h index b0aa79cfe20b6..73a6daaa4e387 100644 --- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h +++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h @@ -84,11 +84,7 @@ TrackingRecHit2DHeterogeneous::TrackingRecHit2DHeterogeneous(uint32_t nH // if empy do not bother if (0 == nHits) { - if -#ifndef __CUDACC__ - constexpr -#endif - (std::is_same::value) { + if constexpr (std::is_same::value) { cms::cuda::copyAsync(m_view, view, stream); } else { m_view.reset(view.release()); // NOLINT: std::move() breaks CUDA version From f14d7b72592725fc03e564be81bf8107d900a333 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Sat, 21 Nov 2020 01:37:28 +0100 Subject: [PATCH 09/12] Move dictionary of HostProduct to CUDADataFormats/Common --- CUDADataFormats/Common/BuildFile.xml | 2 ++ CUDADataFormats/Common/src/classes.h | 7 +++++++ CUDADataFormats/Common/src/classes_def.xml | 4 ++++ CUDADataFormats/TrackingRecHit/src/classes.h | 1 - CUDADataFormats/TrackingRecHit/src/classes_def.xml | 2 -- 5 files changed, 13 insertions(+), 3 deletions(-) create mode 100644 CUDADataFormats/Common/src/classes.h create mode 100644 CUDADataFormats/Common/src/classes_def.xml diff --git a/CUDADataFormats/Common/BuildFile.xml b/CUDADataFormats/Common/BuildFile.xml index b990c1295e31a..c524cada33060 100644 --- a/CUDADataFormats/Common/BuildFile.xml +++ b/CUDADataFormats/Common/BuildFile.xml @@ -1,5 +1,7 @@ + + diff --git a/CUDADataFormats/Common/src/classes.h b/CUDADataFormats/Common/src/classes.h new file mode 100644 index 0000000000000..239e071d513a2 --- /dev/null +++ b/CUDADataFormats/Common/src/classes.h @@ -0,0 +1,7 @@ +#ifndef CUDADataFormats_Common_src_classes_h +#define CUDADataFormats_Common_src_classes_h + +#include "CUDADataFormats/Common/interface/HostProduct.h" +#include "DataFormats/Common/interface/Wrapper.h" + +#endif // CUDADataFormats_Common_src_classes_h diff --git a/CUDADataFormats/Common/src/classes_def.xml b/CUDADataFormats/Common/src/classes_def.xml new file mode 100644 index 0000000000000..024d927595914 --- /dev/null +++ b/CUDADataFormats/Common/src/classes_def.xml @@ -0,0 +1,4 @@ + + + + diff --git a/CUDADataFormats/TrackingRecHit/src/classes.h b/CUDADataFormats/TrackingRecHit/src/classes.h index d80226ec7a14b..3d40821493c5b 100644 --- a/CUDADataFormats/TrackingRecHit/src/classes.h +++ b/CUDADataFormats/TrackingRecHit/src/classes.h @@ -2,7 +2,6 @@ #define CUDADataFormats_SiPixelCluster_src_classes_h #include "CUDADataFormats/Common/interface/Product.h" -#include "CUDADataFormats/Common/interface/HostProduct.h" #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h" #include "DataFormats/Common/interface/Wrapper.h" diff --git a/CUDADataFormats/TrackingRecHit/src/classes_def.xml b/CUDADataFormats/TrackingRecHit/src/classes_def.xml index 02b0eb37d157b..7e1919de510b3 100644 --- a/CUDADataFormats/TrackingRecHit/src/classes_def.xml +++ b/CUDADataFormats/TrackingRecHit/src/classes_def.xml @@ -5,6 +5,4 @@ - - From 711af1870d7748a03e510acdbd69d6662f48242d Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Sat, 21 Nov 2020 02:45:18 +0100 Subject: [PATCH 10/12] Drop c_ prefix of const column accessors of SiPixelClustersCUDA --- .../SiPixelCluster/interface/SiPixelClustersCUDA.h | 8 ++++---- .../plugins/SiPixelRawToClusterGPUKernel.cu | 8 ++++---- RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu | 4 ++-- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h index 47dd32e50ec32..d5d009aaffeb5 100644 --- a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h +++ b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h @@ -27,10 +27,10 @@ class SiPixelClustersCUDA { uint32_t *moduleId() { return moduleId_d.get(); } uint32_t *clusModuleStart() { return clusModuleStart_d.get(); } - uint32_t const *c_moduleStart() const { return moduleStart_d.get(); } - uint32_t const *c_clusInModule() const { return clusInModule_d.get(); } - uint32_t const *c_moduleId() const { return moduleId_d.get(); } - uint32_t const *c_clusModuleStart() const { return clusModuleStart_d.get(); } + uint32_t const *moduleStart() const { return moduleStart_d.get(); } + uint32_t const *clusInModule() const { return clusInModule_d.get(); } + uint32_t const *moduleId() const { return moduleId_d.get(); } + uint32_t const *clusModuleStart() const { return clusModuleStart_d.get(); } class DeviceConstView { public: diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index a522fe0bd966a..d45fdd94961ea 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -637,7 +637,7 @@ namespace pixelgpudetails { findClus<<>>(digis_d.c_moduleInd(), digis_d.c_xx(), digis_d.c_yy(), - clusters_d.c_moduleStart(), + clusters_d.moduleStart(), clusters_d.clusInModule(), clusters_d.moduleId(), digis_d.clus(), @@ -651,9 +651,9 @@ namespace pixelgpudetails { // apply charge cut clusterChargeCut<<>>(digis_d.moduleInd(), digis_d.c_adc(), - clusters_d.c_moduleStart(), + clusters_d.moduleStart(), clusters_d.clusInModule(), - clusters_d.c_moduleId(), + clusters_d.moduleId(), digis_d.clus(), wordCounter); cudaCheck(cudaGetLastError()); @@ -664,7 +664,7 @@ namespace pixelgpudetails { // synchronization/ExternalWork // MUST be ONE block - fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d.c_clusInModule(), clusters_d.clusModuleStart()); + fillHitsModuleStart<<<1, 1024, 0, stream>>>(clusters_d.clusInModule(), clusters_d.clusModuleStart()); // last element holds the number of all clusters cudaCheck(cudaMemcpyAsync(&(nModules_Clusters_h[1]), diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu index 9abb10897961b..aae20b54e150d 100644 --- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu +++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu @@ -39,7 +39,7 @@ namespace pixelgpudetails { pixelCPEforGPU::ParamsOnGPU const* cpeParams, cudaStream_t stream) const { auto nHits = clusters_d.nClusters(); - TrackingRecHit2DCUDA hits_d(nHits, cpeParams, clusters_d.c_clusModuleStart(), stream); + TrackingRecHit2DCUDA hits_d(nHits, cpeParams, clusters_d.clusModuleStart(), stream); int threadsPerBlock = 128; int blocks = digis_d.nModules(); // active modules (with digis) @@ -58,7 +58,7 @@ namespace pixelgpudetails { // assuming full warp of threads is better than a smaller number... if (nHits) { - setHitsLayerStart<<<1, 32, 0, stream>>>(clusters_d.c_clusModuleStart(), cpeParams, hits_d.hitsLayerStart()); + setHitsLayerStart<<<1, 32, 0, stream>>>(clusters_d.clusModuleStart(), cpeParams, hits_d.hitsLayerStart()); cudaCheck(cudaGetLastError()); } From fdab09baf0059697eb299cb2534ed6788a508927 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Sat, 21 Nov 2020 02:52:35 +0100 Subject: [PATCH 11/12] Drop c_ prefix of const column accessor of SiPixelDigiErrorsCUDA --- CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h index fe4ee7d228161..85e8883fa1bd4 100644 --- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h +++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h @@ -25,7 +25,7 @@ class SiPixelDigiErrorsCUDA { const SiPixelFormatterErrors& formatterErrors() const { return formatterErrors_h; } SiPixelErrorCompactVector* error() { return error_d.get(); } - SiPixelErrorCompactVector const* c_error() const { return error_d.get(); } + SiPixelErrorCompactVector const* error() const { return error_d.get(); } using HostDataError = std::pair>; HostDataError dataErrorToHostAsync(cudaStream_t stream) const; From a15f89e0f87fc34933fbfe24f63320303a2b6767 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Sat, 21 Nov 2020 02:58:00 +0100 Subject: [PATCH 12/12] Drop c_ prefix of const column accessor of SiPixelDigisCUDA --- .../SiPixelDigi/interface/SiPixelDigisCUDA.h | 14 +++++++------- .../plugins/SiPixelRawToClusterGPUKernel.cu | 14 +++++++------- 2 files changed, 14 insertions(+), 14 deletions(-) diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h index f07f000e4f387..2dc1f628bf426 100644 --- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h +++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h @@ -34,13 +34,13 @@ class SiPixelDigisCUDA { uint32_t *pdigi() { return pdigi_d.get(); } uint32_t *rawIdArr() { return rawIdArr_d.get(); } - uint16_t const *c_xx() const { return xx_d.get(); } - uint16_t const *c_yy() const { return yy_d.get(); } - uint16_t const *c_adc() const { return adc_d.get(); } - uint16_t const *c_moduleInd() const { return moduleInd_d.get(); } - int32_t const *c_clus() const { return clus_d.get(); } - uint32_t const *c_pdigi() const { return pdigi_d.get(); } - uint32_t const *c_rawIdArr() const { return rawIdArr_d.get(); } + uint16_t const *xx() const { return xx_d.get(); } + uint16_t const *yy() const { return yy_d.get(); } + uint16_t const *adc() const { return adc_d.get(); } + uint16_t const *moduleInd() const { return moduleInd_d.get(); } + int32_t const *clus() const { return clus_d.get(); } + uint32_t const *pdigi() const { return pdigi_d.get(); } + uint32_t const *rawIdArr() const { return rawIdArr_d.get(); } cms::cuda::host::unique_ptr adcToHostAsync(cudaStream_t stream) const; cms::cuda::host::unique_ptr clusToHostAsync(cudaStream_t stream) const; diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index d45fdd94961ea..3eeed5e4cfd97 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -602,8 +602,8 @@ namespace pixelgpudetails { gpuCalibPixel::calibDigis<<>>(isRun2, digis_d.moduleInd(), - digis_d.c_xx(), - digis_d.c_yy(), + digis_d.xx(), + digis_d.yy(), digis_d.adc(), gains, wordCounter, @@ -622,7 +622,7 @@ namespace pixelgpudetails { #endif countModules<<>>( - digis_d.c_moduleInd(), clusters_d.moduleStart(), digis_d.clus(), wordCounter); + digis_d.moduleInd(), clusters_d.moduleStart(), digis_d.clus(), wordCounter); cudaCheck(cudaGetLastError()); // read the number of modules into a data member, used by getProduct()) @@ -634,9 +634,9 @@ namespace pixelgpudetails { #ifdef GPU_DEBUG std::cout << "CUDA findClus kernel launch with " << blocks << " blocks of " << threadsPerBlock << " threads\n"; #endif - findClus<<>>(digis_d.c_moduleInd(), - digis_d.c_xx(), - digis_d.c_yy(), + findClus<<>>(digis_d.moduleInd(), + digis_d.xx(), + digis_d.yy(), clusters_d.moduleStart(), clusters_d.clusInModule(), clusters_d.moduleId(), @@ -650,7 +650,7 @@ namespace pixelgpudetails { // apply charge cut clusterChargeCut<<>>(digis_d.moduleInd(), - digis_d.c_adc(), + digis_d.adc(), clusters_d.moduleStart(), clusters_d.clusInModule(), clusters_d.moduleId(),