implement RecHit SOA and move to new framework

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 alignas(128) 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,13 @@
+#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;
+
+  assert(std::abs(data_h->x)<1.f);
+  data_d_ = cs->make_device_unique<Data>(stream);
+  cuda::memory::async::copy(data_d_.get(), data_h, sizeof(Data), stream.id());
+//  cudaStreamSynchronize(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

@@ -45,6 +45,12 @@ class CUDAProduct: public CUDAProductBase {
     data_(std::move(data))
   {}
 
+  template <typename... Args>
+  explicit CUDAProduct(int device, std::shared_ptr<cuda::stream_t<>> stream, std::shared_ptr<cuda::event_t> event, Args&&... args):
+    CUDAProductBase(device, std::move(stream), std::move(event)),
+    data_(std::forward<Args>(args)...)
+  {}
+
   T data_; //!
 };
 

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,8 @@
+<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

CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc

@@ -0,0 +1,87 @@
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+// #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+
+
+TrackingRecHit2DCUDA::TrackingRecHit2DCUDA(
+                      uint32_t nHits,
+                      pixelCPEforGPU::ParamsOnGPU const * cpeParams,
+                      uint32_t const * hitsModuleStart,
+                      cuda::stream_t<>& stream) : m_nHits(nHits), m_hitsModuleStart(hitsModuleStart){
+
+  edm::Service<CUDAService> cs;
+
+  auto view = cs->make_host_unique<TrackingRecHit2DSOAView>(stream);
+  view->m_nHits = nHits;
+  m_view = cs->make_device_unique<TrackingRecHit2DSOAView>(stream);
+
+  // if empy do not bother
+  if (0==nHits) {
+    cudautils::copyAsync(m_view, view, stream);
+    return;
+  }
+
+
+  // the single arrays are not 128 bit alligned...
+  // the hits are actually accessed in order only in building
+  // if ordering is relevant they may have to be stored phi-ordered by layer or so
+  // this will break 1to1 correspondence with cluster and module locality
+  // so unless proven VERY inefficient we keep it ordered as generated
+  m_store16 = cs->make_device_unique<uint16_t[]>(nHits*n16,stream);
+  m_store32 = cs->make_device_unique<float[]>(nHits*n32+11,stream);
+  m_HistStore = cs->make_device_unique<TrackingRecHit2DSOAView::Hist>(stream);
+
+  auto get16 = [&](int i) { return m_store16.get()+i*nHits;};
+  auto get32 = [&](int i) { return m_store32.get()+i*nHits;};
+
+
+ // copy all the pointers
+  m_hist = view->m_hist = m_HistStore.get();
+
+  view->m_cpeParams = cpeParams;
+  view->m_hitsModuleStart = hitsModuleStart;
+
+  view->m_xl = get32(0);
+  view->m_yl = get32(1);
+  view->m_xerr = get32(2);
+  view->m_yerr = get32(3);
+
+  view->m_xg = get32(4);
+  view->m_yg = get32(5);
+  view->m_zg = get32(6);
+  view->m_rg = get32(7);
+
+  m_iphi = view->m_iphi = reinterpret_cast<int16_t *>(get16(0));
+
+  view->m_charge = reinterpret_cast<int32_t *>(get32(8));
+  view->m_xsize = reinterpret_cast<int16_t *>(get16(2));
+  view->m_ysize = reinterpret_cast<int16_t *>(get16(3));
+  view->m_detInd = get16(1);
+
+  m_hitsLayerStart = view->m_hitsLayerStart = reinterpret_cast<uint32_t *>(get32(n32));
+
+  // transfer view
+  cudautils::copyAsync(m_view, view, stream);
+  // cudaStreamSynchronize(stream.id());
+
+}
+
+
+cudautils::host::unique_ptr<float[]> TrackingRecHit2DCUDA::localCoordToHostAsync(cuda::stream_t<>& stream) const {
+  edm::Service<CUDAService> cs;
+  auto ret = cs->make_host_unique<float[]>(4*nHits(), stream);
+  cudautils::copyAsync(ret, m_store32, 4*nHits(), stream);
+  return ret;
+}
+
+cudautils::host::unique_ptr<uint32_t[]> TrackingRecHit2DCUDA::hitsModuleStartToHostAsync(cuda::stream_t<>& stream) const{
+  edm::Service<CUDAService> cs;
+  auto ret = cs->make_host_unique<uint32_t[]>(2001, stream);
+  cudaMemcpyAsync(ret.get(), m_hitsModuleStart, 4*2001, cudaMemcpyDefault,stream.id());
+  //cudautils::copyAsync(ret, m_hitsModuleStart, 2001, stream);
+  return ret;
+
+}

CUDADataFormats/TrackingRecHit/src/classes.h

@@ -0,0 +1,8 @@
+#ifndef CUDADataFormats_SiPixelCluster_classes_h
+#define CUDADataFormats_SiPixelCluster_classes_h
+
+#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif

CUDADataFormats/TrackingRecHit/src/classes_def.xml

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

CUDADataFormats/TrackingRecHit/test/BuildFile.xml

@@ -0,0 +1,3 @@
+<use name="CUDADataFormats/TrackingRecHit"/>
+<flags CUDA_FLAGS="-g -DGPU_DEBUG"/>
+<bin file="TrackingRecHit2DCUDA_t.cpp TrackingRecHit2DCUDA_t.cu" name="TrackingRecHit2DCUDA_t"/>