diff --git a/CalibTracker/SiPixelESProducers/BuildFile.xml b/CalibTracker/SiPixelESProducers/BuildFile.xml
index ae04d111f1024..6e64a5b4b94ee 100644
--- a/CalibTracker/SiPixelESProducers/BuildFile.xml
+++ b/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"/>
diff --git a/CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h b/CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h
new file mode 100644
index 0000000000000..d86aa93700297
--- /dev/null
+++ b/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
diff --git a/CalibTracker/SiPixelESProducers/plugins/SiPixelROCsStatusAndMappingWrapperESProducer.cc b/CalibTracker/SiPixelESProducers/plugins/SiPixelROCsStatusAndMappingWrapperESProducer.cc
new file mode 100644
index 0000000000000..2c77560a5058e
--- /dev/null
+++ b/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);
diff --git a/CalibTracker/SiPixelESProducers/src/ES_SiPixelROCsStatusAndMappingWrapper.cc b/CalibTracker/SiPixelESProducers/src/ES_SiPixelROCsStatusAndMappingWrapper.cc
new file mode 100644
index 0000000000000..45767102b5958
--- /dev/null
+++ b/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);
diff --git a/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc b/CalibTracker/SiPixelESProducers/src/SiPixelROCsStatusAndMappingWrapper.cc
new file mode 100644
index 0000000000000..1657be1725842
--- /dev/null
+++ b/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)); }
diff --git a/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h b/CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h
new file mode 100644
index 0000000000000..df5b8b24b70dc
--- /dev/null
+++ b/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
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
index df16276520523..5e97610d92286 100644
--- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
+++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
@@ -1,8 +1,15 @@
+// C++ includes
+#include <memory>
+#include <string>
+#include <vector>
+
+// CMSSW includes
 #include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
-#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
+#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "CalibTracker/Records/interface/SiPixelGainCalibrationForHLTGPURcd.h"
+#include "CalibTracker/SiPixelESProducers/interface/SiPixelROCsStatusAndMappingWrapper.h"
 #include "CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h"
 #include "CondFormats/DataRecord/interface/SiPixelFedCablingMapRcd.h"
 #include "CondFormats/SiPixelObjects/interface/SiPixelFedCablingMap.h"
@@ -16,26 +23,22 @@
 #include "FWCore/Framework/interface/ESHandle.h"
 #include "FWCore/Framework/interface/ESTransientHandle.h"
 #include "FWCore/Framework/interface/ESWatcher.h"
-#include "FWCore/Framework/interface/EventSetup.h"
 #include "FWCore/Framework/interface/Event.h"
+#include "FWCore/Framework/interface/EventSetup.h"
 #include "FWCore/Framework/interface/MakerMacros.h"
 #include "FWCore/Framework/interface/stream/EDProducer.h"
 #include "FWCore/MessageLogger/interface/MessageLogger.h"
 #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
-#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
 #include "FWCore/ParameterSet/interface/ParameterSet.h"
+#include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
 #include "FWCore/ServiceRegistry/interface/Service.h"
 #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
 #include "HeterogeneousCore/CUDAServices/interface/CUDAService.h"
-#include "RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPUWrapper.h"
 #include "RecoTracker/Record/interface/CkfComponentsRecord.h"
 
+// local includes
 #include "SiPixelRawToClusterGPUKernel.h"
 
-#include <memory>
-#include <string>
-#include <vector>
-
 class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork> {
 public:
   explicit SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfig);
@@ -58,7 +61,7 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork
   cms::cuda::ContextState ctxState_;
 
   edm::ESWatcher<SiPixelFedCablingMapRcd> recordWatcher_;
-  edm::ESGetToken<SiPixelFedCablingMapGPUWrapper, CkfComponentsRecord> gpuMapToken_;
+  edm::ESGetToken<SiPixelROCsStatusAndMappingWrapper, CkfComponentsRecord> gpuMapToken_;
   edm::ESGetToken<SiPixelGainCalibrationForHLTGPU, SiPixelGainCalibrationForHLTGPURcd> gainsToken_;
   edm::ESGetToken<SiPixelFedCablingMap, SiPixelFedCablingMapRcd> cablingMapToken_;
 
@@ -80,7 +83,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi
     : rawGetToken_(consumes<FEDRawDataCollection>(iConfig.getParameter<edm::InputTag>("InputLabel"))),
       digiPutToken_(produces<cms::cuda::Product<SiPixelDigisCUDA>>()),
       clusterPutToken_(produces<cms::cuda::Product<SiPixelClustersCUDA>>()),
-      gpuMapToken_(esConsumes<SiPixelFedCablingMapGPUWrapper, CkfComponentsRecord>()),
+      gpuMapToken_(esConsumes<SiPixelROCsStatusAndMappingWrapper, CkfComponentsRecord>()),
       gainsToken_(esConsumes<SiPixelGainCalibrationForHLTGPU, SiPixelGainCalibrationForHLTGPURcd>()),
       cablingMapToken_(esConsumes<SiPixelFedCablingMap, SiPixelFedCablingMapRcd>(
           edm::ESInputTag("", iConfig.getParameter<std::string>("CablingMapLabel")))),
@@ -130,7 +133,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent,
   if (hgpuMap->hasQuality() != useQuality_) {
     throw cms::Exception("LogicError")
         << "UseQuality of the module (" << useQuality_
-        << ") differs the one from SiPixelFedCablingMapGPUWrapper. Please fix your configuration.";
+        << ") differs the one from SiPixelROCsStatusAndMappingWrapper. Please fix your configuration.";
   }
   // get the GPU product already here so that the async transfer can begin
   const auto* gpuMap = hgpuMap->getGPUProductAsync(ctx.stream());
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
index f14808dda1e2b..04072943bf0f8 100644
--- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
+++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
@@ -8,29 +8,22 @@
 
 // C++ includes
 #include <cassert>
-#include <chrono>
 #include <cstdio>
 #include <cstdlib>
+#include <cstring>
 #include <fstream>
 #include <iomanip>
 #include <iostream>
-#include <string>
 
 // CUDA includes
-#include <cuda.h>
 #include <cuda_runtime.h>
-#include <thrust/device_vector.h>
-#include <thrust/execution_policy.h>
-#include <thrust/host_vector.h>
-#include <thrust/sort.h>
-#include <thrust/unique.h>
 
 // CMSSW includes
 #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h"
+#include "CondFormats/SiPixelObjects/interface/SiPixelROCsStatusAndMapping.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
-#include "RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPU.h"
 #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h"
 #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h"
 #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h"
@@ -68,7 +61,7 @@ namespace pixelgpudetails {
 
   __device__ bool isBarrel(uint32_t rawId) { return (1 == ((rawId >> 25) & 0x7)); }
 
-  __device__ pixelgpudetails::DetIdGPU getRawId(const SiPixelFedCablingMapGPU *cablingMap,
+  __device__ pixelgpudetails::DetIdGPU getRawId(const SiPixelROCsStatusAndMapping *cablingMap,
                                                 uint8_t fed,
                                                 uint32_t link,
                                                 uint32_t roc) {
@@ -198,7 +191,7 @@ namespace pixelgpudetails {
   __device__ bool dcolIsValid(uint32_t dcol, uint32_t pxid) { return ((dcol < 26) & (2 <= pxid) & (pxid < 162)); }
 
   __device__ uint8_t checkROC(
-      uint32_t errorWord, uint8_t fedId, uint32_t link, const SiPixelFedCablingMapGPU *cablingMap, bool debug = false) {
+      uint32_t errorWord, uint8_t fedId, uint32_t link, const SiPixelROCsStatusAndMapping *cablingMap, bool debug = false) {
     uint8_t errorType = (errorWord >> pixelgpudetails::ROC_shift) & pixelgpudetails::ERROR_mask;
     if (errorType < 25)
       return 0;
@@ -276,7 +269,7 @@ namespace pixelgpudetails {
   __device__ uint32_t getErrRawID(uint8_t fedId,
                                   uint32_t errWord,
                                   uint32_t errorType,
-                                  const SiPixelFedCablingMapGPU *cablingMap,
+                                  const SiPixelROCsStatusAndMapping *cablingMap,
                                   bool debug = false) {
     uint32_t rID = 0xffffffff;
 
@@ -351,7 +344,7 @@ namespace pixelgpudetails {
   }
 
   // Kernel to perform Raw to Digi conversion
-  __global__ void RawToDigi_kernel(const SiPixelFedCablingMapGPU *cablingMap,
+  __global__ void RawToDigi_kernel(const SiPixelROCsStatusAndMapping *cablingMap,
                                    const unsigned char *modToUnp,
                                    const uint32_t wordCounter,
                                    const uint32_t *word,
@@ -524,7 +517,7 @@ namespace pixelgpudetails {
 
   // Interface to outside
   void SiPixelRawToClusterGPUKernel::makeClustersAsync(bool isRun2,
-                                                       const SiPixelFedCablingMapGPU *cablingMap,
+                                                       const SiPixelROCsStatusAndMapping *cablingMap,
                                                        const unsigned char *modToUnp,
                                                        const SiPixelGainForHLTonGPU *gains,
                                                        const WordFedAppender &wordFed,
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
index d214e7784af48..2f52316aa2e78 100644
--- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
+++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
@@ -13,7 +13,7 @@
 #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h"
 #include "DataFormats/SiPixelDigi/interface/PixelErrors.h"
 
-struct SiPixelFedCablingMapGPU;
+struct SiPixelROCsStatusAndMapping;
 class SiPixelGainForHLTonGPU;
 
 namespace pixelgpudetails {
@@ -169,7 +169,7 @@ namespace pixelgpudetails {
     SiPixelRawToClusterGPUKernel& operator=(SiPixelRawToClusterGPUKernel&&) = delete;
 
     void makeClustersAsync(bool isRun2,
-                           const SiPixelFedCablingMapGPU* cablingMap,
+                           const SiPixelROCsStatusAndMapping* cablingMap,
                            const unsigned char* modToUnp,
                            const SiPixelGainForHLTonGPU* gains,
                            const WordFedAppender& wordFed,
diff --git a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py
index 3f8cf314ec2e2..6839e4582bb2b 100644
--- a/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py
+++ b/RecoLocalTracker/SiPixelClusterizer/python/siPixelClustersPreSplitting_cff.py
@@ -3,7 +3,7 @@
 from RecoLocalTracker.SiPixelClusterizer.SiPixelClusterizerPreSplitting_cfi import siPixelClustersPreSplitting
 from RecoLocalTracker.SiPixelClusterizer.siPixelRawToClusterCUDA_cfi import siPixelRawToClusterCUDA as _siPixelRawToClusterCUDA
 from RecoLocalTracker.SiPixelClusterizer.siPixelDigisClustersFromSoA_cfi import siPixelDigisClustersFromSoA as _siPixelDigisClustersFromSoA
-from RecoLocalTracker.SiPixelClusterizer.siPixelFedCablingMapGPUWrapper_cfi import *
+from CalibTracker.SiPixelESProducers.siPixelROCsStatusAndMappingWrapperESProducer_cfi import *
 from CalibTracker.SiPixelESProducers.siPixelGainCalibrationForHLTGPU_cfi import *
 
 siPixelClustersPreSplittingTask = cms.Task(siPixelClustersPreSplitting)