From 148a5cac39e030f75a81929bf130d427d52f8741 Mon Sep 17 00:00:00 2001 From: German Andryeyev Date: Wed, 26 May 2021 01:34:13 -0400 Subject: [PATCH] SWDEV-287137 - Fix regression with dependent signals - Make sure barrier with dependent signals issues before queue index reservation - Don't issue extra barrier if it's already a barrier command with dependent signals Change-Id: I179a8b7adac79eed698f4a4d9eca2606d8e913aa --- rocclr/device/rocm/rocvirtual.cpp | 78 +++++++++++++++++-------------- rocclr/device/rocm/rocvirtual.hpp | 3 +- 2 files changed, 45 insertions(+), 36 deletions(-) diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index fcfb7fd8a5..8e0ffa2bcc 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -798,20 +798,13 @@ bool VirtualGPU::dispatchGenericAqlPacket( void VirtualGPU::dispatchBlockingWait() { auto wait_signals = Barriers().WaitingSignal(); // AQL dispatch doesn't support dependent signals and extra barrier packet must be generated - if (wait_signals.size() != 0) { - for (uint32_t i = 0; i < wait_signals.size(); ++i) { - uint32_t j = i % 5; - barrier_packet_.dep_signal[j] = wait_signals[i]; - constexpr bool kSkipSignal = true; - // If runtime reached the packet limit or the count limit, then flush the barrier - if ((j == 4) || ((i + 1) == wait_signals.size())) { - dispatchBarrierPacket(&barrier_packet_, kNopPacketHeader, kSkipSignal); - barrier_packet_.dep_signal[0] = hsa_signal_t{}; - barrier_packet_.dep_signal[1] = hsa_signal_t{}; - barrier_packet_.dep_signal[2] = hsa_signal_t{}; - barrier_packet_.dep_signal[3] = hsa_signal_t{}; - barrier_packet_.dep_signal[4] = hsa_signal_t{}; - } + for (uint32_t i = 0; i < wait_signals.size(); ++i) { + uint32_t j = i % 5; + barrier_packet_.dep_signal[j] = wait_signals[i]; + constexpr bool kSkipSignal = true; + // If runtime reached the packet limit or the count limit, then flush the barrier + if ((j == 4) || ((i + 1) == wait_signals.size())) { + dispatchBarrierPacket(kNopPacketHeader, kSkipSignal); } } } @@ -860,32 +853,50 @@ bool VirtualGPU::dispatchCounterAqlPacket(hsa_ext_amd_aql_pm4_packet_t* packet, } // ================================================================================================ -void VirtualGPU::dispatchBarrierPacket(hsa_barrier_and_packet_t* packet, - uint16_t packetHeader, bool skipSignal) { +void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal) { const uint32_t queueSize = gpu_queue_->size; const uint32_t queueMask = queueSize - 1; - uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1); - uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); - packet->completion_signal.handle = 0; if (!skipSignal) { - dispatchBlockingWait(); + // Make sure the wait is issued before queue index reservation + auto wait_signals = Barriers().WaitingSignal(); + for (uint32_t i = 0; i < wait_signals.size(); ++i) { + uint32_t j = i % 5; + barrier_packet_.dep_signal[j] = wait_signals[i]; + constexpr bool kSkipSignal = true; + // If runtime reached the packet limit and signals left, then flush the barrier + if ((j == 4) && ((i + 1) < wait_signals.size())) { + dispatchBarrierPacket(kNopPacketHeader, kSkipSignal); + } + } + } + uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1); + uint64_t read = hsa_queue_load_read_index_relaxed(gpu_queue_); + barrier_packet_.completion_signal.handle = 0; + + if (!skipSignal) { // Pool size must grow to the size of pending AQL packets const uint32_t pool_size = index - read; // Get active signal for current dispatch if profiling is necessary - packet->completion_signal = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_, - pool_size); + barrier_packet_.completion_signal = + Barriers().ActiveSignal(kInitSignalValueOne, timestamp_, pool_size); } while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask); hsa_barrier_and_packet_t* aql_loc = &(reinterpret_cast(gpu_queue_->base_address))[index & queueMask]; - *aql_loc = *packet; + *aql_loc = barrier_packet_; __atomic_store_n(reinterpret_cast(aql_loc), packetHeader, __ATOMIC_RELEASE); hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index); + // Clear dependent signals for the next packet + barrier_packet_.dep_signal[0] = hsa_signal_t{}; + barrier_packet_.dep_signal[1] = hsa_signal_t{}; + barrier_packet_.dep_signal[2] = hsa_signal_t{}; + barrier_packet_.dep_signal[3] = hsa_signal_t{}; + barrier_packet_.dep_signal[4] = hsa_signal_t{}; ClPrint(amd::LOG_DEBUG, amd::LOG_AQL, "[%zx] HWq=0x%zx, BarrierAND Header = 0x%x (type=%d, barrier=%d, acquire=%d," " release=%d), " @@ -894,13 +905,14 @@ void VirtualGPU::dispatchBarrierPacket(hsa_barrier_and_packet_t* packet, extractAqlBits(packetHeader, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE), extractAqlBits(packetHeader, HSA_PACKET_HEADER_BARRIER, - HSA_PACKET_HEADER_WIDTH_BARRIER), + HSA_PACKET_HEADER_WIDTH_BARRIER), extractAqlBits(packetHeader, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE, - HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE), + HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE), extractAqlBits(packetHeader, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE, - HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE), - packet->dep_signal[0], packet->dep_signal[1], packet->dep_signal[2], - packet->dep_signal[3], packet->dep_signal[4], packet->completion_signal); + HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE), + barrier_packet_.dep_signal[0], barrier_packet_.dep_signal[1], + barrier_packet_.dep_signal[2], barrier_packet_.dep_signal[3], + barrier_packet_.dep_signal[4], barrier_packet_.completion_signal); } // ================================================================================================ @@ -914,14 +926,13 @@ void VirtualGPU::ResetQueueStates() { // Release the pool, since runtime just completed a barrier // @note: Runtime can reset kernel arg pool only if the barrier with L2 invalidation was issued resetKernArgPool(); - } // ================================================================================================ bool VirtualGPU::releaseGpuMemoryFence(bool skip_cpu_wait) { if (hasPendingDispatch_) { // Dispatch barrier packet into the queue - dispatchBarrierPacket(&barrier_packet_, kBarrierPacketHeader); + dispatchBarrierPacket(kBarrierPacketHeader); hasPendingDispatch_ = false; } @@ -1134,7 +1145,7 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) { //! We can issue a barrier to avoid expensive extra memory allocations. // Dispatch barrier packet into the queue and wait till it finishes. - dispatchBarrierPacket(&barrier_packet_, kBarrierPacketHeader); + dispatchBarrierPacket(kBarrierPacketHeader); if (!Barriers().WaitCurrent()) { LogError("Kernel arguments reset failed"); } @@ -2767,7 +2778,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const } if (gpuKernel.dynamicParallelism()) { - dispatchBarrierPacket(&barrier_packet_, kBarrierPacketHeader, true); + dispatchBarrierPacket(kBarrierPacketHeader, true); static_cast(blitMgr()).runScheduler( getVQVirtualAddress(), schedulerParam_, schedulerQueue_, schedulerSignal_, schedulerThreads_); } @@ -2881,8 +2892,7 @@ void VirtualGPU::submitMarker(amd::Marker& vcmd) { // If there was a pending dispatch use a Barrier packet // with cache flushes. This saves on additional barrier // for cache flushes explicitly and helps wall time - uint16_t header = kNopPacketHeader; - dispatchBarrierPacket(&barrier_packet_, header); + dispatchBarrierPacket(kNopPacketHeader); // Direct dispatch requires a barrier with callback and hasPendingDispatch_ triggers that if (AMD_DIRECT_DISPATCH) { hasPendingDispatch_ = true; diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index 6545cfbff3..85987083b1 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/rocclr/device/rocm/rocvirtual.hpp @@ -368,8 +368,7 @@ class VirtualGPU : public device::VirtualDevice { template bool dispatchGenericAqlPacket(AqlPacket* packet, uint16_t header, uint16_t rest, bool blocking, size_t size = 1); - void dispatchBarrierPacket(hsa_barrier_and_packet_t* packet, uint16_t packetHeader, - bool skipSignal = false); + void dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal = false); bool dispatchCounterAqlPacket(hsa_ext_amd_aql_pm4_packet_t* packet, const uint32_t gfxVersion, bool blocking, const hsa_ven_amd_aqlprofile_1_00_pfn_t* extApi); void dispatchBarrierValuePacket(const hsa_amd_barrier_value_packet_t* packet,