Add option to skip AQL barrier

The change reuses HSA signals for dispatches as a wait signal.
Skipping the barrier requires to  disable L2 cache for sysmem
allocations and extra tracking for HDP access with the large bar.
ROC_BARRIER_SYNC=0 activates the new logic. Barrier sync is
still used by default.
ROC_ACTIVE_WAIT=1 enables unconditional active wait in ROCr.
The change also consolidated ROCr wait logic under single function.

Change-Id: I6bd1be30aa88258da1b1f9de319ef5a45852afd8

rocclr/device/rocm/rocblit.cpp

@@ -379,6 +379,7 @@ bool DmaBlitManager::copyBuffer(device::Memory& srcMemory, device::Memory& dstMe
   return true;
 }
 
+// ================================================================================================
 bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& dstMemory,
                                     const amd::BufferRect& srcRect, const amd::BufferRect& dstRect,
                                     const amd::Coord3D& size, bool entire) const {
@@ -435,8 +436,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
     }
 
     if (isSubwindowRectCopy ) {
-      const hsa_signal_value_t kInitVal = 1;
-      hsa_signal_store_relaxed(completion_signal_, kInitVal);
+      hsa_signal_store_relaxed(completion_signal_, kInitSignalValueOne);
 
       // Copy memory line by line
       hsa_status_t status =
@@ -447,10 +447,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
         return false;
       }
 
-
-      hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0,
-                                                      uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
-      if (val != 0) {
+      if (!WaitForSignal(completion_signal_)) {
         LogError("Async copy failed");
         return false;
       }
@@ -476,9 +473,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
         }
       }
 
-      hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0,
-                                                      uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
-      if (val != 0) {
+      if (!WaitForSignal(completion_signal_)) {
         LogError("Async copy failed");
         return false;
       }
@@ -488,6 +483,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d
   return true;
 }
 
+// ================================================================================================
 bool DmaBlitManager::copyImageToBuffer(device::Memory& srcMemory, device::Memory& dstMemory,
                                        const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin,
                                        const amd::Coord3D& size, bool entire, size_t rowPitch,
@@ -598,6 +594,7 @@ bool DmaBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dstMem
   return result;
 }
 
+// ================================================================================================
 bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
                              const amd::Coord3D& srcOrigin, const amd::Coord3D& dstOrigin,
                              const amd::Coord3D& size, bool enableCopyRect, bool flushDMA) const {
@@ -639,8 +636,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
     srcAgent = dstAgent = dev().getBackendDevice();
   }
 
-  const hsa_signal_value_t kInitVal = 1;
-  hsa_signal_store_relaxed(completion_signal_, kInitVal);
+  hsa_signal_store_relaxed(completion_signal_, kInitSignalValueOne);
 
   // Use SDMA to transfer the data
   status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, size[0], 0, nullptr,
@@ -649,21 +645,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
   if (status == HSA_STATUS_SUCCESS) {
     hsa_signal_value_t val;
 
-    // Use ACTIVE wait for small transfers.
-    // Might want to be dependent on also having an idle GPU
-    // or, if queue is busy, may want to enqueue a blank barrier
-    // before this and wait BLOCKED on its completion signal, followed
-    // by ACTIVE on this.
-
-    constexpr size_t small_transfer_size = 4 * Mi;
-    if (size[0] < small_transfer_size) {
-      val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0,
-                                    std::numeric_limits<uint64_t>::max(), HSA_WAIT_STATE_ACTIVE);
-    } else {
-      val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0,
-                                    std::numeric_limits<uint64_t>::max(), HSA_WAIT_STATE_BLOCKED);
-    }
-    if (val != (kInitVal - 1)) {
+    if (!WaitForSignal(completion_signal_)) {
       LogError("Async copy failed");
       status = HSA_STATUS_ERROR;
     } else {
@@ -676,6 +658,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
   return (status == HSA_STATUS_SUCCESS);
 }
 
+// ================================================================================================
 bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_t size,
                                    address staging, bool hostToDev) const {
   // No allocation is necessary for Full Profile
@@ -693,12 +676,10 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
 
   address hsaBuffer = staging;
 
-  const hsa_signal_value_t kInitVal = 1;
-
   // Allocate requested size of memory
   while (totalSize > 0) {
     size = std::min(totalSize, dev().settings().stagedXferSize_);
-    hsa_signal_silent_store_relaxed(completion_signal_, kInitVal);
+    hsa_signal_silent_store_relaxed(completion_signal_, kInitSignalValueOne);
 
     // Copy data from Host to Device
     if (hostToDev) {
@@ -712,10 +693,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
       status = hsa_amd_memory_async_copy(hostDst + offset, dev().getBackendDevice(), hsaBuffer,
                                          srcAgent, size, 0, nullptr, completion_signal_);
       if (status == HSA_STATUS_SUCCESS) {
-        hsa_signal_value_t val = hsa_signal_wait_scacquire(
-            completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
-
-        if (val != (kInitVal - 1)) {
+        if (!WaitForSignal(completion_signal_)) {
           LogError("Async copy failed");
           return false;
         }
@@ -739,10 +717,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
         hsa_amd_memory_async_copy(hsaBuffer, dstAgent, hostSrc + offset,
                                   dev().getBackendDevice(), size, 0, nullptr, completion_signal_);
     if (status == HSA_STATUS_SUCCESS) {
-      hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ,
-                                                       0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED);
-
-      if (val != (kInitVal - 1)) {
+      if (!WaitForSignal(completion_signal_)) {
         LogError("Async copy failed");
         return false;
       }
@@ -760,6 +735,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
   return true;
 }
 
+// ================================================================================================
 KernelBlitManager::KernelBlitManager(VirtualGPU& gpu, Setup setup)
     : DmaBlitManager(gpu, setup),
       program_(nullptr),
@@ -1659,6 +1635,7 @@ bool KernelBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory
   return result;
 }
 
+// ================================================================================================
 bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost,
                                    const amd::Coord3D& origin, const amd::Coord3D& size,
                                    bool entire) const {
@@ -1667,12 +1644,13 @@ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost,
 
   if (dev().info().largeBar_ && size[0] <= kMaxD2hMemcpySize) {
     if ((srcMemory.owner()->getHostMem() == nullptr) && (srcMemory.owner()->getSvmPtr() != nullptr)) {
-      // CPU read ahead, hence release GPU memory
-      gpu().releaseGpuMemoryFence();
+      // CPU read ahead, hence release GPU memory and force barrier to make sure L2 flush
+      constexpr bool ForceBarrier = true;
+      gpu().releaseGpuMemoryFence(ForceBarrier);
       char* src = reinterpret_cast<char*>(srcMemory.owner()->getSvmPtr());
       std::memcpy(dstHost, src + origin[0], size[0]);
-      // Set HASPENDINGDISPATCH_ FLAG. That will force L2 invalidation on flush
-      gpu().hasPendingDispatch();
+      // The first dispatch will invalidate L2
+      gpu().addSystemScope();
       return true;
     }
   }
@@ -1717,6 +1695,7 @@ bool KernelBlitManager::readBuffer(device::Memory& srcMemory, void* dstHost,
   return result;
 }
 
+// ================================================================================================
 bool KernelBlitManager::readBufferRect(device::Memory& srcMemory, void* dstHost,
                                        const amd::BufferRect& bufRect,
                                        const amd::BufferRect& hostRect, const amd::Coord3D& size,
@@ -1764,6 +1743,7 @@ bool KernelBlitManager::readBufferRect(device::Memory& srcMemory, void* dstHost,
   return result;
 }
 
+// ================================================================================================
 bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemory,
                                     const amd::Coord3D& origin, const amd::Coord3D& size,
                                     bool entire) const {
@@ -1773,12 +1753,13 @@ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemo
   if (dev().info().largeBar_ && size[0] <= kMaxH2dMemcpySize) {
     if ((dstMemory.owner()->getHostMem() == nullptr) && (dstMemory.owner()->getSvmPtr() != nullptr)) {
       // CPU read ahead, hence release GPU memory
-      gpu().releaseGpuMemoryFence();
+      constexpr bool ForceBarrier = true;
+      gpu().releaseGpuMemoryFence(ForceBarrier);
       char* dst = reinterpret_cast<char*>(dstMemory.owner()->getSvmPtr());
       std::memcpy(dst + origin[0], srcHost, size[0]);
       // Set HASPENDINGDISPATCH_ FLAG. Then releaseGpuMemoryFence() will use barrier to invalidate cache
       gpu().hasPendingDispatch();
-      gpu().releaseGpuMemoryFence();
+      gpu().releaseGpuMemoryFence(ForceBarrier);
       return true;
     }
   }
@@ -1825,6 +1806,7 @@ bool KernelBlitManager::writeBuffer(const void* srcHost, device::Memory& dstMemo
   return result;
 }
 
+// ================================================================================================
 bool KernelBlitManager::writeBufferRect(const void* srcHost, device::Memory& dstMemory,
                                         const amd::BufferRect& hostRect,
                                         const amd::BufferRect& bufRect, const amd::Coord3D& size,
@@ -2284,6 +2266,7 @@ address KernelBlitManager::captureArguments(const amd::Kernel* kernel) const {
 void KernelBlitManager::releaseArguments(address args) const {
 }
 
+// ================================================================================================
 bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam,
                                      hsa_queue_t* schedulerQueue,
                                      hsa_signal_t& schedulerSignal,
@@ -2310,7 +2293,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam,
   sp->child_queue = reinterpret_cast<uint64_t>(schedulerQueue);
   sp->complete_signal = schedulerSignal;
 
-  hsa_signal_store_relaxed(schedulerSignal, 1);
+  hsa_signal_store_relaxed(schedulerSignal, kInitSignalValueOne);
 
   sp->scheduler_aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
                              (1 << HSA_PACKET_HEADER_BARRIER) |
@@ -2346,8 +2329,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam,
   }
   releaseArguments(parameters);
 
-  if (hsa_signal_wait_scacquire(schedulerSignal, HSA_SIGNAL_CONDITION_LT, 1, (-1),
-                                HSA_WAIT_STATE_BLOCKED) != 0) {
+  if (!WaitForSignal(schedulerSignal)) {
     LogWarning("Failed schedulerSignal wait");
     return false;
   }
@@ -2355,6 +2337,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam,
   return true;
 }
 
+// ================================================================================================
 bool KernelBlitManager::RunGwsInit(
   uint32_t value) const {
   amd::ScopedLock k(lockXferOps_);

rocclr/device/rocm/rocdevice.cpp

@@ -718,7 +718,7 @@ bool Device::create() {
   }
 
   // Create signal for HMM prefetch operation on device
-  if (HSA_STATUS_SUCCESS != hsa_signal_create(InitSignalValue, 0, nullptr, &prefetch_signal_)) {
+  if (HSA_STATUS_SUCCESS != hsa_signal_create(kInitSignalValueOne, 0, nullptr, &prefetch_signal_)) {
     return false;
   }
 
@@ -1476,6 +1476,12 @@ bool Device::populateOCLDeviceConstants() {
 device::VirtualDevice* Device::createVirtualDevice(amd::CommandQueue* queue) {
   amd::ScopedLock lock(vgpusAccess());
 
+  // If barrier is disabled, then profiling should be enabled to make sure HSA signal is
+  // attached for every dispatch
+  if (!settings().barrier_sync_) {
+    queue->properties().set(CL_QUEUE_PROFILING_ENABLE);
+  }
+
   bool profiling = (queue != nullptr) && queue->properties().test(CL_QUEUE_PROFILING_ENABLE);
   bool cooperative = false;
 
@@ -1732,9 +1738,11 @@ device::Memory* Device::createMemory(amd::Memory& owner) const {
   return memory;
 }
 
+// ================================================================================================
 void* Device::hostAlloc(size_t size, size_t alignment, bool atomics) const {
   void* ptr = nullptr;
-  const hsa_amd_memory_pool_t segment = (!atomics)
+  // If runtime disables barrier, then all host allocations must have L2 disabled
+  const hsa_amd_memory_pool_t segment = (!atomics && settings().barrier_sync_)
       ? (system_coarse_segment_.handle != 0) ? system_coarse_segment_ : system_segment_
       : system_segment_;
   assert(segment.handle != 0);
@@ -1754,10 +1762,12 @@ void* Device::hostAlloc(size_t size, size_t alignment, bool atomics) const {
   return ptr;
 }
 
+// ================================================================================================
 void* Device::hostAgentAlloc(size_t size, const AgentInfo& agentInfo, bool atomics) const {
   void* ptr = nullptr;
   const hsa_amd_memory_pool_t segment =
-      (!atomics) ?
+      // If runtime disables barrier, then all host allocations must have L2 disabled
+      (!atomics && settings().barrier_sync_) ?
           (agentInfo.coarse_grain_pool.handle != 0) ?
               agentInfo.coarse_grain_pool : agentInfo.fine_grain_pool
           : agentInfo.fine_grain_pool;
@@ -1778,6 +1788,7 @@ void* Device::hostAgentAlloc(size_t size, const AgentInfo& agentInfo, bool atomi
   return ptr;
 }
 
+// ================================================================================================
 void* Device::hostNumaAlloc(size_t size, size_t alignment, bool atomics) const {
   void* ptr = nullptr;
 #ifndef ROCCLR_SUPPORT_NUMA_POLICY
@@ -2215,7 +2226,7 @@ bool Device::SvmAllocInit(void* memory, size_t size) const {
 
 #if AMD_HMM_SUPPORT
   // Initialize signal for the barrier
-  hsa_signal_store_relaxed(prefetch_signal_, InitSignalValue);
+  hsa_signal_store_relaxed(prefetch_signal_, kInitSignalValueOne);
 
   // Initiate a prefetch command which should force memory update in HMM
   hsa_status_t status = hsa_amd_svm_prefetch_async(memory, size, getBackendDevice(),
@@ -2226,8 +2237,7 @@ bool Device::SvmAllocInit(void* memory, size_t size) const {
   }
 
   // Wait for the prefetch
-  if (hsa_signal_wait_scacquire(prefetch_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
-                              HSA_WAIT_STATE_BLOCKED) != 0) {
+  if (!WaitForSignal(prefetch_signal_)) {
     LogError("Barrier packet submission failed");
     return false;
   }

rocclr/device/rocm/rocsettings.cpp

@@ -90,6 +90,7 @@ Settings::Settings() {
   hmmFlags_ = (!flagIsDefault(ROC_HMM_FLAGS)) ? ROC_HMM_FLAGS : Hmm::EnableSystemMemory;
 
   rocr_backend_ = true;
+  barrier_sync_ = (!flagIsDefault(ROC_BARRIER_SYNC)) ? ROC_BARRIER_SYNC : true;
 }
 
 bool Settings::create(bool fullProfile, int gfxipMajor, int gfxipMinor, bool coop_groups) {

rocclr/device/rocm/rocsettings.hpp

@@ -51,7 +51,8 @@ class Settings : public device::Settings {
       uint stagedXferRead_ : 1;         //!< Uses a staged buffer read
       uint stagedXferWrite_ : 1;        //!< Uses a staged buffer write
       uint imageBufferWar_ : 1;         //!< Image buffer workaround for Gfx10
-      uint reserved_ : 24;
+      uint barrier_sync_ : 1;           //!< Use AQL barrier command to sync with CPU
+      uint reserved_ : 23;
     };
     uint value_;
   };
@@ -82,7 +83,7 @@ class Settings : public device::Settings {
 
   size_t sdmaCopyThreshold_;  //!< Use SDMA to copy above this size
 
-  uint32_t      hmmFlags_;    //!< HMM functionality control flags
+  uint32_t  hmmFlags_;        //!< HMM functionality control flags
 
   //! Default constructor
   Settings();

rocclr/device/rocm/rocvirtual.cpp

@@ -436,10 +436,12 @@ bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address para
   return true;
 }
 
+// ================================================================================================
 static inline void packet_store_release(uint32_t* packet, uint16_t header, uint16_t rest) {
   __atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE);
 }
 
+// ================================================================================================
 template <typename AqlPacket>
 bool VirtualGPU::dispatchGenericAqlPacket(
   AqlPacket* packet, uint16_t header, uint16_t rest, bool blocking, size_t size) {
@@ -478,10 +480,19 @@ bool VirtualGPU::dispatchGenericAqlPacket(
     }
     signal = packet->completion_signal;
     // Initialize signal for a wait
-    hsa_signal_store_relaxed(signal, InitSignalValue);
+    hsa_signal_store_relaxed(signal, kInitSignalValueOne);
     blocking = true;
   }
 
+  // If runtime doesn't use the barrier, then make sure it tracks the last submitted command
+  if (!dev().settings().barrier_sync_) {
+    // Initialize signal for a wait
+    assert(packet->completion_signal.handle != 0 &&
+        "There is no HSA signal associated with the last command!");
+    hsa_signal_store_relaxed(packet->completion_signal, kInitSignalValueOne);
+    last_signal_ = packet->completion_signal;
+  }
+
   // Insert packet(s)
   // NOTE: need multiple packets to dispatch the performance counter
   //       packet blob of the legacy devices (gfx8)
@@ -521,8 +532,7 @@ bool VirtualGPU::dispatchGenericAqlPacket(
 
   // Wait on signal ?
   if (blocking) {
-    if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1),
-                                HSA_WAIT_STATE_BLOCKED) != 0) {
+    if (!WaitForSignal(signal)) {
       LogPrintfError("Failed signal [0x%lx] wait", signal.handle);
       return false;
     }
@@ -531,16 +541,19 @@ bool VirtualGPU::dispatchGenericAqlPacket(
   return true;
 }
 
+// ================================================================================================
 bool VirtualGPU::dispatchAqlPacket(
   hsa_kernel_dispatch_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) {
   return dispatchGenericAqlPacket(packet, header, rest, blocking);
 }
 
+// ================================================================================================
 bool VirtualGPU::dispatchAqlPacket(
   hsa_barrier_and_packet_t* packet, uint16_t header, uint16_t rest, bool blocking) {
   return dispatchGenericAqlPacket(packet, header, rest, blocking);
 }
 
+// ================================================================================================
 bool VirtualGPU::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) {
@@ -599,32 +612,8 @@ void VirtualGPU::dispatchBarrierPacket(const hsa_barrier_and_packet_t* packet) {
           packet->dep_signal[3], packet->dep_signal[4], packet->completion_signal);
 }
 
-/**
- * @brief Waits on an outstanding kernel without regard to how
- * it was dispatched - with or without a signal
- *
- * @return bool true if Wait returned successfully, false
- * otherwise
- */
-bool VirtualGPU::releaseGpuMemoryFence() {
-  // Return if there is no pending dispatch
-  if (!hasPendingDispatch_) {
-    return false;
-  }
-
-  // Initialize signal for the barrier packet.
-  hsa_signal_store_relaxed(barrier_signal_, InitSignalValue);
-
-  // Dispatch barrier packet into the queue and wait till it finishes.
-  dispatchBarrierPacket(&barrier_packet_);
-  if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
-                              HSA_WAIT_STATE_BLOCKED) != 0) {
-    LogError("Barrier packet submission failed");
-    return false;
-  }
-
-  hasPendingDispatch_ = false;
-
+// ================================================================================================
+void VirtualGPU::ResetQueueStates() {
   // Release all transfer buffers on this command queue
   releaseXferWrite();
 
@@ -633,10 +622,42 @@ bool VirtualGPU::releaseGpuMemoryFence() {
 
   // Release the pool, since runtime just completed a barrier
   resetKernArgPool();
+}
 
+// ================================================================================================
+bool VirtualGPU::releaseGpuMemoryFence(bool force_barrier) {
+  // Return if there is no pending dispatch
+  if (!hasPendingDispatch_) {
+    return false;
+  }
+  hsa_signal_t wait_signal = barrier_signal_;
+
+  // If barrier sync was requested or runtime didn't provide the last signal
+  if (dev().settings().barrier_sync_ || force_barrier) {
+    // Initialize signal for the barrier packet.
+    hsa_signal_store_relaxed(barrier_signal_, kInitSignalValueOne);
+
+    // Dispatch barrier packet into the queue and wait till it finishes.
+    dispatchBarrierPacket(&barrier_packet_);
+  }
+  else {
+    // Take the signal of the last submitted dispatch
+    wait_signal = last_signal_;
+  }
+
+  // Wait for compute work previously submitted
+  if (!WaitForSignal(wait_signal)) {
+    LogError("Waiting for compute work failed!");
+    return false;
+  }
+
+  hasPendingDispatch_ = false;
+
+  ResetQueueStates();
   return true;
 }
 
+// ================================================================================================
 VirtualGPU::VirtualGPU(Device& device, bool profiling, bool cooperative,
                        const std::vector<uint32_t>& cuMask,
                        amd::CommandQueue::Priority priority)
@@ -772,7 +793,7 @@ bool VirtualGPU::create() {
 
   // Create signal for the barrier packet.
   hsa_signal_t signal = {0};
-  if (HSA_STATUS_SUCCESS != hsa_signal_create(InitSignalValue, 0, nullptr, &signal)) {
+  if (HSA_STATUS_SUCCESS != hsa_signal_create(kInitSignalValueOne, 0, nullptr, &signal)) {
     return false;
   }
   barrier_signal_ = signal;
@@ -851,12 +872,11 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) {
       //! We can issue a barrier to avoid expensive extra memory allocations.
 
       // Initialize signal for the barrier packet.
-      hsa_signal_store_relaxed(barrier_signal_, InitSignalValue);
+      hsa_signal_store_relaxed(barrier_signal_, kInitSignalValueOne);
 
       // Dispatch barrier packet into the queue and wait till it finishes.
       dispatchBarrierPacket(&barrier_packet_);
-      if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
-                                  HSA_WAIT_STATE_BLOCKED) != 0) {
+      if (!WaitForSignal(barrier_signal_)) {
         LogError("Kernel arguments reset failed");
       }
 
@@ -879,7 +899,8 @@ void VirtualGPU::profilingBegin(amd::Command& command, bool drmProfiling) {
                         This could have unintended consequences.");
       return;
     }
-    timestamp_ = new Timestamp;
+    // Without barrier profiling will wait for each individual signal
+    timestamp_ = new Timestamp(!dev().settings().barrier_sync_);
     timestamp_->start();
   }
 }
@@ -1202,7 +1223,7 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) {
 void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
 #if AMD_HMM_SUPPORT
   // Initialize signal for the barrier
-  hsa_signal_store_relaxed(barrier_signal_, InitSignalValue);
+  hsa_signal_store_relaxed(barrier_signal_, kInitSignalValueOne);
 
   // Find the requested agent for the transfer
   hsa_agent_t agent = (cmd.cpu_access() ||
@@ -1214,9 +1235,7 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
       const_cast<void*>(cmd.dev_ptr()), cmd.count(), agent, 0, nullptr, barrier_signal_);
 
   // Wait for the prefetch
-  if ((status != HSA_STATUS_SUCCESS) ||
-      hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1),
-                                HSA_WAIT_STATE_BLOCKED) != 0) {
+  if ((status != HSA_STATUS_SUCCESS) || !WaitForSignal(barrier_signal_)) {
     LogError("hsa_amd_svm_prefetch_async failed");
     cmd.setStatus(CL_INVALID_OPERATION);
   }
@@ -2458,20 +2477,38 @@ void VirtualGPU::submitAcquireExtObjects(amd::AcquireExtObjectsCommand& vcmd) {
   profilingEnd(vcmd);
 }
 
+// ================================================================================================
 void VirtualGPU::submitReleaseExtObjects(amd::ReleaseExtObjectsCommand& vcmd) {
   // Make sure VirtualGPU has an exclusive access to the resources
   amd::ScopedLock lock(execution());
   profilingBegin(vcmd);
+  if (!dev().settings().barrier_sync_) {
+    // Force barrier to make sure L2 flush, since interop can be in sysmem
+    constexpr bool ForceBarrier = true;
+    releaseGpuMemoryFence(ForceBarrier);
+  }
   profilingEnd(vcmd);
 }
 
+// ================================================================================================
 void VirtualGPU::flush(amd::Command* list, bool wait) {
-  releaseGpuMemoryFence();
+  // If barrier is requested, then wait for everything, otherwise
+  // a per disaptch wait will occur later in updateCommandsState()
+  if (dev().settings().barrier_sync_) {
+    releaseGpuMemoryFence();
+  }
   updateCommandsState(list);
+
+  // Add extra clean up for resources if releaseGpuMemoryFence() was skipped
+  if (!dev().settings().barrier_sync_) {
+    ResetQueueStates();
+  }
+
   // Release all pinned memory
   releasePinnedMem();
 }
 
+// ================================================================================================
 void VirtualGPU::addXferWrite(Memory& memory) {
   if (xferWriteBuffers_.size() > 7) {
     dev().xferWrite().release(*this, *xferWriteBuffers_.front());

rocclr/device/rocm/rocvirtual.hpp

@@ -42,6 +42,26 @@ struct ProfilingSignal : public amd::HeapObject {
   ProfilingSignal() : ts_(nullptr) { signal_.handle = 0; }
 };
 
+// Initial HSA signal value
+constexpr hsa_signal_value_t kInitSignalValueOne = 1;
+
+inline bool WaitForSignal(hsa_signal_t signal) {
+  constexpr uint64_t Timeout30us = 30000;
+  constexpr uint64_t UnlimitedWait = std::numeric_limits<uint64_t>::max();
+  uint64_t timeout = (ROC_ACTIVE_WAIT) ? UnlimitedWait : Timeout30us;
+
+  // Active wait with a timeout
+  if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
+                                timeout, HSA_WAIT_STATE_ACTIVE) != 0) {
+    // Wait until the completion with CPU suspend
+    if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
+                                  UnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) {
+      return false;
+    }
+  }
+  return true;
+}
+
 // Timestamp for keeping track of some profiling information for various commands
 // including EnqueueNDRangeKernel and clEnqueueCopyBuffer.
 class Timestamp {
@@ -53,6 +73,7 @@ class Timestamp {
   static double ticksToTime_;
   bool splittedDispatch_;
   std::vector<hsa_signal_t> splittedSignals_;
+  bool wait_for_signal_;      //!< Wait for signal before gathering the timestamp values
 
  public:
   uint64_t getStart() {
@@ -75,7 +96,12 @@ class Timestamp {
 
   void setAgent(hsa_agent_t agent) { agent_ = agent; }
 
-  Timestamp() : start_(0), end_(0), profilingSignal_(nullptr), splittedDispatch_(false) {
+  Timestamp(bool wait_for_signal = false)
+    : start_(0)
+    , end_(0)
+    , profilingSignal_(nullptr)
+    , splittedDispatch_(false)
+    , wait_for_signal_(wait_for_signal) {
     agent_.handle = 0;
   }
 
@@ -90,6 +116,9 @@ class Timestamp {
         uint64_t start = UINT64_MAX;
         uint64_t end = 0;
         for (auto it = splittedSignals_.begin(); it < splittedSignals_.end(); it++) {
+          if (wait_for_signal_) {
+            WaitForSignal(*it);
+          }
           hsa_amd_profiling_get_dispatch_time(agent_, *it, &time);
           if (time.start < start) {
             start = time.start;
@@ -101,6 +130,9 @@ class Timestamp {
         start_ = start * ticksToTime_;
         end_ = end * ticksToTime_;
       } else {
+        if (wait_for_signal_) {
+          WaitForSignal(profilingSignal_->signal_);
+        }
         hsa_amd_profiling_get_dispatch_time(agent_, profilingSignal_->signal_, &time);
         start_ = time.start * ticksToTime_;
         end_ = time.end * ticksToTime_;
@@ -125,9 +157,6 @@ class Timestamp {
 
 class VirtualGPU : public device::VirtualDevice {
  public:
-  //! Initial signal value
-  static const hsa_signal_value_t InitSignalValue = 1;
-
   class MemoryDependency : public amd::EmbeddedObject {
    public:
     //! Default constructor
@@ -227,10 +256,9 @@ class VirtualGPU : public device::VirtualDevice {
    * @brief Waits on an outstanding kernel without regard to how
    * it was dispatched - with or without a signal
    *
-   * @return bool true if Wait returned successfully, false
-   * otherwise
+   * @return bool true if Wait returned successfully, false otherwise
    */
-  bool releaseGpuMemoryFence();
+  bool releaseGpuMemoryFence(bool force_barrier    = false);
 
   hsa_agent_t gpu_device() { return gpu_device_; }
   hsa_queue_t* gpu_queue() { return gpu_queue_; }
@@ -314,6 +342,9 @@ class VirtualGPU : public device::VirtualDevice {
   //! Updates AQL header for the upcomming dispatch
   void setAqlHeader(uint16_t header) { aqlHeader_ = header; }
 
+  //! Resets the current queue state. Note: should be called after AQL queue becomes idle
+  void ResetQueueStates();
+
   std::vector<Memory*> xferWriteBuffers_;  //!< Stage write buffers
   std::vector<amd::Memory*> pinnedMems_;   //!< Pinned memory list
 
@@ -336,6 +367,7 @@ class VirtualGPU : public device::VirtualDevice {
   hsa_queue_t* gpu_queue_;  //!< Queue associated with a gpu
   hsa_barrier_and_packet_t barrier_packet_;
   hsa_signal_t barrier_signal_;
+  hsa_signal_t last_signal_ = {};  //!< Last submitted signal
   uint32_t dispatch_id_;  //!< This variable must be updated atomically.
   Device& roc_device_;    //!< roc device object
   PrintfDbg* printfdbg_;

rocclr/utils/flags.hpp

@@ -236,6 +236,10 @@ release(uint, HIP_HIDDEN_FREE_MEM, 0,                                         \
         "0 = Disable")                                                        \
 release(size_t, GPU_FORCE_BLIT_COPY_SIZE, 0,                                  \
         "Size in KB of the threshold below which to force blit instead for sdma") \
+release(bool, ROC_BARRIER_SYNC, true,                                         \
+        "Enable AQL barrier packet for synchronization")                      \
+release(bool, ROC_ACTIVE_WAIT, false,                                         \
+        "Forces unconditional active wait for GPU")                           \
 release(bool, ROC_ENABLE_LARGE_BAR, true,                                     \
         "Enable Large Bar if supported by the device")                        \
 release(bool, HIP_FORCE_QUEUE_PROFILING, false,                               \