From c4eda858354a4092ca9280d9e7bb704e0bca3c87 Mon Sep 17 00:00:00 2001 From: anusha GodavarthySurya Date: Tue, 28 Sep 2021 12:51:17 -0700 Subject: [PATCH] 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 | 9 ++ hipamd/src/hip_graph.cpp | 93 +++++++++++ hipamd/src/hip_graph_helper.hpp | 53 +++++++ hipamd/src/hip_graph_internal.cpp | 227 +++++++++++++++++++++++++++ hipamd/src/hip_graph_internal.hpp | 109 +++++++++---- hipamd/src/hip_hcc.def.in | 9 ++ hipamd/src/hip_hcc.map.in | 9 ++ hipamd/src/hip_memory.cpp | 247 ++++++++++++++++++++++-------- 8 files changed, 665 insertions(+), 91 deletions(-) diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index a1d58d48de..17c0f4e2bd 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -323,3 +323,12 @@ hipGraphNodeFindInClone hipGraphAddChildGraphNode hipGraphChildGraphNodeGetGraph hipGraphExecChildGraphNodeSetParams +hipGraphAddMemcpyNodeFromSymbol +hipGraphMemcpyNodeSetParamsFromSymbol +hipGraphExecMemcpyNodeSetParamsFromSymbol +hipGraphAddMemcpyNodeToSymbol +hipGraphMemcpyNodeSetParamsToSymbol +hipGraphExecMemcpyNodeSetParamsToSymbol +hipGraphExecMemcpyNodeSetParams +hipGraphMemcpyNodeSetParams1D +hipGraphExecMemcpyNodeSetParams1D diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 5864df5c27..e047ec347a 100644 --- a/hipamd/src/hip_graph.cpp +++ b/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(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(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(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(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(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(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(node)->SetCommandParams( + symbol, src, count, offset, kind)); +} diff --git a/hipamd/src/hip_graph_helper.hpp b/hipamd/src/hip_graph_helper.hpp index b28c14e811..fe8316d2e4 100644 --- a/hipamd/src/hip_graph_helper.hpp +++ b/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& commands, void* dst, in hipError_t ihipMemset3DCommand(std::vector& 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); diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index b414478e85..606b2beb44 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/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(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(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(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(command)->validateMemory()) { + delete command; + return hipErrorInvalidValue; + } + } else { + amd::CopyMemoryCommand* command = reinterpret_cast(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(&pCopy)->srcHost = pCopy.srcDevice; + } + } + hipMemoryType dstMemoryType = pCopy.dstMemoryType; + if (dstMemoryType == hipMemoryTypeUnified) { + dstMemoryType = + amd::MemObjMap::FindMemObj(pCopy.dstDevice) ? hipMemoryTypeDevice : hipMemoryTypeHost; + if (srcMemoryType == hipMemoryTypeHost) { + const_cast(&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(&pCopy)->srcDevice = const_cast(pCopy.srcHost); + } + } + if (dstMemoryType == hipMemoryTypeHost) { + amd::Memory* mem = amd::MemObjMap::FindMemObj(pCopy.dstHost); + dstMemoryType = mem ? hipMemoryTypeDevice : hipMemoryTypeHost; + if (dstMemoryType == hipMemoryTypeDevice) { + const_cast(&pCopy)->dstDevice = const_cast(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(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(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(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(commands_[0]); + command->setParams(*dstImage, dstOrigin, copyRegion, + static_cast(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(commands_[0]); + command->setParams(*srcImage, srcOrigin, copyRegion, + static_cast(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(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(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(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", diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index cfe194aaf1..90f61fba3c 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/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(*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 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 hipGraphMemcpyNodeFromSymbol : public hipGraphNode { static_cast(*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 hipGraphMemcpyNodeFromSymbol : public hipGraphNode { offset_ = offset; kind_ = kind; } -}; -template 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(*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 { diff --git a/hipamd/src/hip_hcc.def.in b/hipamd/src/hip_hcc.def.in index bcb8c1413b..e6ed1a6a9b 100644 --- a/hipamd/src/hip_hcc.def.in +++ b/hipamd/src/hip_hcc.def.in @@ -321,3 +321,12 @@ hipGraphNodeFindInClone hipGraphAddChildGraphNode hipGraphChildGraphNodeGetGraph hipGraphExecChildGraphNodeSetParams +hipGraphAddMemcpyNodeFromSymbol +hipGraphMemcpyNodeSetParamsFromSymbol +hipGraphExecMemcpyNodeSetParamsFromSymbol +hipGraphAddMemcpyNodeToSymbol +hipGraphMemcpyNodeSetParamsToSymbol +hipGraphExecMemcpyNodeSetParamsToSymbol +hipGraphExecMemcpyNodeSetParams +hipGraphMemcpyNodeSetParams1D +hipGraphExecMemcpyNodeSetParams1D diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index 9042024e99..1d8fd2d2b5 100644 --- a/hipamd/src/hip_hcc.map.in +++ b/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; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index e44cc42ed0..33dbea764e 100644 --- a/hipamd/src/hip_memory.cpp +++ b/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(srcOrigin)[0] /= elementSize; static_cast(copyRegion)[0] /= elementSize; - amd::BufferRect srcRect; + if (!srcRect.create(static_cast(srcOrigin), static_cast(copyRegion), srcImage->getRowPitch(), srcImage->getSlicePitch())) { return hipErrorInvalidValue; } - amd::BufferRect dstRect; if (!dstRect.create(static_cast(dstOrigin), static_cast(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(dstOrigin)[0] /= elementSize; static_cast(copyRegion)[0] /= elementSize; - amd::BufferRect srcRect; + if (!srcRect.create(static_cast(srcOrigin), static_cast(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(dstOrigin), static_cast(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(srcOrigin), static_cast(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(dstOrigin), static_cast(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(srcOrigin), static_cast(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(dstOrigin), static_cast(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(srcOrigin), static_cast(copyRegion), srcRowPitch, srcSlicePitch)) { return hipErrorInvalidValue; } - amd::BufferRect dstRect; if (!dstRect.create(static_cast(dstOrigin), static_cast(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(srcOrigin), static_cast(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(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(dstOrigin), static_cast(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(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,