From cf8eeabfe2ea09daee0126f39484497885b94a36 Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Mon, 7 Oct 2024 16:18:30 +0200 Subject: [PATCH] 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: db8527f6554bb2fa01804b559a90139bd35615c4] --- projects/clr/hipamd/src/hip_graph.cpp | 63 ++++++++++++++++--- projects/clr/hipamd/src/hip_graph_helper.hpp | 7 ++- .../clr/hipamd/src/hip_graph_internal.cpp | 6 +- .../clr/hipamd/src/hip_graph_internal.hpp | 8 ++- projects/clr/hipamd/src/hip_memory.cpp | 4 +- 5 files changed, 74 insertions(+), 14 deletions(-) diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index d79b5e860b..ae2da97b81 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/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(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(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(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); diff --git a/projects/clr/hipamd/src/hip_graph_helper.hpp b/projects/clr/hipamd/src/hip_graph_helper.hpp index 7fdb378041..3abbbd0f26 100644 --- a/projects/clr/hipamd/src/hip_graph_helper.hpp +++ b/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 \ No newline at end of file +} // namespace hip diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index 485f9ddc7c..f6c7221724 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/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(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; } // ================================================================================================ diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 0e1cbe415b..32b081fc7d 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/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_; } diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 4aa766aec1..66f4675fcb 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/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;