diff --git a/hipamd/src/hip_graph_internal.cpp b/hipamd/src/hip_graph_internal.cpp index f6c7221724..1205f44470 100644 --- a/hipamd/src/hip_graph_internal.cpp +++ b/hipamd/src/hip_graph_internal.cpp @@ -363,6 +363,7 @@ hipError_t GraphExec::Init() { status = CaptureAQLPackets(); } instantiateDeviceId_ = hip::getCurrentDevice()->deviceId(); + static_cast( hip::getCurrentDevice())->retain(); return status; } diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 3fd3be5f4c..856aa839d4 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -505,7 +505,6 @@ struct Graph { amd::ScopedLock lock(graphSetLock_); graphSet_.insert(this); mem_pool_ = device->GetGraphMemoryPool(); - mem_pool_->retain(); graphInstantiated_ = false; roots_.resize(DEBUG_HIP_FORCE_GRAPH_QUEUES); leafs_.resize(DEBUG_HIP_FORCE_GRAPH_QUEUES); @@ -543,9 +542,6 @@ struct Graph { } } graphUserObj_.clear(); - if (mem_pool_ != nullptr) { - mem_pool_->release(); - } memAllocNodePtrs_.clear(); } @@ -764,6 +760,9 @@ struct GraphExec : public amd::ReferenceCountedObject, public Graph { kernArgManager_->release(); } } + if (instantiateDeviceId_ != -1) { + static_cast(g_devices[instantiateDeviceId_])->release(); + } } Node GetClonedNode(Node node) { diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 778f7f6be1..d1367cae5f 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -235,11 +235,13 @@ bool HsaAmdSignalHandler(hsa_signal_value_t value, void* arg) { // Reset last used SDMA engine mask gpu->setLastUsedSdmaEngine(0); + bool isBlocking = ts->GetBlocking(); + // Update the batch, since signal is complete gpu->updateCommandsState(ts->command().GetBatchHead()); // Reset API callback signal. It will release AQL queue and start commands processing - if (callback_signal.handle != 0 && ts->GetBlocking()) { + if (callback_signal.handle != 0 && isBlocking) { hsa_signal_subtract_relaxed(callback_signal, 1); }