[OpenMP] Improve D2D memcpy to use more efficient driver API

[tianshilei1992]

Summary

In current implementation, D2D memcpy is first to copy data back to host and then
copy from host to device. This is very efficient if the device supports D2D
memcpy, like CUDA.

In this patch, D2D memcpy will first try to use native supported driver API. If
it fails, fall back to original way. It is worth noting that D2D memcpy in this
scenerio contains two ideas:

• Same devices: this is the D2D memcpy in the CUDA context.
• Different devices: this is the PeerToPeer memcpy in the CUDA context.

My implementation merges this two parts. It chooses the best API according to
the source device and destination device.

[tianshilei1992]

Just copy the execution results from Summit.

==22767== NVPROF is profiling process 22767, command: ./d2d_memcpy
==22767== Profiling application: ./d2d_memcpy
PASS
==22767== Profiling result:
   Start  Duration            Grid Size      Block Size     Regs*    SSMem*    DSMem*      Size  Throughput  SrcMemType  DstMemType           Device   Context    Stream          Src Dev   Src Ctx          Dst Dev   Dst Ctx  Name
949.72ms  1.7920us                    -               -         -         -         -        1B  544.96KB/s      Device    Pageable  Tesla V100-SXM2         1         7                -         -                -         -  [CUDA memcpy DtoH]
949.77ms  1.7920us                    -               -         -         -         -        1B  544.96KB/s      Device    Pageable  Tesla V100-SXM2         1         7                -         -                -         -  [CUDA memcpy DtoH]
949.80ms  1.5360us                    -               -         -         -         -        4B  2.4835MB/s    Pageable      Device  Tesla V100-SXM2         1         7                -         -                -         -  [CUDA memcpy HtoD]
949.87ms  457.87ms        (2097152 1 1)       (128 1 1)        44      946B        0B         -           -           -           -  Tesla V100-SXM2         1        19                -         -                -         -  __omp_offloading_32_a7b5d52_main_l34 [128]
1.40840s  22.820ms                    -               -         -         -         -  1.0000GB  43.822GB/s      Device      Device  Tesla V100-SXM2         1        19  Tesla V100-SXM2         1  Tesla V100-SXM2         2  [CUDA memcpy PtoP]
1.46565s  1.7920us                    -               -         -         -         -        1B  544.96KB/s      Device    Pageable  Tesla V100-SXM2         2        52                -         -                -         -  [CUDA memcpy DtoH]
1.46568s  1.7920us                    -               -         -         -         -        1B  544.96KB/s      Device    Pageable  Tesla V100-SXM2         2        52                -         -                -         -  [CUDA memcpy DtoH]
1.46572s  1.5360us                    -               -         -         -         -        4B  2.4835MB/s    Pageable      Device  Tesla V100-SXM2         2        52                -         -                -         -  [CUDA memcpy HtoD]
1.48614s  492.70ms        (2097152 1 1)       (128 1 1)        46      946B        0B         -           -           -           -  Tesla V100-SXM2         2        64                -         -                -         -  __omp_offloading_32_a7b5d52_main_l49 [149]
1.97885s  159.89ms                    -               -         -         -         -  1.0000GB  6.2542GB/s      Device    Pageable  Tesla V100-SXM2         2        64                -         -                -         -  [CUDA memcpy DtoH]

Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy

With PeerToPeer copy, the throughput can reach 43+GB/s.

[jdoerfert]

Some high-level comment. I haven't read the main logic part in depth yet but it looks reasonable to me. Others should chime in too.

openmp/libomptarget/plugins/cuda/src/rtl.cpp

804

Any reason not to make this a function and call it?

openmp/libomptarget/src/api.cpp

178

Please move this into the conditional.

openmp/libomptarget/src/device.cpp

398

We use the name? I somehow feel uneasy about this. Don't we have some form of ID?

openmp/libomptarget/src/device.h

160

Documentation please.

189

Documentation please.

[grokos]

openmp/libomptarget/src/device.cpp

398

I agree. Be careful because RTLInfoTy::RTLName is only available in debug builds, so this piece of code will break if we compile the library in release mode.

You can use a direct pointer comparison RTL == OtherDevice.RTL (devices managed by the same RTL will point to the same RTLInfoTy object).

[tianshilei1992]

openmp/libomptarget/src/device.cpp

398

Yes, that's why I removed the macro.
The pointer comparison may not work considering that there is in fact one case can violate it: OpenCL ICD, although it is not part of OpenMP. Maybe adding a new plugin interface here is more appropriate.

[grokos]

I'm happy with the changes, I think the patch looks good now.

@jdoerfert: Can you accept the patch once you are happy with it as well?

[jdoerfert]

LGTM, one nit below.

openmp/libomptarget/src/device.h

189

^

[tianshilei1992]

Updated documentation

[protze.joachim]

I'm looking through the tests in this directory and try to run them on a GPU.

Are those tests supposed to run on GPU?

openmp/libomptarget/test/offloading/d2d_memcpy.c

24–26

What is this supposed to do?
I guess you intended to use a different condition in the if, right?

Why not unconditionally set the dst_device to num_devices-1?

[tianshilei1992]

In D80649#2172381, @protze.joachim wrote:

I'm looking through the tests in this directory and try to run them on a GPU.

Are those tests supposed to run on GPU?

This patch was to add an interface and only implemented NVIDIA part for now. Other target devices can implement the interface if they support similar thing.

openmp/libomptarget/test/offloading/d2d_memcpy.c

24–26

Yeah, that would be better.

[protze.joachim]

I included this change in D84557, where I enable nvptx64 testing for most libomptarget tests.

openmp/libomptarget/include/omptargetplugin.h

 
 // Return an integer different from zero if the provided device image can be
 // supported by the runtime. The functionality is similar to comparing the
 // result of __tgt__rtl__load__binary to NULL. However, this is meant to be a
 // lightweight query to determine if the RTL is suitable for an image without
 // having to load the library, which can be expensive.
 int32_t __tgt_rtl_is_valid_binary(__tgt_device_image *Image);
 
+// Return an integer other than zero if the data can be exchaned from SrcDevId
+// to DstDevId. If it is data exchangable, the device plugin should provide
+// function to move data from source device to destination device directly.
+int32_t __tgt_rtl_is_data_exchangable(int32_t SrcDevId, int32_t DstDevId);
+
 // Initialize the requires flags for the device.
 int64_t __tgt_rtl_init_requires(int64_t RequiresFlags);
 
 // Initialize the specified device. In case of success return 0; otherwise
 // return an error code.
 int32_t __tgt_rtl_init_device(int32_t ID);
 
 // Pass an executable image section described by image to the specified

 int32_t __tgt_rtl_data_retrieve(int32_t ID, void *HostPtr, void *TargetPtr,
                                 int64_t Size);
 
 // Asynchronous version of __tgt_rtl_data_retrieve
 int32_t __tgt_rtl_data_retrieve_async(int32_t ID, void *HostPtr,
                                       void *TargetPtr, int64_t Size,
                                       __tgt_async_info *AsyncInfoPtr);
 
+// Copy the data content from one target device to another target device using
+// its address. This operation does not need to copy data back to host and then
+// from host to another device. In case of success, return zero. Otherwise,
+// return an error code.
+int32_t __tgt_rtl_data_exchange(int32_t SrcID, void *SrcPtr, int32_t DstID,
+                                void *DstPtr, int64_t Size);
+
+// Asynchronous version of __tgt_rtl_data_exchange
+int32_t __tgt_rtl_data_exchange_async(int32_t SrcID, void *SrcPtr,
+                                      int32_t DesID, void *DstPtr, int64_t Size,
+                                      __tgt_async_info *AsyncInfoPtr);
+
 // De-allocate the data referenced by target ptr on the device. In case of
 // success, return zero. Otherwise, return an error code.
 int32_t __tgt_rtl_data_delete(int32_t ID, void *TargetPtr);
 
 // Transfer control to the offloaded entry Entry on the target device.
 // Args and Offsets are arrays of NumArgs size of target addresses and
 // offsets. An offset should be added to the target address before passing it
 // to the outlined function on device side. If AsyncInfoPtr is nullptr, it is

openmp/libomptarget/plugins/cuda/src/rtl.cpp

bool checkResult(CUresult Err, const char *ErrMsg) {
   if (Err == CUDA_SUCCESS)
     return true;
 
   DP(ErrMsg);
   CUDA_ERR_STRING(Err);
   return false;
 }
 
+int memcpyDtoD(const void *SrcPtr, void *DstPtr, int64_t Size,
+               CUstream Stream) {
+  CUresult Err =
+      cuMemcpyDtoDAsync((CUdeviceptr)DstPtr, (CUdeviceptr)SrcPtr, Size, Stream);
+
+  if (Err != CUDA_SUCCESS) {
+    DP("Error when copying data from device to device. Pointers: src "
+       "= " DPxMOD ", dst = " DPxMOD ", size = %" PRId64 "\n",
+       DPxPTR(SrcPtr), DPxPTR(DstPtr), Size);
+    CUDA_ERR_STRING(Err);
+    return OFFLOAD_FAIL;
+  }
+
+  return OFFLOAD_SUCCESS;
+}
+
 // Structure contains per-device data
 struct DeviceDataTy {
   std::list<FuncOrGblEntryTy> FuncGblEntries;
   CUcontext Context = nullptr;
   // Device properties
   int ThreadsPerBlock = 0;
   int BlocksPerGrid = 0;
   int WarpSize = 0;
if (Err != CUDA_SUCCESS) {
          DPxPTR(HstPtr), DPxPTR(TgtPtr), Size);
       CUDA_ERR_STRING(Err);
       return OFFLOAD_FAIL;
     }
 
     return OFFLOAD_SUCCESS;
   }
 
+  int dataExchange(int SrcDevId, const void *SrcPtr, int DstDevId, void *DstPtr,
+                   int64_t Size, __tgt_async_info *AsyncInfoPtr) const {
+    assert(AsyncInfoPtr && "AsyncInfoPtr is nullptr");
+
+    CUresult Err = cuCtxSetCurrent(DeviceData[SrcDevId].Context);
+    if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
+      return OFFLOAD_FAIL;
+
+    CUstream Stream = getStream(SrcDevId, AsyncInfoPtr);
+
+    // If they are two devices, we try peer to peer copy first
+    if (SrcDevId != DstDevId) {
+      int CanAccessPeer = 0;
+      Err = cuDeviceCanAccessPeer(&CanAccessPeer, SrcDevId, DstDevId);
+      if (Err != CUDA_SUCCESS) {
+        DP("Error returned from cuDeviceCanAccessPeer. src = %" PRId32
+           ", dst = %" PRId32 "\n",
+           SrcDevId, DstDevId);
+        CUDA_ERR_STRING(Err);
+        return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
+      }
+
+      if (!CanAccessPeer) {
+        DP("P2P memcpy not supported so fall back to D2D memcpy");
+        return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
+      }
+
+      Err = cuCtxEnablePeerAccess(DeviceData[DstDevId].Context, 0);
+      if (Err != CUDA_SUCCESS) {
+        DP("Error returned from cuCtxEnablePeerAccess. src = %" PRId32
+           ", dst = %" PRId32 "\n",
+           SrcDevId, DstDevId);
+        CUDA_ERR_STRING(Err);
+        return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
+      }
+
+      Err = cuMemcpyPeerAsync((CUdeviceptr)DstPtr, DeviceData[DstDevId].Context,
+                              (CUdeviceptr)SrcPtr, DeviceData[SrcDevId].Context,
+                              Size, Stream);
+      if (Err == CUDA_SUCCESS)
+        return OFFLOAD_SUCCESS;
+
+      DP("Error returned from cuMemcpyPeerAsync. src_ptr = " DPxMOD
+         ", src_id =%" PRId32 ", dst_ptr = %" DPxMOD ", dst_id =%" PRId32 "\n",
+         SrcPtr, SrcDevId, DstPtr, DstDevId);
+      CUDA_ERR_STRING(Err);
+    }
+
+    return memcpyDtoD(SrcPtr, DstPtr, Size, Stream);
+  }
+
   int dataDelete(const int DeviceId, void *TgtPtr) const {
     CUresult Err = cuCtxSetCurrent(DeviceData[DeviceId].Context);
     if (!checkResult(Err, "Error returned from cuCtxSetCurrent\n"))
       return OFFLOAD_FAIL;
 
     Err = cuMemFree((CUdeviceptr)TgtPtr);
     if (!checkResult(Err, "Error returned from cuMemFree\n"))
       return OFFLOAD_FAIL;

 int32_t __tgt_rtl_number_of_devices() { return DeviceRTL.getNumOfDevices(); }
 
 int64_t __tgt_rtl_init_requires(int64_t RequiresFlags) {
   DP("Init requires flags to %ld\n", RequiresFlags);
   DeviceRTL.setRequiresFlag(RequiresFlags);
   return RequiresFlags;
 }
 
+int32_t __tgt_rtl_is_data_exchangable(int32_t src_dev_id, int dst_dev_id) {
+  if (DeviceRTL.isValidDeviceId(src_dev_id) &&
+      DeviceRTL.isValidDeviceId(dst_dev_id))
+    return 1;
+
+  return 0;
+}
+
 int32_t __tgt_rtl_init_device(int32_t device_id) {
   assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
 
   return DeviceRTL.initDevice(device_id);
 }
 
 __tgt_target_table *__tgt_rtl_load_binary(int32_t device_id,
                                           __tgt_device_image *image) {
int32_t __tgt_rtl_data_retrieve_async(int32_t device_id, void *hst_ptr,
                                       __tgt_async_info *async_info_ptr) {
   assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
   assert(async_info_ptr && "async_info_ptr is nullptr");
 
   return DeviceRTL.dataRetrieve(device_id, hst_ptr, tgt_ptr, size,
                                 async_info_ptr);
 }
 
+int32_t __tgt_rtl_data_exchange_async(int32_t src_dev_id, void *src_ptr,
+                                      int dst_dev_id, void *dst_ptr,
+                                      int64_t size,
+                                      __tgt_async_info *async_info_ptr) {
+  assert(DeviceRTL.isValidDeviceId(src_dev_id) && "src_dev_id is invalid");
+  assert(DeviceRTL.isValidDeviceId(dst_dev_id) && "dst_dev_id is invalid");
+  assert(async_info_ptr && "async_info_ptr is nullptr");
+
+  return DeviceRTL.dataExchange(src_dev_id, src_ptr, dst_dev_id, dst_ptr, size,
+                                async_info_ptr);
+}
+
+int32_t __tgt_rtl_data_exchange(int32_t src_dev_id, void *src_ptr,
+                                int32_t dst_dev_id, void *dst_ptr,
+                                int64_t size) {
+  assert(DeviceRTL.isValidDeviceId(src_dev_id) && "src_dev_id is invalid");
+  assert(DeviceRTL.isValidDeviceId(dst_dev_id) && "dst_dev_id is invalid");
+
+  __tgt_async_info async_info;
+  const int32_t rc = __tgt_rtl_data_exchange_async(
+      src_dev_id, src_ptr, dst_dev_id, dst_ptr, size, &async_info);
+  if (rc != OFFLOAD_SUCCESS)
+    return OFFLOAD_FAIL;
+
+  return __tgt_rtl_synchronize(src_dev_id, &async_info);
+}
+
 int32_t __tgt_rtl_data_delete(int32_t device_id, void *tgt_ptr) {
   assert(DeviceRTL.isValidDeviceId(device_id) && "device_id is invalid");
 
   return DeviceRTL.dataDelete(device_id, tgt_ptr);
 }
 
 int32_t __tgt_rtl_run_target_team_region(int32_t device_id, void *tgt_entry_ptr,
                                          void **tgt_args,

openmp/libomptarget/plugins/exports

 VERS1.0 {
   global:
     __tgt_rtl_is_valid_binary;
+    __tgt_rtl_is_data_exchangable;
     __tgt_rtl_number_of_devices;
     __tgt_rtl_init_requires;
     __tgt_rtl_init_device;
     __tgt_rtl_load_binary;
     __tgt_rtl_data_alloc;
     __tgt_rtl_data_submit;
     __tgt_rtl_data_submit_async;
     __tgt_rtl_data_retrieve;
     __tgt_rtl_data_retrieve_async;
+    __tgt_rtl_data_exchange;
+    __tgt_rtl_data_exchange_async;
     __tgt_rtl_data_delete;
     __tgt_rtl_run_target_team_region;
     __tgt_rtl_run_target_team_region_async;
     __tgt_rtl_run_target_region;
     __tgt_rtl_run_target_region_async;
     __tgt_rtl_synchronize;
   local:
     *;
 };

openmp/libomptarget/src/api.cpp

if (src_device == omp_get_initial_device() &&
     DeviceTy& DstDev = Devices[dst_device];
     rc = DstDev.data_submit(dstAddr, srcAddr, length, nullptr);
   } else if (dst_device == omp_get_initial_device()) {
     DP("copy from device to host\n");
     DeviceTy& SrcDev = Devices[src_device];
     rc = SrcDev.data_retrieve(dstAddr, srcAddr, length, nullptr);
   } else {
     DP("copy from device to device\n");
-    void *buffer = malloc(length);
-    DeviceTy& SrcDev = Devices[src_device];
-    DeviceTy& DstDev = Devices[dst_device];
+    DeviceTy &SrcDev = Devices[src_device];
+    DeviceTy &DstDev = Devices[dst_device];
+    // First try to use D2D memcpy which is more efficient. If fails, fall back
+    // to unefficient way.
+    if (SrcDev.isDataExchangable(DstDev)) {
+      rc = SrcDev.data_exchange(srcAddr, DstDev, dstAddr, length, nullptr);
+      if (rc == OFFLOAD_SUCCESS)
+        return OFFLOAD_SUCCESS;
+    }
+
+    void *buffer = malloc(length);
     rc = SrcDev.data_retrieve(buffer, srcAddr, length, nullptr);
     if (rc == OFFLOAD_SUCCESS)
       rc = DstDev.data_submit(dstAddr, buffer, length, nullptr);
     free(buffer);
   }
 
   DP("omp_target_memcpy returns %d\n", rc);
   return rc;

openmp/libomptarget/src/device.h

DeviceTy& operator=(const DeviceTy &d) {
     HostDataToTargetMap = d.HostDataToTargetMap;
     PendingCtorsDtors = d.PendingCtorsDtors;
     ShadowPtrMap = d.ShadowPtrMap;
     LoopTripCnt = d.LoopTripCnt;
 
     return *this;
   }
 
+  // Return true if data can be copied to DstDevice directly
+  bool isDataExchangable(const DeviceTy& DstDevice);
+
   uint64_t getMapEntryRefCnt(void *HstPtrBegin);
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
       bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true,
       bool HasCloseModifier = false);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
       bool UpdateRefCount, bool &IsHostPtr);
   int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
                     bool HasCloseModifier = false);
   int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
   int disassociatePtr(void *HstPtrBegin);
 
   // calls to RTL
   int32_t initOnce();
   __tgt_target_table *load_binary(void *Img);
 
   // Data transfer. When AsyncInfoPtr is nullptr, the transfer will be
   // synchronous.
+  // Copy data from host to device
   int32_t data_submit(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
                       __tgt_async_info *AsyncInfoPtr);
+  // Copy data from device back to host
   int32_t data_retrieve(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size,
                         __tgt_async_info *AsyncInfoPtr);
+  // Copy data from current device to destination device directly
+  int32_t data_exchange(void *SrcPtr, DeviceTy DstDev, void *DstPtr,
+                        int64_t Size, __tgt_async_info *AsyncInfoPtr);
 
   int32_t run_region(void *TgtEntryPtr, void **TgtVarsPtr,
                      ptrdiff_t *TgtOffsets, int32_t TgtVarsSize,
                      __tgt_async_info *AsyncInfoPtr);
   int32_t run_team_region(void *TgtEntryPtr, void **TgtVarsPtr,
                           ptrdiff_t *TgtOffsets, int32_t TgtVarsSize,
                           int32_t NumTeams, int32_t ThreadLimit,
                           uint64_t LoopTripCount,

openmp/libomptarget/src/device.cpp

int32_t DeviceTy::data_retrieve(void *HstPtrBegin, void *TgtPtrBegin,
                                 int64_t Size, __tgt_async_info *AsyncInfoPtr) {
   if (!AsyncInfoPtr || !RTL->data_retrieve_async || !RTL->synchronize)
     return RTL->data_retrieve(RTLDeviceID, HstPtrBegin, TgtPtrBegin, Size);
   else
     return RTL->data_retrieve_async(RTLDeviceID, HstPtrBegin, TgtPtrBegin, Size,
                                     AsyncInfoPtr);
 }
 
+// Copy data from current device to destination device directly
+int32_t DeviceTy::data_exchange(void *SrcPtr, DeviceTy DstDev, void *DstPtr,
+                                int64_t Size, __tgt_async_info *AsyncInfoPtr) {
+  if (!AsyncInfoPtr || !RTL->data_exchange_async || !RTL->synchronize) {
+    assert(RTL->data_exchange && "RTL->data_exchange is nullptr");
+    return RTL->data_exchange(RTLDeviceID, SrcPtr, DstDev.RTLDeviceID, DstPtr,
+                              Size);
+  } else
+    return RTL->data_exchange_async(RTLDeviceID, SrcPtr, DstDev.RTLDeviceID,
+                                    DstPtr, Size, AsyncInfoPtr);
+}
+
 // Run region on device
 int32_t DeviceTy::run_region(void *TgtEntryPtr, void **TgtVarsPtr,
                              ptrdiff_t *TgtOffsets, int32_t TgtVarsSize,
                              __tgt_async_info *AsyncInfoPtr) {
   if (!AsyncInfoPtr || !RTL->run_region || !RTL->synchronize)
     return RTL->run_region(RTLDeviceID, TgtEntryPtr, TgtVarsPtr, TgtOffsets,
                            TgtVarsSize);
   else
return RTL->run_team_region(RTLDeviceID, TgtEntryPtr, TgtVarsPtr,
                                 TgtOffsets, TgtVarsSize, NumTeams, ThreadLimit,
                                 LoopTripCount);
   else
     return RTL->run_team_region_async(RTLDeviceID, TgtEntryPtr, TgtVarsPtr,
                                       TgtOffsets, TgtVarsSize, NumTeams,
                                       ThreadLimit, LoopTripCount, AsyncInfoPtr);
 }
 
+// Whether data can be copied to DstDevice directly
+bool DeviceTy::isDataExchangable(const DeviceTy &DstDevice) {
+  if (RTL != DstDevice.RTL || !RTL->is_data_exchangable)
+    return false;
+
+  if (RTL->is_data_exchangable(RTLDeviceID, DstDevice.RTLDeviceID))
+    return (RTL->data_exchange != nullptr) ||
+           (RTL->data_exchange_async != nullptr);
+
+  return false;
+}
+
 /// Check whether a device has an associated RTL and initialize it if it's not
 /// already initialized.
 bool device_is_ready(int device_num) {
   DP("Checking whether device %d is ready.\n", device_num);
   // Devices.size() can only change while registering a new
   // library, so try to acquire the lock of RTLs' mutex.
   RTLsMtx->lock();
   size_t Devices_size = Devices.size();

openmp/libomptarget/src/rtl.h

 #include <vector>
 
 // Forward declarations.
 struct DeviceTy;
 struct __tgt_bin_desc;
 
 struct RTLInfoTy {
   typedef int32_t(is_valid_binary_ty)(void *);
+  typedef int32_t(is_data_exchangable_ty)(int32_t, int32_t);
   typedef int32_t(number_of_devices_ty)();
   typedef int32_t(init_device_ty)(int32_t);
   typedef __tgt_target_table *(load_binary_ty)(int32_t, void *);
   typedef void *(data_alloc_ty)(int32_t, int64_t, void *);
   typedef int32_t(data_submit_ty)(int32_t, void *, void *, int64_t);
   typedef int32_t(data_submit_async_ty)(int32_t, void *, void *, int64_t,
                                         __tgt_async_info *);
   typedef int32_t(data_retrieve_ty)(int32_t, void *, void *, int64_t);
   typedef int32_t(data_retrieve_async_ty)(int32_t, void *, void *, int64_t,
                                           __tgt_async_info *);
+  typedef int32_t(data_exchange_ty)(int32_t, void *, int32_t, void *, int64_t);
+  typedef int32_t(data_exchange_async_ty)(int32_t, void *, int32_t, void *,
+                                          int64_t, __tgt_async_info *);
   typedef int32_t(data_delete_ty)(int32_t, void *);
   typedef int32_t(run_region_ty)(int32_t, void *, void **, ptrdiff_t *,
                                  int32_t);
   typedef int32_t(run_region_async_ty)(int32_t, void *, void **, ptrdiff_t *,
                                        int32_t, __tgt_async_info *);
   typedef int32_t(run_team_region_ty)(int32_t, void *, void **, ptrdiff_t *,
                                       int32_t, int32_t, int32_t, uint64_t);
   typedef int32_t(run_team_region_async_ty)(int32_t, void *, void **,
struct RTLInfoTy {
   void *LibraryHandler = nullptr;
 
 #ifdef OMPTARGET_DEBUG
   std::string RTLName;
 #endif
 
   // Functions implemented in the RTL.
   is_valid_binary_ty *is_valid_binary = nullptr;
+  is_data_exchangable_ty *is_data_exchangable = nullptr;
   number_of_devices_ty *number_of_devices = nullptr;
   init_device_ty *init_device = nullptr;
   load_binary_ty *load_binary = nullptr;
   data_alloc_ty *data_alloc = nullptr;
   data_submit_ty *data_submit = nullptr;
   data_submit_async_ty *data_submit_async = nullptr;
   data_retrieve_ty *data_retrieve = nullptr;
   data_retrieve_async_ty *data_retrieve_async = nullptr;
+  data_exchange_ty *data_exchange = nullptr;
+  data_exchange_async_ty *data_exchange_async = nullptr;
   data_delete_ty *data_delete = nullptr;
   run_region_ty *run_region = nullptr;
   run_region_async_ty *run_region_async = nullptr;
   run_team_region_ty *run_team_region = nullptr;
   run_team_region_async_ty *run_team_region_async = nullptr;
   init_requires_ty *init_requires = nullptr;
   synchronize_ty *synchronize = nullptr;
 
#endif
   RTLInfoTy(const RTLInfoTy &r) {
     Idx = r.Idx;
     NumberOfDevices = r.NumberOfDevices;
     LibraryHandler = r.LibraryHandler;
 #ifdef OMPTARGET_DEBUG
     RTLName = r.RTLName;
 #endif
     is_valid_binary = r.is_valid_binary;
+    is_data_exchangable = r.is_data_exchangable;
     number_of_devices = r.number_of_devices;
     init_device = r.init_device;
     load_binary = r.load_binary;
     data_alloc = r.data_alloc;
     data_submit = r.data_submit;
     data_submit_async = r.data_submit_async;
     data_retrieve = r.data_retrieve;
     data_retrieve_async = r.data_retrieve_async;
+    data_exchange = r.data_exchange;
+    data_exchange_async = r.data_exchange_async;
     data_delete = r.data_delete;
     run_region = r.run_region;
     run_region_async = r.run_region_async;
     run_team_region = r.run_team_region;
     run_team_region_async = r.run_team_region_async;
     init_requires = r.init_requires;
     isUsed = r.isUsed;
     synchronize = r.synchronize;

openmp/libomptarget/src/rtl.cpp

*((void **)&R.data_submit_async) =
         dlsym(dynlib_handle, "__tgt_rtl_data_submit_async");
     *((void **)&R.data_retrieve_async) =
         dlsym(dynlib_handle, "__tgt_rtl_data_retrieve_async");
     *((void **)&R.run_region_async) =
         dlsym(dynlib_handle, "__tgt_rtl_run_target_region_async");
     *((void **)&R.run_team_region_async) =
         dlsym(dynlib_handle, "__tgt_rtl_run_target_team_region_async");
     *((void **)&R.synchronize) = dlsym(dynlib_handle, "__tgt_rtl_synchronize");
+    *((void **)&R.data_exchange) =
+        dlsym(dynlib_handle, "__tgt_rtl_data_exchange");
+    *((void **)&R.data_exchange_async) =
+        dlsym(dynlib_handle, "__tgt_rtl_data_exchange_async");
+    *((void **)&R.is_data_exchangable) =
+        dlsym(dynlib_handle, "__tgt_rtl_is_data_exchangable");
 
     // No devices are supported by this RTL?
     if (!(R.NumberOfDevices = R.number_of_devices())) {
       DP("No devices supported in this RTL\n");
       continue;
     }
 
     DP("Registering RTL %s supporting %d devices!\n", R.RTLName.c_str(),

openmp/libomptarget/test/offloading/d2d_memcpy.c

+// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-aarch64-unknown-linux-gnu | %fcheck-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64-ibm-linux-gnu | %fcheck-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-powerpc64le-ibm-linux-gnu | %fcheck-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-x86_64-pc-linux-gnu | %fcheck-x86_64-pc-linux-gnu -allow-empty
+
+#include <assert.h>
+#include <omp.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+const int magic_num = 7;
+
+int main(int argc, char *argv[]) {
+  const int N = 128;
+  const int num_devices = omp_get_num_devices();
+
+  // No target device, just return
+  if (num_devices == 0) {
+    printf("PASS\n");
+    return 0;
+  }
+
+  const int src_device = 0;
+  int dst_device = 1;
+  if (dst_device >= num_devices)
+    dst_device = num_devices - 1;
+
+  int length = N * sizeof(int);
+  int *src_ptr = omp_target_alloc(length, src_device);
+  int *dst_ptr = omp_target_alloc(length, dst_device);
+
+  assert(src_ptr && "src_ptr is NULL");
+  assert(dst_ptr && "dst_ptr is NULL");
+
+#pragma omp target teams distribute parallel for device(src_device) \
+                   is_device_ptr(src_ptr)
+  for (int i = 0; i < N; ++i) {
+    src_ptr[i] = magic_num;
+  }
+
+  int rc =
+      omp_target_memcpy(dst_ptr, src_ptr, length, 0, 0, dst_device, src_device);
+
+  assert(rc == 0 && "error in omp_target_memcpy");
+
+  int *buffer = malloc(length);
+
+  assert(buffer && "failed to allocate host buffer");
+
+#pragma omp target teams distribute parallel for device(dst_device) \
+                   map(from: buffer[0:N]) is_device_ptr(dst_ptr)
+  for (int i = 0; i < N; ++i) {
+    buffer[i] = dst_ptr[i] + magic_num;
+  }
+
+  for (int i = 0; i < N; ++i)
+    assert(buffer[i] == 2 * magic_num);
+
+  printf("PASS\n");
+
+  // Free host and device memory
+  free(buffer);
+  omp_target_free(src_ptr, src_device);
+  omp_target_free(dst_ptr, dst_device);
+
+  return 0;
+}
+
+// CHECK: PASS