Synchronise with CMSSW_11_3_0_pre1

CUDADataFormats/Common/BuildFile.xml

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

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/EcalDigi/BuildFile.xml

@@ -0,0 +1,8 @@
+<use name="CUDADataFormats/Common"/>
+<use name="CUDADataFormats/CaloCommon"/>
+<use name="DataFormats/Common"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+  <lib   name="1"/>
+</export>

CUDADataFormats/EcalDigi/interface/DigisCollection.h

@@ -0,0 +1,24 @@
+#ifndef CUDADataFormats_EcalDigi_interface_DigisCollection_h
+#define CUDADataFormats_EcalDigi_interface_DigisCollection_h
+
+#include "CUDADataFormats/CaloCommon/interface/Common.h"
+
+namespace ecal {
+
+  template <typename StoragePolicy>
+  struct DigisCollection : public ::calo::common::AddSize<typename StoragePolicy::TagType> {
+    DigisCollection() = default;
+    DigisCollection(DigisCollection const &) = default;
+    DigisCollection &operator=(DigisCollection const &) = default;
+
+    DigisCollection(DigisCollection &&) = default;
+    DigisCollection &operator=(DigisCollection &&) = default;
+
+    // stride is statically known
+    typename StoragePolicy::template StorageSelector<uint32_t>::type ids;
+    typename StoragePolicy::template StorageSelector<uint16_t>::type data;
+  };
+
+}  // namespace ecal
+
+#endif  // CUDADataFormats_EcalDigi_interface_DigisCollection_h

CUDADataFormats/EcalDigi/src/classes.h

@@ -0,0 +1,3 @@
+#include "CUDADataFormats/Common/interface/Product.h"
+#include "CUDADataFormats/EcalDigi/interface/DigisCollection.h"
+#include "DataFormats/Common/interface/Wrapper.h"

CUDADataFormats/EcalDigi/src/classes_def.xml

@@ -0,0 +1,6 @@
+<lcgdict>
+    <class name="cms::cuda::Product<ecal::DigisCollection<calo::common::ViewStoragePolicy>>" persistent="false" />
+    <class name="cms::cuda::Product<ecal::DigisCollection<calo::common::DevStoragePolicy>>" persistent="false" />
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::DigisCollection<calo::common::ViewStoragePolicy>>>" persistent="false"/>
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::DigisCollection<calo::common::DevStoragePolicy>>>" persistent="false"/>
+</lcgdict>

CUDADataFormats/EcalRecHitSoA/BuildFile.xml

@@ -0,0 +1,10 @@
+<use name="cuda"/>
+<use name="CUDADataFormats/Common"/>
+<use name="CUDADataFormats/CaloCommon"/>
+<use name="DataFormats/Common"/>
+<use name="DataFormats/EcalDigi"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+  <lib   name="1"/>
+</export>

CUDADataFormats/EcalRecHitSoA/interface/EcalRecHit.h

@@ -0,0 +1,45 @@
+#ifndef CUDADataFormats_EcalRecHitSoA_interface_EcalRecHit_h
+#define CUDADataFormats_EcalRecHitSoA_interface_EcalRecHit_h
+
+#include <array>
+#include <vector>
+
+#include "CUDADataFormats/CaloCommon/interface/Common.h"
+#include "CUDADataFormats/EcalRecHitSoA/interface/RecoTypes.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
+
+namespace ecal {
+
+  template <typename StoragePolicy>
+  struct RecHit : public ::calo::common::AddSize<typename StoragePolicy::TagType> {
+    RecHit() = default;
+    RecHit(const RecHit&) = default;
+    RecHit& operator=(const RecHit&) = default;
+
+    RecHit(RecHit&&) = default;
+    RecHit& operator=(RecHit&&) = default;
+
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type energy;
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type time;
+    // should we remove the following, since already included in "extra" ?
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type chi2;
+    typename StoragePolicy::template StorageSelector<uint32_t>::type
+        extra;  // packed uint32_t for timeError, chi2, energyError
+    typename StoragePolicy::template StorageSelector<uint32_t>::type
+        flagBits;  // store rechit condition (see Flags enum) in a bit-wise way
+    typename StoragePolicy::template StorageSelector<uint32_t>::type did;
+
+    template <typename U = typename StoragePolicy::TagType>
+    typename std::enable_if<std::is_same<U, ::calo::common::tags::Vec>::value, void>::type resize(size_t size) {
+      energy.resize(size);
+      time.resize(size);
+      chi2.resize(size);
+      extra.resize(size);
+      flagBits.resize(size);
+      did.resize(size);
+    }
+  };
+
+}  // namespace ecal
+
+#endif  // CUDADataFormats_EcalRecHitSoA_interface_EcalRecHit_h

CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h

@@ -0,0 +1,46 @@
+#ifndef CUDADataFormats_EcalRecHitSoA_interface_EcalUncalibratedRecHit_h
+#define CUDADataFormats_EcalRecHitSoA_interface_EcalUncalibratedRecHit_h
+
+#include <array>
+#include <vector>
+
+#include "CUDADataFormats/CaloCommon/interface/Common.h"
+#include "CUDADataFormats/EcalRecHitSoA/interface/RecoTypes.h"
+#include "DataFormats/EcalDigi/interface/EcalDataFrame.h"
+
+namespace ecal {
+
+  template <typename StoragePolicy>
+  struct UncalibratedRecHit : public ::calo::common::AddSize<typename StoragePolicy::TagType> {
+    UncalibratedRecHit() = default;
+    UncalibratedRecHit(const UncalibratedRecHit&) = default;
+    UncalibratedRecHit& operator=(const UncalibratedRecHit&) = default;
+
+    UncalibratedRecHit(UncalibratedRecHit&&) = default;
+    UncalibratedRecHit& operator=(UncalibratedRecHit&&) = default;
+
+    typename StoragePolicy::template StorageSelector<reco::ComputationScalarType>::type amplitudesAll;
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type amplitude;
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type chi2;
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type pedestal;
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type jitter;
+    typename StoragePolicy::template StorageSelector<reco::StorageScalarType>::type jitterError;
+    typename StoragePolicy::template StorageSelector<uint32_t>::type did;
+    typename StoragePolicy::template StorageSelector<uint32_t>::type flags;
+
+    template <typename U = typename StoragePolicy::TagType>
+    typename std::enable_if<std::is_same<U, ::calo::common::tags::Vec>::value, void>::type resize(size_t size) {
+      amplitudesAll.resize(size * EcalDataFrame::MAXSAMPLES);
+      amplitude.resize(size);
+      pedestal.resize(size);
+      chi2.resize(size);
+      did.resize(size);
+      flags.resize(size);
+      jitter.resize(size);
+      jitterError.resize(size);
+    }
+  };
+
+}  // namespace ecal
+
+#endif  // CUDADataFormats_EcalRecHitSoA_interface_EcalUncalibratedRecHit_h

CUDADataFormats/EcalRecHitSoA/interface/RecoTypes.h

@@ -0,0 +1,13 @@
+#ifndef CUDADataFormats_EcalRecHitSoA_interface_RecoTypes_h
+#define CUDADataFormats_EcalRecHitSoA_interface_RecoTypes_h
+
+namespace ecal {
+  namespace reco {
+
+    using ComputationScalarType = float;
+    using StorageScalarType = float;
+
+  }  // namespace reco
+}  // namespace ecal
+
+#endif  // CUDADataFormats_EcalRecHitSoA_interface_RecoTypes_h

CUDADataFormats/EcalRecHitSoA/src/classes.h

@@ -0,0 +1,4 @@
+#include "CUDADataFormats/Common/interface/Product.h"
+#include "CUDADataFormats/EcalRecHitSoA/interface/EcalRecHit.h"
+#include "CUDADataFormats/EcalRecHitSoA/interface/EcalUncalibratedRecHit.h"
+#include "DataFormats/Common/interface/Wrapper.h"

CUDADataFormats/EcalRecHitSoA/src/classes_def.xml

@@ -0,0 +1,20 @@
+<lcgdict>
+    <class name="cms::cuda::Product<ecal::UncalibratedRecHit<calo::common::ViewStoragePolicy>>" persistent="false"/>
+    <class name="cms::cuda::Product<ecal::UncalibratedRecHit<calo::common::DevStoragePolicy>>" persistent="false"/>
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::UncalibratedRecHit<calo::common::ViewStoragePolicy>>>" persistent="false"/>
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::UncalibratedRecHit<calo::common::DevStoragePolicy>>>" persistent="false"/>
+
+    <class name="cms::cuda::Product<ecal::RecHit<calo::common::ViewStoragePolicy>>" persistent="false"/>
+    <class name="cms::cuda::Product<ecal::RecHit<calo::common::DevStoragePolicy>>" persistent="false"/>
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::RecHit<calo::common::ViewStoragePolicy>>>" persistent="false"/>
+    <class name="edm::Wrapper<cms::cuda::Product<ecal::RecHit<calo::common::DevStoragePolicy>>>" persistent="false"/>
+
+    <class name="ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>" persistent="false"/>
+    <class name="edm::Wrapper<ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>>" persistent="false"/>
+    <class name="ecal::RecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>" persistent="false"/>
+    <class name="edm::Wrapper<ecal::RecHit<calo::common::VecStoragePolicy<calo::common::CUDAHostAllocatorAlias>>>" persistent="false"/>
+    <class name="ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<std::allocator>>" persistent="false"/>
+    <class name="edm::Wrapper<ecal::UncalibratedRecHit<calo::common::VecStoragePolicy<std::allocator>>>" persistent="false"/>
+    <class name="ecal::RecHit<calo::common::VecStoragePolicy<std::allocator>>" persistent="false"/>
+    <class name="edm::Wrapper<ecal::RecHit<calo::common::VecStoragePolicy<std::allocator>>>" persistent="false"/>
+</lcgdict>

CUDADataFormats/SiPixelCluster/BuildFile.xml

@@ -0,0 +1,9 @@
+<use name="cuda"/>
+<use name="rootcore"/>
+<use name="CUDADataFormats/Common"/>
+<use name="DataFormats/Common"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+    <lib name="1"/>
+</export>

CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h

@@ -0,0 +1,63 @@
+#ifndef CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h
+#define CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h
+
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
+
+#include <cuda_runtime.h>
+
+class SiPixelClustersCUDA {
+public:
+  SiPixelClustersCUDA() = default;
+  explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream);
+  ~SiPixelClustersCUDA() = default;
+
+  SiPixelClustersCUDA(const SiPixelClustersCUDA &) = delete;
+  SiPixelClustersCUDA &operator=(const SiPixelClustersCUDA &) = delete;
+  SiPixelClustersCUDA(SiPixelClustersCUDA &&) = default;
+  SiPixelClustersCUDA &operator=(SiPixelClustersCUDA &&) = default;
+
+  void setNClusters(uint32_t nClusters) { nClusters_h = nClusters; }
+
+  uint32_t nClusters() const { return nClusters_h; }
+
+  uint32_t *moduleStart() { return moduleStart_d.get(); }
+  uint32_t *clusInModule() { return clusInModule_d.get(); }
+  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(); }
+
+  class DeviceConstView {
+  public:
+    __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); }
+
+    uint32_t const *moduleStart_;
+    uint32_t const *clusInModule_;
+    uint32_t const *moduleId_;
+    uint32_t const *clusModuleStart_;
+  };
+
+  DeviceConstView *view() const { return view_d.get(); }
+
+private:
+  cms::cuda::device::unique_ptr<uint32_t[]> moduleStart_d;   // index of the first pixel of each module
+  cms::cuda::device::unique_ptr<uint32_t[]> clusInModule_d;  // number of clusters found in each module
+  cms::cuda::device::unique_ptr<uint32_t[]> moduleId_d;      // module id of each module
+
+  // originally from rechits
+  cms::cuda::device::unique_ptr<uint32_t[]> clusModuleStart_d;  // index of the first cluster of each module
+
+  cms::cuda::device::unique_ptr<DeviceConstView> view_d;  // "me" pointer
+
+  uint32_t nClusters_h = 0;
+};
+
+#endif  // CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h

CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h

@@ -0,0 +1,37 @@
+#ifndef CUDADataFormats_SiPixelCluster_interface_gpuClusteringConstants_h
+#define CUDADataFormats_SiPixelCluster_interface_gpuClusteringConstants_h
+
+#include <cstdint>
+#include <limits>
+
+namespace pixelGPUConstants {
+#ifdef GPU_SMALL_EVENTS
+  // kept for testing and debugging
+  constexpr uint32_t maxNumberOfHits = 24 * 1024;
+#else
+  // data at pileup 50 has 18300 +/- 3500 hits; 40000 is around 6 sigma away
+  // tested on MC events with 55-75 pileup events
+  constexpr uint32_t maxNumberOfHits = 48 * 1024;
+#endif
+}  // namespace pixelGPUConstants
+
+namespace gpuClustering {
+#ifdef GPU_SMALL_EVENTS
+  // kept for testing and debugging
+  constexpr uint32_t maxHitsInIter() { return 64; }
+#else
+  // optimized for real data PU 50
+  // tested on MC events with 55-75 pileup events
+  constexpr uint32_t maxHitsInIter() { return 160; }
+#endif
+  constexpr uint32_t maxHitsInModule() { return 1024; }
+
+  constexpr uint16_t maxNumModules = 2000;
+  constexpr int32_t maxNumClustersPerModules = maxHitsInModule();
+  constexpr uint32_t maxNumClusters = pixelGPUConstants::maxNumberOfHits;
+  constexpr uint16_t invalidModuleId = std::numeric_limits<uint16_t>::max() - 1;
+  static_assert(invalidModuleId > maxNumModules);  // invalidModuleId must be > maxNumModules
+
+}  // namespace gpuClustering
+
+#endif  // CUDADataFormats_SiPixelCluster_interface_gpuClusteringConstants_h

CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc

@@ -0,0 +1,19 @@
+#include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+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();
+  view->moduleId_ = moduleId_d.get();
+  view->clusModuleStart_ = clusModuleStart_d.get();
+
+  view_d = cms::cuda::make_device_unique<DeviceConstView>(stream);
+  cms::cuda::copyAsync(view_d, view, stream);
+}

CUDADataFormats/SiPixelCluster/src/classes.h

@@ -0,0 +1,8 @@
+#ifndef CUDADataFormats_SiPixelCluster_src_classes_h
+#define CUDADataFormats_SiPixelCluster_src_classes_h
+
+#include "CUDADataFormats/Common/interface/Product.h"
+#include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif  // CUDADataFormats_SiPixelCluster_src_classes_h

CUDADataFormats/SiPixelCluster/src/classes_def.xml

@@ -0,0 +1,4 @@
+<lcgdict>
+  <class name="cms::cuda::Product<SiPixelClustersCUDA>" persistent="false"/>
+  <class name="edm::Wrapper<cms::cuda::Product<SiPixelClustersCUDA>>" persistent="false"/>
+</lcgdict>

CUDADataFormats/SiPixelDigi/BuildFile.xml

@@ -0,0 +1,10 @@
+<use name="cuda"/>
+<use name="rootcore"/>
+<use name="CUDADataFormats/Common"/>
+<use name="DataFormats/Common"/>
+<use name="DataFormats/SiPixelRawData"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+
+<export>
+    <lib name="1"/>
+</export>