From 3c3c0ca4c53d01b83784c78a439cd1e8f114b95e Mon Sep 17 00:00:00 2001 From: Saleel Kudchadker Date: Tue, 8 Mar 2022 14:15:37 -0800 Subject: [PATCH] SWDEV-301667 - Selectively queue handler - Queue handler for hipEventRecord(aka marker_ts_) only if there is a callback associated with it. Change-Id: I8a9877ae0e342556053abbaacc9510744a8e772a --- rocclr/device/device.hpp | 2 ++ rocclr/device/rocm/rocdevice.cpp | 12 ++++++++++++ rocclr/device/rocm/rocdevice.hpp | 2 ++ rocclr/device/rocm/rocvirtual.cpp | 24 +++++++++++++++++------- rocclr/device/rocm/rocvirtual.hpp | 20 +++++++++++++------- rocclr/platform/command.cpp | 7 ++++++- 6 files changed, 52 insertions(+), 15 deletions(-) diff --git a/rocclr/device/device.hpp b/rocclr/device/device.hpp index 71721689eb..4908a58b98 100644 --- a/rocclr/device/device.hpp +++ b/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 { diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index de73e7e9c7..faaa62d037 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -2631,6 +2631,18 @@ bool Device::IsHwEventReady(const amd::Event& event, bool wait) const { return WaitForSignal(reinterpret_cast(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(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) { diff --git a/rocclr/device/rocm/rocdevice.hpp b/rocclr/device/rocm/rocdevice.hpp index cbaa952ee7..5294d7780f 100644 --- a/rocclr/device/rocm/rocdevice.hpp +++ b/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 diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index c23452612e..7dc7d525fc 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/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(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(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()); diff --git a/rocclr/device/rocm/rocvirtual.hpp b/rocclr/device/rocm/rocvirtual.hpp index b780052ac4..1610b4d7ef 100644 --- a/rocclr/device/rocm/rocvirtual.hpp +++ b/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); } diff --git a/rocclr/platform/command.cpp b/rocclr/platform/command.cpp index 4ce6fa776e..eec70f483d 100644 --- a/rocclr/platform/command.cpp +++ b/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());