Squash Patatrack developments on top of CMSSW_10_4_0

CUDADataFormats/Common/interface/device_unique_ptr.h

@@ -0,0 +1,16 @@
+#ifndef CUDADataFormats_Common_interface_device_unique_ptr_h
+#define CUDADataFormats_Common_interface_device_unique_ptr_h
+
+#include <memory>
+#include <functional>
+
+namespace edm {
+  namespace cuda {
+    namespace device {
+      template <typename T>
+      using unique_ptr = std::unique_ptr<T, std::function<void(void *)>>;
+    }
+  }
+}
+
+#endif

CUDADataFormats/Common/interface/host_unique_ptr.h

@@ -0,0 +1,16 @@
+#ifndef CUDADataFormats_Common_interface_host_unique_ptr_h
+#define CUDADataFormats_Common_interface_host_unique_ptr_h
+
+#include <memory>
+#include <functional>
+
+namespace edm {
+  namespace cuda {
+    namespace host {
+      template <typename T>
+      using unique_ptr = std::unique_ptr<T, std::function<void(void *)>>;
+    }
+  }
+}
+
+#endif

CUDADataFormats/SiPixelCluster/BuildFile.xml

@@ -0,0 +1,8 @@
+<use name="FWCore/ServiceRegistry"/>
+<use name="HeterogeneousCore/CUDAServices"/>
+<use name="cuda-api-wrappers"/>
+
+<export>
+    <lib name="1"/>
+</export>
+

CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h

@@ -0,0 +1,73 @@
+#ifndef CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h
+#define CUDADataFormats_SiPixelCluster_interface_SiPixelClustersCUDA_h
+
+#include "CUDADataFormats/Common/interface/device_unique_ptr.h"
+
+#include <cuda/api_wrappers.h>
+
+class SiPixelClustersCUDA {
+public:
+  SiPixelClustersCUDA() = default;
+  explicit SiPixelClustersCUDA(size_t feds, size_t nelements, cuda::stream_t<>& stream);
+  ~SiPixelClustersCUDA() = default;
+
+  SiPixelClustersCUDA(const SiPixelClustersCUDA&) = delete;
+  SiPixelClustersCUDA& operator=(const SiPixelClustersCUDA&) = delete;
+  SiPixelClustersCUDA(SiPixelClustersCUDA&&) = default;
+  SiPixelClustersCUDA& operator=(SiPixelClustersCUDA&&) = default;
+
+  uint32_t *moduleStart() { return moduleStart_d.get(); }
+  int32_t  *clus() { return clus_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(); }
+  int32_t  const *clus() const { return clus_d.get(); }
+  uint32_t const *clusInModule() const { return clusInModule_d.get(); }
+  uint32_t const *moduleId() const { return moduleId_d.get(); }
+  uint32_t const *clusModuleStart() const { return clusModuleStart_d.get(); }
+
+  uint32_t const *c_moduleStart() const { return moduleStart_d.get(); }
+  int32_t  const *c_clus() const { return clus_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;
+
+#ifdef __CUDACC__
+    __device__ __forceinline__ uint32_t moduleStart(int i) const { return __ldg(moduleStart_+i); }
+    __device__ __forceinline__ int32_t  clus(int i) const { return __ldg(clus_+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); }
+#endif
+
+    friend SiPixelClustersCUDA;
+
+  private:
+    uint32_t const *moduleStart_;
+    int32_t  const *clus_;
+    uint32_t const *clusInModule_;
+    uint32_t const *moduleId_;
+    uint32_t const *clusModuleStart_;
+  };
+
+  DeviceConstView *view() const { return view_d.get(); }
+
+private:
+  edm::cuda::device::unique_ptr<uint32_t[]> moduleStart_d;   // index of the first pixel of each module
+  edm::cuda::device::unique_ptr<int32_t[]>  clus_d;          // cluster id of each pixel
+  edm::cuda::device::unique_ptr<uint32_t[]> clusInModule_d;  // number of clusters found in each module
+  edm::cuda::device::unique_ptr<uint32_t[]> moduleId_d;      // module id of each module
+
+  // originally from rechits
+  edm::cuda::device::unique_ptr<uint32_t[]> clusModuleStart_d;
+
+  edm::cuda::device::unique_ptr<DeviceConstView> view_d;    // "me" pointer
+};
+
+#endif

CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc

@@ -0,0 +1,24 @@
+#include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+
+SiPixelClustersCUDA::SiPixelClustersCUDA(size_t feds, size_t nelements, cuda::stream_t<>& stream) {
+  edm::Service<CUDAService> cs;
+
+  moduleStart_d     = cs->make_device_unique<uint32_t[]>(nelements+1, stream);
+  clus_d            = cs->make_device_unique< int32_t[]>(feds, stream);
+  clusInModule_d    = cs->make_device_unique<uint32_t[]>(nelements, stream);
+  moduleId_d        = cs->make_device_unique<uint32_t[]>(nelements, stream);
+  clusModuleStart_d = cs->make_device_unique<uint32_t[]>(nelements+1, stream);
+
+  auto view = cs->make_host_unique<DeviceConstView>(stream);
+  view->moduleStart_ = moduleStart_d.get();
+  view->clus_ = clus_d.get();
+  view->clusInModule_ = clusInModule_d.get();
+  view->moduleId_ = moduleId_d.get();
+  view->clusModuleStart_ = clusModuleStart_d.get();
+
+  view_d = cs->make_device_unique<DeviceConstView>(stream);
+  cudaMemcpyAsync(view_d.get(), view.get(), sizeof(DeviceConstView), cudaMemcpyDefault, stream.id());
+}

CUDADataFormats/SiPixelDigi/BuildFile.xml

@@ -0,0 +1,7 @@
+<use name="FWCore/ServiceRegistry"/>
+<use name="HeterogeneousCore/CUDAServices"/>
+<use name="cuda-api-wrappers"/>
+
+<export>
+    <lib name="1"/>
+</export>

CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h

@@ -0,0 +1,65 @@
+#ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
+#define CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
+
+#include "CUDADataFormats/Common/interface/device_unique_ptr.h"
+#include "FWCore/Utilities/interface/propagate_const.h"
+
+#include <cuda/api_wrappers.h>
+
+class SiPixelDigisCUDA {
+public:
+  SiPixelDigisCUDA() = default;
+  explicit SiPixelDigisCUDA(size_t nelements, cuda::stream_t<>& stream);
+  ~SiPixelDigisCUDA() = default;
+
+  SiPixelDigisCUDA(const SiPixelDigisCUDA&) = delete;
+  SiPixelDigisCUDA& operator=(const SiPixelDigisCUDA&) = delete;
+  SiPixelDigisCUDA(SiPixelDigisCUDA&&) = default;
+  SiPixelDigisCUDA& operator=(SiPixelDigisCUDA&&) = default;
+
+  uint16_t * xx() { return xx_d.get(); }
+  uint16_t * yy() { return yy_d.get(); }
+  uint16_t * adc() { return adc_d.get(); }
+  uint16_t * moduleInd() { return moduleInd_d.get(); }
+
+  uint16_t const *xx() const { return xx_d.get(); }
+  uint16_t const *yy() const { return yy_d.get(); }
+  uint16_t const *adc() const { return adc_d.get(); }
+  uint16_t const *moduleInd() const { return moduleInd_d.get(); }
+
+  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(); }
+
+  class DeviceConstView {
+  public:
+    DeviceConstView() = default;
+
+#ifdef __CUDACC__
+    __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); }
+#endif
+
+    friend class SiPixelDigisCUDA;
+
+  private:
+    uint16_t const *xx_;
+    uint16_t const *yy_;
+    uint16_t const *adc_;
+    uint16_t const *moduleInd_;
+  };
+
+  const DeviceConstView *view() const { return view_d.get(); }
+
+private:
+  edm::cuda::device::unique_ptr<uint16_t[]> xx_d;        // local coordinates of each pixel
+  edm::cuda::device::unique_ptr<uint16_t[]> yy_d;        //
+  edm::cuda::device::unique_ptr<uint16_t[]> adc_d;       // ADC of each pixel
+  edm::cuda::device::unique_ptr<uint16_t[]> moduleInd_d; // module id of each pixel
+  edm::cuda::device::unique_ptr<DeviceConstView> view_d; // "me" pointer
+};
+
+#endif

CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc

@@ -0,0 +1,25 @@
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+
+#include <cuda_runtime.h>
+
+SiPixelDigisCUDA::SiPixelDigisCUDA(size_t nelements, cuda::stream_t<>& stream) {
+  edm::Service<CUDAService> cs;
+
+  xx_d              = cs->make_device_unique<uint16_t[]>(nelements, stream);
+  yy_d              = cs->make_device_unique<uint16_t[]>(nelements, stream);
+  adc_d             = cs->make_device_unique<uint16_t[]>(nelements, stream);
+  moduleInd_d       = cs->make_device_unique<uint16_t[]>(nelements, stream);
+
+  auto view = cs->make_host_unique<DeviceConstView>(stream);
+  view->xx_ = xx_d.get();
+  view->yy_ = yy_d.get();
+  view->adc_ = adc_d.get();
+  view->moduleInd_ = moduleInd_d.get();
+
+  view_d = cs->make_device_unique<DeviceConstView>(stream);
+  cudaCheck(cudaMemcpyAsync(view_d.get(), view.get(), sizeof(DeviceConstView), cudaMemcpyDefault, stream.id()));
+}

CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h

@@ -0,0 +1,14 @@
+#ifndef CalibTracker_Records_SiPixelGainCalibrationForHLTGPURcd_h
+#define CalibTracker_Records_SiPixelGainCalibrationForHLTGPURcd_h
+
+#include "FWCore/Framework/interface/EventSetupRecordImplementation.h"
+#include "FWCore/Framework/interface/DependentRecordImplementation.h"
+
+#include "CondFormats/DataRecord/interface/SiPixelGainCalibrationForHLTRcd.h"
+#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
+
+#include "boost/mpl/vector.hpp"
+
+class SiPixelGainCalibrationForHLTGPURcd : public edm::eventsetup::DependentRecordImplementation<SiPixelGainCalibrationForHLTGPURcd, boost::mpl::vector<SiPixelGainCalibrationForHLTRcd, TrackerDigiGeometryRecord> > {};
+
+#endif

CalibTracker/Records/src/SiPixelGainCalibrationForHLTGPURcd.cc

@@ -0,0 +1,5 @@
+#include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+
+EVENTSETUP_RECORD_REG(SiPixelGainCalibrationForHLTGPURcd);

CalibTracker/SiPixelESProducers/BuildFile.xml

@@ -7,7 +7,9 @@
 <use   name="DataFormats/SiPixelDigi"/>
 <use   name="CalibTracker/Records"/>
 <use   name="MagneticField/VolumeBasedEngine"/>
+<use   name="HeterogeneousCore/CUDACore"/>
 <use   name="boost"/>
+<use   name="cuda-api-wrappers"/>
 <export>
   <lib   name="1"/>
 </export>

CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h

@@ -0,0 +1,32 @@
+#ifndef CalibTracker_SiPixelESProducers_SiPixelGainCalibrationForHLTGPU_H
+#define CalibTracker_SiPixelESProducers_SiPixelGainCalibrationForHLTGPU_H
+
+#include "HeterogeneousCore/CUDACore/interface/CUDAESProduct.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+
+#include <cuda/api_wrappers.h>
+
+class SiPixelGainCalibrationForHLT;
+class SiPixelGainForHLTonGPU;
+struct SiPixelGainForHLTonGPU_DecodingStructure;
+class TrackerGeometry;
+
+class SiPixelGainCalibrationForHLTGPU {
+public:
+  explicit SiPixelGainCalibrationForHLTGPU(const SiPixelGainCalibrationForHLT& gains, const TrackerGeometry& geom);
+  ~SiPixelGainCalibrationForHLTGPU();
+
+  const SiPixelGainForHLTonGPU *getGPUProductAsync(cuda::stream_t<>& cudaStream) const;
+
+private:
+  const SiPixelGainCalibrationForHLT *gains_ = nullptr;
+  SiPixelGainForHLTonGPU *gainForHLTonHost_ = nullptr;
+  struct GPUData {
+    ~GPUData();
+    SiPixelGainForHLTonGPU *gainForHLTonGPU = nullptr;
+    SiPixelGainForHLTonGPU_DecodingStructure *gainDataOnGPU = nullptr;
+  };
+  CUDAESProduct<GPUData> gpuData_;
+};
+
+#endif

CalibTracker/SiPixelESProducers/plugins/BuildFile.xml

@@ -6,6 +6,8 @@
 <use   name="Geometry/Records"/>
 <use   name="Geometry/TrackerGeometryBuilder"/>
 <use   name="CalibTracker/SiPixelESProducers"/>
+<use   name="HeterogeneousCore/CUDACore"/>
+<use   name="cuda-api-wrappers"/>
 <library   file="*.cc" name="CalibTrackerSiPixelESProducersPlugins">
   <flags   EDM_PLUGIN="1"/>
 </library>

CalibTracker/SiPixelESProducers/plugins/SiPixelGainCalibrationForHLTGPUESProducer.cc

@@ -0,0 +1,47 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
+#include "CondFormats/DataRecord/interface/SiPixelGainCalibrationForHLTRcd.h"
+#include "FWCore/Framework/interface/ESProducer.h"
+#include "FWCore/Framework/interface/EventSetup.h"
+#include "FWCore/Framework/interface/ESHandle.h"
+#include "FWCore/Framework/interface/ModuleFactory.h"
+#include "FWCore/ParameterSet/interface/ParameterSet.h"
+#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
+#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
+
+#include <memory>
+
+class SiPixelGainCalibrationForHLTGPUESProducer: public edm::ESProducer {
+public:
+  explicit SiPixelGainCalibrationForHLTGPUESProducer(const edm::ParameterSet& iConfig);
+  std::unique_ptr<SiPixelGainCalibrationForHLTGPU> produce(const SiPixelGainCalibrationForHLTGPURcd& iRecord);
+
+  static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
+private:
+};
+
+SiPixelGainCalibrationForHLTGPUESProducer::SiPixelGainCalibrationForHLTGPUESProducer(const edm::ParameterSet& iConfig) {
+  setWhatProduced(this);
+}
+
+void SiPixelGainCalibrationForHLTGPUESProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
+  edm::ParameterSetDescription desc;
+  descriptions.add("siPixelGainCalibrationForHLTGPU", desc);
+}
+
+std::unique_ptr<SiPixelGainCalibrationForHLTGPU> SiPixelGainCalibrationForHLTGPUESProducer::produce(const SiPixelGainCalibrationForHLTGPURcd& iRecord) {
+  edm::ESHandle<SiPixelGainCalibrationForHLT> gains;
+  iRecord.getRecord<SiPixelGainCalibrationForHLTRcd>().get(gains);
+
+  edm::ESHandle<TrackerGeometry> geom;
+  iRecord.getRecord<TrackerDigiGeometryRecord>().get(geom);
+
+  return std::make_unique<SiPixelGainCalibrationForHLTGPU>(*gains, *geom);
+}
+
+#include "FWCore/Framework/interface/MakerMacros.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+
+DEFINE_FWK_EVENTSETUP_MODULE(SiPixelGainCalibrationForHLTGPUESProducer);

CalibTracker/SiPixelESProducers/src/ES_SiPixelGainCalibrationForHLTGPU.cc

@@ -0,0 +1,4 @@
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+
+TYPELOOKUP_DATA_REG(SiPixelGainCalibrationForHLTGPU);