SWDEV-24806 - Added support for memcpy related graph APIs

hipGraphMemcpyNodeSetParams1D, hipGraphExecMemcpyNodeSetParams1D, hipGraphExecMemcpyNodeSetParams, hipGraphAddMemcpyNodeFromSymbol, hipGraphMemcpyNodeSetParamsFromSymbol, hipGraphExecMemcpyNodeSetParamsFromSymbol, hipGraphAddMemcpyNodeToSymbol, hipGraphMemcpyNodeSetParamsToSymbol, hipGraphExecMemcpyNodeSetParamsToSymbol

Change-Id: I6896144f340867d938aa1777880c78c46ca6f893

hipamd/src/amdhip.def

@@ -323,3 +323,12 @@ hipGraphNodeFindInClone
 hipGraphAddChildGraphNode
 hipGraphChildGraphNodeGetGraph
 hipGraphExecChildGraphNodeSetParams
+hipGraphAddMemcpyNodeFromSymbol
+hipGraphMemcpyNodeSetParamsFromSymbol
+hipGraphExecMemcpyNodeSetParamsFromSymbol
+hipGraphAddMemcpyNodeToSymbol
+hipGraphMemcpyNodeSetParamsToSymbol
+hipGraphExecMemcpyNodeSetParamsToSymbol
+hipGraphExecMemcpyNodeSetParams
+hipGraphMemcpyNodeSetParams1D
+hipGraphExecMemcpyNodeSetParams1D

hipamd/src/hip_graph.cpp

@@ -413,6 +413,21 @@ hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                                dst, src, count, kind));
 }
 
+hipError_t hipGraphMemcpyNodeSetParams1D(hipGraphNode_t node, void* dst, const void* src,
+                                         size_t count, hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphMemcpyNodeSetParams1D, node, dst, src, count, kind);
+  reinterpret_cast<hipGraphMemcpyNode1D*>(node)->SetParams(dst, src, count, kind);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphExecMemcpyNodeSetParams1D(hipGraphExec_t hGraphExec, hipGraphNode_t node,
+                                             void* dst, const void* src, size_t count,
+                                             hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphExecMemcpyNodeSetParams1D, hGraphExec, node, dst, src, count, kind);
+  HIP_RETURN(
+      reinterpret_cast<hipGraphMemcpyNode1D*>(node)->SetCommandParams(dst, src, count, kind));
+}
+
 hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                  const hipGraphNode_t* pDependencies, size_t numDependencies,
                                  const hipMemsetParams* pMemsetParams) {
@@ -572,6 +587,12 @@ hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DPar
   HIP_RETURN(hipSuccess);
 }
 
+hipError_t hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
+                                           hipMemcpy3DParms* pNodeParams) {
+  HIP_INIT_API(hipGraphExecMemcpyNodeSetParams, hGraphExec, node, pNodeParams);
+  HIP_RETURN(reinterpret_cast<hipGraphMemcpyNode*>(node)->SetCommandParams(pNodeParams));
+}
+
 hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams) {
   HIP_INIT_API(hipGraphMemsetNodeGetParams, node, pNodeParams);
   if (node == nullptr || pNodeParams == nullptr) {
@@ -832,3 +853,75 @@ hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t origina
   }
   HIP_RETURN(hipErrorInvalidValue);
 }
+
+hipError_t hipGraphAddMemcpyNodeFromSymbol(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                           const hipGraphNode_t* pDependencies,
+                                           size_t numDependencies, void* dst, const void* symbol,
+                                           size_t count, size_t offset, hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphAddMemcpyNodeFromSymbol, pGraphNode, graph, pDependencies, numDependencies,
+               dst, symbol, count, offset, kind);
+  size_t sym_size = 0;
+  hipDeviceptr_t device_ptr = nullptr;
+
+  hipError_t status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
+  if (status != hipSuccess) {
+    return status;
+  }
+  *pGraphNode = new hipGraphMemcpyNodeFromSymbol(dst, symbol, count, offset, kind);
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphMemcpyNodeSetParamsFromSymbol(hipGraphNode_t node, void* dst, const void* symbol,
+                                                 size_t count, size_t offset, hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphMemcpyNodeSetParamsFromSymbol, node, dst, symbol, count, offset, kind);
+  reinterpret_cast<hipGraphMemcpyNodeFromSymbol*>(node)->SetParams(dst, symbol, count, offset,
+                                                                   kind);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node,
+                                                     void* dst, const void* symbol, size_t count,
+                                                     size_t offset, hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphExecMemcpyNodeSetParamsFromSymbol, hGraphExec, node, dst, symbol, count,
+               offset, kind);
+  HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeFromSymbol*>(node)->SetCommandParams(
+      dst, symbol, count, offset, kind));
+}
+
+hipError_t hipGraphAddMemcpyNodeToSymbol(hipGraphNode_t* pGraphNode, hipGraph_t graph,
+                                         const hipGraphNode_t* pDependencies,
+                                         size_t numDependencies, const void* symbol,
+                                         const void* src, size_t count, size_t offset,
+                                         hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphAddMemcpyNodeToSymbol, pGraphNode, graph, pDependencies, numDependencies,
+               symbol, src, count, offset, kind);
+  size_t sym_size = 0;
+  hipDeviceptr_t device_ptr = nullptr;
+  hipError_t status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
+  if (status != hipSuccess) {
+    return status;
+  }
+  *pGraphNode = new hipGraphMemcpyNodeToSymbol(symbol, src, count, offset, kind);
+  ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
+  HIP_RETURN(hipSuccess);
+}
+
+hipError_t hipGraphMemcpyNodeSetParamsToSymbol(hipGraphNode_t node, const void* symbol,
+                                               const void* src, size_t count, size_t offset,
+                                               hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphMemcpyNodeSetParamsToSymbol, symbol, src, count, offset, kind);
+  reinterpret_cast<hipGraphMemcpyNodeToSymbol*>(node)->SetParams(symbol, src, count, offset, kind);
+  HIP_RETURN(hipSuccess);
+}
+
+
+hipError_t hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node,
+                                                   const void* symbol, const void* src,
+                                                   size_t count, size_t offset,
+                                                   hipMemcpyKind kind) {
+  HIP_INIT_API(hipGraphExecMemcpyNodeSetParamsToSymbol, hGraphExec, node, symbol, src, count,
+               offset, kind);
+  HIP_RETURN(reinterpret_cast<hipGraphMemcpyNodeToSymbol*>(node)->SetCommandParams(
+      symbol, src, count, offset, kind));
+}

hipamd/src/hip_graph_helper.hpp

@@ -1,3 +1,5 @@
+#include "hip_conversions.hpp"
+
 hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p);
 
 hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
@@ -33,3 +35,54 @@ hipError_t ihipMemsetCommand(std::vector<amd::Command*>& commands, void* dst, in
 
 hipError_t ihipMemset3DCommand(std::vector<amd::Command*>& commands, hipPitchedPtr pitchedDevPtr,
                                int value, hipExtent extent, amd::HostQueue* queue);
+
+hipError_t ihipMemcpySymbol_validate(const void* symbol, size_t sizeBytes, size_t offset,
+                                     size_t& sym_size, hipDeviceptr_t& device_ptr);
+
+hipError_t ihipMemcpyAtoDValidate(hipArray* srcArray, void* dstDevice, amd::Coord3D srcOrigin,
+                                  amd::Coord3D dstOrigin, amd::Coord3D copyRegion,
+                                  size_t dstRowPitch, size_t dstSlicePitch, amd::Memory*& dstMemory,
+                                  amd::Image*& srcImage, amd::BufferRect& srcRect,
+                                  amd::BufferRect& dstRect);
+
+hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray* dstArray, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage,
+                                  amd::Memory*& srcMemory, amd::BufferRect& dstRect,
+                                  amd::BufferRect& srcRect);
+
+hipError_t ihipMemcpyDtoDValidate(void* srcDevice, void* dstDevice, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, size_t dstRowPitch,
+                                  size_t dstSlicePitch, amd::Memory*& srcMemory,
+                                  amd::Memory*& dstMemory, amd::BufferRect& srcRect,
+                                  amd::BufferRect& dstRect);
+
+
+hipError_t ihipMemcpyDtoHValidate(void* srcDevice, void* dstHost, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, size_t dstRowPitch,
+                                  size_t dstSlicePitch, amd::Memory*& srcMemory,
+                                  amd::BufferRect& srcRect, amd::BufferRect& dstRect);
+
+hipError_t ihipMemcpyHtoDValidate(const void* srcHost, void* dstDevice, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, size_t dstRowPitch,
+                                  size_t dstSlicePitch, amd::Memory*& dstMemory,
+                                  amd::BufferRect& srcRect, amd::BufferRect& dstRect);
+
+
+hipError_t ihipMemcpyAtoAValidate(hipArray* srcArray, hipArray* dstArray, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  amd::Image*& srcImage, amd::Image*& dstImage);
+
+
+hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray* dstArray, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage,
+                                  amd::BufferRect& srcRect);
+
+hipError_t ihipMemcpyAtoHValidate(hipArray* srcArray, void* dstHost, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t dstRowPitch, size_t dstSlicePitch, amd::Image*& srcImage,
+                                  amd::BufferRect& dstRect);

hipamd/src/hip_graph_internal.cpp

@@ -47,6 +47,233 @@ const char* GetGraphNodeTypeString(uint32_t op) {
 
 int hipGraphNode::nextID = 0;
 
+hipError_t hipGraphMemcpyNode1D::SetCommandParams(void* dst, const void* src, size_t count,
+                                                  hipMemcpyKind kind) {
+  hipError_t status = ihipMemcpy_validate(dst, src, count, kind);
+  if (status != hipSuccess) {
+    return status;
+  }
+  size_t sOffsetOrig = 0;
+  amd::Memory* origSrcMemory = getMemoryObject(src_, sOffsetOrig);
+  size_t dOffsetOrig = 0;
+  amd::Memory* origDstMemory = getMemoryObject(dst_, dOffsetOrig);
+
+  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 (origDstMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
+      return hipErrorInvalidValue;
+    }
+    amd::WriteMemoryCommand* command = reinterpret_cast<amd::WriteMemoryCommand*>(commands_[0]);
+    command->setParams(*dstMemory->asBuffer(), dOffset, count, src);
+  } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
+    if (origSrcMemory->getContext().devices()[0] != srcMemory->getContext().devices()[0]) {
+      return hipErrorInvalidValue;
+    }
+    amd::ReadMemoryCommand* command = reinterpret_cast<amd::ReadMemoryCommand*>(commands_[0]);
+    command->setParams(*srcMemory->asBuffer(), sOffset, count, dst);
+  } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
+    if (origDstMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
+      return hipErrorInvalidValue;
+    }
+    if (origSrcMemory->getContext().devices()[0] != srcMemory->getContext().devices()[0]) {
+      return hipErrorInvalidValue;
+    }
+    amd::CopyMemoryP2PCommand* command = reinterpret_cast<amd::CopyMemoryP2PCommand*>(commands_[0]);
+    command->setParams(*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, count);
+    // Make sure runtime has valid memory for the command execution. P2P access
+    // requires page table mapping on the current device to another GPU memory
+    if (!static_cast<amd::CopyMemoryP2PCommand*>(command)->validateMemory()) {
+      delete command;
+      return hipErrorInvalidValue;
+    }
+  } else {
+    amd::CopyMemoryCommand* command = reinterpret_cast<amd::CopyMemoryCommand*>(commands_[0]);
+    command->setParams(*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, count);
+  }
+}
+
+hipError_t hipGraphMemcpyNode::SetCommandParams(const hipMemcpy3DParms* pNodeParams) {
+  hipError_t status = ihipMemcpy3D_validate(pNodeParams);
+  if (status != hipSuccess) {
+    return status;
+  }
+  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) {
+    srcMemoryType =
+        amd::MemObjMap::FindMemObj(pCopy.srcDevice) ? hipMemoryTypeDevice : 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;
+    }
+  }
+  hipMemoryType dstMemoryType = pCopy.dstMemoryType;
+  if (dstMemoryType == hipMemoryTypeUnified) {
+    dstMemoryType =
+        amd::MemObjMap::FindMemObj(pCopy.dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost;
+    if (srcMemoryType == hipMemoryTypeHost) {
+      const_cast<HIP_MEMCPY3D*>(&pCopy)->dstHost = pCopy.dstDevice;
+    }
+  }
+
+  // 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) {
+    amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.srcHost);
+    srcMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost;
+    if (srcMemoryType == hipMemoryTypeDevice) {
+      const_cast<HIP_MEMCPY3D*>(&pCopy)->srcDevice = const_cast<void*>(pCopy.srcHost);
+    }
+  }
+  if (dstMemoryType == hipMemoryTypeHost) {
+    amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.dstHost);
+    dstMemoryType = mem ? 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;
+    }
+    amd::WriteMemoryCommand* command = reinterpret_cast<amd::WriteMemoryCommand*>(commands_[0]);
+    command->setParams(*dstMemory, {dstRect.start_, 0, 0}, copyRegion, pCopy.srcHost, dstRect,
+                       srcRect);
+  } 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;
+    }
+    amd::ReadMemoryCommand* command = reinterpret_cast<amd::ReadMemoryCommand*>(commands_[0]);
+    command->setParams(*srcMemory, {srcRect.start_, 0, 0}, copyRegion, pCopy.dstHost, srcRect,
+                       dstRect);
+    command->setSource(*srcMemory);
+    command->setOrigin({srcRect.start_, 0, 0});
+    command->setSize(copyRegion);
+    command->setDestination(pCopy.dstHost);
+    command->setBufRect(srcRect);
+    command->setHostRect(dstRect);
+  } 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;
+    }
+    amd::CopyMemoryCommand* command = reinterpret_cast<amd::CopyMemoryCommand*>(commands_[0]);
+    command->setParams(*srcMemory, *dstMemory, {srcRect.start_, 0, 0}, {dstRect.start_, 0, 0},
+                       copyRegion, srcRect, dstRect);
+  } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) {
+    amd::Image* dstImage;
+    amd::BufferRect srcRect;
+
+    status =
+        ihipMemcpyHtoAValidate(pCopy.srcHost, pCopy.dstArray, srcOrigin, dstOrigin, copyRegion,
+                               pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, dstImage, srcRect);
+    if (status != hipSuccess) {
+      return status;
+    }
+    amd::WriteMemoryCommand* command = reinterpret_cast<amd::WriteMemoryCommand*>(commands_[0]);
+    command->setParams(*dstImage, dstOrigin, copyRegion,
+                       static_cast<const char*>(pCopy.srcHost) + srcRect.start_, pCopy.srcPitch,
+                       pCopy.srcPitch * pCopy.srcHeight);
+  } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) {
+    // Image to Host.
+    amd::Image* srcImage;
+    amd::BufferRect dstRect;
+
+    status =
+        ihipMemcpyAtoHValidate(pCopy.srcArray, pCopy.dstHost, srcOrigin, dstOrigin, copyRegion,
+                               pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcImage, dstRect);
+    if (status != hipSuccess) {
+      return status;
+    }
+    amd::ReadMemoryCommand* command = reinterpret_cast<amd::ReadMemoryCommand*>(commands_[0]);
+    command->setParams(*srcImage, srcOrigin, copyRegion,
+                       static_cast<char*>(pCopy.dstHost) + dstRect.start_, pCopy.dstPitch,
+                       pCopy.dstPitch * pCopy.dstHeight);
+  } 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;
+    }
+    amd::CopyMemoryCommand* command = reinterpret_cast<amd::CopyMemoryCommand*>(commands_[0]);
+    command->setParams(*srcMemory, *dstImage, srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
+  } 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;
+    }
+    amd::CopyMemoryCommand* command = reinterpret_cast<amd::CopyMemoryCommand*>(commands_[0]);
+    command->setParams(*srcImage, *dstMemory, srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
+  } 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;
+    }
+    amd::CopyMemoryCommand* command = reinterpret_cast<amd::CopyMemoryCommand*>(commands_[0]);
+    command->setParams(*srcImage, *dstImage, srcOrigin, dstOrigin, copyRegion);
+  } else {
+    ShouldNotReachHere();
+  }
+  return hipSuccess;
+}
+
 void ihipGraph::AddNode(const Node& node) {
   vertices_.emplace_back(node);
   ClPrint(amd::LOG_INFO, amd::LOG_CODE, "[hipGraph] Add %s(%p)\n",

hipamd/src/hip_graph_internal.hpp

@@ -391,25 +391,28 @@ class hipGraphMemcpyNode : public hipGraphNode {
   void SetParams(const hipMemcpy3DParms* params) {
     std::memcpy(pCopyParams_, params, sizeof(hipMemcpy3DParms));
   }
+  hipError_t SetCommandParams(const hipMemcpy3DParms* pNodeParams);
 };
 
 class hipGraphMemcpyNode1D : public hipGraphNode {
+ protected:
   void* dst_;
   const void* src_;
   size_t count_;
   hipMemcpyKind kind_;
 
  public:
-  hipGraphMemcpyNode1D(void* dst, const void* src, size_t count, hipMemcpyKind kind)
-      : hipGraphNode(hipGraphNodeTypeMemcpy1D), dst_(dst), src_(src), count_(count), kind_(kind) {}
+  hipGraphMemcpyNode1D(void* dst, const void* src, size_t count, hipMemcpyKind kind,
+                       hipGraphNodeType type = hipGraphNodeTypeMemcpy1D)
+      : hipGraphNode(type), dst_(dst), src_(src), count_(count), kind_(kind) {}
 
   ~hipGraphMemcpyNode1D() {}
 
   hipGraphNode* clone() const {
     return new hipGraphMemcpyNode1D(static_cast<hipGraphMemcpyNode1D const&>(*this));
   }
-
-  hipError_t CreateCommand(amd::HostQueue* queue) {
+  
+  virtual hipError_t CreateCommand(amd::HostQueue* queue) {
     commands_.reserve(1);
     amd::Command* command = nullptr;
     hipError_t status = ihipMemcpyCommand(command, dst_, src_, count_, kind_, *queue);
@@ -423,24 +426,21 @@ class hipGraphMemcpyNode1D : public hipGraphNode {
     count_ = count;
     kind_ = kind;
   }
+
+  hipError_t SetCommandParams(void* dst, const void* src, size_t count, hipMemcpyKind kind);
 };
 
-template <class T> class hipGraphMemcpyNodeFromSymbol : public hipGraphNode {
-  void* dst_;
-  const T& symbol_;
-  size_t count_;
+class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D {
+  const void* symbol_;
   size_t offset_;
-  hipMemcpyKind kind_;
 
  public:
   hipGraphMemcpyNodeFromSymbol(void* dst, const void* symbol, size_t count, size_t offset,
                                hipMemcpyKind kind)
-      : hipGraphNode(hipGraphNodeTypeMemcpyFromSymbol),
-        dst_(dst),
+      : hipGraphMemcpyNode1D(dst, nullptr, count, kind, hipGraphNodeTypeMemcpyFromSymbol),
         symbol_(symbol),
-        count_(count),
-        offset_(offset),
-        kind_(kind) {}
+        offset_(offset) {}
+
   ~hipGraphMemcpyNodeFromSymbol() {}
 
   hipGraphNode* clone() const {
@@ -448,7 +448,23 @@ template <class T> class hipGraphMemcpyNodeFromSymbol : public hipGraphNode {
         static_cast<hipGraphMemcpyNodeFromSymbol const&>(*this));
   }
 
-  hipError_t CreateCommand(amd::HostQueue* queue);
+  hipError_t CreateCommand(amd::HostQueue* queue) {
+    commands_.reserve(1);
+    amd::Command* command = nullptr;
+    size_t sym_size = 0;
+    hipDeviceptr_t device_ptr = nullptr;
+
+    hipError_t status = ihipMemcpySymbol_validate(symbol_, count_, offset_, sym_size, device_ptr);
+    if (status != hipSuccess) {
+      return status;
+    }
+    status = ihipMemcpyCommand(command, dst_, device_ptr, count_, kind_, *queue);
+    if (status != hipSuccess) {
+      return status;
+    }
+    commands_.emplace_back(command);
+    return status;
+  }
 
   void SetParams(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind) {
     dst_ = dst;
@@ -457,39 +473,74 @@ template <class T> class hipGraphMemcpyNodeFromSymbol : public hipGraphNode {
     offset_ = offset;
     kind_ = kind;
   }
-};
 
-template <class T> class hipGraphMemcpyNodeToSymbol : public hipGraphNode {
-  const T& symbol_;
-  const void* src_;
-  size_t count_;
+  hipError_t SetCommandParams(void* dst, const void* symbol, size_t count, size_t offset,
+                              hipMemcpyKind kind) {
+    size_t sym_size = 0;
+    hipDeviceptr_t device_ptr = nullptr;
+
+    hipError_t status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
+    if (status != hipSuccess) {
+      return status;
+    }
+    return hipGraphMemcpyNode1D::SetCommandParams(dst, device_ptr, count, kind);
+  }
+};
+class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D {
+  const void* symbol_;
   size_t offset_;
-  hipMemcpyKind kind_;
 
  public:
-  hipGraphMemcpyNodeToSymbol(const T& symbol, void* src, size_t count, size_t offset,
+  hipGraphMemcpyNodeToSymbol(const void* symbol, const void* src, size_t count, size_t offset,
                              hipMemcpyKind kind)
-      : hipGraphNode(hipGraphNodeTypeMemcpyToSymbol),
+      : hipGraphMemcpyNode1D(nullptr, src, count, kind, hipGraphNodeTypeMemcpyToSymbol),
         symbol_(symbol),
-        src_(src),
-        count_(count),
-        offset_(offset),
-        kind_(kind) {}
+        offset_(offset) {}
+
   ~hipGraphMemcpyNodeToSymbol() {}
 
   hipGraphNode* clone() const {
     return new hipGraphMemcpyNodeToSymbol(static_cast<hipGraphMemcpyNodeToSymbol const&>(*this));
   }
 
-  hipError_t CreateCommand(amd::HostQueue* queue);
+  hipError_t CreateCommand(amd::HostQueue* queue) {
+    commands_.reserve(1);
+    amd::Command* command = nullptr;
+    size_t sym_size = 0;
+    hipDeviceptr_t device_ptr = nullptr;
 
-  void SetParams(const T& symbol, void* src, size_t count, size_t offset, hipMemcpyKind kind) {
+    hipError_t status = ihipMemcpySymbol_validate(symbol_, count_, offset_, sym_size, device_ptr);
+    if (status != hipSuccess) {
+      return status;
+    }
+    status = ihipMemcpyCommand(command, device_ptr, src_, count_, kind_, *queue);
+    if (status != hipSuccess) {
+      return status;
+    }
+    commands_.emplace_back(command);
+    return status;
+  }
+
+  void SetParams(const void* symbol, const void* src, size_t count, size_t offset,
+                 hipMemcpyKind kind) {
     symbol_ = symbol;
     src_ = src;
     count_ = count;
     offset_ = offset;
     kind_ = kind;
   }
+
+  hipError_t SetCommandParams(const void* symbol, const void* src, size_t count, size_t offset,
+                              hipMemcpyKind kind) {
+    size_t sym_size = 0;
+    hipDeviceptr_t device_ptr = nullptr;
+
+    hipError_t status = ihipMemcpySymbol_validate(symbol, count, offset, sym_size, device_ptr);
+    if (status != hipSuccess) {
+      return status;
+    }
+    return hipGraphMemcpyNode1D::SetCommandParams(device_ptr, src, count, kind);
+  }
 };
 
 class hipGraphMemsetNode : public hipGraphNode {

hipamd/src/hip_hcc.def.in

@@ -321,3 +321,12 @@ hipGraphNodeFindInClone
 hipGraphAddChildGraphNode
 hipGraphChildGraphNodeGetGraph
 hipGraphExecChildGraphNodeSetParams
+hipGraphAddMemcpyNodeFromSymbol
+hipGraphMemcpyNodeSetParamsFromSymbol
+hipGraphExecMemcpyNodeSetParamsFromSymbol
+hipGraphAddMemcpyNodeToSymbol
+hipGraphMemcpyNodeSetParamsToSymbol
+hipGraphExecMemcpyNodeSetParamsToSymbol
+hipGraphExecMemcpyNodeSetParams
+hipGraphMemcpyNodeSetParams1D
+hipGraphExecMemcpyNodeSetParams1D

hipamd/src/hip_hcc.map.in

@@ -347,6 +347,15 @@ global:
     hipGraphAddChildGraphNode;
     hipGraphChildGraphNodeGetGraph;
     hipGraphExecChildGraphNodeSetParams;
+    hipGraphAddMemcpyNodeFromSymbol;
+    hipGraphMemcpyNodeSetParamsFromSymbol;
+    hipGraphExecMemcpyNodeSetParamsFromSymbol;
+    hipGraphAddMemcpyNodeToSymbol;
+    hipGraphMemcpyNodeSetParamsToSymbol;
+    hipGraphExecMemcpyNodeSetParamsToSymbol;
+    hipGraphExecMemcpyNodeSetParams;
+    hipGraphMemcpyNodeSetParams1D;
+    hipGraphExecMemcpyNodeSetParams1D;
 local:
     *;
 } hip_4.4;

hipamd/src/hip_memory.cpp

@@ -1183,12 +1183,13 @@ hipError_t hipMemcpyDtoHAsync(void* dstHost,
   HIP_RETURN_DURATION(ihipMemcpy(dstHost, srcDevice, ByteCount, hipMemcpyDeviceToHost, *hip::getQueue(stream), true));
 }
 
-hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, void* dstDevice,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch,
-                                 amd::HostQueue* queue) {
+hipError_t ihipMemcpyAtoDValidate(hipArray* srcArray, void* dstDevice, amd::Coord3D srcOrigin,
+                                  amd::Coord3D dstOrigin, amd::Coord3D copyRegion,
+                                  size_t dstRowPitch, size_t dstSlicePitch, amd::Memory*& dstMemory,
+                                  amd::Image*& srcImage, amd::BufferRect& srcRect,
+                                  amd::BufferRect& dstRect) {
   size_t dstOffset = 0;
-  amd::Memory* dstMemory = getMemoryObject(dstDevice, dstOffset);
+  dstMemory = getMemoryObject(dstDevice, dstOffset);
   if (srcArray == nullptr || (dstMemory == nullptr)) {
     return hipErrorInvalidValue;
   }
@@ -1197,19 +1198,18 @@ hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, voi
     return hipErrorInvalidValue;
   }
 
-  amd::Image* srcImage = as_amd(srcMemObj)->asImage();
+  srcImage = as_amd(srcMemObj)->asImage();
   // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
   const size_t elementSize = srcImage->getImageFormat().getElementSize();
   static_cast<size_t*>(srcOrigin)[0] /= elementSize;
   static_cast<size_t*>(copyRegion)[0] /= elementSize;
 
-  amd::BufferRect srcRect;
+
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
                       srcImage->getRowPitch(), srcImage->getSlicePitch())) {
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect dstRect;
   if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
                       dstRowPitch, dstSlicePitch)) {
     return hipErrorInvalidValue;
@@ -1223,10 +1223,26 @@ hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, voi
       !dstMemory->validateRegion(dstOrigin, {copySizeInBytes, 0, 0})) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
 
+hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, void* dstDevice,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch,
+                                 amd::HostQueue* queue) {
+  amd::BufferRect srcRect;
+  amd::BufferRect dstRect;
+  amd::Memory* dstMemory;
+  amd::Image* srcImage;
+  hipError_t status =
+      ihipMemcpyAtoDValidate(srcArray, dstDevice, srcOrigin, dstOrigin, copyRegion, dstRowPitch,
+                             dstSlicePitch, dstMemory, srcImage, srcRect, dstRect);
+  if (status != hipSuccess) {
+    return status;
+  }
   command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_IMAGE_TO_BUFFER,
-                                        amd::Command::EventWaitList{}, *srcImage, *dstMemory,
-                                        srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
+                                       amd::Command::EventWaitList{}, *srcImage, *dstMemory,
+                                       srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
 
   if (command == nullptr) {
     return hipErrorOutOfMemory;
@@ -1234,12 +1250,13 @@ hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, voi
   return hipSuccess;
 }
 
-hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArray* dstArray,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
-                                 amd::HostQueue* queue) {
+hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray* dstArray, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage,
+                                  amd::Memory*& srcMemory, amd::BufferRect& dstRect,
+                                  amd::BufferRect& srcRect) {
   size_t srcOffset = 0;
-  amd::Memory* srcMemory = getMemoryObject(srcDevice, srcOffset);
+  srcMemory = getMemoryObject(srcDevice, srcOffset);
   if ((srcMemory == nullptr) || dstArray == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -1248,13 +1265,13 @@ hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArr
     return hipErrorInvalidValue;
   }
 
-  amd::Image* dstImage = as_amd(dstMemObj)->asImage();
+  dstImage = as_amd(dstMemObj)->asImage();
   // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
   const size_t elementSize = dstImage->getImageFormat().getElementSize();
   static_cast<size_t*>(dstOrigin)[0] /= elementSize;
   static_cast<size_t*>(copyRegion)[0] /= elementSize;
 
-  amd::BufferRect srcRect;
+
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
                       srcRowPitch, srcSlicePitch)) {
     return hipErrorInvalidValue;
@@ -1262,7 +1279,6 @@ hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArr
   srcRect.start_ += srcOffset;
   srcRect.end_ += srcOffset;
 
-  amd::BufferRect dstRect;
   if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
                       dstImage->getRowPitch(), dstImage->getSlicePitch())) {
     return hipErrorInvalidValue;
@@ -1274,10 +1290,26 @@ hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArr
       !dstImage->validateRegion(dstOrigin, copyRegion)) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
 
+hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArray* dstArray,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
+                                 amd::HostQueue* queue) {
+  amd::Image* dstImage;
+  amd::Memory* srcMemory;
+  amd::BufferRect dstRect;
+  amd::BufferRect srcRect;
+  hipError_t status =
+      ihipMemcpyDtoAValidate(srcDevice, dstArray, srcOrigin, dstOrigin, copyRegion, srcRowPitch,
+                             srcSlicePitch, dstImage, srcMemory, dstRect, srcRect);
+  if (status != hipSuccess) {
+    return status;
+  }
   command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER_TO_IMAGE,
-                                        amd::Command::EventWaitList{}, *srcMemory, *dstImage,
-                                        srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
+                                       amd::Command::EventWaitList{}, *srcMemory, *dstImage,
+                                       srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
 
   if (command == nullptr) {
     return hipErrorOutOfMemory;
@@ -1285,20 +1317,21 @@ hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArr
   return hipSuccess;
 }
 
-hipError_t ihipMemcpyDtoDCommand(amd::Command*& command, void* srcDevice, void* dstDevice,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
-                                 size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue) {
+hipError_t ihipMemcpyDtoDValidate(void* srcDevice, void* dstDevice, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, size_t dstRowPitch,
+                                  size_t dstSlicePitch, amd::Memory*& srcMemory,
+                                  amd::Memory*& dstMemory, amd::BufferRect& srcRect,
+                                  amd::BufferRect& dstRect) {
   size_t srcOffset = 0;
-  amd::Memory* srcMemory = getMemoryObject(srcDevice, srcOffset);
+  srcMemory = getMemoryObject(srcDevice, srcOffset);
   size_t dstOffset = 0;
-  amd::Memory* dstMemory = getMemoryObject(dstDevice, dstOffset);
+  dstMemory = getMemoryObject(dstDevice, dstOffset);
 
   if ((srcMemory == nullptr) || (dstMemory == nullptr)) {
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect srcRect;
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
                       srcRowPitch, srcSlicePitch)) {
     return hipErrorInvalidValue;
@@ -1312,7 +1345,6 @@ hipError_t ihipMemcpyDtoDCommand(amd::Command*& command, void* srcDevice, void*
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect dstRect;
   if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
                       dstRowPitch, dstSlicePitch)) {
     return hipErrorInvalidValue;
@@ -1325,7 +1357,26 @@ hipError_t ihipMemcpyDtoDCommand(amd::Command*& command, void* srcDevice, void*
   if (!dstMemory->validateRegion(dstStart, dstSize)) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
 
+hipError_t ihipMemcpyDtoDCommand(amd::Command*& command, void* srcDevice, void* dstDevice,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
+                                 size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue) {
+  amd::Memory* srcMemory;
+  amd::Memory* dstMemory;
+  amd::BufferRect srcRect;
+  amd::BufferRect dstRect;
+
+  hipError_t status = ihipMemcpyDtoDValidate(srcDevice, dstDevice, srcOrigin, dstOrigin, copyRegion,
+                                             srcRowPitch, srcSlicePitch, dstRowPitch, dstSlicePitch,
+                                             srcMemory, dstMemory, srcRect, dstRect);
+  if (status != hipSuccess) {
+    return status;
+  }
+  amd::Coord3D srcStart(srcRect.start_, 0, 0);
+  amd::Coord3D dstStart(dstRect.start_, 0, 0);
   amd::CopyMemoryCommand* copyCommand = new amd::CopyMemoryCommand(
       *queue, CL_COMMAND_COPY_BUFFER_RECT, amd::Command::EventWaitList{}, *srcMemory, *dstMemory,
       srcStart, dstStart, copyRegion, srcRect, dstRect);
@@ -1342,18 +1393,18 @@ hipError_t ihipMemcpyDtoDCommand(amd::Command*& command, void* srcDevice, void*
   return hipSuccess;
 }
 
-hipError_t ihipMemcpyDtoHCommand(amd::Command*& command, void* srcDevice, void* dstHost,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
-                                 size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue) {
+hipError_t ihipMemcpyDtoHValidate(void* srcDevice, void* dstHost, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, size_t dstRowPitch,
+                                  size_t dstSlicePitch, amd::Memory*& srcMemory,
+                                  amd::BufferRect& srcRect, amd::BufferRect& dstRect) {
   size_t srcOffset = 0;
-  amd::Memory* srcMemory = getMemoryObject(srcDevice, srcOffset);
+  srcMemory = getMemoryObject(srcDevice, srcOffset);
 
   if ((srcMemory == nullptr) || (dstHost == nullptr)) {
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect srcRect;
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
                       srcRowPitch, srcSlicePitch)) {
     return hipErrorInvalidValue;
@@ -1367,12 +1418,27 @@ hipError_t ihipMemcpyDtoHCommand(amd::Command*& command, void* srcDevice, void*
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect dstRect;
   if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
                       dstRowPitch, dstSlicePitch)) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
 
+hipError_t ihipMemcpyDtoHCommand(amd::Command*& command, void* srcDevice, void* dstHost,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
+                                 size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue) {
+  amd::Memory* srcMemory;
+  amd::BufferRect srcRect;
+  amd::BufferRect dstRect;
+  hipError_t status = ihipMemcpyDtoHValidate(srcDevice, dstHost, srcOrigin, dstOrigin, copyRegion,
+                                             srcRowPitch, srcSlicePitch, dstRowPitch, dstSlicePitch,
+                                             srcMemory, srcRect, dstRect);
+  if (status != hipSuccess) {
+    return status;
+  }
+  amd::Coord3D srcStart(srcRect.start_, 0, 0);
   amd::ReadMemoryCommand* readCommand =
       new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER_RECT, amd::Command::EventWaitList{},
                                  *srcMemory, srcStart, copyRegion, dstHost, srcRect, dstRect);
@@ -1389,24 +1455,23 @@ hipError_t ihipMemcpyDtoHCommand(amd::Command*& command, void* srcDevice, void*
   return hipSuccess;
 }
 
-hipError_t ihipMemcpyHtoDCommand(amd::Command*& command, const void* srcHost, void* dstDevice,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
-                                 size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue) {
+hipError_t ihipMemcpyHtoDValidate(const void* srcHost, void* dstDevice, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, size_t dstRowPitch,
+                                  size_t dstSlicePitch, amd::Memory*& dstMemory,
+                                  amd::BufferRect& srcRect, amd::BufferRect& dstRect) {
   size_t dstOffset = 0;
-  amd::Memory* dstMemory = getMemoryObject(dstDevice, dstOffset);
+  dstMemory = getMemoryObject(dstDevice, dstOffset);
 
   if ((srcHost == nullptr) || (dstMemory == nullptr)) {
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect srcRect;
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
                       srcRowPitch, srcSlicePitch)) {
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect dstRect;
   if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
                       dstRowPitch, dstSlicePitch)) {
     return hipErrorInvalidValue;
@@ -1419,7 +1484,24 @@ hipError_t ihipMemcpyHtoDCommand(amd::Command*& command, const void* srcHost, vo
   if (!dstMemory->validateRegion(dstStart, dstSize)) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
 
+hipError_t ihipMemcpyHtoDCommand(amd::Command*& command, const void* srcHost, void* dstDevice,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
+                                 size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue) {
+  amd::Memory* dstMemory;
+  amd::BufferRect srcRect;
+  amd::BufferRect dstRect;
+
+  hipError_t status = ihipMemcpyHtoDValidate(srcHost, dstDevice, srcOrigin, dstOrigin, copyRegion,
+                                             srcRowPitch, srcSlicePitch, dstRowPitch, dstSlicePitch,
+                                             dstMemory, srcRect, dstRect);
+  if (status != hipSuccess) {
+    return status;
+  }
+  amd::Coord3D dstStart(dstRect.start_, 0, 0);
   amd::WriteMemoryCommand* writeCommand = new amd::WriteMemoryCommand(
       *queue, CL_COMMAND_WRITE_BUFFER_RECT, amd::Command::EventWaitList{}, *dstMemory, dstStart,
       copyRegion, srcHost, dstRect, srcRect);
@@ -1468,9 +1550,9 @@ hipError_t ihipMemcpyHtoH(const void* srcHost, void* dstHost, amd::Coord3D srcOr
   return hipSuccess;
 }
 
-hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hipArray* dstArray,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, amd::HostQueue* queue) {
+hipError_t ihipMemcpyAtoAValidate(hipArray* srcArray, hipArray* dstArray, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  amd::Image*& srcImage, amd::Image*& dstImage) {
   if (dstArray == nullptr || srcArray == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -1480,8 +1562,8 @@ hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hip
     return hipErrorInvalidValue;
   }
 
-  amd::Image* srcImage = as_amd(srcMemObj)->asImage();
-  amd::Image* dstImage = as_amd(dstMemObj)->asImage();
+  srcImage = as_amd(srcMemObj)->asImage();
+  dstImage = as_amd(dstMemObj)->asImage();
 
   // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
   // Note that src and dst should have the same element size.
@@ -1496,10 +1578,23 @@ hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hip
       !dstImage->validateRegion(dstOrigin, copyRegion)) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
 
-  command =
-      new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_IMAGE, amd::Command::EventWaitList{},
-                                 *srcImage, *dstImage, srcOrigin, dstOrigin, copyRegion);
+hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hipArray* dstArray,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, amd::HostQueue* queue) {
+  amd::Image* srcImage;
+  amd::Image* dstImage;
+
+  hipError_t status = ihipMemcpyAtoAValidate(srcArray, dstArray, srcOrigin, dstOrigin, copyRegion,
+                                             srcImage, dstImage);
+  if (status != hipSuccess) {
+    return status;
+  }
+
+  command = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_IMAGE, amd::Command::EventWaitList{},
+                                       *srcImage, *dstImage, srcOrigin, dstOrigin, copyRegion);
 
   if (command == nullptr) {
     return hipErrorOutOfMemory;
@@ -1507,10 +1602,10 @@ hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hip
   return hipSuccess;
 }
 
-hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hipArray* dstArray,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
-                                 amd::HostQueue* queue) {
+hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray* dstArray, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t srcRowPitch, size_t srcSlicePitch, amd::Image*& dstImage,
+                                  amd::BufferRect& srcRect) {
   if ((srcHost == nullptr) || dstArray == nullptr) {
     return hipErrorInvalidValue;
   }
@@ -1519,13 +1614,12 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect srcRect;
   if (!srcRect.create(static_cast<size_t*>(srcOrigin), static_cast<size_t*>(copyRegion),
                       srcRowPitch, srcSlicePitch)) {
     return hipErrorInvalidValue;
   }
 
-  amd::Image* dstImage = as_amd(dstMemObj)->asImage();
+  dstImage = as_amd(dstMemObj)->asImage();
   // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
   const size_t elementSize = dstImage->getImageFormat().getElementSize();
   static_cast<size_t*>(dstOrigin)[0] /= elementSize;
@@ -1534,6 +1628,21 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi
   if (!dstImage->validateRegion(dstOrigin, copyRegion)) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
+
+hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hipArray* dstArray,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
+                                 amd::HostQueue* queue) {
+  amd::Image* dstImage;
+  amd::BufferRect srcRect;
+
+  hipError_t status = ihipMemcpyHtoAValidate(srcHost, dstArray, srcOrigin, dstOrigin, copyRegion,
+                                             srcRowPitch, srcSlicePitch, dstImage, srcRect);
+  if (status != hipSuccess) {
+    return status;
+  }
 
   command = new amd::WriteMemoryCommand(
       *queue, CL_COMMAND_WRITE_IMAGE, amd::Command::EventWaitList{}, *dstImage, dstOrigin,
@@ -1545,10 +1654,10 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi
   return hipSuccess;
 }
 
-hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, void* dstHost,
-                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
-                                 amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch,
-                                 amd::HostQueue* queue) {
+hipError_t ihipMemcpyAtoHValidate(hipArray* srcArray, void* dstHost, amd::Coord3D& srcOrigin,
+                                  amd::Coord3D& dstOrigin, amd::Coord3D& copyRegion,
+                                  size_t dstRowPitch, size_t dstSlicePitch, amd::Image*& srcImage,
+                                  amd::BufferRect& dstRect) {
   if (srcArray == nullptr || (dstHost == nullptr)) {
     return hipErrorInvalidValue;
   }
@@ -1557,13 +1666,12 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, voi
     return hipErrorInvalidValue;
   }
 
-  amd::BufferRect dstRect;
   if (!dstRect.create(static_cast<size_t*>(dstOrigin), static_cast<size_t*>(copyRegion),
                       dstRowPitch, dstSlicePitch)) {
     return hipErrorInvalidValue;
   }
 
-  amd::Image* srcImage = as_amd(srcMemObj)->asImage();
+  srcImage = as_amd(srcMemObj)->asImage();
   // HIP assumes the width is in bytes, but OCL assumes it's in pixels.
   const size_t elementSize = srcImage->getImageFormat().getElementSize();
   static_cast<size_t*>(srcOrigin)[0] /= elementSize;
@@ -1573,6 +1681,21 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, voi
       !srcImage->isRowSliceValid(dstRowPitch, dstSlicePitch, copyRegion[0], copyRegion[1])) {
     return hipErrorInvalidValue;
   }
+  return hipSuccess;
+}
+
+hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, void* dstHost,
+                                 amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
+                                 amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch,
+                                 amd::HostQueue* queue) {
+  amd::Image* srcImage;
+  amd::BufferRect dstRect;
+
+  hipError_t status = ihipMemcpyAtoHValidate(srcArray, dstHost, srcOrigin, dstOrigin, copyRegion,
+                                             dstRowPitch, dstSlicePitch, srcImage, dstRect);
+  if (status != hipSuccess) {
+    return status;
+  }
 
   command = new amd::ReadMemoryCommand(
       *queue, CL_COMMAND_READ_IMAGE, amd::Command::EventWaitList{}, *srcImage, srcOrigin,