From 239cc6b437918e8544d46ecfcc9f612fac7ddb7f Mon Sep 17 00:00:00 2001 From: anusha GodavarthySurya Date: Wed, 6 Oct 2021 23:09:18 -0700 Subject: [PATCH] SWDEV-240806 - Added support for host graph node Change-Id: I8b3088d34d0b4e8080b55b5da80d821c5f29d8a4 [ROCm/clr commit: ce58dae7048739cd8271efb64e3e7d0f81a7c181] --- projects/clr/hipamd/src/amdhip.def | 4 ++ projects/clr/hipamd/src/hip_graph.cpp | 63 +++++++++++++++++++ .../clr/hipamd/src/hip_graph_internal.hpp | 34 +++++++++- projects/clr/hipamd/src/hip_hcc.def.in | 4 ++ projects/clr/hipamd/src/hip_hcc.map.in | 4 ++ 5 files changed, 107 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/src/amdhip.def b/projects/clr/hipamd/src/amdhip.def index 37e28a22b5..a0becb5fc8 100644 --- a/projects/clr/hipamd/src/amdhip.def +++ b/projects/clr/hipamd/src/amdhip.def @@ -340,3 +340,7 @@ hipGraphAddEventWaitNode hipGraphEventWaitNodeGetEvent hipGraphEventWaitNodeSetEvent hipGraphExecEventWaitNodeSetEvent +hipGraphAddHostNode +hipGraphHostNodeGetParams +hipGraphHostNodeSetParams +hipGraphExecHostNodeSetParams diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 84df06827e..a24d71b793 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -617,6 +617,23 @@ hipError_t capturehipStreamWaitEvent(hipEvent_t& event, hipStream_t& stream, uns return hipSuccess; } +hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*& userData) { + ClPrint(amd::LOG_INFO, amd::LOG_API, "[hipGraph] current capture node Memset2D on stream : %p", + stream); + if (fn == nullptr || userData == nullptr) { + return hipErrorInvalidValue; + } + hipHostNodeParams hostParams = {0}; + hostParams.fn = fn; + hostParams.userData = userData; + hip::Stream* s = reinterpret_cast(stream); + hipGraphNode_t pGraphNode = new hipGraphHostNode(&hostParams); + ihipGraphAddNode(pGraphNode, s->GetCaptureGraph(), s->GetLastCapturedNodes().data(), + s->GetLastCapturedNodes().size()); + s->SetLastCapturedNode(pGraphNode); + return hipSuccess; +} + hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus) { HIP_INIT_API(hipStreamIsCapturing, stream, pCaptureStatus); if (stream == nullptr) { @@ -1345,3 +1362,49 @@ hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraph } HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(event)); } + +hipError_t hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, + const hipGraphNode_t* pDependencies, size_t numDependencies, + const hipHostNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphAddHostNode, pGraphNode, graph, pDependencies, numDependencies, pNodeParams); + if (pGraphNode == nullptr || graph == nullptr || + (numDependencies > 0 && pDependencies == nullptr)) { + HIP_RETURN(hipErrorInvalidValue); + } + if (pNodeParams->fn == nullptr || pNodeParams->userData == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + *pGraphNode = new hipGraphHostNode(pNodeParams); + ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphHostNodeGetParams, node, pNodeParams); + if (node == nullptr || pNodeParams == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + reinterpret_cast(node)->GetParams(pNodeParams); + HIP_RETURN(hipSuccess); +} + +hipError_t hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphHostNodeSetParams, node, pNodeParams); + if (pNodeParams->fn == nullptr || pNodeParams->userData == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(node)->SetParams(pNodeParams)); +} + +hipError_t hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node, + const hipHostNodeParams* pNodeParams) { + HIP_INIT_API(hipGraphExecHostNodeSetParams, hGraphExec, node, pNodeParams); + if (pNodeParams->fn == nullptr || pNodeParams->userData == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node); + if (clonedNode == nullptr) { + HIP_RETURN(hipErrorInvalidValue); + } + HIP_RETURN(reinterpret_cast(clonedNode)->SetParams(pNodeParams)); +} diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index f22eebb572..2ca6046149 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -842,12 +842,42 @@ class hipGraphHostNode : public hipGraphNode { return new hipGraphHostNode(static_cast(*this)); } - hipError_t CreateCommand(amd::HostQueue* queue); + 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; + } + + static void Callback(cl_event event, cl_int command_exec_status, void* user_data) { + hipHostNodeParams* pNodeParams = reinterpret_cast(user_data); + pNodeParams->fn(pNodeParams->userData); + } + + void EnqueueCommands(hipStream_t stream) { + if (!commands_.empty()) { + if (!commands_[0]->setCallback(CL_COMPLETE, hipGraphHostNode::Callback, pNodeParams_)) { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed during setCallback"); + } + commands_[0]->enqueue(); + // Add the new barrier to stall the stream, until the callback is done + amd::Command::EventWaitList eventWaitList; + eventWaitList.push_back(commands_[0]); + amd::Command* block_command = + new amd::Marker(*commands_[0]->queue(), !kMarkerDisableFlush, eventWaitList); + if (block_command == nullptr) { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed during block command creation"); + } + block_command->enqueue(); + block_command->release(); + } + } void GetParams(hipHostNodeParams* params) { std::memcpy(params, pNodeParams_, sizeof(hipHostNodeParams)); } - hipError_t SetParams(hipHostNodeParams* params) { + hipError_t SetParams(const hipHostNodeParams* params) { std::memcpy(pNodeParams_, params, sizeof(hipHostNodeParams)); return hipSuccess; } diff --git a/projects/clr/hipamd/src/hip_hcc.def.in b/projects/clr/hipamd/src/hip_hcc.def.in index e318c9edff..b92267fe79 100644 --- a/projects/clr/hipamd/src/hip_hcc.def.in +++ b/projects/clr/hipamd/src/hip_hcc.def.in @@ -338,3 +338,7 @@ hipGraphAddEventWaitNode hipGraphEventWaitNodeGetEvent hipGraphEventWaitNodeSetEvent hipGraphExecEventWaitNodeSetEvent +hipGraphAddHostNode +hipGraphHostNodeGetParams +hipGraphHostNodeSetParams +hipGraphExecHostNodeSetParams diff --git a/projects/clr/hipamd/src/hip_hcc.map.in b/projects/clr/hipamd/src/hip_hcc.map.in index 89449c00d5..d9491e7d85 100644 --- a/projects/clr/hipamd/src/hip_hcc.map.in +++ b/projects/clr/hipamd/src/hip_hcc.map.in @@ -370,6 +370,10 @@ global: hipGraphEventWaitNodeGetEvent; hipGraphEventWaitNodeSetEvent; hipGraphExecEventWaitNodeSetEvent; + hipGraphAddHostNode; + hipGraphHostNodeGetParams; + hipGraphHostNodeSetParams; + hipGraphExecHostNodeSetParams; local: *; } hip_4.4;