diff --git a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
index d3650e164d44e..dbfb5ff5e1761 100644
--- a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
+++ b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
@@ -58,14 +58,14 @@ class SiPixelClustersCUDA {
   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
+  cms::cuda::device::unique_ptr<uint32_t[]> moduleStart_d;   // index of the first pixel of each module
+  cms::cuda::device::unique_ptr<uint32_t[]> clusInModule_d;  // number of clusters found in each module
+  cms::cuda::device::unique_ptr<uint32_t[]> moduleId_d;      // module id of each module
 
   // originally from rechits
-  cudautils::device::unique_ptr<uint32_t[]> clusModuleStart_d;  // index of the first cluster of each module
+  cms::cuda::device::unique_ptr<uint32_t[]> clusModuleStart_d;  // index of the first cluster of each module
 
-  cudautils::device::unique_ptr<DeviceConstView> view_d;  // "me" pointer
+  cms::cuda::device::unique_ptr<DeviceConstView> view_d;  // "me" pointer
 
   uint32_t nClusters_h;
 };
diff --git a/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc b/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc
index c814cd4a2e131..7bef9d0d8a52f 100644
--- a/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc
+++ b/CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc
@@ -5,17 +5,17 @@
 #include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 
 SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxClusters, cudaStream_t stream) {
-  moduleStart_d = cudautils::make_device_unique<uint32_t[]>(maxClusters + 1, stream);
-  clusInModule_d = cudautils::make_device_unique<uint32_t[]>(maxClusters, stream);
-  moduleId_d = cudautils::make_device_unique<uint32_t[]>(maxClusters, stream);
-  clusModuleStart_d = cudautils::make_device_unique<uint32_t[]>(maxClusters + 1, stream);
+  moduleStart_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters + 1, stream);
+  clusInModule_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters, stream);
+  moduleId_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters, stream);
+  clusModuleStart_d = cms::cuda::make_device_unique<uint32_t[]>(maxClusters + 1, stream);
 
-  auto view = cudautils::make_host_unique<DeviceConstView>(stream);
+  auto view = cms::cuda::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 = cudautils::make_device_unique<DeviceConstView>(stream);
-  cudautils::copyAsync(view_d, view, stream);
+  view_d = cms::cuda::make_device_unique<DeviceConstView>(stream);
+  cms::cuda::copyAsync(view_d, view, stream);
 }
diff --git a/CUDADataFormats/SiPixelCluster/src/classes.h b/CUDADataFormats/SiPixelCluster/src/classes.h
index 08d46244adc7d..0698cb103dab9 100644
--- a/CUDADataFormats/SiPixelCluster/src/classes.h
+++ b/CUDADataFormats/SiPixelCluster/src/classes.h
@@ -1,7 +1,7 @@
 #ifndef CUDADataFormats_SiPixelCluster_classes_h
 #define CUDADataFormats_SiPixelCluster_classes_h
 
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
 #include "DataFormats/Common/interface/Wrapper.h"
 
diff --git a/CUDADataFormats/SiPixelCluster/src/classes_def.xml b/CUDADataFormats/SiPixelCluster/src/classes_def.xml
index ba0706ac4b8aa..70decb9f27df7 100644
--- a/CUDADataFormats/SiPixelCluster/src/classes_def.xml
+++ b/CUDADataFormats/SiPixelCluster/src/classes_def.xml
@@ -1,4 +1,4 @@
 <lcgdict>
-  <class name="CUDAProduct<SiPixelClustersCUDA>" persistent="false"/>
-  <class name="edm::Wrapper<CUDAProduct<SiPixelClustersCUDA>>" persistent="false"/>
+  <class name="cms::cuda::Product<SiPixelClustersCUDA>" persistent="false"/>
+  <class name="edm::Wrapper<cms::cuda::Product<SiPixelClustersCUDA>>" persistent="false"/>
 </lcgdict>
diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h
index 7c18d58a3fc12..1557fd64750e7 100644
--- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h
+++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h
@@ -26,15 +26,15 @@ class SiPixelDigiErrorsCUDA {
   GPU::SimpleVector<PixelErrorCompact> const* c_error() const { return error_d.get(); }
 
   using HostDataError =
-      std::pair<GPU::SimpleVector<PixelErrorCompact>, cudautils::host::unique_ptr<PixelErrorCompact[]>>;
+      std::pair<GPU::SimpleVector<PixelErrorCompact>, cms::cuda::host::unique_ptr<PixelErrorCompact[]>>;
   HostDataError dataErrorToHostAsync(cudaStream_t stream) const;
 
   void copyErrorToHostAsync(cudaStream_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;
+  cms::cuda::device::unique_ptr<PixelErrorCompact[]> data_d;
+  cms::cuda::device::unique_ptr<GPU::SimpleVector<PixelErrorCompact>> error_d;
+  cms::cuda::host::unique_ptr<GPU::SimpleVector<PixelErrorCompact>> error_h;
   PixelFormatterErrors formatterErrors_h;
 };
 
diff --git a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h
index 47efe634ad93d..04207f3e0b385 100644
--- a/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h
+++ b/CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h
@@ -50,10 +50,10 @@ class SiPixelDigisCUDA {
   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(cudaStream_t stream) const;
-  cudautils::host::unique_ptr<int32_t[]> clusToHostAsync(cudaStream_t stream) const;
-  cudautils::host::unique_ptr<uint32_t[]> pdigiToHostAsync(cudaStream_t stream) const;
-  cudautils::host::unique_ptr<uint32_t[]> rawIdArrToHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<uint16_t[]> adcToHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<int32_t[]> clusToHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<uint32_t[]> pdigiToHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<uint32_t[]> rawIdArrToHostAsync(cudaStream_t stream) const;
 
   class DeviceConstView {
   public:
@@ -79,17 +79,17 @@ class SiPixelDigisCUDA {
 
 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
+  cms::cuda::device::unique_ptr<uint16_t[]> xx_d;         // local coordinates of each pixel
+  cms::cuda::device::unique_ptr<uint16_t[]> yy_d;         //
+  cms::cuda::device::unique_ptr<uint16_t[]> adc_d;        // ADC of each pixel
+  cms::cuda::device::unique_ptr<uint16_t[]> moduleInd_d;  // module id of each pixel
+  cms::cuda::device::unique_ptr<int32_t[]> clus_d;        // cluster id of each pixel
+  cms::cuda::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;
+  cms::cuda::device::unique_ptr<uint32_t[]> pdigi_d;
+  cms::cuda::device::unique_ptr<uint32_t[]> rawIdArr_d;
 
   uint32_t nModules_h = 0;
   uint32_t nDigis_h = 0;
diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc
index 7640348c15f08..ffef71092f6c9 100644
--- a/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc
+++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigiErrorsCUDA.cc
@@ -9,32 +9,32 @@
 
 SiPixelDigiErrorsCUDA::SiPixelDigiErrorsCUDA(size_t maxFedWords, PixelFormatterErrors errors, cudaStream_t stream)
     : formatterErrors_h(std::move(errors)) {
-  error_d = cudautils::make_device_unique<GPU::SimpleVector<PixelErrorCompact>>(stream);
-  data_d = cudautils::make_device_unique<PixelErrorCompact[]>(maxFedWords, stream);
+  error_d = cms::cuda::make_device_unique<GPU::SimpleVector<PixelErrorCompact>>(stream);
+  data_d = cms::cuda::make_device_unique<PixelErrorCompact[]>(maxFedWords, stream);
 
-  cudautils::memsetAsync(data_d, 0x00, maxFedWords, stream);
+  cms::cuda::memsetAsync(data_d, 0x00, maxFedWords, stream);
 
-  error_h = cudautils::make_host_unique<GPU::SimpleVector<PixelErrorCompact>>(stream);
+  error_h = cms::cuda::make_host_unique<GPU::SimpleVector<PixelErrorCompact>>(stream);
   GPU::make_SimpleVector(error_h.get(), maxFedWords, data_d.get());
   assert(error_h->empty());
   assert(error_h->capacity() == static_cast<int>(maxFedWords));
 
-  cudautils::copyAsync(error_d, error_h, stream);
+  cms::cuda::copyAsync(error_d, error_h, stream);
 }
 
 void SiPixelDigiErrorsCUDA::copyErrorToHostAsync(cudaStream_t stream) {
-  cudautils::copyAsync(error_h, error_d, stream);
+  cms::cuda::copyAsync(error_h, error_d, stream);
 }
 
 SiPixelDigiErrorsCUDA::HostDataError SiPixelDigiErrorsCUDA::dataErrorToHostAsync(cudaStream_t stream) const {
   // 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 = cudautils::make_host_unique<PixelErrorCompact[]>(error_h->capacity(), stream);
+  auto data = cms::cuda::make_host_unique<PixelErrorCompact[]>(error_h->capacity(), stream);
 
   // but transfer only the required amount
   if (not error_h->empty()) {
-    cudautils::copyAsync(data, data_d, error_h->size(), stream);
+    cms::cuda::copyAsync(data, data_d, error_h->size(), stream);
   }
   auto err = *error_h;
   err.set_data(data.get());
diff --git a/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc b/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc
index a8aab7ab5a4b8..664364b6ff25a 100644
--- a/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc
+++ b/CUDADataFormats/SiPixelDigi/src/SiPixelDigisCUDA.cc
@@ -5,46 +5,46 @@
 #include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
 
 SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream) {
-  xx_d = cudautils::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  yy_d = cudautils::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  adc_d = cudautils::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  moduleInd_d = cudautils::make_device_unique<uint16_t[]>(maxFedWords, stream);
-  clus_d = cudautils::make_device_unique<int32_t[]>(maxFedWords, stream);
+  xx_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
+  yy_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
+  adc_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
+  moduleInd_d = cms::cuda::make_device_unique<uint16_t[]>(maxFedWords, stream);
+  clus_d = cms::cuda::make_device_unique<int32_t[]>(maxFedWords, stream);
 
-  pdigi_d = cudautils::make_device_unique<uint32_t[]>(maxFedWords, stream);
-  rawIdArr_d = cudautils::make_device_unique<uint32_t[]>(maxFedWords, stream);
+  pdigi_d = cms::cuda::make_device_unique<uint32_t[]>(maxFedWords, stream);
+  rawIdArr_d = cms::cuda::make_device_unique<uint32_t[]>(maxFedWords, stream);
 
-  auto view = cudautils::make_host_unique<DeviceConstView>(stream);
+  auto view = cms::cuda::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 = cudautils::make_device_unique<DeviceConstView>(stream);
-  cudautils::copyAsync(view_d, view, stream);
+  view_d = cms::cuda::make_device_unique<DeviceConstView>(stream);
+  cms::cuda::copyAsync(view_d, view, stream);
 }
 
-cudautils::host::unique_ptr<uint16_t[]> SiPixelDigisCUDA::adcToHostAsync(cudaStream_t stream) const {
-  auto ret = cudautils::make_host_unique<uint16_t[]>(nDigis(), stream);
-  cudautils::copyAsync(ret, adc_d, nDigis(), stream);
+cms::cuda::host::unique_ptr<uint16_t[]> SiPixelDigisCUDA::adcToHostAsync(cudaStream_t stream) const {
+  auto ret = cms::cuda::make_host_unique<uint16_t[]>(nDigis(), stream);
+  cms::cuda::copyAsync(ret, adc_d, nDigis(), stream);
   return ret;
 }
 
-cudautils::host::unique_ptr<int32_t[]> SiPixelDigisCUDA::clusToHostAsync(cudaStream_t stream) const {
-  auto ret = cudautils::make_host_unique<int32_t[]>(nDigis(), stream);
-  cudautils::copyAsync(ret, clus_d, nDigis(), stream);
+cms::cuda::host::unique_ptr<int32_t[]> SiPixelDigisCUDA::clusToHostAsync(cudaStream_t stream) const {
+  auto ret = cms::cuda::make_host_unique<int32_t[]>(nDigis(), stream);
+  cms::cuda::copyAsync(ret, clus_d, nDigis(), stream);
   return ret;
 }
 
-cudautils::host::unique_ptr<uint32_t[]> SiPixelDigisCUDA::pdigiToHostAsync(cudaStream_t stream) const {
-  auto ret = cudautils::make_host_unique<uint32_t[]>(nDigis(), stream);
-  cudautils::copyAsync(ret, pdigi_d, nDigis(), stream);
+cms::cuda::host::unique_ptr<uint32_t[]> SiPixelDigisCUDA::pdigiToHostAsync(cudaStream_t stream) const {
+  auto ret = cms::cuda::make_host_unique<uint32_t[]>(nDigis(), stream);
+  cms::cuda::copyAsync(ret, pdigi_d, nDigis(), stream);
   return ret;
 }
 
-cudautils::host::unique_ptr<uint32_t[]> SiPixelDigisCUDA::rawIdArrToHostAsync(cudaStream_t stream) const {
-  auto ret = cudautils::make_host_unique<uint32_t[]>(nDigis(), stream);
-  cudautils::copyAsync(ret, rawIdArr_d, nDigis(), stream);
+cms::cuda::host::unique_ptr<uint32_t[]> SiPixelDigisCUDA::rawIdArrToHostAsync(cudaStream_t stream) const {
+  auto ret = cms::cuda::make_host_unique<uint32_t[]>(nDigis(), stream);
+  cms::cuda::copyAsync(ret, rawIdArr_d, nDigis(), stream);
   return ret;
 }
diff --git a/CUDADataFormats/SiPixelDigi/src/classes.h b/CUDADataFormats/SiPixelDigi/src/classes.h
index 41b135640b883..fca0811e4650f 100644
--- a/CUDADataFormats/SiPixelDigi/src/classes.h
+++ b/CUDADataFormats/SiPixelDigi/src/classes.h
@@ -1,7 +1,7 @@
 #ifndef CUDADataFormats_SiPixelDigi_classes_h
 #define CUDADataFormats_SiPixelDigi_classes_h
 
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
 #include "DataFormats/Common/interface/Wrapper.h"
diff --git a/CUDADataFormats/SiPixelDigi/src/classes_def.xml b/CUDADataFormats/SiPixelDigi/src/classes_def.xml
index 9d6816ed3b14c..ff775afdc2046 100644
--- a/CUDADataFormats/SiPixelDigi/src/classes_def.xml
+++ b/CUDADataFormats/SiPixelDigi/src/classes_def.xml
@@ -1,7 +1,7 @@
 <lcgdict>
-  <class name="CUDAProduct<SiPixelDigisCUDA>" persistent="false"/>
-  <class name="edm::Wrapper<CUDAProduct<SiPixelDigisCUDA>>" persistent="false"/>
+  <class name="cms::cuda::Product<SiPixelDigisCUDA>" persistent="false"/>
+  <class name="edm::Wrapper<cms::cuda::Product<SiPixelDigisCUDA>>" persistent="false"/>
 
-  <class name="CUDAProduct<SiPixelDigiErrorsCUDA>" persistent="false"/>
-  <class name="edm::Wrapper<CUDAProduct<SiPixelDigiErrorsCUDA>>" persistent="false"/>
+  <class name="cms::cuda::Product<SiPixelDigiErrorsCUDA>" persistent="false"/>
+  <class name="edm::Wrapper<cms::cuda::Product<SiPixelDigiErrorsCUDA>>" persistent="false"/>
 </lcgdict>
diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h
index aa551f21b4aad..955f97ca6bd54 100644
--- a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h
+++ b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DHeterogeneous.h
@@ -37,9 +37,9 @@ class TrackingRecHit2DHeterogeneous {
   auto iphi() { return m_iphi; }
 
   // only the local coord and detector index
-  cudautils::host::unique_ptr<float[]> localCoordToHostAsync(cudaStream_t stream) const;
-  cudautils::host::unique_ptr<uint16_t[]> detIndexToHostAsync(cudaStream_t stream) const;
-  cudautils::host::unique_ptr<uint32_t[]> hitsModuleStartToHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<float[]> localCoordToHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<uint16_t[]> detIndexToHostAsync(cudaStream_t stream) const;
+  cms::cuda::host::unique_ptr<uint32_t[]> hitsModuleStartToHostAsync(cudaStream_t stream) const;
 
 private:
   static constexpr uint32_t n16 = 4;
@@ -89,7 +89,7 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(uint32_t nH
         constexpr
 #endif
         (std::is_same<Traits, cudaCompat::GPUTraits>::value) {
-      cudautils::copyAsync(m_view, view, stream);
+      cms::cuda::copyAsync(m_view, view, stream);
     } else {
       m_view.reset(view.release());  // NOLINT: std::move() breaks CUDA version
     }
@@ -136,7 +136,7 @@ TrackingRecHit2DHeterogeneous<Traits>::TrackingRecHit2DHeterogeneous(uint32_t nH
       constexpr
 #endif
       (std::is_same<Traits, cudaCompat::GPUTraits>::value) {
-    cudautils::copyAsync(m_view, view, stream);
+    cms::cuda::copyAsync(m_view, view, stream);
   } else {
     m_view.reset(view.release());  // NOLINT: std::move() breaks CUDA version
   }
diff --git a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc
index e6f223bfec4e3..7b04ed2d530a0 100644
--- a/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc
+++ b/CUDADataFormats/TrackingRecHit/src/TrackingRecHit2DCUDA.cc
@@ -5,15 +5,15 @@
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
 
 template <>
-cudautils::host::unique_ptr<float[]> TrackingRecHit2DCUDA::localCoordToHostAsync(cudaStream_t stream) const {
-  auto ret = cudautils::make_host_unique<float[]>(4 * nHits(), stream);
-  cudautils::copyAsync(ret, m_store32, 4 * nHits(), stream);
+cms::cuda::host::unique_ptr<float[]> TrackingRecHit2DCUDA::localCoordToHostAsync(cudaStream_t stream) const {
+  auto ret = cms::cuda::make_host_unique<float[]>(4 * nHits(), stream);
+  cms::cuda::copyAsync(ret, m_store32, 4 * nHits(), stream);
   return ret;
 }
 
 template <>
-cudautils::host::unique_ptr<uint32_t[]> TrackingRecHit2DCUDA::hitsModuleStartToHostAsync(cudaStream_t stream) const {
-  auto ret = cudautils::make_host_unique<uint32_t[]>(2001, stream);
+cms::cuda::host::unique_ptr<uint32_t[]> TrackingRecHit2DCUDA::hitsModuleStartToHostAsync(cudaStream_t stream) const {
+  auto ret = cms::cuda::make_host_unique<uint32_t[]>(2001, stream);
   cudaCheck(cudaMemcpyAsync(ret.get(), m_hitsModuleStart, 4 * 2001, cudaMemcpyDefault, stream));
   return ret;
 }
diff --git a/CUDADataFormats/TrackingRecHit/src/classes.h b/CUDADataFormats/TrackingRecHit/src/classes.h
index 90cfd0945d76e..d80226ec7a14b 100644
--- a/CUDADataFormats/TrackingRecHit/src/classes.h
+++ b/CUDADataFormats/TrackingRecHit/src/classes.h
@@ -1,7 +1,7 @@
 #ifndef CUDADataFormats_SiPixelCluster_src_classes_h
 #define CUDADataFormats_SiPixelCluster_src_classes_h
 
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/Common/interface/HostProduct.h"
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
 #include "DataFormats/Common/interface/Wrapper.h"
diff --git a/CUDADataFormats/TrackingRecHit/src/classes_def.xml b/CUDADataFormats/TrackingRecHit/src/classes_def.xml
index 4e8325ddce87e..02b0eb37d157b 100644
--- a/CUDADataFormats/TrackingRecHit/src/classes_def.xml
+++ b/CUDADataFormats/TrackingRecHit/src/classes_def.xml
@@ -1,10 +1,10 @@
 <lcgdict>
   <class name="TrackingRecHit2DCPU" persistent="false"/>
   <class name="TrackingRecHit2DHost" persistent="false"/>
-  <class name="CUDAProduct<TrackingRecHit2DGPU>" persistent="false"/>
+  <class name="cms::cuda::Product<TrackingRecHit2DGPU>" persistent="false"/>
   <class name="edm::Wrapper<TrackingRecHit2DCPU>" persistent="false"/>
   <class name="edm::Wrapper<TrackingRecHit2DHost>" persistent="false"/>
-  <class name="edm::Wrapper<CUDAProduct<TrackingRecHit2DGPU>>" persistent="false"/>
+  <class name="edm::Wrapper<cms::cuda::Product<TrackingRecHit2DGPU>>" persistent="false"/>
   <class name="HostProduct<unsigned int[]>" persistent="false"/>
   <class name="edm::Wrapper<HostProduct<unsigned int[]>>" persistent="false"/>
 </lcgdict>
diff --git a/CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp b/CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp
index 592f0267c2f7d..32af6c181ae68 100644
--- a/CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp
+++ b/CUDADataFormats/TrackingRecHit/test/TrackingRecHit2DCUDA_t.cpp
@@ -1,6 +1,6 @@
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/requireCUDADevices.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/requireDevices.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
 
 namespace testTrackingRecHit2D {
@@ -10,7 +10,7 @@ namespace testTrackingRecHit2D {
 }
 
 int main() {
-  requireCUDADevices();
+  cms::cudatest::requireDevices();
 
   cudaStream_t stream;
   cudaCheck(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
diff --git a/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h b/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h
index 8bfefee5c3387..6fb487a244e71 100644
--- a/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h
+++ b/CalibTracker/SiPixelESProducers/interface/SiPixelGainCalibrationForHLTGPU.h
@@ -2,7 +2,7 @@
 #define CalibTracker_SiPixelESProducers_interface_SiPixelGainCalibrationForHLTGPU_h
 
 #include "CondFormats/SiPixelObjects/interface/SiPixelGainCalibrationForHLT.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAESProduct.h"
+#include "HeterogeneousCore/CUDACore/interface/ESProduct.h"
 
 class SiPixelGainCalibrationForHLT;
 class SiPixelGainForHLTonGPU;
@@ -26,7 +26,7 @@ class SiPixelGainCalibrationForHLTGPU {
     SiPixelGainForHLTonGPU *gainForHLTonGPU = nullptr;
     SiPixelGainForHLTonGPU_DecodingStructure *gainDataOnGPU = nullptr;
   };
-  CUDAESProduct<GPUData> gpuData_;
+  cms::cuda::ESProduct<GPUData> gpuData_;
 };
 
 #endif  // CalibTracker_SiPixelESProducers_interface_SiPixelGainCalibrationForHLTGPU_h
diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc
index ad6c46082be8b..be4cc5d9a3336 100644
--- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc
+++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigiErrorsSoAFromCUDA.cc
@@ -1,4 +1,4 @@
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
 #include "DataFormats/SiPixelDigi/interface/SiPixelDigiErrorsSoA.h"
 #include "FWCore/Framework/interface/EventSetup.h"
@@ -8,7 +8,7 @@
 #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
 #include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
 #include "FWCore/ParameterSet/interface/ParameterSet.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h"
+#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
 
 class SiPixelDigiErrorsSoAFromCUDA : public edm::stream::EDProducer<edm::ExternalWork> {
@@ -24,16 +24,17 @@ class SiPixelDigiErrorsSoAFromCUDA : public edm::stream::EDProducer<edm::Externa
                edm::WaitingTaskWithArenaHolder waitingTaskHolder) override;
   void produce(edm::Event& iEvent, const edm::EventSetup& iSetup) override;
 
-  edm::EDGetTokenT<CUDAProduct<SiPixelDigiErrorsCUDA>> digiErrorGetToken_;
+  edm::EDGetTokenT<cms::cuda::Product<SiPixelDigiErrorsCUDA>> digiErrorGetToken_;
   edm::EDPutTokenT<SiPixelDigiErrorsSoA> digiErrorPutToken_;
 
-  cudautils::host::unique_ptr<PixelErrorCompact[]> data_;
+  cms::cuda::host::unique_ptr<PixelErrorCompact[]> data_;
   GPU::SimpleVector<PixelErrorCompact> error_;
   const PixelFormatterErrors* formatterErrors_ = nullptr;
 };
 
 SiPixelDigiErrorsSoAFromCUDA::SiPixelDigiErrorsSoAFromCUDA(const edm::ParameterSet& iConfig)
-    : digiErrorGetToken_(consumes<CUDAProduct<SiPixelDigiErrorsCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
+    : digiErrorGetToken_(
+          consumes<cms::cuda::Product<SiPixelDigiErrorsCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
       digiErrorPutToken_(produces<SiPixelDigiErrorsSoA>()) {}
 
 void SiPixelDigiErrorsSoAFromCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
@@ -46,7 +47,7 @@ void SiPixelDigiErrorsSoAFromCUDA::acquire(const edm::Event& iEvent,
                                            const edm::EventSetup& iSetup,
                                            edm::WaitingTaskWithArenaHolder waitingTaskHolder) {
   // Do the transfer in a CUDA stream parallel to the computation CUDA stream
-  CUDAScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder)};
+  cms::cuda::ScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder)};
 
   const auto& gpuDigiErrors = ctx.get(iEvent, digiErrorGetToken_);
 
diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigisSoAFromCUDA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigisSoAFromCUDA.cc
index 7794032154e98..dbec74585998f 100644
--- a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigisSoAFromCUDA.cc
+++ b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigisSoAFromCUDA.cc
@@ -1,4 +1,4 @@
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h"
 #include "FWCore/Framework/interface/EventSetup.h"
@@ -8,7 +8,7 @@
 #include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
 #include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
 #include "FWCore/ParameterSet/interface/ParameterSet.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h"
+#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
 
 class SiPixelDigisSoAFromCUDA : public edm::stream::EDProducer<edm::ExternalWork> {
@@ -24,19 +24,19 @@ class SiPixelDigisSoAFromCUDA : public edm::stream::EDProducer<edm::ExternalWork
                edm::WaitingTaskWithArenaHolder waitingTaskHolder) override;
   void produce(edm::Event& iEvent, const edm::EventSetup& iSetup) override;
 
-  edm::EDGetTokenT<CUDAProduct<SiPixelDigisCUDA>> digiGetToken_;
+  edm::EDGetTokenT<cms::cuda::Product<SiPixelDigisCUDA>> digiGetToken_;
   edm::EDPutTokenT<SiPixelDigisSoA> digiPutToken_;
 
-  cudautils::host::unique_ptr<uint32_t[]> pdigi_;
-  cudautils::host::unique_ptr<uint32_t[]> rawIdArr_;
-  cudautils::host::unique_ptr<uint16_t[]> adc_;
-  cudautils::host::unique_ptr<int32_t[]> clus_;
+  cms::cuda::host::unique_ptr<uint32_t[]> pdigi_;
+  cms::cuda::host::unique_ptr<uint32_t[]> rawIdArr_;
+  cms::cuda::host::unique_ptr<uint16_t[]> adc_;
+  cms::cuda::host::unique_ptr<int32_t[]> clus_;
 
   int nDigis_;
 };
 
 SiPixelDigisSoAFromCUDA::SiPixelDigisSoAFromCUDA(const edm::ParameterSet& iConfig)
-    : digiGetToken_(consumes<CUDAProduct<SiPixelDigisCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
+    : digiGetToken_(consumes<cms::cuda::Product<SiPixelDigisCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
       digiPutToken_(produces<SiPixelDigisSoA>()) {}
 
 void SiPixelDigisSoAFromCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
@@ -49,7 +49,7 @@ void SiPixelDigisSoAFromCUDA::acquire(const edm::Event& iEvent,
                                       const edm::EventSetup& iSetup,
                                       edm::WaitingTaskWithArenaHolder waitingTaskHolder) {
   // Do the transfer in a CUDA stream parallel to the computation CUDA stream
-  CUDAScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder)};
+  cms::cuda::ScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder)};
 
   const auto& gpuDigis = ctx.get(iEvent, digiGetToken_);
 
diff --git a/RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPUWrapper.h b/RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPUWrapper.h
index 14a5d25504479..2f9eb092bc648 100644
--- a/RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPUWrapper.h
+++ b/RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPUWrapper.h
@@ -1,7 +1,7 @@
 #ifndef RecoLocalTracker_SiPixelClusterizer_SiPixelFedCablingMapGPUWrapper_h
 #define RecoLocalTracker_SiPixelClusterizer_SiPixelFedCablingMapGPUWrapper_h
 
-#include "HeterogeneousCore/CUDACore/interface/CUDAESProduct.h"
+#include "HeterogeneousCore/CUDACore/interface/ESProduct.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/CUDAHostAllocator.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "RecoLocalTracker/SiPixelClusterizer/interface/SiPixelFedCablingMapGPU.h"
@@ -29,7 +29,7 @@ class SiPixelFedCablingMapGPUWrapper {
 
   // returns pointer to GPU memory
   const unsigned char *getModToUnpAllAsync(cudaStream_t cudaStream) const;
-  cudautils::device::unique_ptr<unsigned char[]> getModToUnpRegionalAsync(std::set<unsigned int> const &modules,
+  cms::cuda::device::unique_ptr<unsigned char[]> getModToUnpRegionalAsync(std::set<unsigned int> const &modules,
                                                                           cudaStream_t cudaStream) const;
 
 private:
@@ -44,13 +44,13 @@ class SiPixelFedCablingMapGPUWrapper {
     ~GPUData();
     SiPixelFedCablingMapGPU *cablingMapDevice = nullptr;  // pointer to struct in GPU
   };
-  CUDAESProduct<GPUData> gpuData_;
+  cms::cuda::ESProduct<GPUData> gpuData_;
 
   struct ModulesToUnpack {
     ~ModulesToUnpack();
     unsigned char *modToUnpDefault = nullptr;  // pointer to GPU
   };
-  CUDAESProduct<ModulesToUnpack> modToUnp_;
+  cms::cuda::ESProduct<ModulesToUnpack> modToUnp_;
 };
 
 #endif
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
index 62004d385577d..95aac36dbd197 100644
--- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
+++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc
@@ -1,4 +1,4 @@
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigiErrorsCUDA.h"
@@ -25,7 +25,7 @@
 #include "FWCore/ParameterSet/interface/ParameterSetDescription.h"
 #include "FWCore/ParameterSet/interface/ParameterSet.h"
 #include "FWCore/ServiceRegistry/interface/Service.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.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"
@@ -51,11 +51,11 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork
 
   edm::EDGetTokenT<FEDRawDataCollection> rawGetToken_;
 
-  edm::EDPutTokenT<CUDAProduct<SiPixelDigisCUDA>> digiPutToken_;
-  edm::EDPutTokenT<CUDAProduct<SiPixelDigiErrorsCUDA>> digiErrorPutToken_;
-  edm::EDPutTokenT<CUDAProduct<SiPixelClustersCUDA>> clusterPutToken_;
+  edm::EDPutTokenT<cms::cuda::Product<SiPixelDigisCUDA>> digiPutToken_;
+  edm::EDPutTokenT<cms::cuda::Product<SiPixelDigiErrorsCUDA>> digiErrorPutToken_;
+  edm::EDPutTokenT<cms::cuda::Product<SiPixelClustersCUDA>> clusterPutToken_;
 
-  CUDAContextState ctxState_;
+  cms::cuda::ContextState ctxState_;
 
   edm::ESWatcher<SiPixelFedCablingMapRcd> recordWatcher_;
   edm::ESGetToken<SiPixelFedCablingMapGPUWrapper, CkfComponentsRecord> gpuMapToken_;
@@ -78,8 +78,8 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer<edm::ExternalWork
 
 SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfig)
     : rawGetToken_(consumes<FEDRawDataCollection>(iConfig.getParameter<edm::InputTag>("InputLabel"))),
-      digiPutToken_(produces<CUDAProduct<SiPixelDigisCUDA>>()),
-      clusterPutToken_(produces<CUDAProduct<SiPixelClustersCUDA>>()),
+      digiPutToken_(produces<cms::cuda::Product<SiPixelDigisCUDA>>()),
+      clusterPutToken_(produces<cms::cuda::Product<SiPixelClustersCUDA>>()),
       gpuMapToken_(esConsumes<SiPixelFedCablingMapGPUWrapper, CkfComponentsRecord>()),
       gainsToken_(esConsumes<SiPixelGainCalibrationForHLTGPU, SiPixelGainCalibrationForHLTGPURcd>()),
       cablingMapToken_(esConsumes<SiPixelFedCablingMap, SiPixelFedCablingMapRcd>(
@@ -89,7 +89,7 @@ SiPixelRawToClusterCUDA::SiPixelRawToClusterCUDA(const edm::ParameterSet& iConfi
       usePilotBlade_(iConfig.getParameter<bool>("UsePilotBlade"))  // Control the usage of pilot-blade data, FED=40
 {
   if (includeErrors_) {
-    digiErrorPutToken_ = produces<CUDAProduct<SiPixelDigiErrorsCUDA>>();
+    digiErrorPutToken_ = produces<cms::cuda::Product<SiPixelDigiErrorsCUDA>>();
   }
 
   // regions
@@ -128,7 +128,7 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d
 void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent,
                                       const edm::EventSetup& iSetup,
                                       edm::WaitingTaskWithArenaHolder waitingTaskHolder) {
-  CUDAScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder), ctxState_};
+  cms::cuda::ScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder), ctxState_};
 
   auto hgpuMap = iSetup.getHandle(gpuMapToken_);
   if (hgpuMap->hasQuality() != useQuality_) {
@@ -143,7 +143,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent,
   // get the GPU product already here so that the async transfer can begin
   const auto* gpuGains = hgains->getGPUProductAsync(ctx.stream());
 
-  cudautils::device::unique_ptr<unsigned char[]> modulesToUnpackRegional;
+  cms::cuda::device::unique_ptr<unsigned char[]> modulesToUnpackRegional;
   const unsigned char* gpuModulesToUnpack;
 
   if (regions_) {
@@ -247,7 +247,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent,
 }
 
 void SiPixelRawToClusterCUDA::produce(edm::Event& iEvent, const edm::EventSetup& iSetup) {
-  CUDAScopedContextProduce ctx{ctxState_};
+  cms::cuda::ScopedContextProduce ctx{ctxState_};
 
   auto tmp = gpuAlgo_.getResults();
   ctx.emplace(iEvent, digiPutToken_, std::move(tmp.first));
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
index 8e0d5123e6ecc..53af26ac7527d 100644
--- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
+++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu
@@ -47,8 +47,8 @@ namespace pixelgpudetails {
   constexpr uint32_t MAX_FED_WORDS = pixelgpudetails::MAX_FED * pixelgpudetails::MAX_WORD;
 
   SiPixelRawToClusterGPUKernel::WordFedAppender::WordFedAppender() {
-    word_ = cudautils::make_host_noncached_unique<unsigned int[]>(MAX_FED_WORDS, cudaHostAllocWriteCombined);
-    fedId_ = cudautils::make_host_noncached_unique<unsigned char[]>(MAX_FED_WORDS, cudaHostAllocWriteCombined);
+    word_ = cms::cuda::make_host_noncached_unique<unsigned int[]>(MAX_FED_WORDS, cudaHostAllocWriteCombined);
+    fedId_ = cms::cuda::make_host_noncached_unique<unsigned char[]>(MAX_FED_WORDS, cudaHostAllocWriteCombined);
   }
 
   void SiPixelRawToClusterGPUKernel::WordFedAppender::initializeWordFed(int fedId,
@@ -549,7 +549,7 @@ namespace pixelgpudetails {
     }
     clusters_d = SiPixelClustersCUDA(gpuClustering::MaxNumModules, stream);
 
-    nModules_Clusters_h = cudautils::make_host_unique<uint32_t[]>(2, stream);
+    nModules_Clusters_h = cms::cuda::make_host_unique<uint32_t[]>(2, stream);
 
     if (wordCounter)  // protect in case of empty event....
     {
@@ -558,8 +558,8 @@ namespace pixelgpudetails {
 
       assert(0 == wordCounter % 2);
       // wordCounter is the total no of words in each event to be trasfered on device
-      auto word_d = cudautils::make_device_unique<uint32_t[]>(wordCounter, stream);
-      auto fedId_d = cudautils::make_device_unique<uint8_t[]>(wordCounter, stream);
+      auto word_d = cms::cuda::make_device_unique<uint32_t[]>(wordCounter, stream);
+      auto fedId_d = cms::cuda::make_device_unique<uint8_t[]>(wordCounter, stream);
 
       cudaCheck(
           cudaMemcpyAsync(word_d.get(), wordFed.word(), wordCounter * sizeof(uint32_t), cudaMemcpyDefault, stream));
diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
index 888fc07953d9d..767c5a1e92ad0 100644
--- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
+++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h
@@ -156,8 +156,8 @@ namespace pixelgpudetails {
       const unsigned char* fedId() const { return fedId_.get(); }
 
     private:
-      cudautils::host::noncached::unique_ptr<unsigned int[]> word_;
-      cudautils::host::noncached::unique_ptr<unsigned char[]> fedId_;
+      cms::cuda::host::noncached::unique_ptr<unsigned int[]> word_;
+      cms::cuda::host::noncached::unique_ptr<unsigned char[]> fedId_;
     };
 
     SiPixelRawToClusterGPUKernel() = default;
@@ -187,7 +187,7 @@ namespace pixelgpudetails {
       // stream is still alive
       //
       // technically the statement above is not true anymore now that
-      // the CUDA streams are cached within the CUDAStreamCache, but it is
+      // the CUDA streams are cached within the cms::cuda::StreamCache, but it is
       // still better to release as early as possible
       nModules_Clusters_h.reset();
       return std::make_pair(std::move(digis_d), std::move(clusters_d));
@@ -199,7 +199,7 @@ namespace pixelgpudetails {
     uint32_t nDigis = 0;
 
     // Data to be put in the event
-    cudautils::host::unique_ptr<uint32_t[]> nModules_Clusters_h;
+    cms::cuda::host::unique_ptr<uint32_t[]> nModules_Clusters_h;
     SiPixelDigisCUDA digis_d;
     SiPixelClustersCUDA clusters_d;
     SiPixelDigiErrorsCUDA digiErrors_d;
diff --git a/RecoLocalTracker/SiPixelClusterizer/src/SiPixelFedCablingMapGPUWrapper.cc b/RecoLocalTracker/SiPixelClusterizer/src/SiPixelFedCablingMapGPUWrapper.cc
index d4b8e40dea76b..7d3a9aa8d9b07 100644
--- a/RecoLocalTracker/SiPixelClusterizer/src/SiPixelFedCablingMapGPUWrapper.cc
+++ b/RecoLocalTracker/SiPixelClusterizer/src/SiPixelFedCablingMapGPUWrapper.cc
@@ -127,10 +127,10 @@ const unsigned char* SiPixelFedCablingMapGPUWrapper::getModToUnpAllAsync(cudaStr
   return data.modToUnpDefault;
 }
 
-cudautils::device::unique_ptr<unsigned char[]> SiPixelFedCablingMapGPUWrapper::getModToUnpRegionalAsync(
+cms::cuda::device::unique_ptr<unsigned char[]> SiPixelFedCablingMapGPUWrapper::getModToUnpRegionalAsync(
     std::set<unsigned int> const& modules, cudaStream_t cudaStream) const {
-  auto modToUnpDevice = cudautils::make_device_unique<unsigned char[]>(pixelgpudetails::MAX_SIZE, cudaStream);
-  auto modToUnpHost = cudautils::make_host_unique<unsigned char[]>(pixelgpudetails::MAX_SIZE, cudaStream);
+  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();
diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h
index b22e7a35a6ac7..8ec665f8960b6 100644
--- a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h
+++ b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h
@@ -13,9 +13,8 @@
 
 #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/requireCUDADevices.h"
+#include "HeterogeneousCore/CUDAUtilities/interface/requireDevices.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/launch.h"
-#include "HeterogeneousCore/CUDAUtilities/interface/cudaDeviceCount.h"
 #endif
 
 #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h"
@@ -23,13 +22,7 @@
 
 int main(void) {
 #ifdef __CUDACC__
-  requireCUDADevices();
-
-  if (cudautils::cudaDeviceCount() == 0) {
-    std::cerr << "No CUDA devices on this system"
-              << "\n";
-    exit(EXIT_FAILURE);
-  }
+  cms::cudatest::requireDevices();
 #endif
 
   using namespace gpuClustering;
@@ -44,14 +37,14 @@ int main(void) {
   auto h_clus = std::make_unique<int[]>(numElements);
 
 #ifdef __CUDACC__
-  auto d_id = cudautils::make_device_unique<uint16_t[]>(numElements, nullptr);
-  auto d_x = cudautils::make_device_unique<uint16_t[]>(numElements, nullptr);
-  auto d_y = cudautils::make_device_unique<uint16_t[]>(numElements, nullptr);
-  auto d_adc = cudautils::make_device_unique<uint16_t[]>(numElements, nullptr);
-  auto d_clus = cudautils::make_device_unique<int[]>(numElements, nullptr);
-  auto d_moduleStart = cudautils::make_device_unique<uint32_t[]>(MaxNumModules + 1, nullptr);
-  auto d_clusInModule = cudautils::make_device_unique<uint32_t[]>(MaxNumModules, nullptr);
-  auto d_moduleId = cudautils::make_device_unique<uint32_t[]>(MaxNumModules, nullptr);
+  auto d_id = cms::cuda::make_device_unique<uint16_t[]>(numElements, nullptr);
+  auto d_x = cms::cuda::make_device_unique<uint16_t[]>(numElements, nullptr);
+  auto d_y = cms::cuda::make_device_unique<uint16_t[]>(numElements, nullptr);
+  auto d_adc = cms::cuda::make_device_unique<uint16_t[]>(numElements, nullptr);
+  auto d_clus = cms::cuda::make_device_unique<int[]>(numElements, nullptr);
+  auto d_moduleStart = cms::cuda::make_device_unique<uint32_t[]>(MaxNumModules + 1, nullptr);
+  auto d_clusInModule = cms::cuda::make_device_unique<uint32_t[]>(MaxNumModules, nullptr);
+  auto d_moduleId = cms::cuda::make_device_unique<uint32_t[]>(MaxNumModules, nullptr);
 #else
 
   auto h_moduleStart = std::make_unique<uint32_t[]>(MaxNumModules + 1);
@@ -262,7 +255,7 @@ int main(void) {
     std::cout << "CUDA countModules kernel launch with " << blocksPerGrid << " blocks of " << threadsPerBlock
               << " threads\n";
 
-    cudautils::launch(countModules, {blocksPerGrid, threadsPerBlock}, d_id.get(), d_moduleStart.get(), d_clus.get(), n);
+    cms::cuda::launch(countModules, {blocksPerGrid, threadsPerBlock}, d_id.get(), d_moduleStart.get(), d_clus.get(), n);
 
     blocksPerGrid = MaxNumModules;  //nModules;
 
@@ -270,7 +263,7 @@ int main(void) {
               << " threads\n";
     cudaCheck(cudaMemset(d_clusInModule.get(), 0, MaxNumModules * sizeof(uint32_t)));
 
-    cudautils::launch(findClus,
+    cms::cuda::launch(findClus,
                       {blocksPerGrid, threadsPerBlock},
                       d_id.get(),
                       d_x.get(),
@@ -296,7 +289,7 @@ int main(void) {
     if (ncl != std::accumulate(nclus, nclus + MaxNumModules, 0))
       std::cout << "ERROR!!!!! wrong number of cluster found" << std::endl;
 
-    cudautils::launch(clusterChargeCut,
+    cms::cuda::launch(clusterChargeCut,
                       {blocksPerGrid, threadsPerBlock},
                       d_id.get(),
                       d_adc.get(),
diff --git a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
index e43c45f90523f..70e30563c66c3 100644
--- a/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
+++ b/RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h
@@ -6,7 +6,7 @@
 #include "CalibTracker/SiPixelESProducers/interface/SiPixelCPEGenericDBErrorParametrization.h"
 #include "CondFormats/SiPixelTransient/interface/SiPixelGenError.h"
 #include "CondFormats/SiPixelTransient/interface/SiPixelTemplate.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAESProduct.h"
+#include "HeterogeneousCore/CUDACore/interface/ESProduct.h"
 #include "HeterogeneousCore/CUDAUtilities/interface/CUDAHostAllocator.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
@@ -95,7 +95,7 @@ class PixelCPEFast final : public PixelCPEBase {
     pixelCPEforGPU::ParamsOnGPU h_paramsOnGPU;
     pixelCPEforGPU::ParamsOnGPU *d_paramsOnGPU = nullptr;  // copy of the above on the Device
   };
-  CUDAESProduct<GPUData> gpuData_;
+  cms::cuda::ESProduct<GPUData> gpuData_;
 
   void fillParamsForGpu();
 };
diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu
index 1342ab916e472..4e4f38f329d01 100644
--- a/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu
+++ b/RecoLocalTracker/SiPixelRecHits/plugins/PixelRecHits.cu
@@ -63,8 +63,8 @@ namespace pixelgpudetails {
     }
 
     if (nHits) {
-      auto hws = cudautils::make_device_unique<uint8_t[]>(TrackingRecHit2DSOAView::Hist::wsSize(), stream);
-      cudautils::fillManyFromVector(
+      auto hws = cms::cuda::make_device_unique<uint8_t[]>(TrackingRecHit2DSOAView::Hist::wsSize(), stream);
+      cms::cuda::fillManyFromVector(
           hits_d.phiBinner(), hws.get(), 10, hits_d.iphi(), hits_d.hitsLayerStart(), nHits, 256, stream);
       cudaCheck(cudaGetLastError());
     }
diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc
index 1641719d0537d..4d85c41339020 100644
--- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc
+++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitCUDA.cc
@@ -1,8 +1,7 @@
 #include <cuda_runtime.h>
 
 #include "CUDADataFormats/BeamSpot/interface/BeamSpotCUDA.h"
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
@@ -18,7 +17,7 @@
 #include "FWCore/Utilities/interface/InputTag.h"
 #include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
 #include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h"
+#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
 #include "RecoLocalTracker/Records/interface/TkPixelCPERecord.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h"
@@ -36,11 +35,11 @@ class SiPixelRecHitCUDA : public edm::global::EDProducer<> {
   void produce(edm::StreamID streamID, edm::Event& iEvent, const edm::EventSetup& iSetup) const override;
 
   // The mess with inputs will be cleaned up when migrating to the new framework
-  edm::EDGetTokenT<CUDAProduct<BeamSpotCUDA>> tBeamSpot;
-  edm::EDGetTokenT<CUDAProduct<SiPixelClustersCUDA>> token_;
-  edm::EDGetTokenT<CUDAProduct<SiPixelDigisCUDA>> tokenDigi_;
+  edm::EDGetTokenT<cms::cuda::Product<BeamSpotCUDA>> tBeamSpot;
+  edm::EDGetTokenT<cms::cuda::Product<SiPixelClustersCUDA>> token_;
+  edm::EDGetTokenT<cms::cuda::Product<SiPixelDigisCUDA>> tokenDigi_;
 
-  edm::EDPutTokenT<CUDAProduct<TrackingRecHit2DCUDA>> tokenHit_;
+  edm::EDPutTokenT<cms::cuda::Product<TrackingRecHit2DCUDA>> tokenHit_;
 
   std::string cpeName_;
 
@@ -48,10 +47,10 @@ class SiPixelRecHitCUDA : public edm::global::EDProducer<> {
 };
 
 SiPixelRecHitCUDA::SiPixelRecHitCUDA(const edm::ParameterSet& iConfig)
-    : tBeamSpot(consumes<CUDAProduct<BeamSpotCUDA>>(iConfig.getParameter<edm::InputTag>("beamSpot"))),
-      token_(consumes<CUDAProduct<SiPixelClustersCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
-      tokenDigi_(consumes<CUDAProduct<SiPixelDigisCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
-      tokenHit_(produces<CUDAProduct<TrackingRecHit2DCUDA>>()),
+    : tBeamSpot(consumes<cms::cuda::Product<BeamSpotCUDA>>(iConfig.getParameter<edm::InputTag>("beamSpot"))),
+      token_(consumes<cms::cuda::Product<SiPixelClustersCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
+      tokenDigi_(consumes<cms::cuda::Product<SiPixelDigisCUDA>>(iConfig.getParameter<edm::InputTag>("src"))),
+      tokenHit_(produces<cms::cuda::Product<TrackingRecHit2DCUDA>>()),
       cpeName_(iConfig.getParameter<std::string>("CPE")) {}
 
 void SiPixelRecHitCUDA::fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
@@ -82,17 +81,17 @@ void SiPixelRecHitCUDA::produce(edm::StreamID streamID, edm::Event& iEvent, cons
     throw cms::Exception("Configuration") << "too bad, not a fast cpe gpu processing not possible....";
   }
 
-  edm::Handle<CUDAProduct<SiPixelClustersCUDA>> hclusters;
+  edm::Handle<cms::cuda::Product<SiPixelClustersCUDA>> hclusters;
   iEvent.getByToken(token_, hclusters);
 
-  CUDAScopedContextProduce ctx{*hclusters};
+  cms::cuda::ScopedContextProduce ctx{*hclusters};
   auto const& clusters = ctx.get(*hclusters);
 
-  edm::Handle<CUDAProduct<SiPixelDigisCUDA>> hdigis;
+  edm::Handle<cms::cuda::Product<SiPixelDigisCUDA>> hdigis;
   iEvent.getByToken(tokenDigi_, hdigis);
   auto const& digis = ctx.get(*hdigis);
 
-  edm::Handle<CUDAProduct<BeamSpotCUDA>> hbs;
+  edm::Handle<cms::cuda::Product<BeamSpotCUDA>> hbs;
   iEvent.getByToken(tBeamSpot, hbs);
   auto const& bs = ctx.get(*hbs);
 
diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc
index a4f19ac276a7a..7b072abc1dd47 100644
--- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc
+++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitFromSOA.cc
@@ -1,6 +1,6 @@
 #include <cuda_runtime.h>
 
-#include "CUDADataFormats/Common/interface/CUDAProduct.h"
+#include "CUDADataFormats/Common/interface/Product.h"
 #include "CUDADataFormats/Common/interface/HostProduct.h"
 #include "CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DCUDA.h"
 #include "DataFormats/Common/interface/DetSetVectorNew.h"
@@ -19,7 +19,7 @@
 #include "Geometry/CommonDetUnit/interface/PixelGeomDetUnit.h"
 #include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
 #include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h"
+#include "HeterogeneousCore/CUDACore/interface/ScopedContext.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/pixelCPEforGPU.h"
 
 class SiPixelRecHitFromSOA : public edm::stream::EDProducer<edm::ExternalWork> {
@@ -37,17 +37,18 @@ class SiPixelRecHitFromSOA : public edm::stream::EDProducer<edm::ExternalWork> {
                edm::WaitingTaskWithArenaHolder waitingTaskHolder) override;
   void produce(edm::Event& iEvent, edm::EventSetup const& iSetup) override;
 
-  edm::EDGetTokenT<CUDAProduct<TrackingRecHit2DCUDA>> tokenHit_;  // CUDA hits
-  edm::EDGetTokenT<SiPixelClusterCollectionNew> clusterToken_;    // Legacy Clusters
+  edm::EDGetTokenT<cms::cuda::Product<TrackingRecHit2DCUDA>> tokenHit_;  // CUDA hits
+  edm::EDGetTokenT<SiPixelClusterCollectionNew> clusterToken_;           // Legacy Clusters
 
   uint32_t m_nHits;
-  cudautils::host::unique_ptr<uint16_t[]> m_store16;
-  cudautils::host::unique_ptr<float[]> m_store32;
-  cudautils::host::unique_ptr<uint32_t[]> m_hitsModuleStart;
+  cms::cuda::host::unique_ptr<uint16_t[]> m_store16;
+  cms::cuda::host::unique_ptr<float[]> m_store32;
+  cms::cuda::host::unique_ptr<uint32_t[]> m_hitsModuleStart;
 };
 
 SiPixelRecHitFromSOA::SiPixelRecHitFromSOA(const edm::ParameterSet& iConfig)
-    : tokenHit_(consumes<CUDAProduct<TrackingRecHit2DCUDA>>(iConfig.getParameter<edm::InputTag>("pixelRecHitSrc"))),
+    : tokenHit_(
+          consumes<cms::cuda::Product<TrackingRecHit2DCUDA>>(iConfig.getParameter<edm::InputTag>("pixelRecHitSrc"))),
       clusterToken_(consumes<SiPixelClusterCollectionNew>(iConfig.getParameter<edm::InputTag>("src"))) {
   produces<SiPixelRecHitCollectionNew>();
   produces<HMSstorage>();
@@ -63,8 +64,8 @@ void SiPixelRecHitFromSOA::fillDescriptions(edm::ConfigurationDescriptions& desc
 void SiPixelRecHitFromSOA::acquire(edm::Event const& iEvent,
                                    edm::EventSetup const& iSetup,
                                    edm::WaitingTaskWithArenaHolder waitingTaskHolder) {
-  CUDAProduct<TrackingRecHit2DCUDA> const& inputDataWrapped = iEvent.get(tokenHit_);
-  CUDAScopedContextAcquire ctx{inputDataWrapped, std::move(waitingTaskHolder)};
+  cms::cuda::Product<TrackingRecHit2DCUDA> const& inputDataWrapped = iEvent.get(tokenHit_);
+  cms::cuda::ScopedContextAcquire ctx{inputDataWrapped, std::move(waitingTaskHolder)};
   auto const& inputData = ctx.get(inputDataWrapped);
 
   m_nHits = inputData.nHits();
diff --git a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc
index 7900cf8b2289a..fbe0fd13b84a4 100644
--- a/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc
+++ b/RecoLocalTracker/SiPixelRecHits/plugins/SiPixelRecHitSoAFromLegacy.cc
@@ -22,7 +22,6 @@
 #include "FWCore/Utilities/interface/InputTag.h"
 #include "Geometry/Records/interface/TrackerDigiGeometryRecord.h"
 #include "Geometry/TrackerGeometryBuilder/interface/TrackerGeometry.h"
-#include "HeterogeneousCore/CUDACore/interface/CUDAScopedContext.h"
 #include "RecoLocalTracker/Records/interface/TkPixelCPERecord.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEBase.h"
 #include "RecoLocalTracker/SiPixelRecHits/interface/PixelCPEFast.h"
@@ -251,7 +250,7 @@ void SiPixelRecHitSoAFromLegacy::produce(edm::StreamID streamID, edm::Event& iEv
   for (auto i = 0; i < 11; ++i) {
     output->hitsLayerStart()[i] = hitsModuleStart[cpeView.layerGeometry().layerStart[i]];
   }
-  cudautils::fillManyFromVector(
+  cms::cuda::fillManyFromVector(
       output->phiBinner(), nullptr, 10, output->iphi(), output->hitsLayerStart(), numberOfHits, 256, nullptr);
 
   // std::cout << "created HitSoa for " <<  numberOfClusters << " clusters in " << numberOfDetUnits << " Dets" << std::endl;