From c116411e002661bba45754b5e14669c0319b6dda Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Wed, 17 Nov 2021 15:56:36 -0500 Subject: [PATCH] SWDEV-294669 - Avoid queue drain Use slot wait logic for direct dispatch Change-Id: I431ba1418eb4aa066b9881934f4055b3d338ce3a [ROCm/clr commit: 8e4101b4fd099899f09c888055f8f24087c2b598] --- projects/clr/rocclr/device/rocm/rocvirtual.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 7b71914f9b..620e0ad48e 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -775,8 +775,7 @@ bool VirtualGPU::dispatchGenericAqlPacket( AqlPacket* packet, uint16_t header, uint16_t rest, bool blocking, size_t size) { const uint32_t queueSize = gpu_queue_->size; const uint32_t queueMask = queueSize - 1; - // @note: Reserve extra slot for HW processing. There is unknown race condition in some apps. - const uint32_t sw_queue_size = queueMask - 1; + const uint32_t sw_queue_size = queueMask; // Check for queue full and wait if needed. uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, size); @@ -794,8 +793,10 @@ bool VirtualGPU::dispatchGenericAqlPacket( amd::Os::yield(); } - // Add blocking command if the original value of read index was behind of the queue size - if (blocking || (index - read) >= sw_queue_size) { + // Add blocking command if the original value of read index was behind of the queue size. + // Note: direct dispatch relies on the slot stall above to keep the forward progress + // of the app if a dispatched kernel requires some CPU input for completion + if (blocking || (!AMD_DIRECT_DISPATCH && (index - read) >= sw_queue_size)) { if (packet->completion_signal.handle == 0) { packet->completion_signal = Barriers().ActiveSignal(); }