AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

P4 to Git Change 1517959 by gandryey@gera-lnx-rcf on 2018/02/21 17:24:16

Explorar el Código 
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: dfb5e1b844]
Este commit está contenido en:
foreman
2018-02-21 17:35:38 -05:00
padre 63e460b3bf
commit a9e75064ed
Se han modificado 10 ficheros con 1132 adiciones y 224 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
projects/clr/rocclr/runtime/device/pal/paldefs.hpp
+7 -1
Ver fichero
@@ -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
projects/clr/rocclr/runtime/device/pal/paldevice.cpp
+10 -2
Ver fichero
@@ -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;
projects/clr/rocclr/runtime/device/pal/paldevice.hpp
+4
Ver fichero
@@ -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<Pal::gpusize> freeMem[Pal::GpuHeap::GpuHeapCount]; //!< Free memory counter
amd::Monitor* lockResourceOps_; //!< Lock to serialise resource access
std::list<GpuMemoryReference*>* resourceList_; //!< Active resource list
RgpCaptureMgr* rgpCaptureMgr_; //!< RGP capture manager
};
/*@}*/} // namespace pal
projects/clr/rocclr/runtime/device/pal/palgpuopen.cpp
+623
Ver fichero
@@ -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<Pal::uint8*>(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
projects/clr/rocclr/runtime/device/pal/palgpuopen.hpp
+172
Ver fichero
@@ -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 <queue>
#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));
}
};
projects/clr/rocclr/runtime/device/pal/palsettings.cpp
+7 -2
Ver fichero
@@ -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
projects/clr/rocclr/runtime/device/pal/palsettings.hpp
+4 -1
Ver fichero
@@ -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
projects/clr/rocclr/runtime/device/pal/palvirtual.cpp
+280 -217
Ver fichero
@@ -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<size_t*>(gpuDefQueue->virtualQueue_->data() + offsEvents);
for (j = 0; j < wraps[i].wait_num; ++j) {
uint offs =
static_cast<uint64_t>(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<Memory*>(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<Memory*>(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<HSAILKernel*>(it->second)->index()) {
child = static_cast<HSAILKernel*>(it->second);
}
}
if (child == nullptr) {
printf("Error: couldn't find child kernel!\n");
continue;
}
const uint64_t kernarg_address =
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(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<VirtualGPU*>(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<KernelBlitManager&>(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<SchedulerParam*>(
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<const HSAILKernel&>(*(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<VirtualGPU*>(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<size_t*>(gpuDefQueue->virtualQueue_->data() + offsEvents);
for (j = 0; j < wraps[i].wait_num; ++j) {
uint offs =
static_cast<uint64_t>(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<Memory*>(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<Memory*>(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<HSAILKernel*>(it->second)->index()) {
child = static_cast<HSAILKernel*>(it->second);
}
}
if (child == nullptr) {
printf("Error: couldn't find child kernel!\n");
continue;
}
const uint64_t kernarg_address =
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(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<KernelBlitManager&>(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<SchedulerParam*>(
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;
}
projects/clr/rocclr/runtime/device/pal/palvirtual.hpp
+23 -1
Ver fichero
@@ -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
projects/clr/rocclr/runtime/utils/flags.hpp
+2
Ver fichero
@@ -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 {