From 6405b6cdba0ec509fa71bc9fff069d15fad16249 Mon Sep 17 00:00:00 2001
From: Saleel Kudchadker <Saleel.Kudchadker@amd.com>
Date: Mon, 3 Oct 2022 09:08:40 -0700
Subject: [PATCH] SWDEV-359548 - Detect fence optimization

- Use a dirty flag to determine fence optimization
- If fence is dirty submit a marker at top level to sync.

Change-Id: I53fb19b5bb05b7c7b37c41637a6c7aaf870b639a
---
 rocclr/device/device.hpp          | 3 +++
 rocclr/device/pal/palvirtual.hpp  | 2 ++
 rocclr/device/rocm/rocvirtual.cpp | 7 ++++++-
 rocclr/device/rocm/rocvirtual.hpp | 2 ++
 4 files changed, 13 insertions(+), 1 deletion(-)

diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp
index d846bfac45..8d20b3ad17 100644
--- a/rocclr/device/device.hpp
+++ b/rocclr/device/device.hpp
@@ -1263,6 +1263,9 @@ class VirtualDevice : public amd::HeapObject {
   //! Returns the status of queue handler callback
   virtual bool isHandlerPending() const = 0;
 
+  //! Returns fence state of the VirtualGPU
+  virtual bool isFenceDirty() const = 0;
+
  private:
   //! Disable default copy constructor
   VirtualDevice& operator=(const VirtualDevice&);
diff --git a/rocclr/device/pal/palvirtual.hpp b/rocclr/device/pal/palvirtual.hpp
index 9483456b95..e648869180 100644
--- a/rocclr/device/pal/palvirtual.hpp
+++ b/rocclr/device/pal/palvirtual.hpp
@@ -360,6 +360,8 @@ class VirtualGPU : public device::VirtualDevice {
 
   bool isHandlerPending() const { return false; }
 
+  bool isFenceDirty() const { return false; }
+
   //! Returns GPU device object associated with this kernel
   const Device& dev() const { return gpuDevice_; }
 
diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp
index 25ad112466..9f0d2b9c75 100644
--- a/rocclr/device/rocm/rocvirtual.cpp
+++ b/rocclr/device/rocm/rocvirtual.cpp
@@ -825,6 +825,7 @@ bool VirtualGPU::dispatchGenericAqlPacket(
   if (fence_state_ == amd::Device::kCacheStateSystem &&
       expected_fence_state == amd::Device::kCacheStateSystem) {
     header = dispatchPacketHeader_;
+    fence_dirty_ = true;
   }
 
   fence_state_ = static_cast<Device::CacheState>(expected_fence_state);
@@ -991,6 +992,9 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal,
     barrier_packet_.completion_signal = signal;
   }
 
+  // Reset fence_dirty_ flag if we submit a barrier
+  fence_dirty_ = false;
+
   while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask);
   hsa_barrier_and_packet_t* aql_loc =
     &(reinterpret_cast<hsa_barrier_and_packet_t*>(gpu_queue_->base_address))[index & queueMask];
@@ -1074,7 +1078,8 @@ VirtualGPU::VirtualGPU(Device& device, bool profiling, bool cooperative,
       cuMask_(cuMask),
       priority_(priority),
       copy_command_type_(0),
-      fence_state_(Device::CacheState::kCacheStateInvalid)
+      fence_state_(Device::CacheState::kCacheStateInvalid),
+      fence_dirty_(false)
 {
   index_ = device.numOfVgpus_++;
   gpu_device_ = device.getBackendDevice();
diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp
index 21c5c12951..cffc4a329e 100644
--- a/rocclr/device/rocm/rocvirtual.hpp
+++ b/rocclr/device/rocm/rocvirtual.hpp
@@ -402,6 +402,7 @@ class VirtualGPU : public device::VirtualDevice {
   bool isHandlerPending() const { return barriers_.IsHandlerPending(); }
 
   void* allocKernArg(size_t size, size_t alignment);
+  bool isFenceDirty() const { return fence_dirty_; }
   // } roc OpenCL integration
  private:
   //! Dispatches a barrier with blocking HSA signals
@@ -535,5 +536,6 @@ class VirtualGPU : public device::VirtualDevice {
                                         //!< but ROC profiler expects D2H or H2D detection
   int fence_state_;                     //!< Fence scope
                                         //!< kUnknown/kFlushedToDevice/kFlushedToSystem
+  bool fence_dirty_;                    //!< Fence modified flag
 };
 }