diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 876e727ca3..eb8e66b74a 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -211,12 +211,12 @@ bool HsaAmdSignalHandler(hsa_signal_value_t value, void* arg) { // Save callback signal hsa_signal_t callback_signal = ts->GetCallbackSignal(); - // Update the batch, since signal is complete - ts->gpu()->updateCommandsState(ts->command().GetBatchHead()); - // Reset last used SDMA engine mask ts->gpu()->setLastUsedSdmaEngine(0); + // Update the batch, since signal is complete + ts->gpu()->updateCommandsState(ts->command().GetBatchHead()); + // Reset API callback signal. It will release AQL queue and start commands processing if (callback_signal.handle != 0) { hsa_signal_subtract_relaxed(callback_signal, 1);