diff --git a/CUDADataFormats/Common/BuildFile.xml b/CUDADataFormats/Common/BuildFile.xml index e7a5ba74d80be..c524cada33060 100644 --- a/CUDADataFormats/Common/BuildFile.xml +++ b/CUDADataFormats/Common/BuildFile.xml @@ -1,7 +1,8 @@ - - - - + + + + + - + 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/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) {} 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/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,