SWDEV-546295 - Add new APIs hipMemcpy3DBatchAsync and hipMemcpy3DPeer (#818)

Co-authored-by: Rahul Manocha <rmanocha@amd.com>

[ROCm/clr commit: e3cf50c96d]

projects/clr/hipamd/src/hip_peer.cpp

@@ -21,6 +21,8 @@
 #include <hip/hip_runtime.h>
 
 #include "hip_internal.hpp"
+#include "hip_conversions.hpp"
+
 
 namespace hip {
 
@@ -217,6 +219,39 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src
   HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice, *hip_stream, true, true));
 }
 
+hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms *p) {
+  HIP_INIT_API(hipMemcpy3DPeer, p);
+  if (p == NULL) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  if (p->srcDevice >= static_cast<int>(g_devices.size()) ||
+      p->dstDevice >= static_cast<int>(g_devices.size()) || p->srcDevice < 0 || p->dstDevice < 0) {
+    HIP_RETURN(hipErrorInvalidDevice);
+  }
+  hipMemcpy3DParms copyParms = getMemcpy3DParms(*p);
+  HIP_RETURN(ihipMemcpy3D(&copyParms, nullptr));
+}
+
+hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms *p, hipStream_t stream) {
+  HIP_INIT_API(hipMemcpy3DPeerAsync, p, stream);
+  if (p == NULL) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  if (p->srcDevice >= static_cast<int>(g_devices.size()) ||
+      p->dstDevice >= static_cast<int>(g_devices.size()) || p->srcDevice < 0 || p->dstDevice < 0) {
+    HIP_RETURN(hipErrorInvalidDevice);
+  }
+
+  getStreamPerThread(stream);
+  hip::Stream* hip_stream = hip::getStream(stream);
+  if (hip_stream == nullptr) {
+    return hipErrorInvalidValue;
+  }
+
+  hipMemcpy3DParms copyParms = getMemcpy3DParms(*p);
+  HIP_RETURN(ihipMemcpy3D(&copyParms), stream, true);
+}
+
 hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
   HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags);