diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index a2fafdb35..57ab724d3 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -1182,6 +1182,10 @@ class GraphKernelNode : public GraphNode { kernelParams_.blockDim.y, kernelParams_.blockDim.z, kernelParams_.sharedMemBytes, stream, kernelParams_.kernelParams, kernelParams_.extra, kernelEvents_.startEvent_, kernelEvents_.stopEvent_, flags, coopKernel_, 0, 0, 0, 0, 0); + if (status != hipSuccess) { + return status; + } + if (signal_is_required_) { // Optimize the barriers by adding a signal into the dispatch packet directly command->SetProfiling(); @@ -1372,7 +1376,10 @@ class GraphMemcpyNode : public GraphNode { commands_.reserve(1); amd::Command* command; status = ihipMemcpy3DCommand(command, ©Params_, stream); - commands_.emplace_back(command); + if (status == hipSuccess) { + commands_.emplace_back(command); + } + return status; } @@ -1547,6 +1554,10 @@ class GraphMemcpyNode1D : public GraphMemcpyNode { WorkerThreadLock_.lock(); } status = ihipMemcpyCommand(command, dst_, src_, count_, kind_, *stream); + if (status != hipSuccess) { + return status; + } + hip::MemcpyType type = ihipGetMemcpyType(src_, dst_, kind_); if (type == hipCopyBuffer) { amd::CopyMemoryCommand* cpycmd = reinterpret_cast(command); @@ -1746,10 +1757,10 @@ class GraphMemcpyNodeFromSymbol : public GraphMemcpyNode1D { return status; } status = ihipMemcpyCommand(command, dst_, device_ptr, count_, kind_, *stream); - if (status != hipSuccess) { - return status; + if (status == hipSuccess) { + commands_.emplace_back(command); } - commands_.emplace_back(command); + return status; } @@ -1841,10 +1852,10 @@ class GraphMemcpyNodeToSymbol : public GraphMemcpyNode1D { return status; } status = ihipMemcpyCommand(command, device_ptr, src_, count_, kind_, *stream); - if (status != hipSuccess) { - return status; + if (status == hipSuccess) { + commands_.emplace_back(command); } - commands_.emplace_back(command); + return status; } @@ -2116,7 +2127,10 @@ class GraphEventRecordNode : public GraphNode { commands_.reserve(1); amd::Command* command = nullptr; status = e->recordCommand(command, stream); - commands_.emplace_back(command); + if (status == hipSuccess) { + commands_.emplace_back(command); + } + return status; } @@ -2173,7 +2187,10 @@ class GraphEventWaitNode : public GraphNode { commands_.reserve(1); amd::Command* command; status = e->streamWaitCommand(command, stream); - commands_.emplace_back(command); + if (status == hipSuccess) { + commands_.emplace_back(command); + } + return status; } @@ -2222,6 +2239,9 @@ class GraphHostNode : public GraphNode { amd::Command::EventWaitList waitList; commands_.reserve(1); amd::Command* command = new amd::Marker(*stream, !kMarkerDisableFlush, waitList); + if (command == nullptr) { + return hipErrorOutOfMemory; + } commands_.emplace_back(command); return hipSuccess; } @@ -2283,6 +2303,9 @@ class GraphEmptyNode : public GraphNode { amd::Command::EventWaitList waitList; commands_.reserve(1); amd::Command* command = new amd::Marker(*stream, !kMarkerDisableFlush, waitList); + if (command == nullptr) { + return hipErrorOutOfMemory; + } commands_.emplace_back(command); } return hipSuccess; @@ -2564,7 +2587,10 @@ class GraphDrvMemcpyNode : public GraphNode { commands_.reserve(1); amd::Command* command; status = ihipGetMemcpyParam3DCommand(command, ©Params_, stream); - commands_.emplace_back(command); + if (status == hipSuccess) { + commands_.emplace_back(command); + } + return status; }