From ffcff763d2ab6322e72df585c073670cdfe8c42b Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Mon, 2 Jan 2023 13:08:13 +0000 Subject: [PATCH] SWDEV-375366 - SWDEV-375351 - Handle HtoH case for graph mem cpy impl. Change-Id: I5a8c3c3c22db045f714b0443b8d70a8c6b4a8cea [ROCm/clr commit: fdecc1b3517dc4bcf3881ee8f5045120511a79b8] --- projects/clr/hipamd/src/hip_graph_helper.hpp | 4 ++ .../clr/hipamd/src/hip_graph_internal.hpp | 30 +++++++++-- projects/clr/hipamd/src/hip_memory.cpp | 50 +++++++++++++++---- 3 files changed, 70 insertions(+), 14 deletions(-) diff --git a/projects/clr/hipamd/src/hip_graph_helper.hpp b/projects/clr/hipamd/src/hip_graph_helper.hpp index 4af57c6900..69780338b2 100644 --- a/projects/clr/hipamd/src/hip_graph_helper.hpp +++ b/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, diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index ec016d1ff1..65ccfdc341 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/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); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index e333dce84e..37a53e269e 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/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; }