Rename SiPixelFedCablingMapGPU to SiPixelROCsStatusAndMapping (#582)

CalibTracker/SiPixelESProducers/BuildFile.xml

@@ -1,14 +1,14 @@
-<use name="FWCore/Framework"/>
-<use name="FWCore/ParameterSet"/>
-<use name="FWCore/MessageLogger"/>
+<use name="cuda"/>
+<use name="CalibTracker/Records"/>
 <use name="CondFormats/DataRecord"/>
 <use name="CondFormats/SiPixelObjects"/>
 <use name="DataFormats/Common"/>
 <use name="DataFormats/SiPixelDigi"/>
-<use name="CalibTracker/Records"/>
-<use name="MagneticField/VolumeBasedEngine"/>
+<use name="FWCore/Framework"/>
+<use name="FWCore/MessageLogger"/>
+<use name="FWCore/ParameterSet"/>
 <use name="HeterogeneousCore/CUDACore"/>
-<use name="cuda"/>
+<use name="MagneticField/VolumeBasedEngine"/>
 <use name="RecoTracker/Record"/>
 <export>
   <lib name="1"/>

CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h

@@ -0,0 +1,56 @@
+#ifndef CalibTracker_SiPixelESProducers_interface_SiPixelROCsStatusAndMappingWrapper_h
+#define CalibTracker_SiPixelESProducers_interface_SiPixelROCsStatusAndMappingWrapper_h
+
+#include <set>
+
+#include <cuda_runtime.h>
+
+#include "CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h"
+#include "HeterogeneousCore/CUDACore/interface/ESProduct.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+
+class SiPixelFedCablingMap;
+class TrackerGeometry;
+class SiPixelQuality;
+
+// TODO: since this has more information than just cabling map, maybe we should invent a better name?
+class SiPixelROCsStatusAndMappingWrapper {
+public:
+  SiPixelROCsStatusAndMappingWrapper(SiPixelFedCablingMap const &cablingMap,
+                                 TrackerGeometry const &trackerGeom,
+                                 SiPixelQuality const *badPixelInfo);
+  ~SiPixelROCsStatusAndMappingWrapper();
+
+  bool hasQuality() const { return hasQuality_; }
+
+  // returns pointer to GPU memory
+  const SiPixelROCsStatusAndMapping *getGPUProductAsync(cudaStream_t cudaStream) const;
+
+  // returns pointer to GPU memory
+  const unsigned char *getModToUnpAllAsync(cudaStream_t cudaStream) const;
+  cms::cuda::device::unique_ptr<unsigned char[]> getModToUnpRegionalAsync(std::set<unsigned int> const &modules,
+                                                                          cudaStream_t cudaStream) const;
+
+private:
+  const SiPixelFedCablingMap *cablingMap_;
+  std::vector<unsigned char, cms::cuda::HostAllocator<unsigned char>> modToUnpDefault;
+  unsigned int size;
+  bool hasQuality_;
+
+  SiPixelROCsStatusAndMapping *cablingMapHost = nullptr;  // pointer to struct in CPU
+
+  struct GPUData {
+    ~GPUData();
+    SiPixelROCsStatusAndMapping *cablingMapDevice = nullptr;  // pointer to struct in GPU
+  };
+  cms::cuda::ESProduct<GPUData> gpuData_;
+
+  struct ModulesToUnpack {
+    ~ModulesToUnpack();
+    unsigned char *modToUnpDefault = nullptr;  // pointer to GPU
+  };
+  cms::cuda::ESProduct<ModulesToUnpack> modToUnp_;
+};
+
+#endif  // CalibTracker_SiPixelESProducers_interface_SiPixelROCsStatusAndMappingWrapper_h

CalibTracker/SiPixelESProducers/plugins/SiPixelROCsStatusAndMappingWrapperESProducer.cc

@@ -0,0 +1,67 @@
+#include <memory>
+
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h"
+#include "CondFormats/DataRecord/interface/SiPixelFedCablingMapRcd.h"
+#include "CondFormats/DataRecord/interface/SiPixelQualityRcd.h"
+#include "FWCore/Framework/interface/ESProducer.h"
+#include "FWCore/Framework/interface/ESTransientHandle.h"
+#include "FWCore/Framework/interface/EventSetup.h"
+#include "FWCore/Framework/interface/ModuleFactory.h"
+#include "FWCore/ParameterSet/interface/ParameterSet.h"
+#include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
+#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
+#include "RecoTracker/Record/interface/CkfComponentsRecord.h"  // TODO: eventually use something more limited
+
+class SiPixelROCsStatusAndMappingWrapperESProducer : public edm::ESProducer {
+public:
+  explicit SiPixelROCsStatusAndMappingWrapperESProducer(const edm::ParameterSet& iConfig);
+  std::unique_ptr<SiPixelROCsStatusAndMappingWrapper> produce(const CkfComponentsRecord& iRecord);
+
+  static void fillDescriptions(edm::ConfigurationDescriptions& descriptions);
+
+private:
+  edm::ESGetToken<SiPixelFedCablingMap, SiPixelFedCablingMapRcd> cablingMapToken_;
+  edm::ESGetToken<SiPixelQuality, SiPixelQualityRcd> qualityToken_;
+  edm::ESGetToken<TrackerGeometry, TrackerDigiGeometryRecord> geometryToken_;
+  bool useQuality_;
+};
+
+SiPixelROCsStatusAndMappingWrapperESProducer::SiPixelROCsStatusAndMappingWrapperESProducer(const edm::ParameterSet& iConfig)
+    : useQuality_(iConfig.getParameter<bool>("UseQualityInfo")) {
+  auto const& component = iConfig.getParameter<std::string>("ComponentName");
+  auto cc = setWhatProduced(this, component);
+  cablingMapToken_ = cc.consumes(edm::ESInputTag{"", iConfig.getParameter<std::string>("CablingMapLabel")});
+  if (useQuality_) {
+    qualityToken_ = cc.consumes();
+  }
+  geometryToken_ = cc.consumes();
+}
+
+void SiPixelROCsStatusAndMappingWrapperESProducer::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
+  edm::ParameterSetDescription desc;
+  desc.add<std::string>("ComponentName", "");
+  desc.add<std::string>("CablingMapLabel", "")->setComment("CablingMap label");
+  desc.add<bool>("UseQualityInfo", false);
+  descriptions.addWithDefaultLabel(desc);
+}
+
+std::unique_ptr<SiPixelROCsStatusAndMappingWrapper> SiPixelROCsStatusAndMappingWrapperESProducer::produce(
+    const CkfComponentsRecord& iRecord) {
+  auto cablingMap = iRecord.getTransientHandle(cablingMapToken_);
+
+  const SiPixelQuality* quality = nullptr;
+  if (useQuality_) {
+    auto qualityInfo = iRecord.getTransientHandle(qualityToken_);
+    quality = qualityInfo.product();
+  }
+
+  auto geom = iRecord.getTransientHandle(geometryToken_);
+
+  return std::make_unique<SiPixelROCsStatusAndMappingWrapper>(*cablingMap, *geom, quality);
+}
+
+#include "FWCore/Framework/interface/MakerMacros.h"
+#include "FWCore/Utilities/interface/typelookup.h"
+#include "FWCore/Framework/interface/eventsetuprecord_registration_macro.h"
+
+DEFINE_FWK_EVENTSETUP_MODULE(SiPixelROCsStatusAndMappingWrapperESProducer);

CalibTracker/SiPixelESProducers/src/ES_SiPixelROCsStatusAndMappingWrapper.cc

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

CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc

@@ -0,0 +1,169 @@
+// C++ includes
+#include <algorithm>
+#include <iomanip>
+#include <iostream>
+#include <vector>
+
+// CUDA includes
+#include <cuda_runtime.h>
+
+// CMSSW includes
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelFedCablingMap.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelFedCablingTree.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelQuality.h"
+#include "FWCore/MessageLogger/interface/MessageLogger.h"
+#include "Geometry/CommonDetUnit/interface/GeomDetType.h"
+#include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
+
+SiPixelROCsStatusAndMappingWrapper::SiPixelROCsStatusAndMappingWrapper(SiPixelFedCablingMap const& cablingMap,
+                                                               TrackerGeometry const& trackerGeom,
+                                                               SiPixelQuality const* badPixelInfo)
+    : cablingMap_(&cablingMap), modToUnpDefault(pixelgpudetails::MAX_SIZE), hasQuality_(badPixelInfo != nullptr) {
+  cudaCheck(cudaMallocHost(&cablingMapHost, sizeof(SiPixelROCsStatusAndMapping)));
+
+  std::vector<unsigned int> const& fedIds = cablingMap.fedIds();
+  std::unique_ptr<SiPixelFedCablingTree> const& cabling = cablingMap.cablingTree();
+
+  unsigned int startFed = *(fedIds.begin());
+  unsigned int endFed = *(fedIds.end() - 1);
+
+  sipixelobjects::CablingPathToDetUnit path;
+  int index = 1;
+
+  for (unsigned int fed = startFed; fed <= endFed; fed++) {
+    for (unsigned int link = 1; link <= pixelgpudetails::MAX_LINK; link++) {
+      for (unsigned int roc = 1; roc <= pixelgpudetails::MAX_ROC; roc++) {
+        path = {fed, link, roc};
+        const sipixelobjects::PixelROC* pixelRoc = cabling->findItem(path);
+        cablingMapHost->fed[index] = fed;
+        cablingMapHost->link[index] = link;
+        cablingMapHost->roc[index] = roc;
+        if (pixelRoc != nullptr) {
+          cablingMapHost->RawId[index] = pixelRoc->rawId();
+          cablingMapHost->rocInDet[index] = pixelRoc->idInDetUnit();
+          modToUnpDefault[index] = false;
+          if (badPixelInfo != nullptr)
+            cablingMapHost->badRocs[index] = badPixelInfo->IsRocBad(pixelRoc->rawId(), pixelRoc->idInDetUnit());
+          else
+            cablingMapHost->badRocs[index] = false;
+        } else {  // store some dummy number
+          cablingMapHost->RawId[index] = 9999;
+          cablingMapHost->rocInDet[index] = 9999;
+          cablingMapHost->badRocs[index] = true;
+          modToUnpDefault[index] = true;
+        }
+        index++;
+      }
+    }
+  }  // end of FED loop
+
+  // Given FedId, Link and idinLnk; use the following formula
+  // to get the RawId and idinDU
+  // index = (FedID-1200) * MAX_LINK* MAX_ROC + (Link-1)* MAX_ROC + idinLnk;
+  // where, MAX_LINK = 48, MAX_ROC = 8 for Phase1 as mentioned Danek's email
+  // FedID varies between 1200 to 1338 (In total 108 FED's)
+  // Link varies between 1 to 48
+  // idinLnk varies between 1 to 8
+
+  for (int i = 1; i < index; i++) {
+    if (cablingMapHost->RawId[i] == 9999) {
+      cablingMapHost->moduleId[i] = 9999;
+    } else {
+      /*
+      std::cout << cablingMapHost->RawId[i] << std::endl;
+      */
+      auto gdet = trackerGeom.idToDetUnit(cablingMapHost->RawId[i]);
+      if (!gdet) {
+        LogDebug("SiPixelROCsStatusAndMapping") << " Not found: " << cablingMapHost->RawId[i] << std::endl;
+        continue;
+      }
+      cablingMapHost->moduleId[i] = gdet->index();
+    }
+    LogDebug("SiPixelROCsStatusAndMapping")
+        << "----------------------------------------------------------------------------" << std::endl;
+    LogDebug("SiPixelROCsStatusAndMapping") << i << std::setw(20) << cablingMapHost->fed[i] << std::setw(20)
+                                        << cablingMapHost->link[i] << std::setw(20) << cablingMapHost->roc[i]
+                                        << std::endl;
+    LogDebug("SiPixelROCsStatusAndMapping") << i << std::setw(20) << cablingMapHost->RawId[i] << std::setw(20)
+                                        << cablingMapHost->rocInDet[i] << std::setw(20) << cablingMapHost->moduleId[i]
+                                        << std::endl;
+    LogDebug("SiPixelROCsStatusAndMapping") << i << std::setw(20) << (bool)cablingMapHost->badRocs[i] << std::setw(20)
+                                        << std::endl;
+    LogDebug("SiPixelROCsStatusAndMapping")
+        << "----------------------------------------------------------------------------" << std::endl;
+  }
+
+  cablingMapHost->size = index - 1;
+}
+
+SiPixelROCsStatusAndMappingWrapper::~SiPixelROCsStatusAndMappingWrapper() { cudaCheck(cudaFreeHost(cablingMapHost)); }
+
+const SiPixelROCsStatusAndMapping* SiPixelROCsStatusAndMappingWrapper::getGPUProductAsync(cudaStream_t cudaStream) const {
+  const auto& data = gpuData_.dataForCurrentDeviceAsync(cudaStream, [this](GPUData& data, cudaStream_t stream) {
+    // allocate
+    cudaCheck(cudaMalloc(&data.cablingMapDevice, sizeof(SiPixelROCsStatusAndMapping)));
+
+    // transfer
+    cudaCheck(cudaMemcpyAsync(
+        data.cablingMapDevice, this->cablingMapHost, sizeof(SiPixelROCsStatusAndMapping), cudaMemcpyDefault, stream));
+  });
+  return data.cablingMapDevice;
+}
+
+const unsigned char* SiPixelROCsStatusAndMappingWrapper::getModToUnpAllAsync(cudaStream_t cudaStream) const {
+  const auto& data =
+      modToUnp_.dataForCurrentDeviceAsync(cudaStream, [this](ModulesToUnpack& data, cudaStream_t stream) {
+        cudaCheck(cudaMalloc((void**)&data.modToUnpDefault, pixelgpudetails::MAX_SIZE_BYTE_BOOL));
+        cudaCheck(cudaMemcpyAsync(data.modToUnpDefault,
+                                  this->modToUnpDefault.data(),
+                                  this->modToUnpDefault.size() * sizeof(unsigned char),
+                                  cudaMemcpyDefault,
+                                  stream));
+      });
+  return data.modToUnpDefault;
+}
+
+cms::cuda::device::unique_ptr<unsigned char[]> SiPixelROCsStatusAndMappingWrapper::getModToUnpRegionalAsync(
+    std::set<unsigned int> const& modules, cudaStream_t cudaStream) const {
+  auto modToUnpDevice = cms::cuda::make_device_unique<unsigned char[]>(pixelgpudetails::MAX_SIZE, cudaStream);
+  auto modToUnpHost = cms::cuda::make_host_unique<unsigned char[]>(pixelgpudetails::MAX_SIZE, cudaStream);
+
+  std::vector<unsigned int> const& fedIds = cablingMap_->fedIds();
+  std::unique_ptr<SiPixelFedCablingTree> const& cabling = cablingMap_->cablingTree();
+
+  unsigned int startFed = *(fedIds.begin());
+  unsigned int endFed = *(fedIds.end() - 1);
+
+  sipixelobjects::CablingPathToDetUnit path;
+  int index = 1;
+
+  for (unsigned int fed = startFed; fed <= endFed; fed++) {
+    for (unsigned int link = 1; link <= pixelgpudetails::MAX_LINK; link++) {
+      for (unsigned int roc = 1; roc <= pixelgpudetails::MAX_ROC; roc++) {
+        path = {fed, link, roc};
+        const sipixelobjects::PixelROC* pixelRoc = cabling->findItem(path);
+        if (pixelRoc != nullptr) {
+          modToUnpHost[index] = (not modules.empty()) and (modules.find(pixelRoc->rawId()) == modules.end());
+        } else {  // store some dummy number
+          modToUnpHost[index] = true;
+        }
+        index++;
+      }
+    }
+  }
+
+  cudaCheck(cudaMemcpyAsync(modToUnpDevice.get(),
+                            modToUnpHost.get(),
+                            pixelgpudetails::MAX_SIZE * sizeof(unsigned char),
+                            cudaMemcpyHostToDevice,
+                            cudaStream));
+  return modToUnpDevice;
+}
+
+SiPixelROCsStatusAndMappingWrapper::GPUData::~GPUData() { cudaCheck(cudaFree(cablingMapDevice)); }
+
+SiPixelROCsStatusAndMappingWrapper::ModulesToUnpack::~ModulesToUnpack() { cudaCheck(cudaFree(modToUnpDefault)); }

CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h

@@ -0,0 +1,26 @@
+#ifndef CondFormats_SiPixelObjects_interface_SiPixelROCsStatusAndMapping_h
+#define CondFormats_SiPixelObjects_interface_SiPixelROCsStatusAndMapping_h
+
+namespace pixelgpudetails {
+  // Maximum fed for phase1 is 150 but not all of them are filled
+  // Update the number FED based on maximum fed found in the cabling map
+  constexpr unsigned int MAX_FED = 150;
+  constexpr unsigned int MAX_LINK = 48;  // maximum links/channels for Phase 1
+  constexpr unsigned int MAX_ROC = 8;
+  constexpr unsigned int MAX_SIZE = MAX_FED * MAX_LINK * MAX_ROC;
+  constexpr unsigned int MAX_SIZE_BYTE_BOOL = MAX_SIZE * sizeof(unsigned char);
+}  // namespace pixelgpudetails
+
+// TODO: since this has more information than just cabling map, maybe we should invent a better name?
+struct SiPixelROCsStatusAndMapping {
+  alignas(128) unsigned int fed[pixelgpudetails::MAX_SIZE];
+  alignas(128) unsigned int link[pixelgpudetails::MAX_SIZE];
+  alignas(128) unsigned int roc[pixelgpudetails::MAX_SIZE];
+  alignas(128) unsigned int RawId[pixelgpudetails::MAX_SIZE];
+  alignas(128) unsigned int rocInDet[pixelgpudetails::MAX_SIZE];
+  alignas(128) unsigned int moduleId[pixelgpudetails::MAX_SIZE];
+  alignas(128) unsigned char badRocs[pixelgpudetails::MAX_SIZE];
+  alignas(128) unsigned int size = 0;
+};
+
+#endif  // CondFormats_SiPixelObjects_interface_SiPixelROCsStatusAndMapping_h