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


[ROCm/clr commit: 6a5af4056e]

projects/clr/rocclr/device/rocm/rocblit.cpp

@@ -661,7 +661,7 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory,
       LogError("Async copy failed");
       status = HSA_STATUS_ERROR;
     } else {
-      gpu().hasPendingDispatch();
+      gpu().addSystemScope();
     }
   } else {
     LogPrintfError("Hsa copy from host to device failed with code %d", status);
@@ -749,7 +749,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_
     offset += size;
   }
 
-  gpu().hasPendingDispatch();
+  gpu().addSystemScope();
 
   return true;
 }

projects/clr/rocclr/device/rocm/rocvirtual.cpp

@@ -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

projects/clr/rocclr/device/rocm/rocvirtual.hpp

@@ -261,7 +261,8 @@ class VirtualGPU : public device::VirtualDevice {
 
   void enableSyncBlit() const;
 
-  void hasPendingDispatch() { hasPendingDispatch_ = true;}
+  void hasPendingDispatch() { hasPendingDispatch_ = true; }
+  void addSystemScope() { addSystemScope_ = true; }
 
   // } roc OpenCL integration
  private:
@@ -319,6 +320,7 @@ class VirtualGPU : public device::VirtualDevice {
       uint32_t imageBufferWrtBack_ : 1; //!< Image buffer write back is required
       uint32_t profiling_          : 1; //!< Profiling is enabled
       uint32_t cooperative_        : 1; //!< Cooperative launch is enabled
+      uint32_t addSystemScope_     : 1; //!< Insert a system scope to the next aql
     };
     uint32_t  state_;
   };