From 79d75e7aaf941f2ec90bf1019e92eae463ca54a8 Mon Sep 17 00:00:00 2001 From: Matti Kortelainen Date: Fri, 27 Nov 2020 08:51:39 -0600 Subject: [PATCH] Address pixel local reco PR review comments (#575) Remove SiPixelDigiHeterogeneousConverter as obsolete, should have been removed as part of cms-patatrack#100. 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 with c_ prefix - improve constructor parameter name - use more initializer list - initialize nClusters_h Address review comments for SiPixelDigiErrorsCUDA: - use type alias - remove const getters with c_ prefix and other unnecessary methods - use more initializer list Address review comments for SiPixelDigisCUDA: - remove const getters with c_ prefix and other unnecessary methods - remove commented out default constructor and private: from DeviceConstView - add comments for remaining SiPixelDigisCUDA member arrays Move PixelErrorsCompact and SiPixelDigiErrorsSoa to DataFormats/SiPixelRawData, rename classes Address review comments for SiPixelErrorsSoA - remove redundant assert - move constructor inline Address review comments for SiPixelDigisSoA - remove redundant assert - add comments Enable if constexpr also for CUDA in TrackingRecHit2DHeterogeneous 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 + .../interface/SiPixelClustersCUDA.h | 14 +-- .../SiPixelCluster/src/SiPixelClustersCUDA.cc | 11 +-- .../interface/SiPixelDigiErrorsCUDA.h | 25 ++--- .../SiPixelDigi/interface/SiPixelDigisCUDA.h | 17 +--- .../SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc | 13 ++- .../interface/TrackingRecHit2DHeterogeneous.h | 6 +- CUDADataFormats/TrackingRecHit/src/classes.h | 1 - .../TrackingRecHit/src/classes_def.xml | 2 - .../SiPixelDigi/interface/PixelErrors.h | 21 ---- .../interface/SiPixelDigiErrorsSoA.h | 28 ------ .../SiPixelDigi/interface/SiPixelDigisSoA.h | 14 ++- .../SiPixelDigi/src/SiPixelDigiErrorsSoA.cc | 10 -- .../SiPixelDigi/src/SiPixelDigisSoA.cc | 4 +- DataFormats/SiPixelDigi/src/classes.h | 1 - DataFormats/SiPixelDigi/src/classes_def.xml | 3 - .../interface/SiPixelErrorCompact.h | 13 +++ .../interface/SiPixelErrorsSoA.h | 30 ++++++ .../interface/SiPixelFormatterErrors.h | 12 +++ DataFormats/SiPixelRawData/src/classes.h | 1 + .../SiPixelRawData/src/classes_def.xml | 3 + .../plugins/SiPixelDigiErrorsFromSoA.cc | 8 +- .../plugins/SiPixelDigiErrorsSoAFromCUDA.cc | 14 +-- .../SiPixelDigiHeterogeneousConverter.cc | 95 ------------------- .../plugins/SiPixelRawToClusterGPUKernel.cu | 32 +++---- .../plugins/SiPixelRawToClusterGPUKernel.h | 5 +- 28 files changed, 142 insertions(+), 254 deletions(-) create mode 100644 CUDADataFormats/Common/src/classes.h create mode 100644 CUDADataFormats/Common/src/classes_def.xml 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 delete mode 100644 RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigiHeterogeneousConverter.cc 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/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h index dbfb5ff5e1761..d5d009aaffeb5 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; @@ -32,23 +32,13 @@ class SiPixelClustersCUDA { 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(); } - uint32_t const *c_clusModuleStart() const { return clusModuleStart_d.get(); } - 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/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h index aa06e8dbbd57d..85e8883fa1bd4 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,8 +11,10 @@ class SiPixelDigiErrorsCUDA { public: + 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; @@ -19,23 +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; } - 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(); } + SiPixelErrorCompactVector* error() { return error_d.get(); } + SiPixelErrorCompactVector const* 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; - 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/interface/SiPixelDigisCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h index 04207f3e0b385..2dc1f628bf426 100644 --- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h +++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h @@ -42,14 +42,6 @@ class SiPixelDigisCUDA { 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(); } - 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(); } - cms::cuda::host::unique_ptr adcToHostAsync(cudaStream_t stream) const; cms::cuda::host::unique_ptr clusToHostAsync(cudaStream_t stream) const; cms::cuda::host::unique_ptr pdigiToHostAsync(cudaStream_t stream) const; @@ -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; diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc index 075d408a6f6fc..70bf2e8aa19f5 100644 --- a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc +++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc @@ -7,14 +7,13 @@ #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); - +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); - 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)); @@ -30,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/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 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 @@ - - 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/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/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/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) {} 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..c72c19ce5fda4 --- /dev/null +++ b/DataFormats/SiPixelRawData/interface/SiPixelErrorsSoA.h @@ -0,0 +1,30 @@ +#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) + : error_(error, error + nErrors), formatterErrors_(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/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/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); diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 04072943bf0f8..9a37ec2100661 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -355,7 +355,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) { @@ -390,7 +390,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; } @@ -434,7 +434,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; @@ -450,7 +450,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; @@ -521,7 +521,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, @@ -595,8 +595,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, @@ -615,7 +615,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()) @@ -627,10 +627,10 @@ 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(), - clusters_d.c_moduleStart(), + findClus<<>>(digis_d.moduleInd(), + digis_d.xx(), + digis_d.yy(), + clusters_d.moduleStart(), clusters_d.clusInModule(), clusters_d.moduleId(), digis_d.clus(), @@ -643,10 +643,10 @@ namespace pixelgpudetails { // apply charge cut clusterChargeCut<<>>(digis_d.moduleInd(), - digis_d.c_adc(), - clusters_d.c_moduleStart(), + digis_d.adc(), + clusters_d.moduleStart(), clusters_d.clusInModule(), - clusters_d.c_moduleId(), + clusters_d.moduleId(), digis_d.clus(), wordCounter); cudaCheck(cudaGetLastError()); @@ -657,7 +657,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/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h index 2f52316aa2e78..e06ba8ce735aa 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 SiPixelROCsStatusAndMapping; 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,