SWDEV-417075 - add hipDrvAddMemCpyNode

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: Ie631d7b1788f10171a29d463759a3cba3b2b2007

SWDEV-417075 - add hipDrvGraphAddMemcpyNode

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I6bab3310919643e119cd0004276907e223641cfb

hipamd/include/hip/amd_detail/hip_prof_str.h

@@ -377,7 +377,8 @@ enum hip_api_id_t {
   HIP_API_ID_hipArrayGetInfo = 362,
   HIP_API_ID_hipStreamGetDevice = 363,
   HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray = 364,
-  HIP_API_ID_LAST = 364,
+  HIP_API_ID_hipDrvGraphAddMemcpyNode = 365,
+  HIP_API_ID_LAST = 365,
 
   HIP_API_ID_hipBindTexture = HIP_API_ID_NONE,
   HIP_API_ID_hipBindTexture2D = HIP_API_ID_NONE,
@@ -534,6 +535,7 @@ static inline const char* hip_api_name(const uint32_t id) {
     case HIP_API_ID_hipGraphAddMemAllocNode: return "hipGraphAddMemAllocNode";
     case HIP_API_ID_hipGraphAddMemFreeNode: return "hipGraphAddMemFreeNode";
     case HIP_API_ID_hipGraphAddMemcpyNode: return "hipGraphAddMemcpyNode";
+    case HIP_API_ID_hipDrvGraphAddMemcpyNode: return "hipDrvGraphAddMemcpyNode";
     case HIP_API_ID_hipGraphAddMemcpyNode1D: return "hipGraphAddMemcpyNode1D";
     case HIP_API_ID_hipGraphAddMemcpyNodeFromSymbol: return "hipGraphAddMemcpyNodeFromSymbol";
     case HIP_API_ID_hipGraphAddMemcpyNodeToSymbol: return "hipGraphAddMemcpyNodeToSymbol";
@@ -1148,6 +1150,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
   if (strcmp("hipUserObjectRetain", name) == 0) return HIP_API_ID_hipUserObjectRetain;
   if (strcmp("hipWaitExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipWaitExternalSemaphoresAsync;
   if (strcmp("hipExternalMemoryGetMappedMipmappedArray", name) == 0)  return HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray;
+  if (strcmp("hipDrvGraphAddMemcpyNode", name) == 0) return HIP_API_ID_hipDrvGraphAddMemcpyNode;
   return HIP_API_ID_NONE;
 }
 
@@ -3265,6 +3268,18 @@ typedef struct hip_api_data_s {
       const hipExternalMemoryMipmappedArrayDesc* mipmapDesc;
       hipExternalMemoryMipmappedArrayDesc mipmapDesc__val;
     } hipExternalMemoryGetMappedMipmappedArray;
+
+    struct {
+      hipGraphNode_t* phGraphNode;
+      hipGraphNode_t phGraphNode__val;
+      hipGraph_t hGraph;
+      const hipGraphNode_t* dependencies;
+      hipGraphNode_t dependencies__val;
+      size_t numDependencies;
+      const HIP_MEMCPY3D* copyParams;
+      HIP_MEMCPY3D copyParams__val;
+      hipCtx_t ctx;
+    } hipDrvGraphAddMemcpyNode;
   } args;
   uint64_t *phase_data;
 } hip_api_data_t;
@@ -5438,6 +5453,15 @@ typedef struct hip_api_data_s {
   cb_data.args.hipWaitExternalSemaphoresAsync.numExtSems = (unsigned int)numExtSems; \
   cb_data.args.hipWaitExternalSemaphoresAsync.stream = (hipStream_t)stream; \
 };
+// hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const hipMemcpy3DParms*', 'copyParams'), ('hipCtx_t', 'ctx')]
+#define INIT_hipDrvGraphAddMemcpyNode_CB_ARGS_DATA(cb_data) { \
+  cb_data.args.hipDrvGraphAddMemcpyNode.phGraphNode = (hipGraphNode_t*)phGraphNode; \
+  cb_data.args.hipDrvGraphAddMemcpyNode.hGraph = (hipGraph_t)hGraph; \
+  cb_data.args.hipDrvGraphAddMemcpyNode.dependencies = (const hipGraphNode_t*)dependencies; \
+  cb_data.args.hipDrvGraphAddMemcpyNode.numDependencies = (size_t)numDependencies; \
+  cb_data.args.hipDrvGraphAddMemcpyNode.copyParams = (const HIP_MEMCPY3D*)copyParams; \
+  cb_data.args.hipDrvGraphAddMemcpyNode.ctx = (hipCtx_t)ctx; \
+};
 #define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)
 
 // Macros for non-public API primitives
@@ -6917,6 +6941,12 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
       if (data->args.hipWaitExternalSemaphoresAsync.extSemArray) data->args.hipWaitExternalSemaphoresAsync.extSemArray__val = *(data->args.hipWaitExternalSemaphoresAsync.extSemArray);
       if (data->args.hipWaitExternalSemaphoresAsync.paramsArray) data->args.hipWaitExternalSemaphoresAsync.paramsArray__val = *(data->args.hipWaitExternalSemaphoresAsync.paramsArray);
       break;
+// hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const HIP_MEMCPY3D*', 'copyParams'), ('hipCtx_t', 'ctx')]
+    case HIP_API_ID_hipDrvGraphAddMemcpyNode:
+      if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode) data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val = *(data->args.hipDrvGraphAddMemcpyNode.phGraphNode);
+      if (data->args.hipDrvGraphAddMemcpyNode.dependencies) data->args.hipDrvGraphAddMemcpyNode.dependencies__val = *(data->args.hipDrvGraphAddMemcpyNode.dependencies);
+      if (data->args.hipDrvGraphAddMemcpyNode.copyParams) data->args.hipDrvGraphAddMemcpyNode.copyParams__val = *(data->args.hipDrvGraphAddMemcpyNode.copyParams);
+      break;
     default: break;
   };
 }
@@ -9773,6 +9803,18 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
       oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipWaitExternalSemaphoresAsync.stream);
       oss << ")";
     break;
+    case HIP_API_ID_hipDrvGraphAddMemcpyNode:
+      oss << "hipDrvGraphAddMemcpyNode(";
+      if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode == NULL) oss << "phGraphNode=NULL";
+      else { oss << "phGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val); }
+      oss << ", hGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.hGraph);
+      if (data->args.hipDrvGraphAddMemcpyNode.dependencies == NULL) oss << ", dependencies=NULL";
+      else { oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.dependencies__val); }
+      oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.numDependencies);
+      if (data->args.hipDrvGraphAddMemcpyNode.copyParams == NULL) oss << ", copyParams=NULL";
+      else { oss << ", copyParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.copyParams__val); }
+      oss << ")";
+    break;
     default: oss << "unknown";
   };
   return strdup(oss.str().c_str());

hipamd/src/amdhip.def

@@ -445,3 +445,4 @@ hipGraphAddMemAllocNode
 hipGraphMemAllocNodeGetParams
 hipGraphAddMemFreeNode
 hipGraphMemFreeNodeGetParams
+hipDrvGraphAddMemcpyNode

hipamd/src/hip_graph.cpp

@@ -111,6 +111,23 @@ hipError_t ihipGraphAddMemcpyNode(hip::GraphNode** pGraphNode, hip::Graph* graph
   return status;
 }
 
+hipError_t ihipDrvGraphAddMemcpyNode(hip::GraphNode** pGraphNode, hip::Graph* graph,
+                                     hip::GraphNode* const* pDependencies, size_t numDependencies,
+                                     const HIP_MEMCPY3D* pCopyParams, hipCtx_t ctx,
+                                     bool capture = true) {
+  if (pGraphNode == nullptr || graph == nullptr ||
+      (numDependencies > 0 && pDependencies == nullptr) || pCopyParams == nullptr) {
+    return hipErrorInvalidValue;
+  }
+  hipError_t status = ihipDrvMemcpy3DParamValidate(pCopyParams);
+  if (status != hipSuccess) {
+    return status;
+  }
+  *pGraphNode = new hip::GraphDrvMemcpyNode(pCopyParams);
+  status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture);
+  return status;
+}
+
 hipError_t ihipGraphAddMemcpyNode1D(hip::GraphNode** pGraphNode, hip::Graph* graph,
                                     hip::GraphNode* const* pDependencies, size_t numDependencies,
                                     void* dst, const void* src, size_t count, hipMemcpyKind kind,
@@ -1036,6 +1053,24 @@ hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
   HIP_RETURN(status);
 }
 
+hipError_t hipDrvGraphAddMemcpyNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph,
+                                    const hipGraphNode_t* dependencies, size_t numDependencies,
+                                    const HIP_MEMCPY3D* copyParams, hipCtx_t ctx) {
+  HIP_INIT_API(hipDrvGraphAddMemcpyNode, phGraphNode, hGraph, dependencies, numDependencies,
+               copyParams, ctx);
+  if (phGraphNode == nullptr || hGraph == nullptr ||
+      (numDependencies > 0 && dependencies == nullptr) || ctx == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  hip::GraphNode* node;
+  hipError_t status = ihipDrvGraphAddMemcpyNode(
+    &node, reinterpret_cast<hip::Graph*>(hGraph),
+    reinterpret_cast<hip::GraphNode* const*>(dependencies),
+    numDependencies, copyParams, ctx, false);
+  *phGraphNode = reinterpret_cast<hipGraphNode_t>(node);
+  HIP_RETURN(status);
+}
+
 hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                    const hipGraphNode_t* pDependencies, size_t numDependencies,
                                    void* dst, const void* src, size_t count, hipMemcpyKind kind) {

hipamd/src/hip_graph_helper.hpp

@@ -24,6 +24,10 @@ THE SOFTWARE.
 
 hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p);
 
+hipError_t ihipDrvMemcpy3D_validate(const HIP_MEMCPY3D* pCopy);
+
+hipError_t ihipDrvMemcpy3DParamValidate(const HIP_MEMCPY3D* pCopy);
+
 hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
 
 hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes,
@@ -31,7 +35,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
 
 void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, hip::Stream& stream);
 
-bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind);
+bool IsHtoHMemcpy(void* dst, const void* src);
 
 hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX,
                                      uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
@@ -56,6 +60,9 @@ hipError_t ihipLaunchKernelCommand(amd::Command*& command, hipFunction_t f,
 hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p,
                                hip::Stream* stream);
 
+hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy,
+                                       hip::Stream* stream);
+
 hipError_t ihipMemsetCommand(std::vector<amd::Command*>& commands, void* dst, int64_t value,
                              size_t valueSize, size_t sizeBytes, hip::Stream* stream);
 

hipamd/src/hip_graph_internal.cpp

@@ -102,158 +102,16 @@ hipError_t GraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size_t
 }
 
 hipError_t GraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParams) {
-  hipError_t status = ihipMemcpy3D_validate(pNodeParams);
+  hipError_t status;
+  status = ihipMemcpy3D_validate(pNodeParams);
   if (status != hipSuccess) {
     return status;
   }
-  size_t offset = 0;
+
   const HIP_MEMCPY3D pCopy = hip::getDrvMemcpy3DDesc(*pNodeParams);
-  // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the
-  // (unified virtual address space) base address of the source data and the bytes per row to apply.
-  // {src/dst}Array is ignored.
-  hipMemoryType srcMemoryType = pCopy.srcMemoryType;
-  if (srcMemoryType == hipMemoryTypeUnified) {
-    amd::Memory* memObj = getMemoryObject(pCopy.srcDevice, offset);
-    srcMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) &
-            memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice;
-    if (srcMemoryType == hipMemoryTypeHost) {
-      // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system
-      // memory.
-      const_cast<HIP_MEMCPY3D*>(&pCopy)->srcHost = pCopy.srcDevice;
-      const_cast<HIP_MEMCPY3D*>(&pCopy)->srcXInBytes += offset;
-    }
-  }
-  offset = 0;
-  hipMemoryType dstMemoryType = pCopy.dstMemoryType;
-  if (dstMemoryType == hipMemoryTypeUnified) {
-    amd::Memory* memObj = getMemoryObject(pCopy.dstDevice, offset);
-    dstMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) &
-            memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice;
-    if (dstMemoryType == hipMemoryTypeHost) {
-      const_cast<HIP_MEMCPY3D*>(&pCopy)->dstHost = pCopy.dstDevice;
-      const_cast<HIP_MEMCPY3D*>(&pCopy)->dstXInBytes += offset;
-    }
-  }
-  offset = 0;
-  // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned.
-  // In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning.
-  if (srcMemoryType == hipMemoryTypeHost) {
-    srcMemoryType = getMemoryObject(pCopy.srcHost, offset) ? hipMemoryTypeDevice :
-                    hipMemoryTypeHost;
-
-    if (srcMemoryType == hipMemoryTypeDevice) {
-      const_cast<HIP_MEMCPY3D*>(&pCopy)->srcDevice = const_cast<void*>(pCopy.srcHost);
-    }
-  }
-  offset = 0;
-  if (dstMemoryType == hipMemoryTypeHost) {
-    dstMemoryType = getMemoryObject(pCopy.dstHost, offset) ? hipMemoryTypeDevice :
-                    hipMemoryTypeHost;
-
-    if (dstMemoryType == hipMemoryTypeDevice) {
-      const_cast<HIP_MEMCPY3D*>(&pCopy)->dstDevice = const_cast<void*>(pCopy.dstDevice);
-    }
-  }
-
-  amd::Coord3D srcOrigin = {pCopy.srcXInBytes, pCopy.srcY, pCopy.srcZ};
-  amd::Coord3D dstOrigin = {pCopy.dstXInBytes, pCopy.dstY, pCopy.dstZ};
-  amd::Coord3D copyRegion = {pCopy.WidthInBytes, pCopy.Height, pCopy.Depth};
-
-  if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeDevice)) {
-    // Host to Device.
-
-    amd::Memory* dstMemory;
-    amd::BufferRect srcRect;
-    amd::BufferRect dstRect;
-
-    status =
-        ihipMemcpyHtoDValidate(pCopy.srcHost, pCopy.dstDevice, srcOrigin, dstOrigin, copyRegion,
-                               pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, pCopy.dstPitch,
-                               pCopy.dstPitch * pCopy.dstHeight, dstMemory, srcRect, dstRect);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeHost)) {
-    // Device to Host.
-    amd::Memory* srcMemory;
-    amd::BufferRect srcRect;
-    amd::BufferRect dstRect;
-    status =
-        ihipMemcpyDtoHValidate(pCopy.srcDevice, pCopy.dstHost, srcOrigin, dstOrigin, copyRegion,
-                               pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, pCopy.dstPitch,
-                               pCopy.dstPitch * pCopy.dstHeight, srcMemory, srcRect, dstRect);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeDevice)) {
-    // Device to Device.
-    amd::Memory* srcMemory;
-    amd::Memory* dstMemory;
-    amd::BufferRect srcRect;
-    amd::BufferRect dstRect;
-
-    status = ihipMemcpyDtoDValidate(pCopy.srcDevice, pCopy.dstDevice, srcOrigin, dstOrigin,
-                                    copyRegion, pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight,
-                                    pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcMemory,
-                                    dstMemory, srcRect, dstRect);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) {
-    amd::Image* dstImage;
-    size_t start = 0;
-    status =
-        ihipMemcpyHtoAValidate(pCopy.srcHost, pCopy.dstArray, srcOrigin, dstOrigin, copyRegion,
-                               pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, dstImage, start);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) {
-    // Image to Host.
-    amd::Image* srcImage;
-    size_t start = 0;
-
-    status =
-        ihipMemcpyAtoHValidate(pCopy.srcArray, pCopy.dstHost, srcOrigin, dstOrigin, copyRegion,
-                               pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcImage, start);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeArray)) {
-    // Device to Image.
-    amd::Image* dstImage;
-    amd::Memory* srcMemory;
-    amd::BufferRect dstRect;
-    amd::BufferRect srcRect;
-    status = ihipMemcpyDtoAValidate(pCopy.srcDevice, pCopy.dstArray, srcOrigin, dstOrigin,
-                                    copyRegion, pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight,
-                                    dstImage, srcMemory, dstRect, srcRect);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeDevice)) {
-    // Image to Device.
-    amd::BufferRect srcRect;
-    amd::BufferRect dstRect;
-    amd::Memory* dstMemory;
-    amd::Image* srcImage;
-    status = ihipMemcpyAtoDValidate(pCopy.srcArray, pCopy.dstDevice, srcOrigin, dstOrigin,
-                                    copyRegion, pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight,
-                                    dstMemory, srcImage, srcRect, dstRect);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeArray)) {
-    amd::Image* srcImage;
-    amd::Image* dstImage;
-
-    status = ihipMemcpyAtoAValidate(pCopy.srcArray, pCopy.dstArray, srcOrigin, dstOrigin,
-                                    copyRegion, srcImage, dstImage);
-    if (status != hipSuccess) {
-      return status;
-    }
-  } else {
-    return hipErrorInvalidValue;
+  status = ihipDrvMemcpy3D_validate(&pCopy);
+  if (status != hipSuccess) {
+    return status;
   }
   return hipSuccess;
 }

hipamd/src/hip_graph_internal.hpp

@@ -1173,7 +1173,8 @@ class GraphMemcpyNode : public GraphNode {
   }
 
   virtual hipError_t CreateCommand(hip::Stream* stream) {
-    if (IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.kind)) {
+    if ((copyParams_.kind == hipMemcpyHostToHost || copyParams_.kind == hipMemcpyDefault)
+      && IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr)) {
       return hipSuccess;
     }
     hipError_t status = GraphNode::CreateCommand(stream);
@@ -1188,7 +1189,8 @@ class GraphMemcpyNode : public GraphNode {
   }
 
   virtual void EnqueueCommands(hipStream_t stream) override {
-    if (isEnabled_ && IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.kind)) {
+    if ( (copyParams_.kind == hipMemcpyHostToHost || copyParams_.kind == hipMemcpyDefault) &&
+          isEnabled_ && IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr)) {
       ihipHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr,
                      copyParams_.extent.width * copyParams_.extent.height *
                      copyParams_.extent.depth, *hip::getStream(stream));
@@ -1326,7 +1328,7 @@ class GraphMemcpyNode1D : public GraphMemcpyNode {
   }
 
   virtual hipError_t CreateCommand(hip::Stream* stream) {
-    if (IsHtoHMemcpy(dst_, src_, kind_)) {
+    if ((kind_ == hipMemcpyHostToHost || kind_ == hipMemcpyDefault) && IsHtoHMemcpy(dst_, src_)) {
       return hipSuccess;
     }
     hipError_t status = GraphNode::CreateCommand(stream);
@@ -1341,7 +1343,10 @@ class GraphMemcpyNode1D : public GraphMemcpyNode {
   }
 
   virtual void EnqueueCommands(hipStream_t stream) {
-    bool isH2H = IsHtoHMemcpy(dst_, src_, kind_);
+    bool isH2H = false;
+    if ((kind_ == hipMemcpyHostToHost || kind_ == hipMemcpyDefault) && IsHtoHMemcpy(dst_, src_)) {
+      isH2H = true;
+    }
     if (!isH2H) {
       if (commands_.empty()) return;
       // commands_ should have just 1 item
@@ -2233,4 +2238,80 @@ class GraphMemFreeNode : public GraphNode {
     *params = device_ptr_;
   }
 };
+
+class GraphDrvMemcpyNode : public GraphNode {
+  HIP_MEMCPY3D copyParams_;
+
+ public:
+  GraphDrvMemcpyNode(const HIP_MEMCPY3D* pCopyParams)
+      : GraphNode(hipGraphNodeTypeMemcpy, "solid", "trapezium", "MEMCPY") {
+    copyParams_ = *pCopyParams;
+  }
+  ~GraphDrvMemcpyNode() {}
+
+  GraphDrvMemcpyNode(const GraphDrvMemcpyNode& rhs) : GraphNode(rhs) {
+    copyParams_ = rhs.copyParams_;
+  }
+
+  GraphNode* clone() const {
+    return new GraphDrvMemcpyNode(static_cast<GraphDrvMemcpyNode const&>(*this));
+  }
+
+  hipError_t CreateCommand(hip::Stream* stream) {
+    if(copyParams_.srcMemoryType == hipMemoryTypeHost &&
+       copyParams_.dstMemoryType == hipMemoryTypeHost &&
+       IsHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost)) {
+      return hipSuccess;
+    }
+    hipError_t status = GraphNode::CreateCommand(stream);
+    if (status != hipSuccess) {
+      return status;
+    }
+    commands_.reserve(1);
+    amd::Command* command;
+    status = ihipGetMemcpyParam3DCommand(command, &copyParams_, stream);
+    commands_.emplace_back(command);
+    return status;
+  }
+
+  void EnqueueCommands(hipStream_t stream) override {
+    bool isHtoH = false;
+    if(copyParams_.srcMemoryType == hipMemoryTypeHost &&
+       copyParams_.dstMemoryType == hipMemoryTypeHost &&
+       IsHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost)) {
+      isHtoH = true;
+    }
+    if (isEnabled_ && isHtoH) {
+      ihipHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost,
+                     copyParams_.WidthInBytes * copyParams_.Height *
+                     copyParams_.Depth, *hip::getStream(stream));
+      return;
+    }
+    GraphNode::EnqueueCommands(stream);
+  }
+
+  void GetParams(HIP_MEMCPY3D* params) {
+    std::memcpy(params, &copyParams_, sizeof(HIP_MEMCPY3D));
+  }
+  hipError_t SetParams(const HIP_MEMCPY3D* params) {
+    hipError_t status = ValidateParams(params);
+    if (status != hipSuccess) {
+      return status;
+    }
+    std::memcpy(&copyParams_, params, sizeof(HIP_MEMCPY3D));
+    return hipSuccess;
+  }
+  hipError_t SetParams(GraphNode* node) {
+    const GraphDrvMemcpyNode* memcpyNode = static_cast<GraphDrvMemcpyNode const*>(node);
+    return SetParams(&memcpyNode->copyParams_);
+  }
+  // ToDo: use this when commands are cloned and command params are to be updated
+  hipError_t ValidateParams(const HIP_MEMCPY3D* pNodeParams) {
+    hipError_t status = ihipDrvMemcpy3D_validate(pNodeParams);
+    if (status != hipSuccess) {
+      return status;
+    }
+    return hipSuccess;
+  }
+};
 }  // namespace hip

hipamd/src/hip_hcc.map.in

@@ -523,6 +523,7 @@ global:
     hipArrayGetInfo;
     hipArrayGetDescriptor;
     hipArray3DGetDescriptor;
+    hipDrvGraphAddMemcpyNode;
 local:
     *;
 } hip_5.5;

hipamd/src/hip_memory.cpp

@@ -470,17 +470,13 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
   }
   return hipSuccess;
 }
-
-// ================================================================================================
-bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind) {
+bool IsHtoHMemcpy(void* dst, const void* src) {
   size_t sOffset = 0;
   amd::Memory* srcMemory = getMemoryObject(src, sOffset);
   size_t dOffset = 0;
   amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
   if (srcMemory == nullptr && dstMemory == nullptr) {
-    if (kind == hipMemcpyHostToHost || kind == hipMemcpyDefault) {
-      return true;
-    }
+    return true;
   }
   return false;
 }
@@ -2123,12 +2119,12 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray_t srcArray, vo
   return hipSuccess;
 }
 
-hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy,
-                                       hip::Stream* stream) {
+void ihipCopyMemParamSet(const HIP_MEMCPY3D* pCopy, hipMemoryType& srcMemType,
+                         hipMemoryType& dstMemType)  {
   size_t offset = 0;
-  // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the
-  // (unified virtual address space) base address of the source data and the bytes per row to
-  // apply. {src/dst}Array is ignored.
+  // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch
+  // specify the (unified virtual address space)
+  // base address of the source data and the bytes per row to apply. {src/dst}Array is ignored.
   hipMemoryType srcMemoryType = pCopy->srcMemoryType;
   if (srcMemoryType == hipMemoryTypeUnified) {
     amd::Memory* memObj = getMemoryObject(pCopy->srcDevice, offset);
@@ -2138,11 +2134,14 @@ hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3
     } else {
       srcMemoryType = hipMemoryTypeHost;
     }
+
     if (srcMemoryType == hipMemoryTypeHost) {
-      // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system
-      // memory.
+      // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we
+      //  detect system memory.
       const_cast<HIP_MEMCPY3D*>(pCopy)->srcHost = pCopy->srcDevice;
       const_cast<HIP_MEMCPY3D*>(pCopy)->srcXInBytes += offset;
+      // We don't need detect memory type again for hipMemoryTypeUnified
+      const_cast<HIP_MEMCPY3D*>(pCopy)->srcMemoryType = srcMemoryType;
     }
   }
   offset = 0;
@@ -2159,12 +2158,13 @@ hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3
     if (dstMemoryType == hipMemoryTypeHost) {
       const_cast<HIP_MEMCPY3D*>(pCopy)->dstHost = pCopy->dstDevice;
       const_cast<HIP_MEMCPY3D*>(pCopy)->dstXInBytes += offset;
+      // We don't need detect memory type again for hipMemoryTypeUnified
+      const_cast<HIP_MEMCPY3D*>(pCopy)->dstMemoryType = dstMemoryType;
     }
   }
-
+  offset = 0;
   // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned.
   // In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning.
-  offset = 0;
   if (srcMemoryType == hipMemoryTypeHost) {
     srcMemoryType = getMemoryObject(pCopy->srcHost, offset) ? hipMemoryTypeDevice :
                     hipMemoryTypeHost;
@@ -2179,9 +2179,18 @@ hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3
                     hipMemoryTypeHost;
 
     if (dstMemoryType == hipMemoryTypeDevice) {
-      const_cast<HIP_MEMCPY3D*>(pCopy)->dstDevice = const_cast<void*>(pCopy->dstHost);
+      const_cast<HIP_MEMCPY3D*>(pCopy)->dstDevice = const_cast<void*>(pCopy->dstDevice);
     }
   }
+  srcMemType = srcMemoryType;
+  dstMemType = dstMemoryType;
+}
+
+hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy,
+                                       hip::Stream* stream) {
+  hipMemoryType srcMemoryType;
+  hipMemoryType dstMemoryType;
+  ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType);
 
   amd::Coord3D srcOrigin = {pCopy->srcXInBytes, pCopy->srcY, pCopy->srcZ};
   amd::Coord3D dstOrigin = {pCopy->dstXInBytes, pCopy->dstY, pCopy->dstZ};
@@ -2250,7 +2259,6 @@ inline hipError_t ihipMemcpyCmdEnqueue(amd::Command* command, bool isAsync = fal
 
 hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool isAsync = false) {
   hipError_t status;
-  size_t offset = 0;
   if (pCopy == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -2262,55 +2270,9 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool
                   pCopy->Height, pCopy->Depth);
     return hipSuccess;
   }
-  // If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the (unified virtual address space)
-  // base address of the source data and the bytes per row to apply. {src/dst}Array is ignored.
-  hipMemoryType srcMemoryType = pCopy->srcMemoryType;
-  if (srcMemoryType == hipMemoryTypeUnified) {
-    amd::Memory* memObj = getMemoryObject(pCopy->srcDevice, offset);
-    if (memObj != nullptr) {
-      srcMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) &
-            memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice;
-    } else {
-      srcMemoryType = hipMemoryTypeHost;
-    }
-
-    if (srcMemoryType == hipMemoryTypeHost) {
-      // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system memory.
-      const_cast<HIP_MEMCPY3D*>(pCopy)->srcHost = pCopy->srcDevice;
-      const_cast<HIP_MEMCPY3D*>(pCopy)->srcXInBytes += offset;
-      // We don't need detect memory type again for hipMemoryTypeUnified
-      const_cast<HIP_MEMCPY3D*>(pCopy)->srcMemoryType = srcMemoryType;
-    }
-  }
-  offset = 0;
-  hipMemoryType dstMemoryType = pCopy->dstMemoryType;
-  if (dstMemoryType == hipMemoryTypeUnified) {
-    amd::Memory* memObj = getMemoryObject(pCopy->dstDevice, offset);
-    if (memObj != nullptr) {
-      dstMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) &
-            memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice;
-    } else {
-      dstMemoryType = hipMemoryTypeHost;
-    }
-
-    if (dstMemoryType == hipMemoryTypeHost) {
-      const_cast<HIP_MEMCPY3D*>(pCopy)->dstHost = pCopy->dstDevice;
-      const_cast<HIP_MEMCPY3D*>(pCopy)->dstXInBytes += offset;
-      // We don't need detect memory type again for hipMemoryTypeUnified
-      const_cast<HIP_MEMCPY3D*>(pCopy)->dstMemoryType = dstMemoryType;
-    }
-  }
-  // If {src/dst}MemoryType is hipMemoryTypeHost, check if the memory was prepinned.
-  // In that case upgrade the copy type to hipMemoryTypeDevice to avoid extra pinning.
-  offset = 0;
-  if (srcMemoryType == hipMemoryTypeHost) {
-    srcMemoryType = getMemoryObject(pCopy->srcHost, offset) ? hipMemoryTypeDevice :
-                    hipMemoryTypeHost;
-  }
-  if (dstMemoryType == hipMemoryTypeHost) {
-    dstMemoryType = getMemoryObject(pCopy->dstHost, offset) ? hipMemoryTypeDevice :
-                    hipMemoryTypeHost;
-  }
+  hipMemoryType srcMemoryType;
+  hipMemoryType dstMemoryType;
+  ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType);
 
   if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeHost)) {
     amd::Coord3D srcOrigin = {pCopy->srcXInBytes, pCopy->srcY, pCopy->srcZ};
@@ -2842,6 +2804,170 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) {
   return hipSuccess;
 }
 
+hipError_t ihipDrvMemcpy3DParamValidate(const HIP_MEMCPY3D* p) {
+  // Passing more than one non-zero source or destination will cause hipMemcpy3D() to
+  // return an error.
+  if (p == nullptr || ((p->srcArray != nullptr) && (p->srcHost != nullptr)) ||
+      ((p->dstArray != nullptr) && (p->dstHost != nullptr))) {
+    return hipErrorInvalidValue;
+  }
+  // The struct passed to hipMemcpy3D() must specify one of srcArray or srcPtr and one of dstArray
+  // or dstPtr.
+  if (((p->srcArray == nullptr) && (p->srcHost == nullptr)) ||
+      ((p->dstArray == nullptr) && (p->dstHost == nullptr))) {
+    return hipErrorInvalidValue;
+  }
+
+  // If the source and destination are both arrays, hipMemcpy3D() will return an error if they do
+  // not have the same element size.
+  if (((p->srcArray != nullptr) && (p->dstArray != nullptr)) &&
+      (hip::getElementSize(p->srcArray) != hip::getElementSize(p->dstArray))) {
+    return hipErrorInvalidValue;
+  }
+
+  // dst/src pitch must be less than max pitch
+  auto* deviceHandle = g_devices[hip::getCurrentDevice()->deviceId()]->devices()[0];
+  const auto& info = deviceHandle->info();
+  constexpr auto int32_max = static_cast<uint64_t>(std::numeric_limits<int32_t>::max());
+  auto maxPitch = std::min(info.maxMemAllocSize_, int32_max);
+
+  // negative pitch cases
+  if (p->srcPitch >= maxPitch || p->dstPitch >= maxPitch) {
+    return hipErrorInvalidValue;
+  }
+
+  if (p->dstArray == nullptr && p->srcArray == nullptr) {
+    if ((p->WidthInBytes + p->srcXInBytes > p->srcPitch) ||
+        (p->WidthInBytes + p->dstXInBytes > p->dstPitch)) {
+      return hipErrorInvalidValue;
+    }
+  }
+  if (p->srcMemoryType < hipMemoryTypeHost || p->srcMemoryType > hipMemoryTypeManaged) {
+    return hipErrorInvalidMemcpyDirection;
+  }
+  if (p->dstMemoryType < hipMemoryTypeHost || p->dstMemoryType > hipMemoryTypeManaged) {
+    return hipErrorInvalidMemcpyDirection;
+  }
+  return hipSuccess;
+}
+
+hipError_t ihipDrvMemcpy3D_validate(const HIP_MEMCPY3D* pCopy) {
+  hipError_t status;
+  if (pCopy->WidthInBytes == 0 || pCopy->Height == 0 || pCopy->Depth == 0) {
+    LogPrintfInfo("Either Width :%d or Height: %d and Depth: %d is zero", pCopy->WidthInBytes,
+                  pCopy->Height, pCopy->Depth);
+    return hipSuccess;
+  }
+  hipMemoryType srcMemoryType;
+  hipMemoryType dstMemoryType;
+  ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType);
+
+  amd::Coord3D srcOrigin = {pCopy->srcXInBytes, pCopy->srcY, pCopy->srcZ};
+  amd::Coord3D dstOrigin = {pCopy->dstXInBytes, pCopy->dstY, pCopy->dstZ};
+  amd::Coord3D copyRegion = {pCopy->WidthInBytes, pCopy->Height, pCopy->Depth};
+
+  if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeDevice)) {
+    // Host to Device.
+
+    amd::Memory* dstMemory;
+    amd::BufferRect srcRect;
+    amd::BufferRect dstRect;
+
+    status =
+        ihipMemcpyHtoDValidate(pCopy->srcHost, pCopy->dstDevice, srcOrigin, dstOrigin, copyRegion,
+                               pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch,
+                               pCopy->dstPitch * pCopy->dstHeight, dstMemory, srcRect, dstRect);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeHost)) {
+    // Device to Host.
+    amd::Memory* srcMemory;
+    amd::BufferRect srcRect;
+    amd::BufferRect dstRect;
+    status =
+        ihipMemcpyDtoHValidate(pCopy->srcDevice, pCopy->dstHost, srcOrigin, dstOrigin, copyRegion,
+                               pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch,
+                               pCopy->dstPitch * pCopy->dstHeight, srcMemory, srcRect, dstRect);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeDevice)) {
+    // Device to Device.
+    amd::Memory* srcMemory;
+    amd::Memory* dstMemory;
+    amd::BufferRect srcRect;
+    amd::BufferRect dstRect;
+
+    status = ihipMemcpyDtoDValidate(pCopy->srcDevice, pCopy->dstDevice, srcOrigin, dstOrigin,
+                                    copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight,
+                                    pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, srcMemory,
+                                    dstMemory, srcRect, dstRect);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) {
+    amd::Image* dstImage;
+    size_t start = 0;
+
+    status =
+        ihipMemcpyHtoAValidate(pCopy->srcHost, pCopy->dstArray, srcOrigin, dstOrigin, copyRegion,
+                               pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, dstImage,
+                               start);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) {
+    // Image to Host.
+    amd::Image* srcImage;
+    size_t start = 0;
+
+    status =
+        ihipMemcpyAtoHValidate(pCopy->srcArray, pCopy->dstHost, srcOrigin, dstOrigin, copyRegion,
+                               pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, srcImage,
+                               start);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeArray)) {
+    // Device to Image.
+    amd::Image* dstImage;
+    amd::Memory* srcMemory;
+    amd::BufferRect dstRect;
+    amd::BufferRect srcRect;
+    status = ihipMemcpyDtoAValidate(pCopy->srcDevice, pCopy->dstArray, srcOrigin, dstOrigin,
+                                    copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight,
+                                    dstImage, srcMemory, dstRect, srcRect);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeDevice)) {
+    // Image to Device.
+    amd::BufferRect srcRect;
+    amd::BufferRect dstRect;
+    amd::Memory* dstMemory;
+    amd::Image* srcImage;
+    status = ihipMemcpyAtoDValidate(pCopy->srcArray, pCopy->dstDevice, srcOrigin, dstOrigin,
+                                    copyRegion, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight,
+                                    dstMemory, srcImage, srcRect, dstRect);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeArray)) {
+    amd::Image* srcImage;
+    amd::Image* dstImage;
+
+    status = ihipMemcpyAtoAValidate(pCopy->srcArray, pCopy->dstArray, srcOrigin, dstOrigin,
+                                    copyRegion, srcImage, dstImage);
+    if (status != hipSuccess) {
+      return status;
+    }
+  } else {
+    return hipErrorInvalidValue;
+  }
+  return hipSuccess;
+}
+
 hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p,
                                hip::Stream* stream) {
   const HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*p);