Next prototype of the framework integration (#100)

Provide a mechanism for a chain of modules to share a resource, that can be e.g. CUDA device memory or a CUDA stream.
Minimize data movements between the CPU and the device, and support multiple devices.
Allow the same job configuration to be used on all hardware combinations.

See HeterogeneousCore/CUDACore/README.md for a more detailed description and examples.

CUDADataFormats/Common/BuildFile.xml

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

CUDADataFormats/SiPixelCluster/BuildFile.xml

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

CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h

@@ -0,0 +1,76 @@
+#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 <cuda/api_wrappers.h>
+
+class SiPixelClustersCUDA {
+public:
+  SiPixelClustersCUDA() = default;
+  explicit SiPixelClustersCUDA(size_t maxClusters, cuda::stream_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(); }
+
+  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;
+
+#ifdef __CUDACC__
+    __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); }
+#endif
+
+    friend SiPixelClustersCUDA;
+
+  private:
+    uint32_t const *moduleStart_;
+    uint32_t const *clusInModule_;
+    uint32_t const *moduleId_;
+    uint32_t const *clusModuleStart_;
+  };
+
+  DeviceConstView *view() const { return view_d.get(); }
+
+private:
+  cudautils::device::unique_ptr<uint32_t[]> moduleStart_d;   // index of the first pixel of each module
+  cudautils::device::unique_ptr<uint32_t[]> clusInModule_d;  // number of clusters found in each module
+  cudautils::device::unique_ptr<uint32_t[]> moduleId_d;      // module id of each module
+
+  // originally from rechits
+  cudautils::device::unique_ptr<uint32_t[]> clusModuleStart_d;
+
+  cudautils::device::unique_ptr<DeviceConstView> view_d;    // "me" pointer
+
+  uint32_t nClusters_h;
+};
+
+#endif

CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc

@@ -0,0 +1,23 @@
+#include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+
+SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxClusters, cuda::stream_t<>& stream) {
+  edm::Service<CUDAService> cs;
+
+  moduleStart_d     = cs->make_device_unique<uint32_t[]>(maxClusters+1, stream);
+  clusInModule_d    = cs->make_device_unique<uint32_t[]>(maxClusters, stream);
+  moduleId_d        = cs->make_device_unique<uint32_t[]>(maxClusters, stream);
+  clusModuleStart_d = cs->make_device_unique<uint32_t[]>(maxClusters+1, stream);
+
+  auto view = cs->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 = cs->make_device_unique<DeviceConstView>(stream);
+  cudautils::copyAsync(view_d, view, stream);
+}

CUDADataFormats/SiPixelCluster/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/SiPixelCluster/interface/SiPixelClustersCUDA.h"
+#include "DataFormats/Common/interface/Wrapper.h"
+
+#endif

CUDADataFormats/SiPixelCluster/src/classes_def.xml

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

CUDADataFormats/SiPixelDigi/BuildFile.xml

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

CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h

@@ -0,0 +1,40 @@
+#ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h
+#define CUDADataFormats_SiPixelDigi_interface_SiPixelDigiErrorsCUDA_h
+
+#include "DataFormats/SiPixelDigi/interface/PixelErrors.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/GPUSimpleVector.h"
+
+#include <cuda/api_wrappers.h>
+
+class SiPixelDigiErrorsCUDA {
+public:
+  SiPixelDigiErrorsCUDA() = default;
+  explicit SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cuda::stream_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; }
+
+  GPU::SimpleVector<PixelErrorCompact> *error() { return error_d.get(); }
+  GPU::SimpleVector<PixelErrorCompact> const *error() const { return error_d.get(); }
+  GPU::SimpleVector<PixelErrorCompact> const *c_error() const { return error_d.get(); }
+
+  using HostDataError = std::pair<GPU::SimpleVector<PixelErrorCompact>, cudautils::host::unique_ptr<PixelErrorCompact[]>>;
+  HostDataError dataErrorToHostAsync(cuda::stream_t<>& stream) const;
+
+  void copyErrorToHostAsync(cuda::stream_t<>& stream);
+
+private:
+  cudautils::device::unique_ptr<PixelErrorCompact[]> data_d;
+  cudautils::device::unique_ptr<GPU::SimpleVector<PixelErrorCompact>> error_d;
+  cudautils::host::unique_ptr<GPU::SimpleVector<PixelErrorCompact>> error_h;
+  PixelFormatterErrors formatterErrors_h;
+};
+
+#endif

CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h

@@ -0,0 +1,99 @@
+#ifndef CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
+#define CUDADataFormats_SiPixelDigi_interface_SiPixelDigisCUDA_h
+
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+#include <cuda/api_wrappers.h>
+
+class SiPixelDigisCUDA {
+public:
+  SiPixelDigisCUDA() = default;
+  explicit SiPixelDigisCUDA(size_t maxFedWords, cuda::stream_t<>& stream);
+  ~SiPixelDigisCUDA() = default;
+
+  SiPixelDigisCUDA(const SiPixelDigisCUDA&) = delete;
+  SiPixelDigisCUDA& operator=(const SiPixelDigisCUDA&) = delete;
+  SiPixelDigisCUDA(SiPixelDigisCUDA&&) = default;
+  SiPixelDigisCUDA& operator=(SiPixelDigisCUDA&&) = default;
+
+  void setNModulesDigis(uint32_t nModules, uint32_t nDigis) {
+    nModules_h = nModules;
+    nDigis_h = nDigis;
+  }
+
+  uint32_t nModules() const { return nModules_h; }
+  uint32_t nDigis() const { return nDigis_h; }
+
+  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(); }
+  int32_t  * clus() { return clus_d.get(); }
+  uint32_t * pdigi() { return pdigi_d.get(); }
+  uint32_t * rawIdArr() { return rawIdArr_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(); }
+  int32_t  const *clus() const { return clus_d.get(); } 
+  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(); }
+
+  cudautils::host::unique_ptr<uint16_t[]> adcToHostAsync(cuda::stream_t<>& stream) const;
+  cudautils::host::unique_ptr< int32_t[]> clusToHostAsync(cuda::stream_t<>& stream) const;
+  cudautils::host::unique_ptr<uint32_t[]> pdigiToHostAsync(cuda::stream_t<>& stream) const;
+  cudautils::host::unique_ptr<uint32_t[]> rawIdArrToHostAsync(cuda::stream_t<>& stream) const;
+
+  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); }
+    __device__ __forceinline__ int32_t  clus(int i) const { return __ldg(clus_+i); }
+#endif
+
+    friend class SiPixelDigisCUDA;
+
+  private:
+    uint16_t const *xx_;
+    uint16_t const *yy_;
+    uint16_t const *adc_;
+    uint16_t const *moduleInd_;
+    int32_t  const *clus_;
+  };
+
+  const DeviceConstView *view() const { return view_d.get(); }
+
+private:
+  // These are consumed by downstream device code
+  cudautils::device::unique_ptr<uint16_t[]> xx_d;        // local coordinates of each pixel
+  cudautils::device::unique_ptr<uint16_t[]> yy_d;        //
+  cudautils::device::unique_ptr<uint16_t[]> adc_d;       // ADC of each pixel
+  cudautils::device::unique_ptr<uint16_t[]> moduleInd_d; // module id of each pixel
+  cudautils::device::unique_ptr<int32_t[]>  clus_d;      // cluster id of each pixel
+  cudautils::device::unique_ptr<DeviceConstView> view_d; // "me" pointer
+
+  // These are for CPU output; should we (eventually) place them to a
+  // separate product?
+  cudautils::device::unique_ptr<uint32_t[]> pdigi_d;
+  cudautils::device::unique_ptr<uint32_t[]> rawIdArr_d;
+
+  uint32_t nModules_h = 0;
+  uint32_t nDigis_h = 0;
+};
+
+#endif

CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc

@@ -0,0 +1,44 @@
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/memsetAsync.h"
+
+SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cuda::stream_t<>& stream):
+  formatterErrors_h(std::move(errors))
+{
+  edm::Service<CUDAService> cs;
+
+  error_d = cs->make_device_unique<GPU::SimpleVector<PixelErrorCompact>>(stream);
+  data_d = cs->make_device_unique<PixelErrorCompact[]>(maxFedWords, stream);
+
+  cudautils::memsetAsync(data_d, 0x00, maxFedWords, stream);
+
+  error_h = cs->make_host_unique<GPU::SimpleVector<PixelErrorCompact>>(stream);
+  GPU::make_SimpleVector(error_h.get(), maxFedWords, data_d.get());
+  assert(error_h->size() == 0);
+  assert(error_h->capacity() == static_cast<int>(maxFedWords));
+
+  cudautils::copyAsync(error_d, error_h, stream);
+}
+
+void SiPixelDigiErrorsCUDA::copyErrorToHostAsync(cuda::stream_t<>& stream) {
+  cudautils::copyAsync(error_h, error_d, stream);
+}
+
+SiPixelDigiErrorsCUDA::HostDataError SiPixelDigiErrorsCUDA::dataErrorToHostAsync(cuda::stream_t<>& stream) const {
+  edm::Service<CUDAService> cs;
+  // 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 = cs->make_host_unique<PixelErrorCompact[]>(error_h->capacity(), stream);
+
+  // but transfer only the required amount
+  if(error_h->size() > 0) {
+    cudautils::copyAsync(data, data_d, error_h->size(), stream);
+  }
+  auto err = *error_h;
+  err.set_data(data.get());
+  return HostDataError(std::move(err), std::move(data));
+}

CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc

@@ -0,0 +1,56 @@
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
+
+#include "FWCore/ServiceRegistry/interface/Service.h"
+#include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
+
+SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cuda::stream_t<>& stream) {
+  edm::Service<CUDAService> cs;
+
+  xx_d              = cs->make_device_unique<uint16_t[]>(maxFedWords, stream);
+  yy_d              = cs->make_device_unique<uint16_t[]>(maxFedWords, stream);
+  adc_d             = cs->make_device_unique<uint16_t[]>(maxFedWords, stream);
+  moduleInd_d       = cs->make_device_unique<uint16_t[]>(maxFedWords, stream);
+  clus_d            = cs->make_device_unique< int32_t[]>(maxFedWords, stream);
+
+  pdigi_d           = cs->make_device_unique<uint32_t[]>(maxFedWords, stream);
+  rawIdArr_d        = cs->make_device_unique<uint32_t[]>(maxFedWords, 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->clus_ = clus_d.get();
+
+  view_d = cs->make_device_unique<DeviceConstView>(stream);
+  cudautils::copyAsync(view_d, view, stream);
+}
+
+cudautils::host::unique_ptr<uint16_t[]> SiPixelDigisCUDA::adcToHostAsync(cuda::stream_t<>& stream) const {
+  edm::Service<CUDAService> cs;
+  auto ret = cs->make_host_unique<uint16_t[]>(nDigis(), stream);
+  cudautils::copyAsync(ret, adc_d, nDigis(), stream);
+  return ret;
+}
+
+cudautils::host::unique_ptr<int32_t[]> SiPixelDigisCUDA::clusToHostAsync(cuda::stream_t<>& stream) const {
+  edm::Service<CUDAService> cs;
+  auto ret = cs->make_host_unique<int32_t[]>(nDigis(), stream);
+  cudautils::copyAsync(ret, clus_d, nDigis(), stream);
+  return ret;
+}
+
+cudautils::host::unique_ptr<uint32_t[]> SiPixelDigisCUDA::pdigiToHostAsync(cuda::stream_t<>& stream) const {
+  edm::Service<CUDAService> cs;
+  auto ret = cs->make_host_unique<uint32_t[]>(nDigis(), stream);
+  cudautils::copyAsync(ret, pdigi_d, nDigis(), stream);
+  return ret;
+}
+
+cudautils::host::unique_ptr<uint32_t[]> SiPixelDigisCUDA::rawIdArrToHostAsync(cuda::stream_t<>& stream) const {
+  edm::Service<CUDAService> cs;
+  auto ret = cs->make_host_unique<uint32_t[]>(nDigis(), stream);
+  cudautils::copyAsync(ret, rawIdArr_d, nDigis(), stream);
+  return ret;
+}