SWDEV-375366 - SWDEV-375351 - Handle HtoH case for graph mem cpy impl.

Change-Id: I5a8c3c3c22db045f714b0443b8d70a8c6b4a8cea


[ROCm/clr commit: fdecc1b351]

projects/clr/hipamd/src/hip_graph_helper.hpp

@@ -7,6 +7,10 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes, hip
 hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes,
                              hipMemcpyKind kind, amd::HostQueue& queue, bool isAsync = false);
 
+void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, amd::HostQueue& queue);
+
+bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind);
+
 hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX,
                                      uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
                                      uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -1037,6 +1037,9 @@ class hipGraphMemcpyNode : public hipGraphNode {
   }
 
   hipError_t CreateCommand(amd::HostQueue* queue) {
+    if (IsHtoHMemcpy(pCopyParams_->dstPtr.ptr, pCopyParams_->srcPtr.ptr, pCopyParams_->kind)) {
+      return hipSuccess;
+    }
     hipError_t status = hipGraphNode::CreateCommand(queue);
     if (status != hipSuccess) {
       return status;
@@ -1048,6 +1051,16 @@ class hipGraphMemcpyNode : public hipGraphNode {
     return status;
   }
 
+  void EnqueueCommands(hipStream_t stream) override {
+    if (isEnabled_ && IsHtoHMemcpy(pCopyParams_->dstPtr.ptr, pCopyParams_->srcPtr.ptr, pCopyParams_->kind)) {
+      ihipHtoHMemcpy(pCopyParams_->dstPtr.ptr, pCopyParams_->srcPtr.ptr,
+                     pCopyParams_->extent.width * pCopyParams_->extent.height *
+                     pCopyParams_->extent.depth, *hip::getQueue(stream));
+      return;
+    }
+    hipGraphNode::EnqueueCommands(stream);
+  }
+
   void GetParams(hipMemcpy3DParms* params) {
     std::memcpy(params, pCopyParams_, sizeof(hipMemcpy3DParms));
   }
@@ -1170,6 +1183,9 @@ class hipGraphMemcpyNode1D : public hipGraphNode {
   }
 
   virtual hipError_t CreateCommand(amd::HostQueue* queue) {
+    if (IsHtoHMemcpy(dst_, src_, kind_)) {
+      return hipSuccess;
+    }
     hipError_t status = hipGraphNode::CreateCommand(queue);
     if (status != hipSuccess) {
       return status;
@@ -1182,10 +1198,18 @@ class hipGraphMemcpyNode1D : public hipGraphNode {
   }
 
   void EnqueueCommands(hipStream_t stream) {
-    if (commands_.empty()) return;
-    // commands_ should have just 1 item
-    assert(commands_.size() == 1 && "Invalid command size in hipGraphMemcpyNode1D");
+    bool isH2H = IsHtoHMemcpy(dst_, src_, kind_);
+    if (!isH2H) {
+      if (commands_.empty()) return;
+      // commands_ should have just 1 item
+      assert(commands_.size() == 1 && "Invalid command size in hipGraphMemcpyNode1D");
+    }
     if (isEnabled_) {
+      //HtoH
+      if (isH2H) {
+        ihipHtoHMemcpy(dst_, src_, count_, *hip::getQueue(stream));
+        return;
+      }
       amd::Command* command = commands_[0];
       amd::HostQueue* cmdQueue = command->queue();
       amd::HostQueue* queue = hip::getQueue(stream);

projects/clr/hipamd/src/hip_memory.cpp

@@ -321,7 +321,18 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
   memObj->getUserData().deviceId = hip::getCurrentDevice()->deviceId();
   return hipSuccess;
 }
-
+bool IsHtoHMemcpyValid(void* dst, const void* src, hipMemcpyKind kind) {
+  size_t sOffset = 0;
+  amd::Memory* srcMemory = getMemoryObject(src, sOffset);
+  size_t dOffset = 0;
+  amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
+  if (src && dst && srcMemory == nullptr && dstMemory == nullptr) {
+    if (kind != hipMemcpyHostToHost && kind != hipMemcpyDefault) {
+      return false;
+    }
+  }
+  return true;
+}
 hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
                                       hipMemcpyKind kind) {
   if (dst == nullptr || src == nullptr) {
@@ -337,6 +348,10 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
       (srcMemory && sizeBytes > (srcMemory->getSize() - sOffset))) {
     return hipErrorInvalidValue;
   }
+  //If src and dst ptr are null then kind must be either h2h or def.
+  if (!IsHtoHMemcpyValid(dst, src, kind)) {
+    return hipErrorInvalidValue;
+  }
   return hipSuccess;
 }
 
@@ -431,7 +446,22 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
   }
   return hipSuccess;
 }
-
+bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind) {
+  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 false;
+}
+void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, amd::HostQueue& queue) {
+  queue.finish();
+  memcpy(dst, src, sizeBytes);
+}
 // ================================================================================================
 hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
                       amd::HostQueue& queue, bool isAsync = false) {
@@ -451,14 +481,9 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   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)) {
-      queue.finish();
-      memcpy(dst, src, sizeBytes);
-      return hipSuccess;
-    } else {
-      return hipErrorInvalidValue;
-    }
+  if (srcMemory == nullptr && dstMemory == nullptr) {
+    ihipHtoHMemcpy(dst, src, sizeBytes, queue);
+    return hipSuccess;
   } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
     isAsync = false;
   } else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
@@ -2631,7 +2656,10 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) {
   if (p->kind < hipMemcpyHostToHost || p->kind > hipMemcpyDefault) {
     return hipErrorInvalidMemcpyDirection;
   }
-
+  //If src and dst ptr are null then kind must be either h2h or def.
+  if (!IsHtoHMemcpyValid(p->dstPtr.ptr, p->srcPtr.ptr, p->kind)) {
+    return hipErrorInvalidValue;
+  }
   return hipSuccess;
 }