SWDEV-489619 - Fix memcpy tests with capture stream enabled

- Added missing validation as graph node should not be created
 if parameters are invalid
 - Fix conversion of input params to graphNode params

Change-Id: I37ab04942b5fb2eb07386850cb7dbbf26f9ca967


[ROCm/clr commit: db8527f655]

projects/clr/hipamd/src/hip_graph.cpp

@@ -415,6 +415,12 @@ hipError_t capturehipMemcpy3DAsync(hipStream_t& stream, const hipMemcpy3DParms*&
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
+
+  // Skip zero-sized copies
+  if (p->extent.width == 0 || p->extent.height == 0 || p->extent.depth == 0) {
+    return hipSuccess;
+  }
+
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   hip::GraphNode* pGraphNode;
   hipError_t status =
@@ -435,6 +441,16 @@ hipError_t capturehipMemcpy2DAsync(hipStream_t& stream, void*& dst, size_t& dpit
   if (dst == nullptr || src == nullptr) {
     return hipErrorInvalidValue;
   }
+
+  // Skip zero-sized copies
+  if (width == 0 || height == 0) {
+    return hipSuccess;
+  }
+
+  if ((width > dpitch) || (width > spitch)) {
+    return hipErrorInvalidPitchValue;
+  }
+
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -470,9 +486,15 @@ hipError_t capturehipMemcpy2DFromArrayAsync(hipStream_t& stream, void*& dst, siz
                                             hipMemcpyKind& kind) {
   ClPrint(amd::LOG_INFO, amd::LOG_API,
           "[hipGraph] Current capture node Memcpy2DFromArray on stream : %p", stream);
-  if (src == nullptr || dst == nullptr) {
-    return hipErrorInvalidValue;
+
+  // Skip zero-sized copies
+  if (width == 0 || height == 0) {
+    return hipSuccess;
   }
+
+  HIP_RETURN_ONFAIL(hipMemcpy2DValidateArray(src, wOffsetSrc, hOffsetSrc, width, height));
+  HIP_RETURN_ONFAIL(hipMemcpy2DValidateBuffer(dst, dpitch, width));
+
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
@@ -541,6 +563,21 @@ hipError_t capturehipMemcpyParam2DAsync(hipStream_t& stream, const hip_Memcpy2D*
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
+
+  if ((pCopy->srcDevice == nullptr && pCopy->srcMemoryType == hipMemoryTypeDevice) ||
+      (pCopy->dstDevice == nullptr && pCopy->dstMemoryType == hipMemoryTypeDevice) ||
+      (pCopy->srcHost == nullptr && pCopy->srcMemoryType == hipMemoryTypeHost) ||
+      (pCopy->dstHost == nullptr && pCopy->dstMemoryType == hipMemoryTypeHost) ||
+      (pCopy->srcArray == nullptr && pCopy->srcMemoryType == hipMemoryTypeArray) ||
+      (pCopy->dstArray == nullptr && pCopy->dstMemoryType == hipMemoryTypeArray)) {
+    return hipErrorInvalidValue;
+  }
+
+  /// Skip zero-sized copies
+  if (pCopy->WidthInBytes == 0 || pCopy->Height == 0) {
+    return hipSuccess;
+  }
+
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   hip::GraphNode* pGraphNode;
   hipMemcpy3DParms p = {};
@@ -563,14 +600,25 @@ hipError_t capturehipMemcpyParam2DAsync(hipStream_t& stream, const hip_Memcpy2D*
   if (pCopy->dstHost != nullptr) {
     p.dstPtr.ptr = const_cast<void*>(pCopy->dstHost);
   }
+
+
+  // If array is participating in the copy, the extent is defined in terms of that array's elements.
+  // If no array is participating in the copy then the extents are defined in elements of unsigned
+  // char.
   p.extent = {pCopy->WidthInBytes, pCopy->Height, 1};
-  if (pCopy->srcMemoryType == hipMemoryTypeHost && pCopy->dstMemoryType == hipMemoryTypeDevice) {
+  if (pCopy->srcArray != nullptr) {
+    p.extent.width /= getElementSize(pCopy->srcArray);
+  } else if (pCopy->dstArray != nullptr) {
+    p.extent.width /= getElementSize(pCopy->dstArray);
+  }
+
+  if (pCopy->srcMemoryType == hipMemoryTypeHost && pCopy->dstMemoryType == hipMemoryTypeHost) {
+    p.kind = hipMemcpyHostToHost;
+  } else if (pCopy->srcMemoryType == hipMemoryTypeHost) {
     p.kind = hipMemcpyHostToDevice;
-  } else if (pCopy->srcMemoryType == hipMemoryTypeDevice &&
-             pCopy->dstMemoryType == hipMemoryTypeHost) {
+  } else if (pCopy->dstMemoryType == hipMemoryTypeHost) {
     p.kind = hipMemcpyDeviceToHost;
-  } else if (pCopy->srcMemoryType == hipMemoryTypeDevice &&
-             pCopy->dstMemoryType == hipMemoryTypeDevice) {
+  } else {
     p.kind = hipMemcpyDeviceToDevice;
   }
   hipError_t status =
@@ -601,6 +649,7 @@ hipError_t capturehipMemcpyAtoHAsync(hipStream_t& stream, void*& dstHost, hipArr
   p.srcPos = {srcOffset, 0, 0};
   p.dstPtr.ptr = dstHost;
   p.extent = {ByteCount / hip::getElementSize(p.srcArray), 1, 1};
+  p.kind = hipMemcpyDeviceToHost;
   hipError_t status =
       ihipGraphAddMemcpyNode(&pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(),
                              s->GetLastCapturedNodes().size(), &p);

projects/clr/hipamd/src/hip_graph_helper.hpp

@@ -27,6 +27,11 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p);
 
 hipError_t ihipDrvMemcpy3D_validate(const HIP_MEMCPY3D* pCopy);
 
+hipError_t hipMemcpy2DValidateArray(hipArray_const_t arr, size_t wOffset, size_t hOffset,
+                                    size_t width, size_t height);
+
+hipError_t hipMemcpy2DValidateBuffer(const void* buf, size_t pitch, size_t width);
+
 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,
@@ -123,4 +128,4 @@ hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord
 hipError_t ihipGraphMemsetParams_validate(const hipMemsetParams* pNodeParams);
 
 hip::MemcpyType ihipGetMemcpyType(const void* src, void* dst, hipMemcpyKind kind);
-}  // namespace hip
+}  // namespace hip

projects/clr/hipamd/src/hip_graph_internal.cpp

@@ -403,7 +403,7 @@ hipError_t GraphExec::AllocKernelArgForGraphNode() {
       }
     }
     if (node->GraphCaptureEnabled()) {
-      node->CaptureAndFormPacket(capture_stream_, GetKernelArgManager());
+      status = node->CaptureAndFormPacket(capture_stream_, GetKernelArgManager());
     } else if (node->GetType() == hipGraphNodeTypeGraph) {
       auto childNode = reinterpret_cast<hip::ChildGraphNode*>(node);
       if (childNode->GetChildGraph()->max_streams_ == 1) {
@@ -444,9 +444,9 @@ hipError_t GraphExec::CaptureAQLPackets() {
 hipError_t GraphExec::UpdateAQLPacket(hip::GraphNode* node) {
   hipError_t status = hipSuccess;
   if (max_streams_ == 1) {
-    node->CaptureAndFormPacket(capture_stream_, kernArgManager_);
+    status = node->CaptureAndFormPacket(capture_stream_, kernArgManager_);
   }
-  return hipSuccess;
+  return status;
 }
 
 // ================================================================================================

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -269,8 +269,12 @@ struct GraphNode : public hipGraphNodeDOTAttribute {
   size_t GetKerArgSize() const { return alignedKernArgSize_; }
   size_t GetKernargSegmentByteSize() const { return kernargSegmentByteSize_; }
   size_t GetKernargSegmentAlignment() const { return kernargSegmentAlignment_; }
-  void CaptureAndFormPacket(hip::Stream* capture_stream, GraphKernelArgManager* kernArgMgr) {
+  hipError_t CaptureAndFormPacket(hip::Stream* capture_stream, GraphKernelArgManager* kernArgMgr) {
     hipError_t status = CreateCommand(capture_stream);
+    if (status != hipSuccess) {
+      return status;
+    }
+
     gpuPackets_.clear();
     for (auto& command : commands_) {
       command->setPktCapturingState(true, &gpuPackets_, kernArgMgr, &capturedKernelName_);
@@ -281,6 +285,8 @@ struct GraphNode : public hipGraphNodeDOTAttribute {
     }
     // Commands are captured and released. Clear them from the object.
     commands_.clear();
+
+    return status;
   }
   hip::Stream* GetQueue() const { return stream_; }
 

projects/clr/hipamd/src/hip_memory.cpp

@@ -2667,8 +2667,8 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) {
   }
 
   if (p->dstArray == nullptr && p->srcArray == nullptr) {
-    if ((p->extent.width + p->dstPos.x > p->dstPtr.pitch) ||
-        (p->extent.width + p->srcPos.x > p->srcPtr.pitch)) {
+    if ((p->dstPtr.pitch != 0 && (p->extent.width + p->dstPos.x > p->dstPtr.pitch)) ||
+        (p->srcPtr.pitch != 0 && (p->extent.width + p->srcPos.x > p->srcPtr.pitch))) {
       return hipErrorInvalidValue;
     }
     auto totalExtentBytes = p->extent.width * p->extent.height * p->extent.depth;