CTD19 developments for review and merging (fixed)

CUDADataFormats/BeamSpot/BuildFile.xml

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

CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h

@@ -0,0 +1,32 @@
+#ifndef CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h
+#define CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h
+
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+
+#include <cuda/api_wrappers.h>
+
+class BeamSpotCUDA {
+public:
+  // alignas(128) doesn't really make sense as there is only one
+  // beamspot per event? 
+  struct Data {
+    float x,y,z;   // position
+    // TODO: add covariance matrix
+
+    float sigmaZ;
+    float beamWidthX, beamWidthY;
+    float dxdz, dydz;
+    float emittanceX, emittanceY;
+    float betaStar;
+  };
+
+  BeamSpotCUDA() = default;
+  BeamSpotCUDA(Data const* data_h, cuda::stream_t<>& stream);
+
+  Data const* data() const { return data_d_.get(); }
+
+private:
+  cudautils::device::unique_ptr<Data> data_d_;
+};
+
+#endif

CUDADataFormats/BeamSpot/src/BeamSpotCUDA.cc

@@ -0,0 +1,11 @@
+#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+
+BeamSpotCUDA::BeamSpotCUDA(Data const* data_h, cuda::stream_t<>& stream) {
+  edm::Service<CUDAService> cs;
+
+  data_d_ = cs->make_device_unique<Data>(stream);
+  cuda::memory::async::copy(data_d_.get(), data_h, sizeof(Data), stream.id());
+}

CUDADataFormats/BeamSpot/src/classes.h

@@ -0,0 +1,8 @@
+#ifndef CUDADataFormats_BeamSpot_classes_h
+#define CUDADataFormats_BeamSpot_classes_h
+
+#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif

CUDADataFormats/BeamSpot/src/classes_def.xml

@@ -0,0 +1,4 @@
+<lcgdict>
+  <class name="CUDAProduct<BeamSpotCUDA>" persistent="false"/>
+  <class name="edm::Wrapper<CUDAProduct<BeamSpotCUDA>>" persistent="false"/>
+</lcgdict>

CUDADataFormats/Common/interface/CUDAProduct.h

@@ -40,11 +40,17 @@ class CUDAProduct: public CUDAProductBase {
   friend class CUDAScopedContext;
   friend class edm::Wrapper<CUDAProduct<T>>;
 
-  explicit CUDAProduct(int device, std::shared_ptr<cuda::stream_t<>> stream, std::shared_ptr<cuda::event_t> event, T data):
-    CUDAProductBase(device, std::move(stream), std::move(event)),
+  explicit CUDAProduct(int device, std::shared_ptr<cuda::stream_t<>> stream, T data):
+    CUDAProductBase(device, std::move(stream)),
     data_(std::move(data))
   {}
 
+  template <typename... Args>
+  explicit CUDAProduct(int device, std::shared_ptr<cuda::stream_t<>> stream, Args&&... args):
+    CUDAProductBase(device, std::move(stream)),
+    data_(std::forward<Args>(args)...)
+  {}
+
   T data_; //!
 };
 

CUDADataFormats/Common/interface/CUDAProductBase.h

@@ -40,12 +40,18 @@ class CUDAProductBase {
   cuda::event_t *event() { return event_.get(); }
 
 protected:
-  explicit CUDAProductBase(int device, std::shared_ptr<cuda::stream_t<>> stream, std::shared_ptr<cuda::event_t> event);
+  explicit CUDAProductBase(int device, std::shared_ptr<cuda::stream_t<>> stream):
+    stream_{std::move(stream)},
+    device_{device}
+  {}
 
 private:
   friend class CUDAScopedContext;
 
-  // Intended to be used only from CUDAScopedContext
+  // The following functions are intended to be used only from CUDAScopedContext
+  void setEvent(std::shared_ptr<cuda::event_t> event) {
+    event_ = std::move(event);
+  }
   const std::shared_ptr<cuda::stream_t<>>& streamPtr() const { return stream_; }
 
   bool mayReuseStream() const {

CUDADataFormats/Common/src/CUDAProductBase.cc

@@ -3,12 +3,6 @@
 #include "FWCore/ServiceRegistry/interface/Service.h"
 #include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
 
-CUDAProductBase::CUDAProductBase(int device, std::shared_ptr<cuda::stream_t<>> stream, std::shared_ptr<cuda::event_t> event):
-  stream_(std::move(stream)),
-  event_(std::move(event)),
-  device_(device)
-{}
-
 bool CUDAProductBase::isAvailable() const {
   // In absence of event, the product was available already at the end
   // of produce() of the producer.

CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h

@@ -0,0 +1,26 @@
+#ifndef CUDADataFormats_SiPixelCluster_gpuClusteringConstants_H
+#define CUDADataFormats_SiPixelCluster_gpuClusteringConstants_H
+
+#include <cstdint>
+
+namespace pixelGPUConstants {
+#ifdef GPU_SMALL_EVENTS
+  constexpr uint32_t maxNumberOfHits = 24*1024;
+#else
+  constexpr uint32_t maxNumberOfHits = 48*1024; // data at pileup 50 has 18300 +/- 3500 hits; 40000 is around 6 sigma away
+#endif
+}
+
+namespace gpuClustering {
+  constexpr uint32_t maxHitsInModule() { return 256;}
+
+  constexpr uint32_t MaxNumModules  = 2000;
+  constexpr uint32_t MaxNumPixels   = 256 * 2000;   // this does not mean maxPixelPerModule == 256!
+  constexpr uint32_t MaxNumClustersPerModules = 1024;
+  constexpr uint32_t MaxHitsInModule = maxHitsInModule();
+  constexpr uint32_t MaxNumClusters = pixelGPUConstants::maxNumberOfHits;
+  constexpr uint16_t InvId          = 9999;         // must be > MaxNumModules
+
+}
+
+#endif 

CUDADataFormats/TrackingRecHit/BuildFile.xml

@@ -0,0 +1,9 @@
+<use name="CUDADataFormats/Common"/>
+<use name="FWCore/ServiceRegistry"/>
+<use name="HeterogeneousCore/CUDAUtilities"/>
+<use name="HeterogeneousCore/CUDAServices"/>
+<use name="cuda-api-wrappers"/>
+<use name="rootcore"/>
+<export>
+    <lib name="1"/>
+</export>

CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h

@@ -0,0 +1,162 @@
+#ifndef CUDADataFormats_TrackingRecHit_TrackingRecHit2DCUDA_H
+#define CUDADataFormats_TrackingRecHit_TrackingRecHit2DCUDA_H
+
+
+#include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/HistoContainer.h"
+
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+#include <cuda/api_wrappers.h>
+
+namespace pixelCPEforGPU {
+  struct ParamsOnGPU;
+}
+
+class TrackingRecHit2DSOAView {
+public:
+
+  static constexpr uint32_t maxHits() { return gpuClustering::MaxNumClusters;}
+  using hindex_type = uint16_t; // if above is <=2^16
+
+  using Hist = HistoContainer<int16_t,128,gpuClustering::MaxNumClusters,8*sizeof(int16_t),uint16_t,10>;
+
+  friend class TrackingRecHit2DCUDA;
+
+  __device__ __forceinline__ uint32_t nHits() const { return m_nHits; }
+
+  __device__ __forceinline__ float & xLocal(int i) { return m_xl[i];}
+  __device__ __forceinline__ float xLocal(int i)  const  { return __ldg(m_xl+i);}
+  __device__ __forceinline__ float & yLocal(int i) { return m_yl[i];}
+  __device__ __forceinline__ float yLocal(int i)  const  { return __ldg(m_yl+i);}
+
+  __device__ __forceinline__ float & xerrLocal(int i) { return m_xerr[i];}
+  __device__ __forceinline__ float xerrLocal(int i)  const  { return __ldg(m_xerr+i);}
+  __device__ __forceinline__ float & yerrLocal(int i) { return m_yerr[i];}
+  __device__ __forceinline__ float yerrLocal(int i)  const  { return __ldg(m_yerr+i);}
+
+  __device__ __forceinline__ float & xGlobal(int i) { return m_xg[i];}
+  __device__ __forceinline__ float xGlobal(int i)  const  { return __ldg(m_xg+i);}
+  __device__ __forceinline__ float & yGlobal(int i) { return m_yg[i];}
+  __device__ __forceinline__ float yGlobal(int i)  const  { return __ldg(m_yg+i);}
+  __device__ __forceinline__ float & zGlobal(int i) { return m_zg[i];}
+  __device__ __forceinline__ float zGlobal(int i)  const  { return __ldg(m_zg+i);}
+  __device__ __forceinline__ float & rGlobal(int i) { return m_rg[i];}
+  __device__ __forceinline__ float rGlobal(int i)  const  { return __ldg(m_rg+i);}
+
+  __device__ __forceinline__ int16_t & iphi(int i) { return m_iphi[i];}
+  __device__ __forceinline__ int16_t  iphi(int i) const { return __ldg(m_iphi+i);}
+
+  __device__ __forceinline__ int32_t & charge(int i) { return m_charge[i];}
+  __device__ __forceinline__ int32_t  charge(int i) const { return __ldg(m_charge+i);}
+  __device__ __forceinline__ int16_t & clusterSizeX(int i) { return m_xsize[i];}
+  __device__ __forceinline__ int16_t  clusterSizeX(int i) const { return __ldg(m_xsize+i);}
+  __device__ __forceinline__ int16_t & clusterSizeY(int i) { return m_ysize[i];}
+  __device__ __forceinline__ int16_t  clusterSizeY(int i) const { return __ldg(m_ysize+i);}
+  __device__ __forceinline__ uint16_t & detectorIndex(int i)  { return m_detInd[i];}
+  __device__ __forceinline__ uint16_t  detectorIndex(int i) const { return __ldg(m_detInd+i);}
+
+
+  __device__ __forceinline__  pixelCPEforGPU::ParamsOnGPU const & cpeParams() const { return  *m_cpeParams; }
+
+  __device__ __forceinline__  uint32_t hitsModuleStart(int i) const { return __ldg(m_hitsModuleStart+i);}
+
+  __device__ __forceinline__  uint32_t * hitsLayerStart() { return  m_hitsLayerStart; }
+  __device__ __forceinline__  uint32_t const * hitsLayerStart() const { return  m_hitsLayerStart; }
+
+
+
+  __device__ __forceinline__ Hist & phiBinner() { return *m_hist; }
+  __device__ __forceinline__ Hist const & phiBinner() const { return  *m_hist; }
+private:
+
+  // local coord
+  float *m_xl, *m_yl;
+  float *m_xerr, *m_yerr;
+
+  // global coord
+  float *m_xg, *m_yg, *m_zg, *m_rg;
+  int16_t * m_iphi;
+
+  // cluster properties
+  int32_t * m_charge;
+  int16_t * m_xsize;
+  int16_t * m_ysize;
+  uint16_t * m_detInd;
+
+
+  // supporting objects
+  pixelCPEforGPU::ParamsOnGPU const * m_cpeParams;    // forwarded from setup, NOT owned
+  uint32_t const * m_hitsModuleStart; // forwarded from clusters
+
+  uint32_t * m_hitsLayerStart;
+
+  Hist * m_hist;
+
+  uint32_t m_nHits;
+};
+
+
+class TrackingRecHit2DCUDA {
+public:
+
+  using Hist = TrackingRecHit2DSOAView::Hist;
+
+  TrackingRecHit2DCUDA() = default;
+
+  explicit 
+  TrackingRecHit2DCUDA(uint32_t nHits, 
+                       pixelCPEforGPU::ParamsOnGPU const * cpeParams,
+                       uint32_t const * hitsModuleStart,
+                      cuda::stream_t<>& stream);
+  ~TrackingRecHit2DCUDA() = default;
+
+  TrackingRecHit2DCUDA(const TrackingRecHit2DCUDA&) = delete;
+  TrackingRecHit2DCUDA& operator=(const TrackingRecHit2DCUDA&) = delete;
+  TrackingRecHit2DCUDA(TrackingRecHit2DCUDA&&) = default;
+  TrackingRecHit2DCUDA& operator=(TrackingRecHit2DCUDA&&) = default;
+
+
+  TrackingRecHit2DSOAView  * view() { return m_view.get(); }
+  TrackingRecHit2DSOAView const * view() const { return m_view.get(); }
+
+  auto  nHits() const { return m_nHits; }
+
+  auto hitsModuleStart() const { return m_hitsModuleStart;}
+  auto hitsLayerStart() { return m_hitsLayerStart; }
+  auto phiBinner()  { return m_hist; }
+  auto iphi() { return m_iphi;}
+
+
+  // only the local coord and detector index
+  cudautils::host::unique_ptr<float[]> localCoordToHostAsync(cuda::stream_t<>& stream) const;
+  cudautils::host::unique_ptr<uint16_t[]> detIndexToHostAsync(cuda::stream_t<>& stream) const;
+  cudautils::host::unique_ptr<uint32_t[]> hitsModuleStartToHostAsync(cuda::stream_t<>& stream) const;
+
+private:
+
+  static constexpr uint32_t n16 = 4;
+  static constexpr uint32_t n32 = 9;
+  static_assert(sizeof(uint32_t) == sizeof(float)); // just stating the obvious
+
+  cudautils::device::unique_ptr<uint16_t[]>  m_store16;
+  cudautils::device::unique_ptr<float[]>  m_store32;
+
+  cudautils::device::unique_ptr<TrackingRecHit2DSOAView::Hist> m_HistStore;
+
+  cudautils::device::unique_ptr<TrackingRecHit2DSOAView> m_view;
+
+  uint32_t m_nHits;
+
+  uint32_t const * m_hitsModuleStart; // needed for legacy  this is on GPU!
+
+  // needed as kernel params...
+  Hist * m_hist;
+  uint32_t * m_hitsLayerStart;
+  int16_t * m_iphi;
+
+};
+
+
+#endif