SWDEV-301667 - Selectively queue handler

- Queue handler for hipEventRecord(aka marker_ts_) only if there is a
callback associated with it.

Change-Id: I8a9877ae0e342556053abbaacc9510744a8e772a


[ROCm/clr commit: 3c3c0ca4c5]

projects/clr/rocclr/device/device.hpp

@@ -1749,6 +1749,8 @@ class Device : public RuntimeObject {
     return false;
   };
 
+  virtual void getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const {};
+
   virtual const uint32_t getPreferredNumaNode() const { return 0; }
   virtual void ReleaseGlobalSignal(void* signal) const {}
   virtual const bool isFineGrainSupported() const {

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

@@ -2631,6 +2631,18 @@ bool Device::IsHwEventReady(const amd::Event& event, bool wait) const {
   return WaitForSignal<Timeout>(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_);
 }
 
+// ================================================================================================
+void Device::getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const {
+  void* hw_event = (event.NotifyEvent() != nullptr) ?
+    event.NotifyEvent()->HwEvent() : event.HwEvent();
+  if (hw_event == nullptr) {
+    ClPrint(amd::LOG_INFO, amd::LOG_SIG, "No HW event to read time");
+    *start = *end = 0;
+  } else {
+    fetchSignalTime(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, getBackendDevice(),
+                    start, end);
+  }
+}
 // ================================================================================================
 static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
   if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {

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

@@ -259,6 +259,7 @@ class NullDevice : public amd::Device {
       cl_set_device_clock_mode_output_amd* pSetClockModeOutput) { return true; }
 
   virtual bool IsHwEventReady(const amd::Event& event, bool wait = false) const { return false; }
+  virtual void getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const {};
   virtual void ReleaseGlobalSignal(void* signal) const {}
 
 #if defined(__clang__)
@@ -438,6 +439,7 @@ class Device : public NullDevice {
                             cl_set_device_clock_mode_output_amd* pSetClockModeOutput);
 
   virtual bool IsHwEventReady(const amd::Event& event, bool wait = false) const;
+  virtual void getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const;
   virtual void ReleaseGlobalSignal(void* signal) const;
 
   //! Allocate host memory in terms of numa policy set by user

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

@@ -413,10 +413,15 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(
     ts->retain();
     prof_signal->ts_ = ts;
     ts->AddProfilingSignal(prof_signal);
+    uint32_t init_value = kInitSignalValueOne;
+    bool enqueHandler= false;
+    enqueHandler = AMD_DIRECT_DISPATCH &&
+                   (ts->command().Callback() != nullptr ||
+                    ts->command().GetBatchHead() != nullptr )  &&
+                   (!ts->command().CpuWaitRequested());
     // If direct dispatch is enabled and the batch head isn't null, then it's a marker and
     // requires the batch update upon HSA signal completion
-    if (AMD_DIRECT_DISPATCH && (ts->command().GetBatchHead() != nullptr) &&
-        !ts->command().CpuWaitRequested()) {
+    if (enqueHandler) {
       uint32_t init_value = kInitSignalValueOne;
       // If API callback is enabled, then use a blocking signal for AQL queue.
       // HSA signal will be acquired in SW and released after HSA signal callback
@@ -434,6 +439,11 @@ hsa_signal_t VirtualGPU::HwQueueTracker::ActiveSignal(
         ClPrint(amd::LOG_INFO, amd::LOG_SIG, "Set Handler: handle(0x%lx), timestamp(%p)",
           prof_signal->signal_.handle, prof_signal);
       }
+
+      // Update the current command/marker with HW event
+      prof_signal->retain();
+      ts->command().SetHwEvent(prof_signal);
+    } else if (ts->command().profilingInfo().marker_ts_ ) {
       // Update the current command/marker with HW event
       prof_signal->retain();
       ts->command().SetHwEvent(prof_signal);
@@ -519,7 +529,7 @@ bool VirtualGPU::HwQueueTracker::CpuWaitForSignal(ProfilingSignal* signal) {
     ts->checkGpuTime();
     ts->release();
     signal->ts_ = nullptr;
-  } else if (!signal->done_) {
+  } else if (hsa_signal_load_relaxed(signal->signal_) > 0) {
     amd::ScopedLock lock(signal->LockSignalOps());
     ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Host wait on completion_signal=0x%zx",
             signal->signal_.handle);
@@ -1344,8 +1354,7 @@ void VirtualGPU::updateCommandsState(amd::Command* list) const {
     while (current != nullptr) {
       if (current->data() != nullptr) {
         ts = reinterpret_cast<Timestamp*>(current->data());
-        startTimeStamp = ts->getStart();
-        endTimeStamp = ts->getStart();
+        ts->getTime(&startTimeStamp, &endTimeStamp);
         break;
       }
       current = current->getNext();
@@ -1370,8 +1379,7 @@ void VirtualGPU::updateCommandsState(amd::Command* list) const {
         // Since this is a valid command to get a timestamp, we use the
         // timestamp provided by the runtime (saved in the data())
         ts = reinterpret_cast<Timestamp*>(current->data());
-        startTimeStamp = ts->getStart();
-        endTimeStamp = ts->getEnd();
+        ts->getTime(&startTimeStamp, &endTimeStamp);
         ts->release();
         current->setData(nullptr);
       } else {
@@ -1397,6 +1405,8 @@ void VirtualGPU::updateCommandsState(amd::Command* list) const {
   }
 }
 
+// ================================================================================================
+
 void VirtualGPU::submitReadMemory(amd::ReadMemoryCommand& cmd) {
   // Make sure VirtualGPU has an exclusive access to the resources
   amd::ScopedLock lock(execution());

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

@@ -81,6 +81,16 @@ inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false) {
   return true;
 }
 
+inline void fetchSignalTime(hsa_signal_t signal, hsa_agent_t gpu_device,
+                            uint64_t* start, uint64_t* end) {
+  if (start != nullptr && end != nullptr) {
+    hsa_amd_profiling_dispatch_time_t time = {};
+    hsa_amd_profiling_get_dispatch_time(gpu_device, signal, &time);
+    *start = time.start;
+    *end = time.end;
+  }
+}
+
 // Timestamp for keeping track of some profiling information for various commands
 // including EnqueueNDRangeKernel and clEnqueueCopyBuffer.
 class Timestamp : public amd::ReferenceCountedObject {
@@ -111,14 +121,10 @@ class Timestamp : public amd::ReferenceCountedObject {
 
   ~Timestamp() {}
 
-  uint64_t getStart() {
+  uint64_t getTime(uint64_t* start, uint64_t* end) {
     checkGpuTime();
-    return start_;
-  }
-
-  uint64_t getEnd() {
-    checkGpuTime();
-    return end_;
+    *start = start_;
+    *end = end_;
   }
 
   void AddProfilingSignal(ProfilingSignal* signal) { signals_.push_back(signal); }

projects/clr/rocclr/platform/command.cpp

@@ -353,9 +353,14 @@ void Command::enqueue() {
     // when multiple threads submit/flush/update the batch at the same time
     ScopedLock sl(queue_->vdev()->execution());
     queue_->FormSubmissionBatch(this);
-    if ((type() == CL_COMMAND_MARKER || type() == 0)) {
+
+    bool isMarker = (type() == CL_COMMAND_MARKER || type() == 0);
+    if (isMarker) {
       // The current HSA signal tracking logic requires profiling enabled for the markers
       EnableProfiling();
+    }
+
+    if (isMarker && !profilingInfo().marker_ts_) {
       // Update batch head for the current marker. Hence the status of all commands can be
       // updated upon the marker completion
       SetBatchHead(queue_->GetSubmittionBatch());