From a9e75064eda44dc1073c6123fe48d2bae3a892a9 Mon Sep 17 00:00:00 2001 From: foreman Date: Wed, 21 Feb 2018 17:35:38 -0500 Subject: [PATCH] P4 to Git Change 1517959 by gandryey@gera-lnx-rcf on 2018/02/21 17:24:16 SWDEV-79445 - Add RGP trace capture capability into runtime. - Initial implementation - PAL_RGP_DISP_COUNT controls the number of captured dispatches(default 10). - RD panel and service are required in order to capture the traces and RD Profiler for viewing and SQTT exports Affected files ... ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/build/Makefile.pal#17 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldefs.hpp#31 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#72 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.hpp#23 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palgpuopen.cpp#1 add ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palgpuopen.hpp#1 add ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palsettings.cpp#43 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palsettings.hpp#15 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#73 edit ... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.hpp#41 edit ... //depot/stg/opencl/drivers/opencl/runtime/utils/flags.hpp#283 edit [ROCm/clr commit: dfb5e1b8449c4f4a654b32dea8eeaec574fa1708] --- .../clr/rocclr/runtime/device/pal/paldefs.hpp | 8 +- .../rocclr/runtime/device/pal/paldevice.cpp | 12 +- .../rocclr/runtime/device/pal/paldevice.hpp | 4 + .../rocclr/runtime/device/pal/palgpuopen.cpp | 623 ++++++++++++++++++ .../rocclr/runtime/device/pal/palgpuopen.hpp | 172 +++++ .../rocclr/runtime/device/pal/palsettings.cpp | 9 +- .../rocclr/runtime/device/pal/palsettings.hpp | 5 +- .../rocclr/runtime/device/pal/palvirtual.cpp | 497 ++++++++------ .../rocclr/runtime/device/pal/palvirtual.hpp | 24 +- projects/clr/rocclr/runtime/utils/flags.hpp | 2 + 10 files changed, 1132 insertions(+), 224 deletions(-) create mode 100644 projects/clr/rocclr/runtime/device/pal/palgpuopen.cpp create mode 100644 projects/clr/rocclr/runtime/device/pal/palgpuopen.hpp diff --git a/projects/clr/rocclr/runtime/device/pal/paldefs.hpp b/projects/clr/rocclr/runtime/device/pal/paldefs.hpp index 2247f429e1..69d79d13ed 100644 --- a/projects/clr/rocclr/runtime/device/pal/paldefs.hpp +++ b/projects/clr/rocclr/runtime/device/pal/paldefs.hpp @@ -170,7 +170,13 @@ static const AMDDeviceInfo Gfx9PlusSubDeviceInfo[] = { }; // Supported OpenCL versions -enum OclVersion { OpenCL10, OpenCL11, OpenCL12, OpenCL20, OpenCL21 }; +enum OclVersion { + OpenCL10 = 0x10, + OpenCL11 = 0x11, + OpenCL12 = 0x12, + OpenCL20 = 0x20, + OpenCL21 = 0x21 +}; struct MemoryFormat { cl_image_format clFormat_; //!< CL image format diff --git a/projects/clr/rocclr/runtime/device/pal/paldevice.cpp b/projects/clr/rocclr/runtime/device/pal/paldevice.cpp index a4a5fe35c8..7621350b92 100644 --- a/projects/clr/rocclr/runtime/device/pal/paldevice.cpp +++ b/projects/clr/rocclr/runtime/device/pal/paldevice.cpp @@ -659,7 +659,8 @@ Device::Device() globalScratchBuf_(nullptr), srdManager_(nullptr), lockResourceOps_(nullptr), - resourceList_(nullptr) + resourceList_(nullptr), + rgpCaptureMgr_(nullptr) {} Device::~Device() { @@ -717,6 +718,9 @@ Device::~Device() { } device_ = nullptr; + + // Delete developer driver manager + delete rgpCaptureMgr_; } extern const char* SchedulerSourceCode; @@ -914,6 +918,8 @@ bool Device::create(Pal::IDevice* device) { return true; } +static Pal::IPlatform* platform; + bool Device::initializeHeapResources() { amd::ScopedLock k(lockForInitHeap_); if (!heapInitComplete_) { @@ -983,6 +989,9 @@ bool Device::initializeHeapResources() { return false; } xferQueue_->enableSyncedBlit(); + + // Create RGP capture manager + rgpCaptureMgr_ = RgpCaptureMgr::Create(platform, *this); } return true; } @@ -1081,7 +1090,6 @@ static int reportHook(int reportType, char* message, int* returnValue) { #endif // _WIN32 & DEBUG static char* platformObj; -static Pal::IPlatform* platform; bool Device::init() { uint32_t numDevices = 0; diff --git a/projects/clr/rocclr/runtime/device/pal/paldevice.hpp b/projects/clr/rocclr/runtime/device/pal/paldevice.hpp index f972afe72e..b92ea73456 100644 --- a/projects/clr/rocclr/runtime/device/pal/paldevice.hpp +++ b/projects/clr/rocclr/runtime/device/pal/paldevice.hpp @@ -18,6 +18,7 @@ #include "device/pal/paldefs.hpp" #include "device/pal/palsettings.hpp" #include "device/pal/palappprofile.hpp" +#include "device/pal/palgpuopen.hpp" #include "acl.h" #include "memory" @@ -469,6 +470,8 @@ class Device : public NullDevice { //! Return private device context for internal allocations amd::Context& context() const { return *context_; } + RgpCaptureMgr* rgpCaptureMgr() const { return rgpCaptureMgr_; } + //! Update free memory for OCL extension void updateFreeMemory(Pal::GpuHeap heap, //!< PAL GPU heap for update Pal::gpusize size, //!< Size of alocated/destroyed memory @@ -590,6 +593,7 @@ class Device : public NullDevice { std::atomic freeMem[Pal::GpuHeap::GpuHeapCount]; //!< Free memory counter amd::Monitor* lockResourceOps_; //!< Lock to serialise resource access std::list* resourceList_; //!< Active resource list + RgpCaptureMgr* rgpCaptureMgr_; //!< RGP capture manager }; /*@}*/} // namespace pal diff --git a/projects/clr/rocclr/runtime/device/pal/palgpuopen.cpp b/projects/clr/rocclr/runtime/device/pal/palgpuopen.cpp new file mode 100644 index 0000000000..8d5b89da99 --- /dev/null +++ b/projects/clr/rocclr/runtime/device/pal/palgpuopen.cpp @@ -0,0 +1,623 @@ +/* + ************************************************************************************************** + * + * Trade secret of Advanced Micro Devices, Inc. + * Copyright (c) 2016, Advanced Micro Devices, Inc., (unpublished) + * + * All rights reserved. This notice is intended as a precaution against inadvertent publication + * and does not imply publication or any waiver of confidentiality. The year included in + * the foregoing notice is the year of creation of the work. + * + ************************************************************************************************** + */ +#include "device/pal/palgpuopen.hpp" +#include "device/pal/paldevice.hpp" +#include "device/pal/palvirtual.hpp" + +// PAL headers +#include "palCmdAllocator.h" +#include "palFence.h" +#include "palQueueSemaphore.h" + +// gpuutil headers +#include "gpuUtil/palGpaSession.h" + +// gpuopen headers +#include "devDriverServer.h" +#include "msgChannel.h" +#include "msgTransport.h" +#include "protocols/rgpServer.h" +#include "protocols/driverControlServer.h" + +namespace pal +{ +// ================================================================================================ +RgpCaptureMgr::RgpCaptureMgr(Pal::IPlatform* platform, const Device& device) + : + device_(device), + dev_driver_server_(platform->GetDevDriverServer()), + num_prep_disp_(0), + trace_gpu_mem_limit_(0), + global_disp_count_(1), // Must start from 1 according to RGP spec + trace_enabled_(false), + inst_tracing_enabled_(false) +{ + memset(&trace_, 0, sizeof(trace_)); +} + +// ================================================================================================ +RgpCaptureMgr::~RgpCaptureMgr() +{ + DestroyRGPTracing(); +} + +// ================================================================================================ +// Creates the GPU Open Developer Mode manager class. +RgpCaptureMgr* RgpCaptureMgr::Create(Pal::IPlatform* platform, const Device& device) +{ + RgpCaptureMgr* mgr = new RgpCaptureMgr(platform, device); + + if (mgr != nullptr && !mgr->Init(platform)) { + delete mgr; + mgr = nullptr; + } + + return mgr; +} + +// ================================================================================================ +bool RgpCaptureMgr::Init(Pal::IPlatform* platform) +{ + if (dev_driver_server_ == nullptr) { + return false; + } + const Settings& settings = device_.settings(); + // Tell RGP that the server (i.e. the driver) supports tracing if requested. + rgp_server_ = dev_driver_server_->GetRGPServer(); + if (rgp_server_ == nullptr) { + return false; + } + + // Finalize RGP settings + Finalize(); + + bool result = true; + + // Fail initialization of trace resources if SQTT tracing has been force-disabled from + // the panel (this will consequently fail the trace), or if the chosen device's gfxip + // does not support SQTT. + // + // It's necessary to check this during RGP tracing init in addition to devmode init because + // during the earlier devmode init we may be in a situation where some enumerated physical + // devices support tracing and others do not. + if (GpuSupportsTracing(device_.properties(), settings) == false) { + result = false; + } + + // Create a GPA session object for this trace session + if (result) { + assert(trace_.gpa_session_ == nullptr); + + const uint32_t api_version = settings.oclVersion_; + + trace_.gpa_session_ = new GpuUtil::GpaSession( + platform, + device_.iDev(), + api_version >> 4, // OCL API version major + api_version & 0xf, // OCL API version minor + RgpSqttInstrumentationSpecVersion, + RgpSqttInstrumentationApiVersion); + + if (trace_.gpa_session_ == nullptr) { + result = false; + } + } + + // Initialize the GPA session + if (result && (trace_.gpa_session_->Init() != Pal::Result::Success)) { + result = false; + } + + // Initialize trace resources required by each queue (and queue family) + bool hasDebugVmid = true; + + if (result) { + //result = InitTraceQueueResources(trace_, &hasDebugVmid); + } + + // If we've failed to acquire the debug VMID, fail to trace + if (hasDebugVmid == false) { + result = false; + } + + if (!result) { + // If we've failed to initialize tracing, permanently disable traces + if (rgp_server_ != nullptr) { + rgp_server_->DisableTraces(); + + trace_enabled_ = false; + } + + // Clean up if we failed + DestroyRGPTracing(); + } else { + PostDeviceCreate(); + } + + return result; +} + +// ================================================================================================ +// Called during initial device enumeration prior to calling Pal::IDevice::CommitSettingsAndInit(). +// +// This finalizes the developer driver manager. +void RgpCaptureMgr::Finalize() +{ + // Figure out if the gfxip supports tracing. We decide tracing if there is at least one + // enumerated GPU that can support tracing. Since we don't yet know if that GPU will be + // picked as the target of an eventual VkDevice, this check is imperfect. + // In mixed-GPU situations where an unsupported GPU is picked for tracing, + // trace capture will fail with an error. + bool hw_support_tracing = false; + + if ((rgp_server_->EnableTraces() == DevDriver::Result::Success)) { + if (GpuSupportsTracing(device_.properties(), device_.settings())) { + hw_support_tracing = true; + } + } + + if (hw_support_tracing == false) { + rgp_server_->DisableTraces(); + } + + // Finalize the devmode manager + dev_driver_server_->Finalize(); + + // Figure out if tracing support should be enabled or not + trace_enabled_ = (rgp_server_ != nullptr) && rgp_server_->TracesEnabled(); +} + + +// ================================================================================================ +// Waits for the driver to be resumed if it's currently paused. +void RgpCaptureMgr::WaitForDriverResume() +{ + auto* pDriverControlServer = dev_driver_server_->GetDriverControlServer(); + + assert(pDriverControlServer != nullptr); + + pDriverControlServer->WaitForDriverResume(); +} + +// ================================================================================================ +// Called before a swap chain presents. This signals a frame-end boundary and +// is used to coordinate RGP trace start/stop. +void RgpCaptureMgr::PostDispatch(VirtualGPU* pQueue) +{ + if (rgp_server_->TracesEnabled()) { + // If there's currently a trace running, submit the trace-end command buffer + if (trace_.status_ == TraceStatus::Running) { + amd::ScopedLock traceLock(&trace_mutex_); + trace_.sqtt_disp_count_++; + if (trace_.sqtt_disp_count_ >= device_.settings().rgpSqttDispCount_) { + if (EndRGPHardwareTrace(pQueue) != Pal::Result::Success) { + FinishRGPTrace(true); + } + } + } + + if (IsQueueTimingActive()) { + // Call TimedQueuePresent() to insert commands that collect GPU timestamp. + Pal::IQueue* pPalQueue = pQueue->queue(MainEngine).iQueue_; + + // Currently nothing in the PresentInfo struct is used for inserting a timed present marker. + GpuUtil::TimedQueuePresentInfo timedPresentInfo = {}; + //Pal::Result result = trace_.gpa_session_->TimedQueuePresent(pPalQueue, timedPresentInfo); + //assert(result == Pal::Result::Success); + } + } +} + +// ================================================================================================ +Pal::Result RgpCaptureMgr::CheckForTraceResults() +{ + assert(trace_.status_ == TraceStatus::WaitingForResults); + + Pal::Result result = Pal::Result::NotReady; + + // Check if trace results are ready + if (trace_.gpa_session_->IsReady() && // GPA session is ready + (trace_.trace_begin_queue_->isDone(&trace_.end_event_))) // "Trace end" cmdbuf has retired + { + bool success = false; + + // Fetch required trace data size from GPA session + size_t traceDataSize = 0; + void* pTraceData = nullptr; + + trace_.gpa_session_->GetResults(trace_.gpa_sample_id_, &traceDataSize, nullptr); + + // Allocate memory for trace data + if (traceDataSize > 0) { + pTraceData = amd::AlignedMemory::allocate(traceDataSize, 256); + } + + if (pTraceData != nullptr) { + // Get trace data from GPA session + if (trace_.gpa_session_->GetResults(trace_.gpa_sample_id_, &traceDataSize, pTraceData) == + Pal::Result::Success) { + // Transmit trace data to anyone who's listening + auto devResult = rgp_server_->WriteTraceData( + static_cast(pTraceData), traceDataSize); + + success = (devResult == DevDriver::Result::Success); + } + + amd::AlignedMemory::deallocate(pTraceData); + } + + if (success) { + result = Pal::Result::Success; + } + } + + return result; +} + +// ================================================================================================ +// Called after a swap chain presents. This signals a (next) frame-begin boundary and is +// used to coordinate RGP trace start/stop. +void RgpCaptureMgr::PreDispatch(VirtualGPU* pQueue) +{ + // Wait for the driver to be resumed in case it's been paused. + WaitForDriverResume(); + + if (rgp_server_->TracesEnabled()) { + amd::ScopedLock traceLock(&trace_mutex_); + + // Check if there's an RGP trace request pending and we're idle + if ((trace_.status_ == TraceStatus::Idle) && rgp_server_->IsTracePending()) { + // Attempt to start preparing for a trace + if (PrepareRGPTrace(pQueue) == Pal::Result::Success) { + // Attempt to start the trace immediately if we do not need to prepare + if (num_prep_disp_ == 0) { + if (BeginRGPTrace(pQueue) != Pal::Result::Success) { + FinishRGPTrace(true); + } + } + } + } + else if (trace_.status_ == TraceStatus::Preparing) { + // Wait some number of "preparation frames" before starting the trace in order to get enough + // timer samples to sync CPU/GPU clock domains. + trace_.prepared_disp_count_++; + + // Take a calibration timing measurement sample for this frame. + trace_.gpa_session_->SampleTimingClocks(); + + // Start the SQTT trace if we've waited a sufficient number of preparation frames + if (trace_.prepared_disp_count_ >= num_prep_disp_) { + Pal::Result result = BeginRGPTrace(pQueue); + + if (result != Pal::Result::Success) { + FinishRGPTrace(true); + } + } + } + // Check if we're ending a trace waiting for SQTT to turn off. + // If SQTT has turned off, end the trace + else if (trace_.status_ == TraceStatus::WaitingForSqtt) { + Pal::Result result = Pal::Result::Success; + + if (trace_.trace_begin_queue_->isDone(&trace_.end_sqtt_event_)) { + result = EndRGPTrace(pQueue); + } else { + // todo: There is a wait inside the trace end for now + result = EndRGPTrace(pQueue); + } + + if (result != Pal::Result::Success) { + FinishRGPTrace(true); + } + } + // Check if we're waiting for final trace results. + else if (trace_.status_ == TraceStatus::WaitingForResults) { + Pal::Result result = CheckForTraceResults(); + + // Results ready: finish trace + if (result == Pal::Result::Success) { + FinishRGPTrace(false); + } + // Error while computing results: abort trace + else if (result != Pal::Result::NotReady) { + FinishRGPTrace(true); + } + } + } + + global_disp_count_++; +} + +// ================================================================================================ +// This function starts preparing for an RGP trace. Preparation involves some N frames of +// lead-up time during which timing samples are accumulated to synchronize CPU and GPU clock domains. +// +// This function transitions from the Idle state to the Preparing state. +Pal::Result RgpCaptureMgr::PrepareRGPTrace(VirtualGPU* pQueue) +{ + assert(trace_.status_ == TraceStatus::Idle); + + // We can only trace using a single device at a time currently, so recreate RGP trace + // resources against this new one if the device is changing. + Pal::Result result = Pal::Result::Success; + + const auto traceParameters = rgp_server_->QueryTraceParameters(); + + num_prep_disp_ = traceParameters.numPreparationFrames; + trace_gpu_mem_limit_ = traceParameters.gpuMemoryLimitInMb * 1024 * 1024; + inst_tracing_enabled_ = traceParameters.flags.enableInstructionTokens; + + // Notify the RGP server that we are starting a trace + if (rgp_server_->BeginTrace() != DevDriver::Result::Success) { + result = Pal::Result::ErrorUnknown; + } + + // Tell the GPA session class we're starting a trace + if (result == Pal::Result::Success) { + GpuUtil::GpaSessionBeginInfo info = {}; + + info.flags.enableQueueTiming = false;// trace_.queueTimingEnabled; + + result = trace_.gpa_session_->Begin(info); + } + + trace_.prepared_disp_count_ = 0; + trace_.sqtt_disp_count_ = 0; + + // Sample the timing clocks prior to starting a trace. + if (result == Pal::Result::Success) { + trace_.gpa_session_->SampleTimingClocks(); + } + + if (result == Pal::Result::Success) { + // Remember which queue started the trace + trace_.trace_prepare_queue_ = pQueue; + trace_.trace_begin_queue_ = nullptr; + + trace_.status_ = TraceStatus::Preparing; + } else { + // We failed to prepare for the trace so abort it. + if (rgp_server_ != nullptr) { + const DevDriver::Result devDriverResult = rgp_server_->AbortTrace(); + + // AbortTrace should always succeed unless we've used the api incorrectly. + assert(devDriverResult == DevDriver::Result::Success); + } + } + + return result; +} + +// ================================================================================================ +// This function begins an RGP trace by initializing all dependent resources and submitting +// the "begin trace" information command buffer. +// +// This function transitions from the Preparing state to the Running state. +Pal::Result RgpCaptureMgr::BeginRGPTrace(VirtualGPU* pQueue) +{ + assert(trace_.status_ == TraceStatus::Preparing); + assert(trace_enabled_); + + // We can only trace using a single device at a time currently, so recreate RGP trace + // resources against this new one if the device is changing. + Pal::Result result = Pal::Result::Success; + + if (result == Pal::Result::Success) { + // Only allow trace to start if the queue family at prep-time matches the queue + // family at begin time because the command buffer engine type must match + if (trace_.trace_prepare_queue_ != pQueue) { + result = Pal::Result::ErrorIncompatibleQueue; + } + } + + // Start a GPA tracing sample with SQTT enabled + if (result == Pal::Result::Success) { + GpuUtil::GpaSampleConfig sampleConfig = {}; + + sampleConfig.type = GpuUtil::GpaSampleType::Trace; + sampleConfig.sqtt.gpuMemoryLimit = trace_gpu_mem_limit_; + sampleConfig.sqtt.flags.enable = true; + sampleConfig.sqtt.flags.supressInstructionTokens = (inst_tracing_enabled_ == false); + + // Fill GPU commands + pQueue->eventBegin(MainEngine); + trace_.gpa_sample_id_ = trace_.gpa_session_->BeginSample( + pQueue->queue(MainEngine).iCmd(), sampleConfig); + pQueue->eventEnd(MainEngine, trace_.begin_sqtt_event_); + } + + // Submit the trace-begin command buffer + if (result == Pal::Result::Success) { + static constexpr bool NeedFlush = true; + // Update the global GPU event + pQueue->setGpuEvent(trace_.begin_sqtt_event_, NeedFlush); + } + + // Make the trace active and remember which queue started it + if (result == Pal::Result::Success) { + trace_.status_ = TraceStatus::Running; + trace_.trace_begin_queue_ = pQueue; + } + + return result; +} + +// ================================================================================================ +// This function submits the command buffer to stop SQTT tracing. Full tracing still continues. +// +// This function transitions from the Running state to the WaitingForSqtt state. +Pal::Result RgpCaptureMgr::EndRGPHardwareTrace(VirtualGPU* pQueue) +{ + assert(trace_.status_ == TraceStatus::Running); + + Pal::Result result = Pal::Result::Success; + + // Only allow SQTT trace to start and end on the same queue because it's critical that these are + // in the same order + if (pQueue != trace_.trace_begin_queue_) { + result = Pal::Result::ErrorIncompatibleQueue; + } + + // Tell the GPA session to insert any necessary commands to end the tracing sample and + // end the session itself + if (result == Pal::Result::Success) { + assert(trace_.gpa_session_ != nullptr); + + // Write CB commands to finish the SQTT + pQueue->eventBegin(MainEngine); + trace_.gpa_session_->EndSample(pQueue->queue(MainEngine).iCmd(), trace_.gpa_sample_id_); + pQueue->eventEnd(MainEngine, trace_.end_sqtt_event_); + + static constexpr bool NeedFlush = true; + // Update the global GPU event + pQueue->setGpuEvent(trace_.end_sqtt_event_, NeedFlush); + + trace_.status_ = TraceStatus::WaitingForSqtt; + + // Execute a device wait idle + if (device_.settings().rgpSqttWaitIdle_) { + // Make sure the trace is done. Note: required for SDMA data write back + pQueue->waitForEvent(&trace_.end_sqtt_event_); + } + } + + return result; +} + +// ================================================================================================ +// This function ends a running RGP trace. +// +// This function transitions from the WaitingForSqtt state to WaitingForResults state. +Pal::Result RgpCaptureMgr::EndRGPTrace(VirtualGPU* pQueue) +{ + assert(trace_.status_ == TraceStatus::WaitingForSqtt); + + Pal::Result result = Pal::Result::Success; + + // Tell the GPA session to insert any necessary commands to end the tracing sample and + // end the session itself + if (result == Pal::Result::Success) { + assert(trace_.gpa_session_ != nullptr); + // Initiate SDMA copy + pQueue->eventBegin(SdmaEngine); + result = trace_.gpa_session_->End(pQueue->queue(SdmaEngine).iCmd()); + pQueue->eventEnd(SdmaEngine, trace_.end_event_); + } + + // Submit the trace-end command buffer + if (result == Pal::Result::Success) { + static constexpr bool NeedFlush = true; + // Update the global GPU event + pQueue->setGpuEvent(trace_.end_event_, NeedFlush); + + trace_.status_ = TraceStatus::WaitingForResults; + + if (device_.settings().rgpSqttWaitIdle_) { + // Make sure the transfer is done + pQueue->waitForEvent(&trace_.end_event_); + } + } + + return result; +} + +// ================================================================================================ +// This function resets and possibly cancels a currently active (between begin/end) RGP trace. +// It frees any dependent resources. +void RgpCaptureMgr::FinishRGPTrace(bool aborted) +{ + if (trace_.trace_prepare_queue_ == nullptr) { + return; + } + + // Inform RGP protocol that we're done with the trace, either by aborting it or finishing normally + if (aborted) { + rgp_server_->AbortTrace(); + } else { + rgp_server_->EndTrace(); + } + + if (trace_.gpa_session_ != nullptr) { + trace_.gpa_session_->Reset(); + } + + // Reset tracing state to idle + trace_.prepared_disp_count_ = 0; + trace_.sqtt_disp_count_ = 0; + trace_.gpa_sample_id_ = 0; + trace_.status_ = TraceStatus::Idle; + trace_.trace_prepare_queue_ = nullptr; + trace_.trace_begin_queue_ = nullptr; +} + +// ================================================================================================ +// Destroys device-persistent RGP resources +void RgpCaptureMgr::DestroyRGPTracing() +{ + if (trace_.status_ != TraceStatus::Idle) { + FinishRGPTrace(true); + } + + // Destroy the GPA session + if (trace_.gpa_session_ != nullptr) { + //Util::Destructor(trace_.gpa_session_); + delete trace_.gpa_session_; + trace_.gpa_session_ = nullptr; + } + + memset(&trace_, 0, sizeof(trace_)); +} + +// ================================================================================================ +// Returns true if the given device properties/settings support tracing. +bool RgpCaptureMgr::GpuSupportsTracing( + const Pal::DeviceProperties& props, + const Settings& settings) +{ + return props.gfxipProperties.flags.supportRgpTraces && !settings.rgpSqttForceDisable_; +} + +// ================================================================================================ +// Called when a new device is created. This will preallocate reusable RGP trace resources +// for that device. +void RgpCaptureMgr::PostDeviceCreate() +{ + amd::ScopedLock traceLock(&trace_mutex_); + + auto* pDriverControlServer = dev_driver_server_->GetDriverControlServer(); + + assert(pDriverControlServer != nullptr); + + // If the driver hasn't been marked as fully initialized yet, mark it now. + // We consider the time after the logical device creation to be the fully initialized driver + // position. This is mainly because PAL is fully initialized at this point and we also know + // whether or not the debug vmid has been acquired. External tools use this information to + // decide when it's reasonable to make certain requests of the driver through protocol functions. + if (pDriverControlServer->IsDriverInitialized() == false) { + pDriverControlServer->FinishDriverInitialization(); + } +} + +// ================================================================================================ +// Called prior to a device's being destroyed. This will free persistent RGP trace resources for +// that device. +void RgpCaptureMgr::PreDeviceDestroy() +{ + amd::ScopedLock traceLock(&trace_mutex_); + // If we are idle, we can re-initialize trace resources based on the new device. + if (trace_.status_ == TraceStatus::Idle) { + DestroyRGPTracing(); + } +} + +}; // namespace vk diff --git a/projects/clr/rocclr/runtime/device/pal/palgpuopen.hpp b/projects/clr/rocclr/runtime/device/pal/palgpuopen.hpp new file mode 100644 index 0000000000..0b3889763f --- /dev/null +++ b/projects/clr/rocclr/runtime/device/pal/palgpuopen.hpp @@ -0,0 +1,172 @@ +/* +*********************************************************************************************************************** +* +* Trade secret of Advanced Micro Devices, Inc. +* Copyright (c) 2016 Advanced Micro Devices, Inc. (unpublished) +* +* All rights reserved. This notice is intended as a precaution against inadvertent publication and +* does not imply publication or any waiver of confidentiality. The year included in the foregoing +* notice is the year of creation of the work. +* +*********************************************************************************************************************** +*/ +#pragma once + +#include +#include "device/pal/paldefs.hpp" +#include "platform/commandqueue.hpp" +#include "protocols/rgpServer.h" +#include "device/blit.hpp" + +// PAL headers +#include "palUtil.h" +#include "palPlatform.h" +#include "palCmdBuffer.h" +#include "palCmdAllocator.h" +#include "palQueue.h" +#include "palFence.h" +#include "palLinearAllocator.h" +#include "palHashMap.h" +#include "palQueue.h" +#include "palUtil.h" + +// gpuopen headers +#include "gpuopen.h" + +// PAL forward declarations +namespace Pal +{ +class ICmdBuffer; +class IFence; +class IQueueSemaphore; +struct PalPublicSettings; +} + +// GpuUtil forward declarations +namespace GpuUtil +{ +class GpaSession; +}; + +// GPUOpen forward declarations +namespace DevDriver +{ +class DevDriverServer; +class IMsgChannel; +struct MessageBuffer; + +namespace DriverControlProtocol +{ +enum struct DeviceClockMode : uint32_t; +class HandlerServer; +} + +namespace SettingsProtocol +{ +class HandlerServer; +} + +} + +namespace pal +{ +class Settings; +class Device; +class VirtualGPU; + +// RGP SQTT Instrumentation Specification version (API-independent) +constexpr uint32_t RgpSqttInstrumentationSpecVersion = 1; + +// RGP SQTT Instrumentation Specification version for Vulkan-specific tables +constexpr uint32_t RgpSqttInstrumentationApiVersion = 0; + +// ===================================================================================================================== +// This class provides functionality to interact with the GPU Open Developer Mode message passing service and the rest +// of the driver. +class RgpCaptureMgr +{ +public: + ~RgpCaptureMgr(); + + static RgpCaptureMgr* Create(Pal::IPlatform* platform, const Device& device); + + void Finalize(); + + void PreDispatch(VirtualGPU* pQueue); + void PostDispatch(VirtualGPU* pQueue); + void WaitForDriverResume(); + + void PostDeviceCreate(); + void PreDeviceDestroy(); + void FinishRGPTrace(bool aborted); + + bool IsQueueTimingActive() const; + +private: + // Steps that an RGP trace goes through + enum class TraceStatus + { + Idle = 0, // No active trace and none requested + Preparing, // A trace has been requested but is not active yet because we are + // currently sampling timing information over some number of lead frames. + Running, // SQTT and queue timing is currently active for all command buffer submits. + WaitingForSqtt, + WaitingForResults // Tracing is no longer active, but all results are not yet ready. + }; + + // All per-device state to support RGP tracing + struct TraceState + { + TraceStatus status_; // Current trace status (idle, running, etc.) + + GpuEvent begin_sqtt_event_; // Event that is signaled when a trace-end cmdbuf retires + GpuEvent end_sqtt_event_; // Event that is signaled when a trace-end cmdbuf retires + GpuEvent end_event_; // Event that is signaled when a trace-end cmdbuf retires + + VirtualGPU* trace_prepare_queue_; // The queue that triggered the full start of a trace + VirtualGPU* trace_begin_queue_; // The queue that triggered starting SQTT + + GpuUtil::GpaSession* gpa_session_; // GPA session helper object for building RGP data + uint32_t gpa_sample_id_; // Sample ID associated with the current trace + bool queue_timing_; // Queue timing is enabled + + uint32_t prepared_disp_count_; // Number of dispatches counted while preparing for a trace + uint32_t sqtt_disp_count_; // Number of dispatches counted while SQTT tracing is active + }; + + RgpCaptureMgr(Pal::IPlatform* platform, const Device& device); + + bool Init(Pal::IPlatform* platform); + Pal::Result PrepareRGPTrace(VirtualGPU* pQueue); + Pal::Result BeginRGPTrace(VirtualGPU* pQueue); + Pal::Result EndRGPHardwareTrace(VirtualGPU* pQueue); + Pal::Result EndRGPTrace(VirtualGPU* pQueue); + void DestroyRGPTracing(); + Pal::Result CheckForTraceResults(); + static bool GpuSupportsTracing(const Pal::DeviceProperties& props, const Settings& settings); + + const Device& device_; + DevDriver::DevDriverServer* dev_driver_server_; + DevDriver::RGPProtocol::RGPServer* rgp_server_; + amd::Monitor trace_mutex_; + TraceState trace_; + uint32_t num_prep_disp_; + uint32_t trace_gpu_mem_limit_; + uint32_t global_disp_count_; + bool trace_enabled_; // True if tracing is currently enabled (master flag) + bool inst_tracing_enabled_; // Enable instruction-level SQTT tokens + + PAL_DISALLOW_DEFAULT_CTOR(RgpCaptureMgr); + PAL_DISALLOW_COPY_AND_ASSIGN(RgpCaptureMgr); +}; + +// ===================================================================================================================== +// Returns true if queue operations are currently being timed by RGP traces. +inline bool RgpCaptureMgr::IsQueueTimingActive() const +{ + return (trace_.queue_timing_ && + (trace_.status_ == TraceStatus::Running || + trace_.status_ == TraceStatus::Preparing || + trace_.status_ == TraceStatus::WaitingForSqtt)); +} +}; diff --git a/projects/clr/rocclr/runtime/device/pal/palsettings.cpp b/projects/clr/rocclr/runtime/device/pal/palsettings.cpp index 90061bce3f..680d16e1fe 100644 --- a/projects/clr/rocclr/runtime/device/pal/palsettings.cpp +++ b/projects/clr/rocclr/runtime/device/pal/palsettings.cpp @@ -133,12 +133,17 @@ Settings::Settings() { // Disable SDMA workaround by default sdamPageFaultWar_ = false; + + // SQTT buffer size in bytes + rgpSqttDispCount_ = PAL_RGP_DISP_COUNT; + rgpSqttWaitIdle_ = true; + rgpSqttForceDisable_ = false; } bool Settings::create(const Pal::DeviceProperties& palProp, const Pal::GpuMemoryHeapProperties* heaps, const Pal::WorkStationCaps& wscaps, - bool reportAsOCL12Device) { - // uint target = calAttr.target; + bool reportAsOCL12Device) +{ uint32_t osVer = 0x0; // Disable thread trace by default for all devices diff --git a/projects/clr/rocclr/runtime/device/pal/palsettings.hpp b/projects/clr/rocclr/runtime/device/pal/palsettings.hpp index 5f0e495e1c..a6755da9d5 100644 --- a/projects/clr/rocclr/runtime/device/pal/palsettings.hpp +++ b/projects/clr/rocclr/runtime/device/pal/palsettings.hpp @@ -64,7 +64,9 @@ class Settings : public device::Settings { uint useDeviceQueue_ : 1; //!< Submit to separate device queue uint singleFpDenorm_ : 1; //!< Support Single FP Denorm uint sdamPageFaultWar_ : 1; //!< SDMA page fault workaround - uint reserved_ : 8; + uint rgpSqttWaitIdle_: 1; //!< Wait for idle after SQTT trace + uint rgpSqttForceDisable_: 1; //!< Disables SQTT + uint reserved_ : 6; }; uint value_; }; @@ -94,6 +96,7 @@ class Settings : public device::Settings { size_t resourceCacheSize_; //!< Resource cache size in MB size_t numMemDependencies_; //!< The array size for memory dependencies tracking uint64_t maxAllocSize_; //!< Maximum single allocation size + uint rgpSqttDispCount_; //!< The number of dispatches captured in SQTT amd::LibrarySelector libSelector_; //!< Select linking libraries for compiler diff --git a/projects/clr/rocclr/runtime/device/pal/palvirtual.cpp b/projects/clr/rocclr/runtime/device/pal/palvirtual.cpp index 014d0e1efa..479f64096e 100644 --- a/projects/clr/rocclr/runtime/device/pal/palvirtual.cpp +++ b/projects/clr/rocclr/runtime/device/pal/palvirtual.cpp @@ -872,6 +872,12 @@ bool VirtualGPU::create(bool profiling, uint deviceQueueSize, uint rtCUs, return false; } + // If the developer mode manager is available and it's not a device queue, + // then enable RGP capturing + if ((index() != 0) && dev().rgpCaptureMgr() != nullptr) { + state_.rgpCaptureEnabled_ = true; + } + return true; } @@ -902,6 +908,11 @@ bool VirtualGPU::allocHsaQueueMem() { } VirtualGPU::~VirtualGPU() { + // Destroy RGP trace + if (rgpCaptureEna()) { + dev().rgpCaptureMgr()->FinishRGPTrace(true); + } + // Not safe to remove a queue. So lock the device amd::ScopedLock k(dev().lockAsyncOps()); amd::ScopedLock lock(dev().vgpusAccess()); @@ -1808,6 +1819,253 @@ void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& vcmd) { profilingEnd(vcmd); } +// ================================================================================================ +void VirtualGPU::PrintChildren(const HSAILKernel& hsaKernel, VirtualGPU* gpuDefQueue) +{ + AmdAqlWrap* wraps = + (AmdAqlWrap*)(&((AmdVQueueHeader*)gpuDefQueue->virtualQueue_->data())[1]); + uint p = 0; + for (uint i = 0; i < gpuDefQueue->vqHeader_->aql_slot_num; ++i) { + if (wraps[i].state != 0) { + uint j; + if (p == GPU_PRINT_CHILD_KERNEL) { + break; + } + p++; + std::stringstream print; + print.flags(std::ios::right | std::ios_base::hex | std::ios_base::uppercase); + print << "Slot#: " << i << "\n"; + print << "\tenqueue_flags: " << wraps[i].enqueue_flags << "\n"; + print << "\tcommand_id: " << wraps[i].command_id << "\n"; + print << "\tchild_counter: " << wraps[i].child_counter << "\n"; + print << "\tcompletion: " << wraps[i].completion << "\n"; + print << "\tparent_wrap: " << wraps[i].parent_wrap << "\n"; + print << "\twait_list: " << wraps[i].wait_list << "\n"; + print << "\twait_num: " << wraps[i].wait_num << "\n"; + uint offsEvents = wraps[i].wait_list - gpuDefQueue->virtualQueue_->vmAddress(); + size_t* events = + reinterpret_cast(gpuDefQueue->virtualQueue_->data() + offsEvents); + for (j = 0; j < wraps[i].wait_num; ++j) { + uint offs = + static_cast(events[j]) - gpuDefQueue->virtualQueue_->vmAddress(); + AmdEvent* eventD = (AmdEvent*)(gpuDefQueue->virtualQueue_->data() + offs); + print << "Wait Event#: " << j << "\n"; + print << "\tState: " << eventD->state << "; Counter: " << eventD->counter << "\n"; + } + print << "WorkGroupSize[ " << wraps[i].aql.workgroup_size_x << ", "; + print << wraps[i].aql.workgroup_size_y << ", "; + print << wraps[i].aql.workgroup_size_z << "]\n"; + print << "GridSize[ " << wraps[i].aql.grid_size_x << ", "; + print << wraps[i].aql.grid_size_y << ", "; + print << wraps[i].aql.grid_size_z << "]\n"; + + uint64_t* kernels = + (uint64_t*)(const_cast(hsaKernel.prog().kernelTable())->map(this)); + for (j = 0; j < hsaKernel.prog().kernels().size(); ++j) { + if (kernels[j] == wraps[i].aql.kernel_object) { + break; + } + } + const_cast(hsaKernel.prog().kernelTable())->unmap(this); + HSAILKernel* child = nullptr; + for (auto it = hsaKernel.prog().kernels().begin(); + it != hsaKernel.prog().kernels().end(); ++it) { + if (j == static_cast(it->second)->index()) { + child = static_cast(it->second); + } + } + if (child == nullptr) { + printf("Error: couldn't find child kernel!\n"); + continue; + } + const uint64_t kernarg_address = + static_cast(reinterpret_cast(wraps[i].aql.kernarg_address)); + uint offsArg = kernarg_address - gpuDefQueue->virtualQueue_->vmAddress(); + address argum = gpuDefQueue->virtualQueue_->data() + offsArg; + print << "Kernel: " << child->name() << "\n"; + for (auto arg : child->arguments()) { + const char* extraArgName = nullptr; + switch (arg->type_) { + case HSAIL_ARGTYPE_HIDDEN_GLOBAL_OFFSET_X: + extraArgName = "Offset0: "; + break; + case HSAIL_ARGTYPE_HIDDEN_GLOBAL_OFFSET_Y: + extraArgName = "Offset1: "; + break; + case HSAIL_ARGTYPE_HIDDEN_GLOBAL_OFFSET_Z: + extraArgName = "Offset2: "; + break; + case HSAIL_ARGTYPE_HIDDEN_PRINTF_BUFFER: + extraArgName = "PrintfBuf: "; + break; + case HSAIL_ARGTYPE_HIDDEN_DEFAULT_QUEUE: + extraArgName = "VqueuePtr: "; + break; + case HSAIL_ARGTYPE_HIDDEN_COMPLETION_ACTION: + extraArgName = "AqlWrap: "; + break; + case HSAIL_ARGTYPE_HIDDEN_NONE: + extraArgName = "Unknown: "; + break; + default: + break; + } + if (extraArgName) { + print << "\t" << extraArgName << *(size_t*)argum; + print << "\n"; + argum += sizeof(size_t); + continue; + } + print << "\t" << arg->name_ << ": "; + for (int s = arg->size_ - 1; s >= 0; --s) { + print.width(2); + print.fill('0'); + print << (uint32_t)(argum[s]); + } + argum += arg->size_; + print << "\n"; + } + printf("%s", print.str().c_str()); + } + } +} + +// ================================================================================================ +bool VirtualGPU::PreDeviceEnqueue( + const amd::Kernel& kernel, + const HSAILKernel& hsaKernel, + VirtualGPU** gpuDefQueue, + uint64_t* vmDefQueue) +{ + amd::DeviceQueue* defQueue = kernel.program().context().defDeviceQueue(dev()); + if (nullptr == defQueue) { + LogError("Default device queue wasn't allocated"); + return false; + } + else { + if (dev().settings().useDeviceQueue_) { + *gpuDefQueue = static_cast(defQueue->vDev()); + if ((*gpuDefQueue)->hwRing() == hwRing()) { + LogError("Can't submit the child kernels to the same HW ring as the host queue!"); + return false; + } + } + else { + createVirtualQueue(defQueue->size()); + *gpuDefQueue = this; + } + } + *vmDefQueue = (*gpuDefQueue)->virtualQueue_->vmAddress(); + + (*gpuDefQueue)->writeVQueueHeader(*this, hsaKernel.prog().kernelTable()->vmAddress()); + // Add memory handles before the actual dispatch + addVmMemory((*gpuDefQueue)->virtualQueue_); + addVmMemory((*gpuDefQueue)->schedParams_); + addVmMemory(hsaKernel.prog().kernelTable()); + return true; +} + +// ================================================================================================ +void VirtualGPU::PostDeviceEnqueue( + const amd::Kernel& kernel, + const HSAILKernel& hsaKernel, + VirtualGPU* gpuDefQueue, + uint64_t vmDefQueue, + uint64_t vmParentWrap, + GpuEvent* gpuEvent) +{ + amd::DeviceQueue* defQueue = kernel.program().context().defDeviceQueue(dev()); + + // Make sure exculsive access to the device queue + amd::ScopedLock(defQueue->lock()); + + if (GPU_PRINT_CHILD_KERNEL != 0) { + waitForEvent(gpuEvent); + PrintChildren(hsaKernel, gpuDefQueue); + } + + if (!dev().settings().useDeviceQueue_) { + // Add the termination handshake to the host queue + eventBegin(MainEngine); + iCmd()->CmdVirtualQueueHandshake(vmParentWrap + offsetof(AmdAqlWrap, state), AQL_WRAP_DONE, + vmParentWrap + offsetof(AmdAqlWrap, child_counter), 0, + dev().settings().useDeviceQueue_); + eventEnd(MainEngine, *gpuEvent); + } + + // Get the global loop start before the scheduler + Pal::gpusize loopStart = gpuDefQueue->iCmd()->CmdVirtualQueueDispatcherStart(); + static_cast(gpuDefQueue->blitMgr()) + .runScheduler(*gpuDefQueue->virtualQueue_, *gpuDefQueue->schedParams_, + gpuDefQueue->schedParamIdx_, + gpuDefQueue->vqHeader_->aql_slot_num / (DeviceQueueMaskSize * maskGroups_)); + const static bool FlushL2 = true; + gpuDefQueue->addBarrier(FlushL2); + + // Get the address of PM4 template and add write it to params + //! @note DMA flush must not occur between patch and the scheduler + Pal::gpusize patchStart = gpuDefQueue->iCmd()->CmdVirtualQueueDispatcherStart(); + // Program parameters for the scheduler + SchedulerParam* param = &reinterpret_cast( + gpuDefQueue->schedParams_->data())[gpuDefQueue->schedParamIdx_]; + param->signal = 1; + // Scale clock to 1024 to avoid 64 bit div in the scheduler + param->eng_clk = (1000 * 1024) / dev().info().maxClockFrequency_; + param->hw_queue = patchStart + sizeof(uint32_t) /* Rewind packet*/; + param->hsa_queue = gpuDefQueue->hsaQueueMem()->vmAddress(); + param->releaseHostCP = 0; + param->parentAQL = vmParentWrap; + param->dedicatedQueue = dev().settings().useDeviceQueue_; + param->useATC = dev().settings().svmFineGrainSystem_; + + // Fill the scratch buffer information + if (hsaKernel.prog().maxScratchRegs() > 0) { + pal::Memory* scratchBuf = dev().scratch(gpuDefQueue->hwRing())->memObj_; + param->scratchSize = scratchBuf->size(); + param->scratch = scratchBuf->vmAddress(); + param->numMaxWaves = 32 * dev().info().maxComputeUnits_; + param->scratchOffset = dev().scratch(gpuDefQueue->hwRing())->offset_; + addVmMemory(scratchBuf); + } + else { + param->numMaxWaves = 0; + param->scratchSize = 0; + param->scratch = 0; + param->scratchOffset = 0; + } + + // Add all kernels in the program to the mem list. + //! \note Runtime doesn't know which one will be called + hsaKernel.prog().fillResListWithKernels(*this); + + Pal::gpusize signalAddr = gpuDefQueue->schedParams_->vmAddress() + + gpuDefQueue->schedParamIdx_ * sizeof(SchedulerParam); + gpuDefQueue->eventBegin(MainEngine); + gpuDefQueue->iCmd()->CmdVirtualQueueDispatcherEnd( + signalAddr, loopStart, + gpuDefQueue->vqHeader_->aql_slot_num / (DeviceQueueMaskSize * maskGroups_)); + // Note: Device enqueue can't have extra commands after INDIRECT_BUFFER call. + // Thus TS command for profiling has to follow in the next CB. + constexpr bool ForceSubmitFirst = true; + gpuDefQueue->eventEnd(MainEngine, *gpuEvent, ForceSubmitFirst); + + if (dev().settings().useDeviceQueue_) { + // Add the termination handshake to the host queue + eventBegin(MainEngine); + iCmd()->CmdVirtualQueueHandshake(vmParentWrap + offsetof(AmdAqlWrap, state), AQL_WRAP_DONE, + vmParentWrap + offsetof(AmdAqlWrap, child_counter), + signalAddr, dev().settings().useDeviceQueue_); + eventEnd(MainEngine, *gpuEvent); + } + + ++gpuDefQueue->schedParamIdx_ %= gpuDefQueue->schedParams_->size() / sizeof(SchedulerParam); + //! \todo optimize the wrap around + if (gpuDefQueue->schedParamIdx_ == 0) { + gpuDefQueue->schedParams_->wait(*gpuDefQueue); + } +} + +// ================================================================================================ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) { // Make sure VirtualGPU has an exclusive access to the resources amd::ScopedLock lock(execution()); @@ -1822,20 +2080,26 @@ void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) { profilingEnd(vcmd); } +// ================================================================================================ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel, const_address parameters, bool nativeMem, - amd::Event* enqueueEvent) { + amd::Event* enqueueEvent) +{ uint64_t vmParentWrap = 0; uint64_t vmDefQueue = 0; - amd::DeviceQueue* defQueue = kernel.program().context().defDeviceQueue(dev()); VirtualGPU* gpuDefQueue = nullptr; amd::HwDebugManager* dbgManager = dev().hwDebugMgr(); + // If RGP capturing is enabled, then start SQTT trace + if (rgpCaptureEna()) { + dev().rgpCaptureMgr()->PreDispatch(this); + } + // Get the HSA kernel object const HSAILKernel& hsaKernel = static_cast(*(kernel.getDeviceKernel(dev()))); bool printfEnabled = (hsaKernel.printfInfo().size() > 0) ? true : false; - if (!printfDbgHSA().init(*this, printfEnabled)) { + if (printfEnabled && !printfDbgHSA().init(*this, printfEnabled)) { LogError("Printf debug buffer initialization failed!"); return false; } @@ -1849,28 +2113,10 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const AddKernel(kernel); if (hsaKernel.dynamicParallelism()) { - if (nullptr == defQueue) { - LogError("Default device queue wasn't allocated"); + // Initialize GPU device queue for execution (gpuDefQueue) + if (!PreDeviceEnqueue(kernel, hsaKernel, &gpuDefQueue, &vmDefQueue)) { return false; - } else { - if (dev().settings().useDeviceQueue_) { - gpuDefQueue = static_cast(defQueue->vDev()); - if (gpuDefQueue->hwRing() == hwRing()) { - LogError("Can't submit the child kernels to the same HW ring as the host queue!"); - return false; - } - } else { - createVirtualQueue(defQueue->size()); - gpuDefQueue = this; - } } - vmDefQueue = gpuDefQueue->virtualQueue_->vmAddress(); - - gpuDefQueue->writeVQueueHeader(*this, hsaKernel.prog().kernelTable()->vmAddress()); - // Add memory handles before the actual dispatch - addVmMemory(gpuDefQueue->virtualQueue_); - addVmMemory(gpuDefQueue->schedParams_); - addVmMemory(hsaKernel.prog().kernelTable()); } // setup the storage for the memory pointers of the kernel parameters @@ -1881,7 +2127,8 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const bool needFlush = false; - // Avoid flushing when PerfCounter is enabled, to make sure PerfStart/dispatch/PerfEnd are in the same cmdBuffer + // Avoid flushing when PerfCounter is enabled, to make sure PerfStart/dispatch/PerfEnd + // are in the same cmdBuffer if (!state_.perfCounterEnabled_) { dmaFlushMgmt_.findSplitSize(dev(), sizes.global().product(), hsaKernel.aqlCodeSize()); if (dmaFlushMgmt().dispatchSplitSize() != 0) { @@ -1911,6 +2158,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const } } } + for (int iter = 0; iter < iteration; ++iter) { GpuEvent gpuEvent(queues_[MainEngine]->cmdBufId()); uint32_t id = gpuEvent.id; @@ -1981,212 +2229,27 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const dbgManager->executePostDispatchCallBack(); } + // Execute scheduler for device enqueue if (hsaKernel.dynamicParallelism()) { - // Make sure exculsive access to the device queue - amd::ScopedLock(defQueue->lock()); - - if (GPU_PRINT_CHILD_KERNEL != 0) { - waitForEvent(&gpuEvent); - - AmdAqlWrap* wraps = - (AmdAqlWrap*)(&((AmdVQueueHeader*)gpuDefQueue->virtualQueue_->data())[1]); - uint p = 0; - for (uint i = 0; i < gpuDefQueue->vqHeader_->aql_slot_num; ++i) { - if (wraps[i].state != 0) { - uint j; - if (p == GPU_PRINT_CHILD_KERNEL) { - break; - } - p++; - std::stringstream print; - print.flags(std::ios::right | std::ios_base::hex | std::ios_base::uppercase); - print << "Slot#: " << i << "\n"; - print << "\tenqueue_flags: " << wraps[i].enqueue_flags << "\n"; - print << "\tcommand_id: " << wraps[i].command_id << "\n"; - print << "\tchild_counter: " << wraps[i].child_counter << "\n"; - print << "\tcompletion: " << wraps[i].completion << "\n"; - print << "\tparent_wrap: " << wraps[i].parent_wrap << "\n"; - print << "\twait_list: " << wraps[i].wait_list << "\n"; - print << "\twait_num: " << wraps[i].wait_num << "\n"; - uint offsEvents = wraps[i].wait_list - gpuDefQueue->virtualQueue_->vmAddress(); - size_t* events = - reinterpret_cast(gpuDefQueue->virtualQueue_->data() + offsEvents); - for (j = 0; j < wraps[i].wait_num; ++j) { - uint offs = - static_cast(events[j]) - gpuDefQueue->virtualQueue_->vmAddress(); - AmdEvent* eventD = (AmdEvent*)(gpuDefQueue->virtualQueue_->data() + offs); - print << "Wait Event#: " << j << "\n"; - print << "\tState: " << eventD->state << "; Counter: " << eventD->counter << "\n"; - } - print << "WorkGroupSize[ " << wraps[i].aql.workgroup_size_x << ", "; - print << wraps[i].aql.workgroup_size_y << ", "; - print << wraps[i].aql.workgroup_size_z << "]\n"; - print << "GridSize[ " << wraps[i].aql.grid_size_x << ", "; - print << wraps[i].aql.grid_size_y << ", "; - print << wraps[i].aql.grid_size_z << "]\n"; - - uint64_t* kernels = - (uint64_t*)(const_cast(hsaKernel.prog().kernelTable())->map(this)); - for (j = 0; j < hsaKernel.prog().kernels().size(); ++j) { - if (kernels[j] == wraps[i].aql.kernel_object) { - break; - } - } - const_cast(hsaKernel.prog().kernelTable())->unmap(this); - HSAILKernel* child = nullptr; - for (auto it = hsaKernel.prog().kernels().begin(); - it != hsaKernel.prog().kernels().end(); ++it) { - if (j == static_cast(it->second)->index()) { - child = static_cast(it->second); - } - } - if (child == nullptr) { - printf("Error: couldn't find child kernel!\n"); - continue; - } - const uint64_t kernarg_address = - static_cast(reinterpret_cast(wraps[i].aql.kernarg_address)); - uint offsArg = kernarg_address - gpuDefQueue->virtualQueue_->vmAddress(); - address argum = gpuDefQueue->virtualQueue_->data() + offsArg; - print << "Kernel: " << child->name() << "\n"; - for (auto arg : child->arguments()) { - const char* extraArgName = nullptr; - switch (arg->type_) { - case HSAIL_ARGTYPE_HIDDEN_GLOBAL_OFFSET_X: - extraArgName = "Offset0: "; - break; - case HSAIL_ARGTYPE_HIDDEN_GLOBAL_OFFSET_Y: - extraArgName = "Offset1: "; - break; - case HSAIL_ARGTYPE_HIDDEN_GLOBAL_OFFSET_Z: - extraArgName = "Offset2: "; - break; - case HSAIL_ARGTYPE_HIDDEN_PRINTF_BUFFER: - extraArgName = "PrintfBuf: "; - break; - case HSAIL_ARGTYPE_HIDDEN_DEFAULT_QUEUE: - extraArgName = "VqueuePtr: "; - break; - case HSAIL_ARGTYPE_HIDDEN_COMPLETION_ACTION: - extraArgName = "AqlWrap: "; - break; - case HSAIL_ARGTYPE_HIDDEN_NONE: - extraArgName = "Unknown: "; - break; - default: - break; - } - if (extraArgName) { - print << "\t" << extraArgName << *(size_t*)argum; - print << "\n"; - argum += sizeof(size_t); - continue; - } - print << "\t" << arg->name_ << ": "; - for (int s = arg->size_ - 1; s >= 0; --s) { - print.width(2); - print.fill('0'); - print << (uint32_t)(argum[s]); - } - argum += arg->size_; - print << "\n"; - } - printf("%s", print.str().c_str()); - } - } - } - - if (!dev().settings().useDeviceQueue_) { - // Add the termination handshake to the host queue - eventBegin(MainEngine); - iCmd()->CmdVirtualQueueHandshake(vmParentWrap + offsetof(AmdAqlWrap, state), AQL_WRAP_DONE, - vmParentWrap + offsetof(AmdAqlWrap, child_counter), 0, - dev().settings().useDeviceQueue_); - eventEnd(MainEngine, gpuEvent); - } - - // Get the global loop start before the scheduler - Pal::gpusize loopStart = gpuDefQueue->iCmd()->CmdVirtualQueueDispatcherStart(); - static_cast(gpuDefQueue->blitMgr()) - .runScheduler(*gpuDefQueue->virtualQueue_, *gpuDefQueue->schedParams_, - gpuDefQueue->schedParamIdx_, - gpuDefQueue->vqHeader_->aql_slot_num / (DeviceQueueMaskSize * maskGroups_)); - const static bool FlushL2 = true; - gpuDefQueue->addBarrier(FlushL2); - - // Get the address of PM4 template and add write it to params - //! @note DMA flush must not occur between patch and the scheduler - Pal::gpusize patchStart = gpuDefQueue->iCmd()->CmdVirtualQueueDispatcherStart(); - // Program parameters for the scheduler - SchedulerParam* param = &reinterpret_cast( - gpuDefQueue->schedParams_->data())[gpuDefQueue->schedParamIdx_]; - param->signal = 1; - // Scale clock to 1024 to avoid 64 bit div in the scheduler - param->eng_clk = (1000 * 1024) / dev().info().maxClockFrequency_; - param->hw_queue = patchStart + sizeof(uint32_t) /* Rewind packet*/; - param->hsa_queue = gpuDefQueue->hsaQueueMem()->vmAddress(); - param->releaseHostCP = 0; - param->parentAQL = vmParentWrap; - param->dedicatedQueue = dev().settings().useDeviceQueue_; - param->useATC = dev().settings().svmFineGrainSystem_; - - // Fill the scratch buffer information - if (hsaKernel.prog().maxScratchRegs() > 0) { - pal::Memory* scratchBuf = dev().scratch(gpuDefQueue->hwRing())->memObj_; - param->scratchSize = scratchBuf->size(); - param->scratch = scratchBuf->vmAddress(); - param->numMaxWaves = 32 * dev().info().maxComputeUnits_; - param->scratchOffset = dev().scratch(gpuDefQueue->hwRing())->offset_; - addVmMemory(scratchBuf); - } else { - param->numMaxWaves = 0; - param->scratchSize = 0; - param->scratch = 0; - param->scratchOffset = 0; - } - - // Add all kernels in the program to the mem list. - //! \note Runtime doesn't know which one will be called - hsaKernel.prog().fillResListWithKernels(*this); - - Pal::gpusize signalAddr = gpuDefQueue->schedParams_->vmAddress() + - gpuDefQueue->schedParamIdx_ * sizeof(SchedulerParam); - gpuDefQueue->eventBegin(MainEngine); - gpuDefQueue->iCmd()->CmdVirtualQueueDispatcherEnd( - signalAddr, loopStart, - gpuDefQueue->vqHeader_->aql_slot_num / (DeviceQueueMaskSize * maskGroups_)); - // Note: Device enqueue can't have extra commands after INDIRECT_BUFFER call. - // Thus TS command for profiling has to follow in the next CB. - constexpr bool ForceSubmitFirst = true; - gpuDefQueue->eventEnd(MainEngine, gpuEvent, ForceSubmitFirst); - - if (dev().settings().useDeviceQueue_) { - // Add the termination handshake to the host queue - eventBegin(MainEngine); - iCmd()->CmdVirtualQueueHandshake(vmParentWrap + offsetof(AmdAqlWrap, state), AQL_WRAP_DONE, - vmParentWrap + offsetof(AmdAqlWrap, child_counter), - signalAddr, dev().settings().useDeviceQueue_); - eventEnd(MainEngine, gpuEvent); - } - - ++gpuDefQueue->schedParamIdx_ %= gpuDefQueue->schedParams_->size() / sizeof(SchedulerParam); - //! \todo optimize the wrap around - if (gpuDefQueue->schedParamIdx_ == 0) { - gpuDefQueue->schedParams_->wait(*gpuDefQueue); - } + PostDeviceEnqueue(kernel, hsaKernel, gpuDefQueue, vmDefQueue, vmParentWrap, &gpuEvent); } + if (id != gpuEvent.id) { LogError("Something is wrong. ID mismatch!\n"); } // Update the global GPU event setGpuEvent(gpuEvent, needFlush); - if (!printfDbgHSA().output(*this, printfEnabled, hsaKernel.printfInfo())) { + if (printfEnabled && !printfDbgHSA().output(*this, printfEnabled, hsaKernel.printfInfo())) { LogError("Couldn't read printf data from the buffer!\n"); return false; } } + if (rgpCaptureEna()) { + dev().rgpCaptureMgr()->PostDispatch(this); + } + return true; } diff --git a/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp b/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp index 07265462e9..43b67b17f1 100644 --- a/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp +++ b/projects/clr/rocclr/runtime/device/pal/palvirtual.hpp @@ -153,7 +153,7 @@ class VirtualGPU : public device::VirtualDevice { Pal::IFence* iCmdFences_[MaxCmdBuffers]; //!< PAL fences, associated with CMD const amd::Kernel* last_kernel_; //!< Last submitted kernel - private: + private: void DumpMemoryReferences() const; Pal::IDevice* iDev_; //!< PAL device uint cmdBufIdSlot_; //!< Command buffer ID slot for submissions @@ -201,6 +201,7 @@ class VirtualGPU : public device::VirtualDevice { uint forceWait_ : 1; //!< Forces wait in flush() uint profileEnabled_ : 1; //!< Profiling is enabled for WaveLimiter uint perfCounterEnabled_ : 1; //!< PerfCounter is enabled + uint rgpCaptureEnabled_ : 1; //!< RGP capture is enabled in the runtime }; uint value_; State() : value_(0) {} @@ -496,6 +497,9 @@ class VirtualGPU : public device::VirtualDevice { const Resource& dst //!< Destination resource for SDMA transfer ); + //! Checks if RGP capture is enabled + bool rgpCaptureEna() const { return state_.rgpCaptureEnabled_; } + protected: void profileEvent(EngineType engine, bool type) const; @@ -561,6 +565,24 @@ class VirtualGPU : public device::VirtualDevice { HwDbgKernelInfo& kernelInfo //!< kernel info for the dispatch ); + void PrintChildren(const HSAILKernel& hsaKernel, //!< The parent HSAIL kernel + VirtualGPU* gpuDefQueue //!< Device queue for children execution + ); + + bool PreDeviceEnqueue(const amd::Kernel& kernel, //!< Parent amd kernel object + const HSAILKernel& hsaKernel, //!< Parent HSAIL object + VirtualGPU** gpuDefQueue, //!< [Return] GPU default queue + uint64_t* vmDefQueue //!< [Return] VM handle to the virtual queue + ); + + void PostDeviceEnqueue(const amd::Kernel& kernel, //!< Parent amd kernel object + const HSAILKernel& hsaKernel, //!< Parent HSAIL object + VirtualGPU* gpuDefQueue, //!< GPU default queue + uint64_t vmDefQueue, //!< VM handle to the virtual queue + uint64_t vmParentWrap, //!< VM handle to the wrapped AQL packet location + GpuEvent* gpuEvent //!< [Return] GPU event associated with the device enqueue + ); + Device& gpuDevice_; //!< physical GPU device amd::Monitor execution_; //!< Lock to serialise access to all device objects uint index_; //!< The virtual device unique index diff --git a/projects/clr/rocclr/runtime/utils/flags.hpp b/projects/clr/rocclr/runtime/utils/flags.hpp index a91245c701..aba52edb5b 100644 --- a/projects/clr/rocclr/runtime/utils/flags.hpp +++ b/projects/clr/rocclr/runtime/utils/flags.hpp @@ -211,6 +211,8 @@ release(bool, GPU_VEGA10_ONLY, VEGA10_ONLY, \ "1 = Report vega10 only on OCL/ROCR") \ release_on_stg(bool, PAL_DISABLE_SDMA, false, \ "1 = Disable SDMA for PAL") \ +release_on_stg(uint, PAL_RGP_DISP_COUNT, 10, \ + "The number of dispatches for RGP capture with SQTT") \ namespace amd {