diff --git a/projects/clr/rocclr/device/rocm/rocdefs.hpp b/projects/clr/rocclr/device/rocm/rocdefs.hpp index 6c7cc776af..a7dca96fe8 100644 --- a/projects/clr/rocclr/device/rocm/rocdefs.hpp +++ b/projects/clr/rocclr/device/rocm/rocdefs.hpp @@ -39,8 +39,7 @@ enum HwQueueEngine : uint32_t { Compute = 0, SdmaRead = 1, SdmaWrite = 2, - Unknown = 3, - External = 4 + Unknown = 3 }; } // namespace amd::roc diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 3b10056c8a..111329e6b0 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -534,28 +534,28 @@ std::vector& VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngi if (!skip_internal_signal) { external_signals_.push_back(signal_list_[current_id_]); } + } - // Validate all signals for the wait and skip already completed - for (uint32_t i = 0; i < external_signals_.size(); ++i) { - // Early signal status check - if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) { - const Settings& settings = gpu_.dev().settings(); - // Actively wait on CPU to avoid extra overheads of signal tracking on GPU. - // For small copies set forced wait - if (!WaitForSignal(external_signals_[i]->signal_, false, - external_signals_[i]->flags_.forceHostWait_)) { - if (settings.cpu_wait_for_signal_) { - // Wait on CPU for completion if requested - CpuWaitForSignal(external_signals_[i]); - } else { - // Add HSA signal for tracking on GPU - waiting_signals_.push_back(external_signals_[i]->signal_); - } + // Validate all signals for the wait and skip already completed + for (uint32_t i = 0; i < external_signals_.size(); ++i) { + // Early signal status check + if (hsa_signal_load_relaxed(external_signals_[i]->signal_) > 0) { + const Settings& settings = gpu_.dev().settings(); + // Actively wait on CPU to avoid extra overheads of signal tracking on GPU. + // For small copies set forced wait + if (!WaitForSignal(external_signals_[i]->signal_, false, + external_signals_[i]->flags_.forceHostWait_)) { + if (settings.cpu_wait_for_signal_) { + // Wait on CPU for completion if requested + CpuWaitForSignal(external_signals_[i]); + } else { + // Add HSA signal for tracking on GPU + waiting_signals_.push_back(external_signals_[i]->signal_); } } } - external_signals_.clear(); } + external_signals_.clear(); // Return the array of waiting HSA signals return waiting_signals_; diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index 8ae66fde31..a904e31ace 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -253,7 +253,6 @@ class VirtualGPU : public device::VirtualDevice { //! Adds an external signal(submission in another queue) for dependency tracking void AddExternalSignal(ProfilingSignal* signal) { external_signals_.push_back(signal); - engine_ = HwQueueEngine::External; } //! Get the last active signal on the queue