Tracks and hits portable (cms-sw#3)

* STUFF that make things work * Use DeviceToDevice copy for cpe params Co-authored-by: Dimitris Papagiannis <[email protected]>
borzari · Nov 24, 2022 · e7f5b13 · e7f5b13
1 parent 65eb849
commit e7f5b13
Show file tree

Hide file tree

Showing 39 changed files with 417 additions and 354 deletions.
diff --git a/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h b/CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
@@ -21,19 +21,11 @@ using SiPixelClustersCUDASoA = SiPixelClustersCUDALayout<>;
 class SiPixelClustersCUDA : public cms::cuda::PortableDeviceCollection<SiPixelClustersCUDALayout<>> {
 public:
   SiPixelClustersCUDA() = default;
-  // explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream);
-  explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
-      : PortableDeviceCollection<SiPixelClustersCUDALayout<>>(maxModules + 1, stream) {}
   ~SiPixelClustersCUDA() = default;
 
-  // // Restrict view
-  // using RestrictConstView =
-  //    Layout::ConstViewTemplate<cms::soa::RestrictQualify::enabled, cms::soa::RangeChecking::disabled>;
-  //
-  // RestrictConstView restrictConstView() const { return RestrictConstView(layout()); }
+  explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
+      : PortableDeviceCollection<SiPixelClustersCUDALayout<>>(maxModules + 1, stream) {}
 
-  // SiPixelClustersCUDA(const SiPixelClustersCUDA &) = delete;
-  // SiPixelClustersCUDA &operator=(const SiPixelClustersCUDA &) = delete;
   SiPixelClustersCUDA(SiPixelClustersCUDA &&) = default;
   SiPixelClustersCUDA &operator=(SiPixelClustersCUDA &&) = default;
 
@@ -45,41 +37,7 @@ class SiPixelClustersCUDA : public cms::cuda::PortableDeviceCollection<SiPixelCl
   uint32_t nClusters() const { return nClusters_h; }
   int32_t offsetBPIX2() const { return offsetBPIX2_h; }
 
-  // uint32_t *moduleStart() { return moduleStart_d.get(); }
-  // uint32_t *clusInModule() { return clusInModule_d.get(); }
-  // uint32_t *moduleId() { return moduleId_d.get(); }
-  // uint32_t *clusModuleStart() { return clusModuleStart_d.get(); }
-  //
-  // uint32_t const *moduleStart() const { return moduleStart_d.get(); }
-  // uint32_t const *clusInModule() const { return clusInModule_d.get(); }
-  // uint32_t const *moduleId() const { return moduleId_d.get(); }
-  // uint32_t const *clusModuleStart() const { return clusModuleStart_d.get(); }
-  //
-  // class SiPixelClustersCUDASOAView {
-  // public:
-  //   __device__ __forceinline__ uint32_t moduleStart(int i) const { return __ldg(moduleStart_ + i); }
-  //   __device__ __forceinline__ uint32_t clusInModule(int i) const { return __ldg(clusInModule_ + i); }
-  //   __device__ __forceinline__ uint32_t moduleId(int i) const { return __ldg(moduleId_ + i); }
-  //   __device__ __forceinline__ uint32_t clusModuleStart(int i) const { return __ldg(clusModuleStart_ + i); }
-  //
-  //   uint32_t const *moduleStart_;
-  //   uint32_t const *clusInModule_;
-  //   uint32_t const *moduleId_;
-  //   uint32_t const *clusModuleStart_;
-  // };
-  //
-  // SiPixelClustersCUDASOAView const *view() const { return view_d.get(); }
-
 private:
-  // 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
-  // cms::cuda::device::unique_ptr<uint32_t[]> clusModuleStart_d;  // index of the first cluster of each module
-  //
-  // cms::cuda::device::unique_ptr<SiPixelClustersCUDASOAView> view_d;  // "me" pointer
-
   uint32_t nClusters_h = 0;
   int32_t offsetBPIX2_h = 0;
 };

diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoADevice.h
@@ -15,34 +15,54 @@ namespace trackingRecHit
     TrackingRecHitSoADevice() = default;  // cms::cuda::Product needs this
 
     // Constructor which specifies the SoA size
-    explicit TrackingRecHitSoADevice(uint32_t nHits, bool isPhase2, int32_t offsetBPIX2, pixelCPEforGPU::ParamsOnGPU const* cpeParams, uint32_t const* hitsModuleStart, cudaStream_t stream)
-        : PortableDeviceCollection<TrackingRecHitSoALayout<>>(nHits, stream), nHits_(nHits), cpeParams_(cpeParams), hitsModuleStart_(hitsModuleStart), offsetBPIX2_(offsetBPIX2)
-        {
-          nModules_ = isPhase2 ? phase2PixelTopology::numberOfModules : phase1PixelTopology::numberOfModules;
-          phiBinner_ = &(view().phiBinner());
-          // phiBinner_ = cms::cuda::make_device_unique<TrackingRecHit2DSOAView::PhiBinner>(stream).get();
-          cudaCheck(cudaMemcpyAsync(&(view().nHits()), &nHits, sizeof(uint32_t),cudaMemcpyHostToDevice,stream));
-          cudaCheck(cudaMemcpyAsync(&(view().nMaxModules()), &nModules_, sizeof(uint32_t),cudaMemcpyHostToDevice,stream));
-          cudaCheck(cudaMemcpyAsync(view().hitsModuleStart().data(), hitsModuleStart, sizeof(uint32_t) * int(nModules_ + 1),cudaMemcpyHostToDevice,stream));
-          // cudaCheck(cudaMemcpyAsync(&(view().cpeParams()), cpeParams, int(sizeof(pixelCPEforGPU::ParamsOnGPU)),cudaMemcpyHostToDevice,stream));
-          cudaCheck(cudaMemcpyAsync(&(view().offsetBPIX2()), &offsetBPIX2, sizeof(int32_t),cudaMemcpyHostToDevice,stream));
+    explicit TrackingRecHitSoADevice(uint32_t nHits,
+                                     bool isPhase2,
+                                     int32_t offsetBPIX2,
+                                     pixelCPEforGPU::ParamsOnGPU const* cpeParams,
+                                     uint32_t const* hitsModuleStart,
+                                     cudaStream_t stream)
+        : PortableDeviceCollection<TrackingRecHitSoALayout<>>(nHits, stream),
+          nHits_(nHits),
+          cpeParams_(cpeParams),
+          hitsModuleStart_(hitsModuleStart),
+          offsetBPIX2_(offsetBPIX2) {
+      cudaCheck(cudaGetLastError());
+      cudaCheck(cudaDeviceSynchronize());
 
-        }
+      nModules_ = isPhase2 ? phase2PixelTopology::numberOfModules : phase1PixelTopology::numberOfModules;
+      phiBinner_ = &(view().phiBinner());
+      // phiBinner_ = cms::cuda::make_device_unique<TrackingRecHit2DSOAView::PhiBinner>(stream).get();
+      cudaCheck(cudaMemcpyAsync(&(view().nHits()), &nHits, sizeof(uint32_t), cudaMemcpyHostToDevice, stream));
+      cudaCheck(cudaMemcpyAsync(&(view().nMaxModules()), &nModules_, sizeof(uint32_t), cudaMemcpyHostToDevice, stream));
+      cudaCheck(cudaMemcpyAsync(view().hitsModuleStart().data(),
+                                hitsModuleStart,
+                                sizeof(uint32_t) * int(nModules_ + 1),
+                                cudaMemcpyHostToDevice,
+                                stream));
+      cudaCheck(
+          cudaMemcpyAsync(&(view().offsetBPIX2()), &offsetBPIX2, sizeof(int32_t), cudaMemcpyHostToDevice, stream));
+
+      // cpeParams argument is a pointer to device memory, copy
+      // its contents into the Layout.
+
+      cudaCheck(cudaMemcpyAsync(
+          &(view().cpeParams()), cpeParams, int(sizeof(pixelCPEforGPU::ParamsOnGPU)), cudaMemcpyDeviceToDevice, stream));
+    }
 
     uint32_t nHits() const { return nHits_; } //go to size of view
     uint32_t nModules() const { return nModules_; }
 
     cms::cuda::host::unique_ptr<float[]> localCoordToHostAsync(cudaStream_t stream) const {
       auto ret = cms::cuda::make_host_unique<float[]>(4 * nHits(), stream);
       size_t rowSize = sizeof(float) * nHits();
-      printf("%d \n",nModules());
-      printf("%d \n",nHits());
-      cudaCheck(cudaMemcpyAsync(ret.get(), view().xLocal() , rowSize * 4, cudaMemcpyDeviceToHost, stream));
+      printf("nModules=%d \n", nModules());
+      printf("nHits=%d \n", nHits());
+      cudaCheck(cudaMemcpyAsync(ret.get(), view().xLocal(), rowSize * 4, cudaMemcpyDeviceToHost, stream));
       // cudaCheck(cudaMemcpyAsync(ret.get() + rowSize , view().yLocal() , rowSize, cudaMemcpyDeviceToHost, stream));
       // cudaCheck(cudaMemcpyAsync(ret.get() + size_t(rowSize * 2), view().xerrLocal() , rowSize, cudaMemcpyDeviceToHost, stream));
       // cudaCheck(cudaMemcpyAsync(ret.get() + size_t(rowSize * 3) , view().yerrLocal() , rowSize, cudaMemcpyDeviceToHost, stream));
       return ret;
-    } //move to utilities
+    }  //move to utilities
 
     cms::cuda::host::unique_ptr<uint32_t[]> hitsModuleStartToHostAsync(cudaStream_t stream) const {
       // printf("%d \n",nModules());
@@ -58,7 +78,8 @@ namespace trackingRecHit
 
     private:
       uint32_t nHits_; //Needed for the host SoA size
-      pixelCPEforGPU::ParamsOnGPU const* cpeParams_; //TODO: this is used not that much from the hits (only once in BrokenLineFit), would make sens to remove it from this class.
+      pixelCPEforGPU::ParamsOnGPU const*
+          cpeParams_;  //TODO: this is used not that much from the hits (only once in BrokenLineFit), would make sens to remove it from this class.
       uint32_t const* hitsModuleStart_;
       uint32_t offsetBPIX2_;
 

diff --git a/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h b/CUDADataFormats/TrackingRecHit/interface/TrackingRecHitSoAHost.h
@@ -20,18 +20,24 @@ namespace trackingRecHit
     explicit TrackingRecHitSoAHost(uint32_t nHits, cudaStream_t stream)
     : PortableHostCollection<TrackingRecHitSoALayout<>>(nHits, stream) {}
 
-    explicit TrackingRecHitSoAHost(uint32_t nHits, bool isPhase2, int32_t offsetBPIX2, pixelCPEforGPU::ParamsOnGPU const* cpeParams, uint32_t const* hitsModuleStart, cudaStream_t stream)
-        : PortableHostCollection<TrackingRecHitSoALayout<>>(nHits, stream), nHits_(nHits), cpeParams_(cpeParams), offsetBPIX2_(offsetBPIX2)
-        {
-          nModules_ = isPhase2 ? phase2PixelTopology::numberOfModules : phase1PixelTopology::numberOfModules;
-
-          view().nHits() = nHits;
-          view().nMaxModules() = nModules_;
-          std::copy(hitsModuleStart,hitsModuleStart+nModules_+1,view().hitsModuleStart().begin());
-
-          view().offsetBPIX2() = offsetBPIX2;
-
-        }
+    explicit TrackingRecHitSoAHost(uint32_t nHits,
+                                   bool isPhase2,
+                                   int32_t offsetBPIX2,
+                                   pixelCPEforGPU::ParamsOnGPU const* cpeParams,
+                                   uint32_t const* hitsModuleStart,
+                                   cudaStream_t stream)
+        : PortableHostCollection<TrackingRecHitSoALayout<>>(nHits, stream),
+          nHits_(nHits),
+          cpeParams_(cpeParams),
+          offsetBPIX2_(offsetBPIX2) {
+      nModules_ = isPhase2 ? phase2PixelTopology::numberOfModules : phase1PixelTopology::numberOfModules;
+      std::cout << "PORCA MADONNA!!!!!!!!!!!!!!!!!" << std::endl;
+      view().nHits() = nHits;
+      view().nMaxModules() = nModules_;
+      std::copy(hitsModuleStart, hitsModuleStart + nModules_ + 1, view().hitsModuleStart().begin());
+      memcpy(&(view().cpeParams()), cpeParams, sizeof(pixelCPEforGPU::ParamsOnGPU));
+      view().offsetBPIX2() = offsetBPIX2;
+    }
 
     uint32_t nHits() const { return nHits_; }
     uint32_t nModules() const { return nModules_; }

diff --git a/CUDADataFormats/TrackingRecHit/src/classes_def.xml b/CUDADataFormats/TrackingRecHit/src/classes_def.xml
@@ -1,15 +1,17 @@
 <lcgdict>
-  <class name="TrackingRecHit2DCPU" persistent="false"/>
-  <class name="edm::Wrapper<TrackingRecHit2DCPU>" persistent="false"/>
-  <class name="TrackingRecHit2DHost" persistent="false"/>
-  <class name="edm::Wrapper<TrackingRecHit2DHost>" persistent="false"/>
-  <class name="cms::cuda::Product<TrackingRecHit2DGPU>" persistent="false"/>
-  <class name="edm::Wrapper<cms::cuda::Product<TrackingRecHit2DGPU>>" persistent="false"/>
-  <class name="TrackingRecHit2DReduced" persistent="false"/>
-  <class name="edm::Wrapper<TrackingRecHit2DReduced>" persistent="false"/>
+  <!-- <class name="TrackingRecHit2DCPU" persistent="false"/> -->
+  <!-- <class name="edm::Wrapper<TrackingRecHit2DCPU>" persistent="false"/> -->
+  <!-- <class name="TrackingRecHit2DHost" persistent="false"/> -->
+  <!-- <class name="edm::Wrapper<TrackingRecHit2DHost>" persistent="false"/> -->
+  <!-- <class name="cms::cuda::Product<TrackingRecHit2DGPU>" persistent="false"/> -->
+  <!-- <class name="edm::Wrapper<cms::cuda::Product<TrackingRecHit2DGPU>>" persistent="false"/> -->
+  <!-- <class name="TrackingRecHit2DReduced" persistent="false"/> -->
+  <!-- <class name="edm::Wrapper<TrackingRecHit2DReduced>" persistent="false"/> -->
 
   <class name="trackingRecHit::TrackingRecHitSoAHost" persistent="false"/>
   <class name="edm::Wrapper<trackingRecHit::TrackingRecHitSoAHost>" persistent="false"/>
+  <class name="trackingRecHit::TrackingRecHitSoADevice" persistent="false"/>  
   <class name="cms::cuda::Product<trackingRecHit::TrackingRecHitSoADevice>" persistent="false"/>
+  <class name="edm::Wrapper<trackingRecHit::TrackingRecHitSoADevice>" persistent="false"/>  
   <class name="edm::Wrapper<cms::cuda::Product<trackingRecHit::TrackingRecHitSoADevice>>" persistent="false"/>
 </lcgdict>
diff --git a/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigisSoAFromCUDA.cc b/EventFilter/SiPixelRawToDigi/plugins/SiPixelDigisSoAFromCUDA.cc
@@ -1,4 +1,5 @@
 #include "CUDADataFormats/Common/interface/Product.h"
+#include "CUDADataFormats/Common/interface/PortableHostCollection.h"
 #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h"
 #include "DataFormats/SiPixelDigi/interface/SiPixelDigisSoA.h"
 #include "FWCore/Framework/interface/EventSetup.h"
@@ -27,7 +28,7 @@ class SiPixelDigisSoAFromCUDA : public edm::stream::EDProducer<edm::ExternalWork
   edm::EDGetTokenT<cms::cuda::Product<SiPixelDigisCUDA>> digiGetToken_;
   edm::EDPutTokenT<SiPixelDigisSoA> digiPutToken_;
 
-  cms::cuda::host::unique_ptr<std::byte[]> store_;
+  cms::cuda::PortableHostCollection<SiPixelDigisSoALayout<>> digis_h;
 
   int nDigis_;
 };
@@ -48,31 +49,22 @@ void SiPixelDigisSoAFromCUDA::acquire(const edm::Event& iEvent,
   // Do the transfer in a CUDA stream parallel to the computation CUDA stream
   cms::cuda::ScopedContextAcquire ctx{iEvent.streamID(), std::move(waitingTaskHolder)};
 
-  const auto& gpuDigis = ctx.get(iEvent, digiGetToken_);
+  const auto& digis_d = ctx.get(iEvent, digiGetToken_);
 
-  nDigis_ = gpuDigis.nDigis();
-  store_ = gpuDigis.copyAllToHostAsync(ctx.stream());
+  nDigis_ = digis_d.nDigis();
+  digis_h = cms::cuda::PortableHostCollection<SiPixelDigisSoALayout<>>(digis_d.view().metadata().size(), ctx.stream());
+  cudaCheck(cudaMemcpyAsync(
+      digis_h.buffer().get(), digis_d.const_buffer().get(), digis_d.bufferSize(), cudaMemcpyDeviceToHost, ctx.stream()));
+  cudaCheck(cudaGetLastError());
 }
 
 void SiPixelDigisSoAFromCUDA::produce(edm::Event& iEvent, const edm::EventSetup& iSetup) {
-  // The following line copies the data from the pinned host memory to
-  // regular host memory. In principle that feels unnecessary (why not
-  // just use the pinned host memory?). There are a few arguments for
-  // doing it though
-  // - Now can release the pinned host memory back to the (caching) allocator
-  //   * if we'd like to keep the pinned memory, we'd need to also
-  //     keep the CUDA stream around as long as that, or allow pinned
-  //     host memory to be allocated without a CUDA stream
-  // - What if a CPU algorithm would produce the same SoA? We can't
-  //   use cudaMallocHost without a GPU...
-
-  // auto tmp_view = SiPixelDigisCUDASOAView(store_, nDigis_, SiPixelDigisCUDASOAView::StorageLocationHost::kMAX);
-  SiPixelDigisCUDASOA tmp_layout(store_.get(), nDigis_);
-  SiPixelDigisCUDASOA::View tmp_view(tmp_layout);
-
-  iEvent.emplace(digiPutToken_, nDigis_, tmp_view.pdigi(), tmp_view.rawIdArr(), tmp_view.adc(), tmp_view.clus());
-
-  store_.reset();
+  iEvent.emplace(digiPutToken_,
+                 nDigis_,
+                 digis_h.view().pdigi(),
+                 digis_h.view().rawIdArr(),
+                 digis_h.view().adc(),
+                 digis_h.view().clus());
 }
 
 // define as framework plugin

diff --git a/HLTrigger/Configuration/python/HLT_2022v15_cff.py b/HLTrigger/Configuration/python/HLT_2022v15_cff.py
@@ -84802,7 +84802,7 @@
 fragment.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/python/HLT_FULL_cff.py b/HLTrigger/Configuration/python/HLT_FULL_cff.py
@@ -121310,7 +121310,7 @@
 fragment.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/python/HLT_GRun_cff.py b/HLTrigger/Configuration/python/HLT_GRun_cff.py
@@ -85517,7 +85517,7 @@
 fragment.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/python/HLT_HIon_cff.py b/HLTrigger/Configuration/python/HLT_HIon_cff.py
@@ -25503,7 +25503,7 @@
 fragment.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/python/HLT_PRef_cff.py b/HLTrigger/Configuration/python/HLT_PRef_cff.py
@@ -9248,7 +9248,7 @@
 fragment.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/test/OnLine_HLT_2022v15.py b/HLTrigger/Configuration/test/OnLine_HLT_2022v15.py
@@ -85070,7 +85070,7 @@
 process.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/test/OnLine_HLT_FULL.py b/HLTrigger/Configuration/test/OnLine_HLT_FULL.py
@@ -121947,7 +121947,7 @@
 process.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/test/OnLine_HLT_GRun.py b/HLTrigger/Configuration/test/OnLine_HLT_GRun.py
@@ -85785,7 +85785,7 @@
 process.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/test/OnLine_HLT_HIon.py b/HLTrigger/Configuration/test/OnLine_HLT_HIon.py
@@ -25771,7 +25771,7 @@
 process.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),

diff --git a/HLTrigger/Configuration/test/OnLine_HLT_PRef.py b/HLTrigger/Configuration/test/OnLine_HLT_PRef.py
@@ -9516,7 +9516,7 @@
 process.hltSiPixelRecHitsSoA = SwitchProducerCUDA(
    cpu = cms.EDAlias(
        hltSiPixelRecHitsFromLegacy = cms.VPSet( 
-         cms.PSet(  type = cms.string( "cmscudacompatCPUTraitsTrackingRecHit2DHeterogeneous" )         ),
+         cms.PSet(  type = cms.string( "trackingRecHitTrackingRecHitSoAHost" )         ),
          cms.PSet(  type = cms.string( "uintAsHostProduct" )         )
        )
    ),