From de2b06a7a72eae8689a2f3194fcc901b53132171 Mon Sep 17 00:00:00 2001
From: Saleel Kudchadker <Saleel.Kudchadker@amd.com>
Date: Fri, 20 Oct 2023 23:51:48 +0000
Subject: [PATCH] SWDEV-301667 - Reset addSystemScope_

If we submit a systemScope Barrier, we should reset this flag as there
is no need for dispatch AQL again to flush caches/HDP

Change-Id: I55710feb4ba6650852e785b5cadfa64c6b9ce14e
---
 rocclr/device/rocm/rocvirtual.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp
index eb8e66b74a..5ab54a02d9 100644
--- a/rocclr/device/rocm/rocvirtual.cpp
+++ b/rocclr/device/rocm/rocvirtual.cpp
@@ -1014,9 +1014,10 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal,
     barrier_packet_.completion_signal = signal;
   }
 
-  // Reset fence_dirty_ flag if we submit a barrier
+  // Reset fence_dirty_ and addSystemScope_ flag if we submit a barrier with system scopes
   if (cache_state == amd::Device::kCacheStateSystem) {
     fence_dirty_ = false;
+    addSystemScope_ = false;
   }
 
   while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask);