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

@@ -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&);

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_; }
 

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();

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
 };
 }