Address some pixel local reco PR review comments

CUDADataFormats/Common/BuildFile.xml

@@ -1,5 +1,7 @@
 <iftool name="cuda">
+  <use name="DataFormats/Common"/>
   <use name="HeterogeneousCore/CUDAUtilities"/>
+  <use name="rootcore"/>
   <export>
     <lib name="1"/>
   </export>

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

CUDADataFormats/Common/src/classes_def.xml

@@ -0,0 +1,4 @@
+<lcgdict>
+  <class name="HostProduct<unsigned int[]>" persistent="false"/>
+  <class name="edm::Wrapper<HostProduct<unsigned int[]>>" persistent="false"/>
+</lcgdict>

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<DeviceConstView> view_d;  // "me" pointer
 
-  uint32_t nClusters_h;
+  uint32_t nClusters_h = 0;
 };
 
 #endif

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<uint32_t[]>(maxClusters + 1, stream);
-  clusInModule_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters, stream);
-  moduleId_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters, stream);
-  clusModuleStart_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters + 1, stream);
-
+SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
+    : moduleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)),
+      clusInModule_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
+      moduleId_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
+      clusModuleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)) {
   auto view = cms::cuda::make_host_unique<DeviceConstView>(stream);
   view->moduleStart_ = moduleStart_d.get();
   view->clusInModule_ = clusInModule_d.get();

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,32 +11,32 @@
 
 class SiPixelDigiErrorsCUDA {
 public:
+  using SiPixelErrorCompactVector = cms::cuda::SimpleVector<SiPixelErrorCompact>;
+
   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;
   SiPixelDigiErrorsCUDA& operator=(const SiPixelDigiErrorsCUDA&) = delete;
   SiPixelDigiErrorsCUDA(SiPixelDigiErrorsCUDA&&) = default;
   SiPixelDigiErrorsCUDA& operator=(SiPixelDigiErrorsCUDA&&) = default;
 
-  const PixelFormatterErrors& formatterErrors() const { return formatterErrors_h; }
+  const SiPixelFormatterErrors& formatterErrors() const { return formatterErrors_h; }
 
-  cms::cuda::SimpleVector<PixelErrorCompact>* error() { return error_d.get(); }
-  cms::cuda::SimpleVector<PixelErrorCompact> const* error() const { return error_d.get(); }
-  cms::cuda::SimpleVector<PixelErrorCompact> 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::SimpleVector<PixelErrorCompact>, cms::cuda::host::unique_ptr<PixelErrorCompact[]>>;
+  using HostDataError = std::pair<SiPixelErrorCompactVector, cms::cuda::host::unique_ptr<SiPixelErrorCompact[]>>;
   HostDataError dataErrorToHostAsync(cudaStream_t stream) const;
 
   void copyErrorToHostAsync(cudaStream_t stream);
 
 private:
-  cms::cuda::device::unique_ptr<PixelErrorCompact[]> data_d;
-  cms::cuda::device::unique_ptr<cms::cuda::SimpleVector<PixelErrorCompact>> error_d;
-  cms::cuda::host::unique_ptr<cms::cuda::SimpleVector<PixelErrorCompact>> error_h;
-  PixelFormatterErrors formatterErrors_h;
+  cms::cuda::device::unique_ptr<SiPixelErrorCompact[]> data_d;
+  cms::cuda::device::unique_ptr<SiPixelErrorCompactVector> error_d;
+  cms::cuda::host::unique_ptr<SiPixelErrorCompactVector> error_h;
+  SiPixelFormatterErrors formatterErrors_h;
 };
 
 #endif

CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h

@@ -42,32 +42,19 @@ 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<uint16_t[]> adcToHostAsync(cudaStream_t stream) const;
   cms::cuda::host::unique_ptr<int32_t[]> clusToHostAsync(cudaStream_t stream) const;
   cms::cuda::host::unique_ptr<uint32_t[]> pdigiToHostAsync(cudaStream_t stream) const;
   cms::cuda::host::unique_ptr<uint32_t[]> rawIdArrToHostAsync(cudaStream_t stream) const;
 
   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<uint32_t[]> pdigi_d;
-  cms::cuda::device::unique_ptr<uint32_t[]> rawIdArr_d;
+  cms::cuda::device::unique_ptr<uint32_t[]> pdigi_d;     // packed digi (row, col, adc) of each pixel
+  cms::cuda::device::unique_ptr<uint32_t[]> rawIdArr_d;  // DetId of each pixel
 
   uint32_t nModules_h = 0;
   uint32_t nDigis_h = 0;

CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc

@@ -7,14 +7,13 @@
 
 #include <cassert>
 
-SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream)
-    : formatterErrors_h(std::move(errors)) {
-  error_d = cms::cuda::make_device_unique<cms::cuda::SimpleVector<PixelErrorCompact>>(stream);
-  data_d = cms::cuda::make_device_unique<PixelErrorCompact[]>(maxFedWords, stream);
-
+SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, SiPixelFormatterErrors errors, cudaStream_t stream)
+    : data_d(cms::cuda::make_device_unique<SiPixelErrorCompact[]>(maxFedWords, stream)),
+      error_d(cms::cuda::make_device_unique<SiPixelErrorCompactVector>(stream)),
+      error_h(cms::cuda::make_host_unique<SiPixelErrorCompactVector>(stream)),
+      formatterErrors_h(std::move(errors)) {
   cms::cuda::memsetAsync(data_d, 0x00, maxFedWords, stream);
 
-  error_h = cms::cuda::make_host_unique<cms::cuda::SimpleVector<PixelErrorCompact>>(stream);
   cms::cuda::make_SimpleVector(error_h.get(), maxFedWords, data_d.get());
   assert(error_h->empty());
   assert(error_h->capacity() == static_cast<int>(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<PixelErrorCompact[]>(error_h->capacity(), stream);
+  auto data = cms::cuda::make_host_unique<SiPixelErrorCompact[]>(error_h->capacity(), stream);
 
   // but transfer only the required amount
   if (not error_h->empty()) {

CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h

@@ -84,11 +84,7 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(uint32_t nH
 
   // if empy do not bother
   if (0 == nHits) {
-    if
-#ifndef __CUDACC__
-        constexpr
-#endif
-        (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
+    if constexpr (std::is_same<Traits, cms::cudacompat::GPUTraits>::value) {
       cms::cuda::copyAsync(m_view, view, stream);
     } else {
       m_view.reset(view.release());  // NOLINT: std::move() breaks CUDA version

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"
 

CUDADataFormats/TrackingRecHit/src/classes_def.xml

@@ -5,6 +5,4 @@
   <class name="edm::Wrapper<TrackingRecHit2DCPU>" persistent="false"/>
   <class name="edm::Wrapper<TrackingRecHit2DHost>" persistent="false"/>
   <class name="edm::Wrapper<cms::cuda::Product<TrackingRecHit2DGPU>>" persistent="false"/>
-  <class name="HostProduct<unsigned int[]>" persistent="false"/>
-  <class name="edm::Wrapper<HostProduct<unsigned int[]>>" persistent="false"/>
 </lcgdict>

DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h

@@ -4,6 +4,12 @@
 #include <cstdint>
 #include <vector>
 
+// 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<int32_t>& clusVector() const { return clus_; }
 
 private:
-  std::vector<uint32_t> pdigi_;
-  std::vector<uint32_t> rawIdArr_;
-  std::vector<uint16_t> adc_;
-  std::vector<int32_t> clus_;
+  std::vector<uint32_t> pdigi_;     // packed digi (row, col, adc) of each pixel
+  std::vector<uint32_t> rawIdArr_;  // DetId of each pixel
+  std::vector<uint16_t> adc_;       // ADC of each pixel
+  std::vector<int32_t> clus_;       // cluster id of each pixel
 };
 
 #endif

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) {}

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"

DataFormats/SiPixelDigi/src/classes_def.xml

@@ -52,7 +52,4 @@
 
    <class name="SiPixelDigisSoA" persistent="false"/>
    <class name="edm::Wrapper<SiPixelDigisSoA>" persistent="false"/>
-
-   <class name="SiPixelDigiErrorsSoA" persistent="false"/>
-   <class name="edm::Wrapper<SiPixelDigiErrorsSoA>" persistent="false"/>
 </lcgdict>

DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h

@@ -0,0 +1,13 @@
+#ifndef DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h
+#define DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h
+
+#include <cstdint>
+
+struct SiPixelErrorCompact {
+  uint32_t rawId;
+  uint32_t word;
+  uint8_t errorType;
+  uint8_t fedId;
+};
+
+#endif  // DataFormats_SiPixelRawData_interface_SiPixelErrorCompact_h

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 <cstdint>
+#include <vector>
+
+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<SiPixelErrorCompact> &errorVector() const { return error_; }
+
+private:
+  std::vector<SiPixelErrorCompact> error_;
+  const SiPixelFormatterErrors *formatterErrors_ = nullptr;
+};
+
+#endif