Use system scope for packet following sdma copies

SWDEV-234947
SWDEV-236298
Instead of forcing a barrier packet, just inject system scope on the next packet.

Change-Id: If9bcee23e08dfe5db731235e2fcb30582cbd4c1c
This commit is contained in:
Christophe Paquot
2020-05-14 13:03:36 -07:00
committed by Christophe Paquot
parent 3624b8df16
commit 6a5af4056e
3 changed files with 11 additions and 3 deletions
+6
View File
@@ -436,6 +436,12 @@ bool VirtualGPU::dispatchGenericAqlPacket(
uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_);
hsa_signal_t signal;
if (addSystemScope_) {
header &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE);
header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE);
addSystemScope_ = false;
}
// TODO: placeholder to setup the kernel to populate start and end timestamp.
if (timestamp_ != nullptr) {
// Find signal slot