AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-470612 - Avoid processing internal signals

Browse Source 
If only external signals were provided, then just process it
without adding internal signals

Change-Id: Iaefd65d0f8b0a64b9f6a864a9bd73de20a29dfa4


[ROCm/clr commit: 18187cd8fe]
This commit is contained in:
German Andryeyev
2024-07-24 17:10:16 -04:00
parent 47dcfbae6b
commit 9d1d3a6493
3 changed files with 18 additions and 20 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+17 -17
View File
@@ -534,28 +534,28 @@ std::vector<hsa_signal_t>& 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<true>(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<true>(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_;