diff --git a/hipamd/include/hip/amd_detail/hip_prof_str.h b/hipamd/include/hip/amd_detail/hip_prof_str.h index 4753028e00..c67a8886e0 100644 --- a/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -377,7 +377,8 @@ enum hip_api_id_t { HIP_API_ID_hipArrayGetInfo = 362, HIP_API_ID_hipStreamGetDevice = 363, HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray = 364, - HIP_API_ID_LAST = 364, + HIP_API_ID_hipDrvGraphAddMemcpyNode = 365, + HIP_API_ID_LAST = 365, HIP_API_ID_hipBindTexture = HIP_API_ID_NONE, HIP_API_ID_hipBindTexture2D = HIP_API_ID_NONE, @@ -534,6 +535,7 @@ static inline const char* hip_api_name(const uint32_t id) { case HIP_API_ID_hipGraphAddMemAllocNode: return "hipGraphAddMemAllocNode"; case HIP_API_ID_hipGraphAddMemFreeNode: return "hipGraphAddMemFreeNode"; case HIP_API_ID_hipGraphAddMemcpyNode: return "hipGraphAddMemcpyNode"; + case HIP_API_ID_hipDrvGraphAddMemcpyNode: return "hipDrvGraphAddMemcpyNode"; case HIP_API_ID_hipGraphAddMemcpyNode1D: return "hipGraphAddMemcpyNode1D"; case HIP_API_ID_hipGraphAddMemcpyNodeFromSymbol: return "hipGraphAddMemcpyNodeFromSymbol"; case HIP_API_ID_hipGraphAddMemcpyNodeToSymbol: return "hipGraphAddMemcpyNodeToSymbol"; @@ -1148,6 +1150,7 @@ static inline uint32_t hipApiIdByName(const char* name) { if (strcmp("hipUserObjectRetain", name) == 0) return HIP_API_ID_hipUserObjectRetain; if (strcmp("hipWaitExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipWaitExternalSemaphoresAsync; if (strcmp("hipExternalMemoryGetMappedMipmappedArray", name) == 0) return HIP_API_ID_hipExternalMemoryGetMappedMipmappedArray; + if (strcmp("hipDrvGraphAddMemcpyNode", name) == 0) return HIP_API_ID_hipDrvGraphAddMemcpyNode; return HIP_API_ID_NONE; } @@ -3265,6 +3268,18 @@ typedef struct hip_api_data_s { const hipExternalMemoryMipmappedArrayDesc* mipmapDesc; hipExternalMemoryMipmappedArrayDesc mipmapDesc__val; } hipExternalMemoryGetMappedMipmappedArray; + + struct { + hipGraphNode_t* phGraphNode; + hipGraphNode_t phGraphNode__val; + hipGraph_t hGraph; + const hipGraphNode_t* dependencies; + hipGraphNode_t dependencies__val; + size_t numDependencies; + const HIP_MEMCPY3D* copyParams; + HIP_MEMCPY3D copyParams__val; + hipCtx_t ctx; + } hipDrvGraphAddMemcpyNode; } args; uint64_t *phase_data; } hip_api_data_t; @@ -5438,6 +5453,15 @@ typedef struct hip_api_data_s { cb_data.args.hipWaitExternalSemaphoresAsync.numExtSems = (unsigned int)numExtSems; \ cb_data.args.hipWaitExternalSemaphoresAsync.stream = (hipStream_t)stream; \ }; +// hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const hipMemcpy3DParms*', 'copyParams'), ('hipCtx_t', 'ctx')] +#define INIT_hipDrvGraphAddMemcpyNode_CB_ARGS_DATA(cb_data) { \ + cb_data.args.hipDrvGraphAddMemcpyNode.phGraphNode = (hipGraphNode_t*)phGraphNode; \ + cb_data.args.hipDrvGraphAddMemcpyNode.hGraph = (hipGraph_t)hGraph; \ + cb_data.args.hipDrvGraphAddMemcpyNode.dependencies = (const hipGraphNode_t*)dependencies; \ + cb_data.args.hipDrvGraphAddMemcpyNode.numDependencies = (size_t)numDependencies; \ + cb_data.args.hipDrvGraphAddMemcpyNode.copyParams = (const HIP_MEMCPY3D*)copyParams; \ + cb_data.args.hipDrvGraphAddMemcpyNode.ctx = (hipCtx_t)ctx; \ +}; #define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data) // Macros for non-public API primitives @@ -6917,6 +6941,12 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) { if (data->args.hipWaitExternalSemaphoresAsync.extSemArray) data->args.hipWaitExternalSemaphoresAsync.extSemArray__val = *(data->args.hipWaitExternalSemaphoresAsync.extSemArray); if (data->args.hipWaitExternalSemaphoresAsync.paramsArray) data->args.hipWaitExternalSemaphoresAsync.paramsArray__val = *(data->args.hipWaitExternalSemaphoresAsync.paramsArray); break; +// hipDrvGraphAddMemcpyNode[('hipGraphNode_t*', 'phGraphNode'), ('hipGraph_t', 'hGraph'), ('const hipGraphNode_t*', 'dependencies'), ('size_t', 'numDependencies'), ('const HIP_MEMCPY3D*', 'copyParams'), ('hipCtx_t', 'ctx')] + case HIP_API_ID_hipDrvGraphAddMemcpyNode: + if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode) data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val = *(data->args.hipDrvGraphAddMemcpyNode.phGraphNode); + if (data->args.hipDrvGraphAddMemcpyNode.dependencies) data->args.hipDrvGraphAddMemcpyNode.dependencies__val = *(data->args.hipDrvGraphAddMemcpyNode.dependencies); + if (data->args.hipDrvGraphAddMemcpyNode.copyParams) data->args.hipDrvGraphAddMemcpyNode.copyParams__val = *(data->args.hipDrvGraphAddMemcpyNode.copyParams); + break; default: break; }; } @@ -9773,6 +9803,18 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipWaitExternalSemaphoresAsync.stream); oss << ")"; break; + case HIP_API_ID_hipDrvGraphAddMemcpyNode: + oss << "hipDrvGraphAddMemcpyNode("; + if (data->args.hipDrvGraphAddMemcpyNode.phGraphNode == NULL) oss << "phGraphNode=NULL"; + else { oss << "phGraphNode="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.phGraphNode__val); } + oss << ", hGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.hGraph); + if (data->args.hipDrvGraphAddMemcpyNode.dependencies == NULL) oss << ", dependencies=NULL"; + else { oss << ", dependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.dependencies__val); } + oss << ", numDependencies="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.numDependencies); + if (data->args.hipDrvGraphAddMemcpyNode.copyParams == NULL) oss << ", copyParams=NULL"; + else { oss << ", copyParams="; roctracer::hip_support::detail::operator<<(oss, data->args.hipDrvGraphAddMemcpyNode.copyParams__val); } + oss << ")"; + break; default: oss << "unknown"; }; return strdup(oss.str().c_str()); diff --git a/hipamd/src/amdhip.def b/hipamd/src/amdhip.def index f3f1a64599..a3bbd32766 100644 --- a/hipamd/src/amdhip.def +++ b/hipamd/src/amdhip.def @@ -445,3 +445,4 @@ hipGraphAddMemAllocNode hipGraphMemAllocNodeGetParams hipGraphAddMemFreeNode hipGraphMemFreeNodeGetParams +hipDrvGraphAddMemcpyNode diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index d87b19e40b..f135c3bf9b 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -111,6 +111,23 @@ hipError_t ihipGraphAddMemcpyNode(hip::GraphNode** pGraphNode, hip::Graph* graph return status; } +hipError_t ihipDrvGraphAddMemcpyNode(hip::GraphNode** pGraphNode, hip::Graph* graph, + hip::GraphNode* const* pDependencies, size_t numDependencies, + const HIP_MEMCPY3D* pCopyParams, hipCtx_t ctx, + bool capture = true) { + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr) || pCopyParams == nullptr) { + return hipErrorInvalidValue; + } + hipError_t status = ihipDrvMemcpy3DParamValidate(pCopyParams); + if (status != hipSuccess) { + return status; + } + *pGraphNode = new hip::GraphDrvMemcpyNode(pCopyParams); + status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies, capture); + return status; +} + hipError_t ihipGraphAddMemcpyNode1D(hip::GraphNode** pGraphNode, hip::Graph* graph, hip::GraphNode* const* pDependencies, size_t numDependencies, void* dst, const void* src, size_t count, hipMemcpyKind kind, @@ -1036,6 +1053,24 @@ hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, HIP_RETURN(status); } +hipError_t hipDrvGraphAddMemcpyNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph, + const hipGraphNode_t* dependencies, size_t numDependencies, + const HIP_MEMCPY3D* copyParams, hipCtx_t ctx) { + HIP_INIT_API(hipDrvGraphAddMemcpyNode, phGraphNode, hGraph, dependencies, numDependencies, + copyParams, ctx); + if (phGraphNode == nullptr || hGraph == nullptr || + (numDependencies > 0 && dependencies == nullptr) || ctx == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + hip::GraphNode* node; + hipError_t status = ihipDrvGraphAddMemcpyNode( + &node, reinterpret_cast(hGraph), + reinterpret_cast(dependencies), + numDependencies, copyParams, ctx, false); + *phGraphNode = reinterpret_cast(node); + HIP_RETURN(status); +} + hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, const hipGraphNode_t* pDependencies, size_t numDependencies, void* dst, const void* src, size_t count, hipMemcpyKind kind) { diff --git a/hipamd/src/hip_graph_helper.hpp b/hipamd/src/hip_graph_helper.hpp index 63f3ae139f..1b8bf34516 100644 --- a/hipamd/src/hip_graph_helper.hpp +++ b/hipamd/src/hip_graph_helper.hpp @@ -24,6 +24,10 @@ THE SOFTWARE. hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p); +hipError_t ihipDrvMemcpy3D_validate(const HIP_MEMCPY3D* pCopy); + +hipError_t ihipDrvMemcpy3DParamValidate(const HIP_MEMCPY3D* pCopy); + 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, @@ -31,7 +35,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, hip::Stream& stream); -bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind); +bool IsHtoHMemcpy(void* dst, const void* src); hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, @@ -56,6 +60,9 @@ hipError_t ihipLaunchKernelCommand(amd::Command*& command, hipFunction_t f, hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p, hip::Stream* stream); +hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy, + hip::Stream* stream); + hipError_t ihipMemsetCommand(std::vector& commands, void* dst, int64_t value, size_t valueSize, size_t sizeBytes, hip::Stream* stream); diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index 077267e917..884b4cdf6b 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -102,158 +102,16 @@ hipError_t GraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size_t } hipError_t GraphMemcpyNode::ValidateParams(const hipMemcpy3DParms* pNodeParams) { - hipError_t status = ihipMemcpy3D_validate(pNodeParams); + hipError_t status; + status = ihipMemcpy3D_validate(pNodeParams); if (status != hipSuccess) { return status; } - size_t offset = 0; + 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) { - amd::Memory* memObj = getMemoryObject(pCopy.srcDevice, offset); - srcMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & - memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice; - 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; - const_cast(&pCopy)->srcXInBytes += offset; - } - } - offset = 0; - hipMemoryType dstMemoryType = pCopy.dstMemoryType; - if (dstMemoryType == hipMemoryTypeUnified) { - amd::Memory* memObj = getMemoryObject(pCopy.dstDevice, offset); - dstMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & - memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice; - if (dstMemoryType == hipMemoryTypeHost) { - const_cast(&pCopy)->dstHost = pCopy.dstDevice; - const_cast(&pCopy)->dstXInBytes += offset; - } - } - offset = 0; - // 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) { - srcMemoryType = getMemoryObject(pCopy.srcHost, offset) ? hipMemoryTypeDevice : - hipMemoryTypeHost; - - if (srcMemoryType == hipMemoryTypeDevice) { - const_cast(&pCopy)->srcDevice = const_cast(pCopy.srcHost); - } - } - offset = 0; - if (dstMemoryType == hipMemoryTypeHost) { - dstMemoryType = getMemoryObject(pCopy.dstHost, offset) ? 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; - } - } 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; - } - } 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; - } - } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) { - amd::Image* dstImage; - size_t start = 0; - status = - ihipMemcpyHtoAValidate(pCopy.srcHost, pCopy.dstArray, srcOrigin, dstOrigin, copyRegion, - pCopy.srcPitch, pCopy.srcPitch * pCopy.srcHeight, dstImage, start); - if (status != hipSuccess) { - return status; - } - } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) { - // Image to Host. - amd::Image* srcImage; - size_t start = 0; - - status = - ihipMemcpyAtoHValidate(pCopy.srcArray, pCopy.dstHost, srcOrigin, dstOrigin, copyRegion, - pCopy.dstPitch, pCopy.dstPitch * pCopy.dstHeight, srcImage, start); - if (status != hipSuccess) { - return status; - } - } 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; - } - } 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; - } - } 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; - } - } else { - return hipErrorInvalidValue; + status = ihipDrvMemcpy3D_validate(&pCopy); + if (status != hipSuccess) { + return status; } return hipSuccess; } diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 86f3544344..27892301c4 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -1173,7 +1173,8 @@ class GraphMemcpyNode : public GraphNode { } virtual hipError_t CreateCommand(hip::Stream* stream) { - if (IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.kind)) { + if ((copyParams_.kind == hipMemcpyHostToHost || copyParams_.kind == hipMemcpyDefault) + && IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr)) { return hipSuccess; } hipError_t status = GraphNode::CreateCommand(stream); @@ -1188,7 +1189,8 @@ class GraphMemcpyNode : public GraphNode { } virtual void EnqueueCommands(hipStream_t stream) override { - if (isEnabled_ && IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.kind)) { + if ( (copyParams_.kind == hipMemcpyHostToHost || copyParams_.kind == hipMemcpyDefault) && + isEnabled_ && IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr)) { ihipHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.extent.width * copyParams_.extent.height * copyParams_.extent.depth, *hip::getStream(stream)); @@ -1326,7 +1328,7 @@ class GraphMemcpyNode1D : public GraphMemcpyNode { } virtual hipError_t CreateCommand(hip::Stream* stream) { - if (IsHtoHMemcpy(dst_, src_, kind_)) { + if ((kind_ == hipMemcpyHostToHost || kind_ == hipMemcpyDefault) && IsHtoHMemcpy(dst_, src_)) { return hipSuccess; } hipError_t status = GraphNode::CreateCommand(stream); @@ -1341,7 +1343,10 @@ class GraphMemcpyNode1D : public GraphMemcpyNode { } virtual void EnqueueCommands(hipStream_t stream) { - bool isH2H = IsHtoHMemcpy(dst_, src_, kind_); + bool isH2H = false; + if ((kind_ == hipMemcpyHostToHost || kind_ == hipMemcpyDefault) && IsHtoHMemcpy(dst_, src_)) { + isH2H = true; + } if (!isH2H) { if (commands_.empty()) return; // commands_ should have just 1 item @@ -2233,4 +2238,80 @@ class GraphMemFreeNode : public GraphNode { *params = device_ptr_; } }; + +class GraphDrvMemcpyNode : public GraphNode { + HIP_MEMCPY3D copyParams_; + + public: + GraphDrvMemcpyNode(const HIP_MEMCPY3D* pCopyParams) + : GraphNode(hipGraphNodeTypeMemcpy, "solid", "trapezium", "MEMCPY") { + copyParams_ = *pCopyParams; + } + ~GraphDrvMemcpyNode() {} + + GraphDrvMemcpyNode(const GraphDrvMemcpyNode& rhs) : GraphNode(rhs) { + copyParams_ = rhs.copyParams_; + } + + GraphNode* clone() const { + return new GraphDrvMemcpyNode(static_cast(*this)); + } + + hipError_t CreateCommand(hip::Stream* stream) { + if(copyParams_.srcMemoryType == hipMemoryTypeHost && + copyParams_.dstMemoryType == hipMemoryTypeHost && + IsHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost)) { + return hipSuccess; + } + hipError_t status = GraphNode::CreateCommand(stream); + if (status != hipSuccess) { + return status; + } + commands_.reserve(1); + amd::Command* command; + status = ihipGetMemcpyParam3DCommand(command, ©Params_, stream); + commands_.emplace_back(command); + return status; + } + + void EnqueueCommands(hipStream_t stream) override { + bool isHtoH = false; + if(copyParams_.srcMemoryType == hipMemoryTypeHost && + copyParams_.dstMemoryType == hipMemoryTypeHost && + IsHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost)) { + isHtoH = true; + } + if (isEnabled_ && isHtoH) { + ihipHtoHMemcpy(copyParams_.dstHost, copyParams_.srcHost, + copyParams_.WidthInBytes * copyParams_.Height * + copyParams_.Depth, *hip::getStream(stream)); + return; + } + GraphNode::EnqueueCommands(stream); + } + + void GetParams(HIP_MEMCPY3D* params) { + std::memcpy(params, ©Params_, sizeof(HIP_MEMCPY3D)); + } + hipError_t SetParams(const HIP_MEMCPY3D* params) { + hipError_t status = ValidateParams(params); + if (status != hipSuccess) { + return status; + } + std::memcpy(©Params_, params, sizeof(HIP_MEMCPY3D)); + return hipSuccess; + } + hipError_t SetParams(GraphNode* node) { + const GraphDrvMemcpyNode* memcpyNode = static_cast(node); + return SetParams(&memcpyNode->copyParams_); + } + // ToDo: use this when commands are cloned and command params are to be updated + hipError_t ValidateParams(const HIP_MEMCPY3D* pNodeParams) { + hipError_t status = ihipDrvMemcpy3D_validate(pNodeParams); + if (status != hipSuccess) { + return status; + } + return hipSuccess; + } +}; } // namespace hip diff --git a/hipamd/src/hip_hcc.map.in b/hipamd/src/hip_hcc.map.in index a0d22286af..3e92924d91 100644 --- a/hipamd/src/hip_hcc.map.in +++ b/hipamd/src/hip_hcc.map.in @@ -523,6 +523,7 @@ global: hipArrayGetInfo; hipArrayGetDescriptor; hipArray3DGetDescriptor; + hipDrvGraphAddMemcpyNode; local: *; } hip_5.5; diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 70add33a5f..785a3a8fff 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -470,17 +470,13 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, } return hipSuccess; } - -// ================================================================================================ -bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind) { +bool IsHtoHMemcpy(void* dst, const void* src) { 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 true; } return false; } @@ -2123,12 +2119,12 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray_t srcArray, vo return hipSuccess; } -hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy, - hip::Stream* stream) { +void ihipCopyMemParamSet(const HIP_MEMCPY3D* pCopy, hipMemoryType& srcMemType, + hipMemoryType& dstMemType) { size_t offset = 0; - // 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. + // 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) { amd::Memory* memObj = getMemoryObject(pCopy->srcDevice, offset); @@ -2138,11 +2134,14 @@ hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3 } else { srcMemoryType = hipMemoryTypeHost; } + if (srcMemoryType == hipMemoryTypeHost) { - // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we detect system - // memory. + // {src/dst}Host may be unitialized. Copy over {src/dst}Device into it if we + // detect system memory. const_cast(pCopy)->srcHost = pCopy->srcDevice; const_cast(pCopy)->srcXInBytes += offset; + // We don't need detect memory type again for hipMemoryTypeUnified + const_cast(pCopy)->srcMemoryType = srcMemoryType; } } offset = 0; @@ -2159,12 +2158,13 @@ hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3 if (dstMemoryType == hipMemoryTypeHost) { const_cast(pCopy)->dstHost = pCopy->dstDevice; const_cast(pCopy)->dstXInBytes += offset; + // We don't need detect memory type again for hipMemoryTypeUnified + const_cast(pCopy)->dstMemoryType = dstMemoryType; } } - + offset = 0; // 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. - offset = 0; if (srcMemoryType == hipMemoryTypeHost) { srcMemoryType = getMemoryObject(pCopy->srcHost, offset) ? hipMemoryTypeDevice : hipMemoryTypeHost; @@ -2179,9 +2179,18 @@ hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3 hipMemoryTypeHost; if (dstMemoryType == hipMemoryTypeDevice) { - const_cast(pCopy)->dstDevice = const_cast(pCopy->dstHost); + const_cast(pCopy)->dstDevice = const_cast(pCopy->dstDevice); } } + srcMemType = srcMemoryType; + dstMemType = dstMemoryType; +} + +hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy, + hip::Stream* stream) { + hipMemoryType srcMemoryType; + hipMemoryType dstMemoryType; + ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType); amd::Coord3D srcOrigin = {pCopy->srcXInBytes, pCopy->srcY, pCopy->srcZ}; amd::Coord3D dstOrigin = {pCopy->dstXInBytes, pCopy->dstY, pCopy->dstZ}; @@ -2250,7 +2259,6 @@ inline hipError_t ihipMemcpyCmdEnqueue(amd::Command* command, bool isAsync = fal hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool isAsync = false) { hipError_t status; - size_t offset = 0; if (pCopy == nullptr) { return hipErrorInvalidValue; } @@ -2262,55 +2270,9 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool pCopy->Height, pCopy->Depth); return hipSuccess; } - // 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) { - amd::Memory* memObj = getMemoryObject(pCopy->srcDevice, offset); - if (memObj != nullptr) { - srcMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & - memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice; - } else { - srcMemoryType = 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; - const_cast(pCopy)->srcXInBytes += offset; - // We don't need detect memory type again for hipMemoryTypeUnified - const_cast(pCopy)->srcMemoryType = srcMemoryType; - } - } - offset = 0; - hipMemoryType dstMemoryType = pCopy->dstMemoryType; - if (dstMemoryType == hipMemoryTypeUnified) { - amd::Memory* memObj = getMemoryObject(pCopy->dstDevice, offset); - if (memObj != nullptr) { - dstMemoryType = ((CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_USE_HOST_PTR) & - memObj->getMemFlags()) ? hipMemoryTypeHost : hipMemoryTypeDevice; - } else { - dstMemoryType = hipMemoryTypeHost; - } - - if (dstMemoryType == hipMemoryTypeHost) { - const_cast(pCopy)->dstHost = pCopy->dstDevice; - const_cast(pCopy)->dstXInBytes += offset; - // We don't need detect memory type again for hipMemoryTypeUnified - const_cast(pCopy)->dstMemoryType = dstMemoryType; - } - } - // 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. - offset = 0; - if (srcMemoryType == hipMemoryTypeHost) { - srcMemoryType = getMemoryObject(pCopy->srcHost, offset) ? hipMemoryTypeDevice : - hipMemoryTypeHost; - } - if (dstMemoryType == hipMemoryTypeHost) { - dstMemoryType = getMemoryObject(pCopy->dstHost, offset) ? hipMemoryTypeDevice : - hipMemoryTypeHost; - } + hipMemoryType srcMemoryType; + hipMemoryType dstMemoryType; + ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType); if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeHost)) { amd::Coord3D srcOrigin = {pCopy->srcXInBytes, pCopy->srcY, pCopy->srcZ}; @@ -2842,6 +2804,170 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) { return hipSuccess; } +hipError_t ihipDrvMemcpy3DParamValidate(const HIP_MEMCPY3D* p) { + // Passing more than one non-zero source or destination will cause hipMemcpy3D() to + // return an error. + if (p == nullptr || ((p->srcArray != nullptr) && (p->srcHost != nullptr)) || + ((p->dstArray != nullptr) && (p->dstHost != nullptr))) { + return hipErrorInvalidValue; + } + // The struct passed to hipMemcpy3D() must specify one of srcArray or srcPtr and one of dstArray + // or dstPtr. + if (((p->srcArray == nullptr) && (p->srcHost == nullptr)) || + ((p->dstArray == nullptr) && (p->dstHost == nullptr))) { + return hipErrorInvalidValue; + } + + // If the source and destination are both arrays, hipMemcpy3D() will return an error if they do + // not have the same element size. + if (((p->srcArray != nullptr) && (p->dstArray != nullptr)) && + (hip::getElementSize(p->srcArray) != hip::getElementSize(p->dstArray))) { + return hipErrorInvalidValue; + } + + // dst/src pitch must be less than max pitch + auto* deviceHandle = g_devices[hip::getCurrentDevice()->deviceId()]->devices()[0]; + const auto& info = deviceHandle->info(); + constexpr auto int32_max = static_cast(std::numeric_limits::max()); + auto maxPitch = std::min(info.maxMemAllocSize_, int32_max); + + // negative pitch cases + if (p->srcPitch >= maxPitch || p->dstPitch >= maxPitch) { + return hipErrorInvalidValue; + } + + if (p->dstArray == nullptr && p->srcArray == nullptr) { + if ((p->WidthInBytes + p->srcXInBytes > p->srcPitch) || + (p->WidthInBytes + p->dstXInBytes > p->dstPitch)) { + return hipErrorInvalidValue; + } + } + if (p->srcMemoryType < hipMemoryTypeHost || p->srcMemoryType > hipMemoryTypeManaged) { + return hipErrorInvalidMemcpyDirection; + } + if (p->dstMemoryType < hipMemoryTypeHost || p->dstMemoryType > hipMemoryTypeManaged) { + return hipErrorInvalidMemcpyDirection; + } + return hipSuccess; +} + +hipError_t ihipDrvMemcpy3D_validate(const HIP_MEMCPY3D* pCopy) { + hipError_t status; + if (pCopy->WidthInBytes == 0 || pCopy->Height == 0 || pCopy->Depth == 0) { + LogPrintfInfo("Either Width :%d or Height: %d and Depth: %d is zero", pCopy->WidthInBytes, + pCopy->Height, pCopy->Depth); + return hipSuccess; + } + hipMemoryType srcMemoryType; + hipMemoryType dstMemoryType; + ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType); + + 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; + } + } 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; + } + } 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; + } + } else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) { + amd::Image* dstImage; + size_t start = 0; + + status = + ihipMemcpyHtoAValidate(pCopy->srcHost, pCopy->dstArray, srcOrigin, dstOrigin, copyRegion, + pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, dstImage, + start); + if (status != hipSuccess) { + return status; + } + } else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) { + // Image to Host. + amd::Image* srcImage; + size_t start = 0; + + status = + ihipMemcpyAtoHValidate(pCopy->srcArray, pCopy->dstHost, srcOrigin, dstOrigin, copyRegion, + pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, srcImage, + start); + if (status != hipSuccess) { + return status; + } + } 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; + } + } 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; + } + } 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; + } + } else { + return hipErrorInvalidValue; + } + return hipSuccess; +} + hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p, hip::Stream* stream) { const HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*p);