diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h index 40bac15efe..893dbbfc34 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h @@ -279,6 +279,7 @@ enum hip_api_id_t { HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureObjectTextureDesc = HIP_API_ID_NONE, HIP_API_ID_hipGetTextureReference = HIP_API_ID_NONE, + HIP_API_ID_hipGraphAddMemcpyNode1D = HIP_API_ID_NONE, HIP_API_ID_hipMemcpy2DArrayToArray = HIP_API_ID_NONE, HIP_API_ID_hipMemcpyArrayToArray = HIP_API_ID_NONE, HIP_API_ID_hipMemcpyAtoA = HIP_API_ID_NONE, @@ -3769,6 +3770,8 @@ typedef struct hip_api_data_s { #define INIT_hipGetTextureObjectTextureDesc_CB_ARGS_DATA(cb_data) {}; // hipGetTextureReference() #define INIT_hipGetTextureReference_CB_ARGS_DATA(cb_data) {}; +// hipGraphAddMemcpyNode1D() +#define INIT_hipGraphAddMemcpyNode1D_CB_ARGS_DATA(cb_data) {}; // hipMemcpy2DArrayToArray() #define INIT_hipMemcpy2DArrayToArray_CB_ARGS_DATA(cb_data) {}; // hipMemcpyArrayToArray() diff --git a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h index fbe70bab8a..5175cebcb5 100644 --- a/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvidia_detail/nvidia_hip_runtime_api.h @@ -2292,6 +2292,13 @@ inline static hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGr cudaGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, pCopyParams)); } +inline static 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) { + return hipCUDAErrorTohipError( + cudaGraphAddMemcpyNode1D(pGraphNode, graph, pDependencies, numDependencies, dst, src, count, kind)); +} + inline static hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const hipGraphNode_t* pDependencies, size_t numDependencies, diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 9a3c43ed39..7731f2aea2 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -281,6 +281,7 @@ hipGraphDestroy hipGraphAddKernelNode hipGraphAddMemsetNode hipGraphAddMemcpyNode +hipGraphAddMemcpyNode1D hipGraphInstantiate hipGraphLaunch hipStreamIsCapturing @@ -300,3 +301,5 @@ hipGraphMemcpyNodeSetParams hipGraphMemsetNodeGetParams hipGraphMemsetNodeSetParams hipGraphAddDependencies +hipGraphExecKernelNodeSetParams +hipGraphAddEmptyNode diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index ed592cbbfb..623e2f024c 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -27,13 +27,7 @@ thread_local std::vector g_captureStreams; std::unordered_map hipGraphExec::activeGraphExec_; -hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, - const hipGraphNode_t* pDependencies, size_t numDependencies, - const hipKernelNodeParams* pNodeParams) { - if (pGraphNode == nullptr || graph == nullptr || - (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr) { - return hipErrorInvalidValue; - } +hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams) { hipFunction_t func = nullptr; hipError_t status = PlatformState::instance().getStatFunc(&func, pNodeParams->func, ihipGetDevice()); @@ -56,6 +50,24 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, if (status != hipSuccess) { return status; } +} + +hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipKernelNodeParams* pNodeParams) { + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr) || pNodeParams == nullptr) { + return hipErrorInvalidValue; + } + hipError_t status = ihipValidateKernelParams(pNodeParams); + if (hipSuccess != status) { + return status; + } + hipFunction_t func = nullptr; + status = PlatformState::instance().getStatFunc(&func, pNodeParams->func, ihipGetDevice()); + if ((status != hipSuccess) || (func == nullptr)) { + return hipErrorInvalidDeviceFunction; + } *pGraphNode = new hipGraphKernelNode(pNodeParams, func); if (numDependencies == 0) { graph->AddNode(*pGraphNode); @@ -88,6 +100,26 @@ hipError_t ihipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, return hipSuccess; } +hipError_t ihipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + void* dst, const void* src, size_t count, hipMemcpyKind kind) { + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr)) { + return hipErrorInvalidValue; + } + ihipMemcpy_validate(dst, src, count, kind); + *pGraphNode = new hipGraphMemcpyNode1D(dst, src, count, kind); + if (numDependencies == 0) { + graph->AddNode(*pGraphNode); + } + for (size_t i = 0; i < numDependencies; i++) { + if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { + return hipErrorInvalidValue; + } + } + return hipSuccess; +} + hipError_t ihipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, const hipGraphNode_t* pDependencies, size_t numDependencies, const hipMemsetParams* pMemsetParams) { @@ -268,6 +300,7 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node StreamWaitEvent on stream : %p, Event %p", stream, event); + hip::Stream* s = reinterpret_cast(stream); hip::Event* e = reinterpret_cast(event); @@ -281,7 +314,7 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns } s->AddCrossCapturedNode(e->GetNodesPrevToRecorded()); g_captureStreams.push_back(stream); - return hipSuccess; + HIP_RETURN(hipSuccess); } hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus** pCaptureStatus) { @@ -355,7 +388,7 @@ hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, HIP_INIT_API(hipGraphAddKernelNode, pGraphNode, graph, pDependencies, numDependencies, pNodeParams); HIP_RETURN_DURATION( - ihipGraphAddKernelNode(pGraphNode, graph, pDependencies, numDependencies, pNodeParams);); + ihipGraphAddKernelNode(pGraphNode, graph, pDependencies, numDependencies, pNodeParams)); } hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -365,7 +398,17 @@ hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, pCopyParams); HIP_RETURN_DURATION( - ihipGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, pCopyParams);); + ihipGraphAddMemcpyNode(pGraphNode, graph, pDependencies, numDependencies, pCopyParams)); +} + +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) { + HIP_INIT_API(hipGraphAddMemcpyNode1D, pGraphNode, graph, pDependencies, numDependencies, dst, src, + count, kind); + + HIP_RETURN_DURATION(ihipGraphAddMemcpyNode1D(pGraphNode, graph, pDependencies, numDependencies, + dst, src, count, kind)); } hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, @@ -375,7 +418,26 @@ hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, pMemsetParams); HIP_RETURN_DURATION( - ihipGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, pMemsetParams);); + ihipGraphAddMemsetNode(pGraphNode, graph, pDependencies, numDependencies, pMemsetParams)); +} + +hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies) { + HIP_INIT_API(hipGraphAddEmptyNode, pGraphNode, graph, pDependencies, numDependencies); + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr)) { + return HIP_RETURN(hipErrorInvalidValue); + } + *pGraphNode = new hipGraphEmptyNode(); + if (numDependencies == 0) { + graph->AddNode(*pGraphNode); + } + for (size_t i = 0; i < numDependencies; i++) { + if (graph->AddEdge(*(pDependencies + i), *pGraphNode) != hipSuccess) { + return hipErrorInvalidValue; + } + } + return hipSuccess; } hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph, @@ -417,8 +479,9 @@ hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) { hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes) { HIP_INIT_API(hipGraphGetNodes, graph, nodes, numNodes); if (graph == nullptr || numNodes == nullptr) { - *numNodes = graph->GetNodeCount(); + return HIP_RETURN(hipErrorInvalidValue); } + *numNodes = graph->GetNodeCount(); if (*numNodes > 0) { nodes = graph->GetNodes().data(); } @@ -510,3 +573,12 @@ hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from, } HIP_RETURN(hipSuccess); } + +hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + const hipKernelNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphExecKernelNodeSetParams, hGraphExec, node, pNodeParams); + if (hGraphExec == nullptr || node == nullptr || pNodeParams == nullptr) { + return HIP_RETURN(hipErrorInvalidValue); + } + return reinterpret_cast(node)->SetCommandParams(pNodeParams); +} diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index fc187c20fd..312aca3c14 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -32,6 +32,7 @@ #include "hip_graph_helper.hpp" typedef hipGraphNode* Node; +hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams); struct hipGraphNode { protected: @@ -123,6 +124,24 @@ class hipGraphKernelNode : public hipGraphNode { void SetParams(const hipKernelNodeParams* params) { std::memcpy(pKernelParams_, params, sizeof(hipKernelNodeParams)); } + hipError_t SetCommandParams(const hipKernelNodeParams* params) { + if (params->func != pKernelParams_->func) { + return hipErrorInvalidValue; + } + // updates kernel params + hipError_t status = ihipValidateKernelParams(params); + if (hipSuccess != status) { + return status; + } + size_t globalWorkOffset[3] = {0}; + size_t globalWorkSize[3] = {params->gridDim.x, params->gridDim.y, params->gridDim.z}; + size_t localWorkSize[3] = {params->blockDim.x, params->blockDim.y, params->blockDim.z}; + reinterpret_cast(commands_[0]) + ->setSizes(globalWorkOffset, globalWorkSize, localWorkSize); + reinterpret_cast(commands_[0]) + ->setSharedMemBytes(params->sharedMemBytes); + return hipSuccess; + } }; class hipGraphMemcpyNode : public hipGraphNode { @@ -150,7 +169,6 @@ class hipGraphMemcpyNode : public hipGraphNode { } }; - class hipGraphMemcpyNode1D : public hipGraphNode { void* dst_; const void* src_; @@ -315,6 +333,20 @@ class hipGraphHostNode : public hipGraphNode { } }; +class hipGraphEmptyNode : public hipGraphNode { + public: + hipGraphEmptyNode() : hipGraphNode(hipGraphNodeTypeEmpty) {} + ~hipGraphEmptyNode() {} + + hipError_t CreateCommand(amd::HostQueue* queue) { + amd::Command::EventWaitList waitList; + commands_.reserve(1); + amd::Command* command = new amd::Marker(*queue, !kMarkerDisableFlush, waitList); + commands_.emplace_back(command); + return hipSuccess; + } +}; + struct hipGraphExec { std::vector> parallelLists_; std::vector levelOrder_; diff --git a/projects/clr/hipamd/src/hip_hcc.def.in b/projects/clr/hipamd/src/hip_hcc.def.in index 73d747e98e..d9cb153016 100755 --- a/projects/clr/hipamd/src/hip_hcc.def.in +++ b/projects/clr/hipamd/src/hip_hcc.def.in @@ -281,6 +281,7 @@ hipGraphDestroy hipGraphAddKernelNode hipGraphAddMemsetNode hipGraphAddMemcpyNode +hipGraphAddMemcpyNode1D hipGraphInstantiate hipGraphLaunch hipStreamIsCapturing @@ -306,3 +307,5 @@ hipGraphMemcpyNodeSetParams hipGraphMemsetNodeGetParams hipGraphMemsetNodeSetParams hipGraphAddDependencies +hipGraphExecKernelNodeSetParams +hipGraphAddEmptyNode diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index e0e37df664..d997d7d71d 100755 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -291,6 +291,7 @@ global: hipGraphAddKernelNode; hipGraphAddMemsetNode; hipGraphAddMemcpyNode; + hipGraphAddMemcpyNode1D; hipGraphInstantiate; hipGraphLaunch; hipStreamIsCapturing; @@ -326,6 +327,8 @@ global: hipStreamWaitValue64; hipStreamWriteValue32; hipStreamWriteValue64; + hipGraphExecKernelNodeSetParams; + hipGraphAddEmptyNode; local: *; } hip_4.3; diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index 0c93cfc0e2..568fd6f818 100755 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -40,7 +40,8 @@ Stream::Stream(hip::Device* dev, Priority p, unsigned int f, bool null_stream, flags_(f), null_(null_stream), cuMask_(cuMask), - captureStatus_(captureStatus) {} + captureStatus_(captureStatus), + originStream_(false) {} // ================================================================================================ Stream::~Stream() {