diff --git a/src/cudauvm/CUDADataFormats/BeamSpotCUDA.cc b/src/cudauvm/CUDADataFormats/BeamSpotCUDA.cc
index 3d146947c..9f5566d97 100644
--- a/src/cudauvm/CUDADataFormats/BeamSpotCUDA.cc
+++ b/src/cudauvm/CUDADataFormats/BeamSpotCUDA.cc
@@ -1,9 +1,15 @@
 #include "CUDADataFormats/BeamSpotCUDA.h"
 
 #include "CUDACore/cudaCheck.h"
-#include "CUDACore/device_unique_ptr.h"
+#include "CUDACore/managed_unique_ptr.h"
 
-BeamSpotCUDA::BeamSpotCUDA(Data const* data_h, cudaStream_t stream) {
-  data_d_ = cms::cuda::make_device_unique<Data>(stream);
-  cudaCheck(cudaMemcpyAsync(data_d_.get(), data_h, sizeof(Data), cudaMemcpyHostToDevice, stream));
+BeamSpotCUDA::BeamSpotCUDA(Data const& data_h, int device, cudaStream_t stream) {
+  data_d_ = cms::cuda::make_managed_unique<Data>(stream);
+  *data_d_ = data_h;
+#ifndef CUDAUVM_DISABLE_ADVICE
+  cudaCheck(cudaMemAdvise(data_d_.get(), sizeof(Data), cudaMemAdviseSetReadMostly, device));
+#endif
+#ifndef CUDAUVM_DISABLE_PREFETCH
+  cudaCheck(cudaMemPrefetchAsync(data_d_.get(), sizeof(Data), device, stream));
+#endif
 }
diff --git a/src/cudauvm/CUDADataFormats/BeamSpotCUDA.h b/src/cudauvm/CUDADataFormats/BeamSpotCUDA.h
index e529f6eeb..1d950cac1 100644
--- a/src/cudauvm/CUDADataFormats/BeamSpotCUDA.h
+++ b/src/cudauvm/CUDADataFormats/BeamSpotCUDA.h
@@ -1,7 +1,7 @@
 #ifndef CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h
 #define CUDADataFormats_BeamSpot_interface_BeamSpotCUDA_h
 
-#include "CUDACore/device_unique_ptr.h"
+#include "CUDACore/managed_unique_ptr.h"
 
 #include <cuda_runtime.h>
 
@@ -21,12 +21,12 @@ class BeamSpotCUDA {
   };
 
   BeamSpotCUDA() = default;
-  BeamSpotCUDA(Data const* data_h, cudaStream_t stream);
+  BeamSpotCUDA(Data const& data_h, int device, cudaStream_t stream);
 
   Data const* data() const { return data_d_.get(); }
 
 private:
-  cms::cuda::device::unique_ptr<Data> data_d_;
+  cms::cuda::managed::unique_ptr<Data> data_d_;
 };
 
 #endif
diff --git a/src/cudauvm/plugin-BeamSpotProducer/BeamSpotToCUDA.cc b/src/cudauvm/plugin-BeamSpotProducer/BeamSpotToCUDA.cc
index df4331be6..9ec2efa1d 100644
--- a/src/cudauvm/plugin-BeamSpotProducer/BeamSpotToCUDA.cc
+++ b/src/cudauvm/plugin-BeamSpotProducer/BeamSpotToCUDA.cc
@@ -5,7 +5,6 @@
 #include "Framework/PluginFactory.h"
 #include "Framework/EDProducer.h"
 #include "CUDACore/ScopedContext.h"
-#include "CUDACore/host_noncached_unique_ptr.h"
 
 #include <cuda_runtime.h>
 
@@ -20,20 +19,17 @@ class BeamSpotToCUDA : public edm::EDProducer {
 
 private:
   edm::EDPutTokenT<cms::cuda::Product<BeamSpotCUDA>> bsPutToken_;
-
-  cms::cuda::host::noncached::unique_ptr<BeamSpotCUDA::Data> bsHost;
 };
 
 BeamSpotToCUDA::BeamSpotToCUDA(edm::ProductRegistry& reg)
-    : bsPutToken_{reg.produces<cms::cuda::Product<BeamSpotCUDA>>()},
-      bsHost{cms::cuda::make_host_noncached_unique<BeamSpotCUDA::Data>(cudaHostAllocWriteCombined)} {}
+    : bsPutToken_{reg.produces<cms::cuda::Product<BeamSpotCUDA>>()} {}
 
 void BeamSpotToCUDA::produce(edm::Event& iEvent, const edm::EventSetup& iSetup) {
-  *bsHost = iSetup.get<BeamSpotCUDA::Data>();
+  auto const& bs = iSetup.get<BeamSpotCUDA::Data>();
 
   cms::cuda::ScopedContextProduce ctx{iEvent.streamID()};
 
-  ctx.emplace(iEvent, bsPutToken_, bsHost.get(), ctx.stream());
+  ctx.emplace(iEvent, bsPutToken_, bs, ctx.device(), ctx.stream());
 }
 
 DEFINE_FWK_MODULE(BeamSpotToCUDA);