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

SWDEV-272496 - Wait on CPU before switching to GPU wait

Browse Source 
GPU waits have noticeable overheads on compute with extra
AQL barrier packet and on SDMA with power saving features. This
change introduces a wait on CPU for 30 us in case the app has tiny
operations.

Change-Id: I761ba3af595f3f48544980058a9077dda15aa5f9
This commit is contained in:
German Andryeyev
2021-02-17 14:48:08 -05:00
committed by Saleel Kudchadker
parent 9900b46c21
commit ac387f9b03
2 changed files with 29 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/device/rocm/rocvirtual.cpp
+9 -5
View File
@@ -388,11 +388,15 @@ hsa_signal_t* VirtualGPU::HwQueueTracker::WaitingSignal(HwQueueEngine engine) {
// Early signal status check
if (hsa_signal_load_relaxed(prof_signal->signal_) > 0) {
const Settings& settings = gpu_.dev().settings();
// Wait on CPU if requested
if (settings.cpu_wait_for_signal_) {
CpuWaitForSignal(prof_signal);
} else {
return &prof_signal->signal_;
// Actively wait on CPU for 30 us to avoid extra overheads of signal tracking on GPU
if (!WaitForSignal<kTimeout30us>(prof_signal->signal_)) {
if (settings.cpu_wait_for_signal_) {
// Wait on CPU for completion if requested
CpuWaitForSignal(prof_signal);
} else {
// Return HSA signal for tracking on GPU
return &prof_signal->signal_;
}
}
}
}