From 5e5b471bb2fd99dcd5b9b9f5b4c2df487b7cbaf5 Mon Sep 17 00:00:00 2001
From: Peter Boyle <paboyle@ph.ed.ac.uk>
Date: Wed, 21 Feb 2024 14:47:06 -0500
Subject: [PATCH] Put/Get and DEviceToDevice

---
 Grid/threads/Accelerator.h | 20 ++++++++++++++++++++
 1 file changed, 20 insertions(+)

diff --git a/Grid/threads/Accelerator.h b/Grid/threads/Accelerator.h
index c52afebc..d26f261e 100644
--- a/Grid/threads/Accelerator.h
+++ b/Grid/threads/Accelerator.h
@@ -275,6 +275,7 @@ inline void acceleratorCopyDeviceToDeviceAsynch(void *from,void *to,size_t bytes
 }
 inline void acceleratorCopySynchronise(void) { cudaStreamSynchronize(copyStream); };
 
+
 inline int  acceleratorIsCommunicable(void *ptr)
 {
   //  int uvm=0;
@@ -533,6 +534,12 @@ inline void acceleratorCopySynchronise(void) { auto discard=hipStreamSynchronize
 
 #endif
 
+inline void acceleratorCopyDeviceToDevice(void *from,void *to,size_t bytes)
+{
+  acceleratorCopyDeviceToDeviceAsynch(from,to,bytes);
+  acceleratorCopySynchronise();
+}
+
 //////////////////////////////////////////////
 // CPU Target - No accelerator just thread instead
 //////////////////////////////////////////////
@@ -644,4 +651,17 @@ accelerator_inline void acceleratorFence(void)
   return;
 }
 
+template<class T> void acceleratorPut(T& dev,T&host)
+{
+  acceleratorCopyToDevice(&host,&dev,sizeof(T));
+}
+template<class T> T acceleratorGet(T& dev)
+{
+  T host;
+  acceleratorCopyFromDevice(&dev,&host,sizeof(T));
+  return host;
+}
+
+
+
 NAMESPACE_END(Grid);