AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Support debugging hw exceptions.

Browse Source 
Change-Id: I9780147294af2e9457fa54693580735452ee2ae6


[ROCm/ROCR-Runtime commit: 206e87d28b]
This commit is contained in:
Sean Keely
2021-05-25 16:32:07 -05:00
parent 183531963c
commit 5c9500d50b
6 changed files with 211 additions and 59 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h
+8 -1
View File
@@ -222,8 +222,12 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
void Suspend();
/// @brief Handler for hardware queue events.
template <bool HandleExceptions>
static bool DynamicScratchHandler(hsa_signal_value_t error_code, void* arg);
/// @brief Handler for KFD exceptions.
static bool ExceptionHandler(hsa_signal_value_t error_code, void* arg);
// AQL packet ring buffer
void* ring_buf_;
@@ -261,7 +265,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
KernelMutex pm4_ib_mutex_;
// Error handler control variable.
std::atomic<uint32_t> dynamicScratchState;
std::atomic<uint32_t> dynamicScratchState, exceptionState;
enum { ERROR_HANDLER_DONE = 1, ERROR_HANDLER_TERMINATE = 2, ERROR_HANDLER_SCRATCH_RETRY = 4 };
// Queue currently suspended or scheduled
@@ -270,6 +274,9 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo
// Thunk dispatch and wavefront scheduling priority
HSA_QUEUE_PRIORITY priority_;
// Exception notification signal
Signal* exception_signal_;
// Shared event used for queue errors
static HsaEvent* queue_event_;
projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h
+12 -3
View File
@@ -106,6 +106,11 @@ class Runtime {
hsa_amd_memory_pool_link_info_t info;
};
struct KfdVersion_t {
HsaVersionInfo version;
bool supports_exception_debugging;
};
/// @brief Open connection to kernel driver and increment reference count.
static hsa_status_t Acquire();
@@ -350,9 +355,13 @@ class Runtime {
uint64_t sys_clock_freq() const { return sys_clock_freq_; }
void KfdVersion(const HsaVersionInfo& version) { kfd_version = version; }
void KfdVersion(const HsaVersionInfo& version) { kfd_version.version = version; }
HsaVersionInfo KfdVersion() const { return kfd_version; }
void KfdVersion(bool exception_debugging) {
kfd_version.supports_exception_debugging = exception_debugging;
}
KfdVersion_t KfdVersion() const { return kfd_version; }
protected:
static void AsyncEventsLoop(void*);
@@ -575,7 +584,7 @@ class Runtime {
InterruptSignal::EventPool EventPool;
// Kfd version
HsaVersionInfo kfd_version;
KfdVersion_t kfd_version;
// Frees runtime memory when the runtime library is unloaded if safe to do so.
// Failure to release the runtime indicates an incorrect application but is
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+159 -53
View File
@@ -96,7 +96,8 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
pm4_ib_size_b_(0x1000),
dynamicScratchState(0),
suspended_(false),
priority_(HSA_QUEUE_PRIORITY_NORMAL) {
priority_(HSA_QUEUE_PRIORITY_NORMAL),
exception_signal_(nullptr) {
// When queue_full_workaround_ is set to 1, the ring buffer is internally
// doubled in size. Virtual addresses in the upper half of the ring allocation
// are mapped to the same set of pages backing the lower half.
@@ -211,28 +212,6 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
assert(amd_queue_.private_segment_aperture_base_hi != 0 && "No private region found.");
}
// Ensure the amd_queue_ is fully initialized before creating the KFD queue.
// This ensures that the debugger can access the fields once it detects there
// is a KFD queue. The debugger may access the aperture addresses, queue
// scratch base, and queue type.
HSAKMT_STATUS kmt_status;
kmt_status = hsaKmtCreateQueue(node_id, HSA_QUEUE_COMPUTE_AQL, 100, priority_, ring_buf_,
ring_buf_alloc_bytes_, NULL, &queue_rsrc);
if (kmt_status != HSAKMT_STATUS_SUCCESS)
throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
"Queue create failed at hsaKmtCreateQueue\n");
// Complete populating the doorbell signal structure.
signal_.legacy_hardware_doorbell_ptr =
(volatile uint32_t*)queue_rsrc.Queue_DoorBell;
// Bind Id of Queue such that is unique i.e. it is not re-used by another
// queue (AQL, HOST) in the same process during its lifetime.
amd_queue_.hsa_queue.id = this->GetQueueId();
queue_id_ = queue_rsrc.QueueId;
MAKE_NAMED_SCOPE_GUARD(QueueGuard, [&]() { hsaKmtDestroyQueue(queue_id_); });
MAKE_NAMED_SCOPE_GUARD(EventGuard, [&]() {
ScopedAcquire<KernelMutex> _lock(&queue_lock_);
queue_count_--;
@@ -243,7 +222,9 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
});
MAKE_NAMED_SCOPE_GUARD(SignalGuard, [&]() {
HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal);
if (amd_queue_.queue_inactive_signal.handle != 0)
HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal);
if (exception_signal_ != nullptr) exception_signal_->DestroySignal();
});
if (core::g_use_interrupt_wait) {
@@ -260,21 +241,67 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr
auto Signal = new core::InterruptSignal(0, queue_event_);
assert(Signal != nullptr && "Should have thrown!\n");
amd_queue_.queue_inactive_signal = core::InterruptSignal::Convert(Signal);
exception_signal_ = new core::InterruptSignal(0, queue_event_);
assert(exception_signal_ != nullptr && "Should have thrown!\n");
} else {
EventGuard.Dismiss();
auto Signal = new core::DefaultSignal(0);
assert(Signal != nullptr && "Should have thrown!\n");
amd_queue_.queue_inactive_signal = core::DefaultSignal::Convert(Signal);
exception_signal_ = new core::DefaultSignal(0);
assert(exception_signal_ != nullptr && "Should have thrown!\n");
}
// Ensure the amd_queue_ is fully initialized before creating the KFD queue.
// This ensures that the debugger can access the fields once it detects there
// is a KFD queue. The debugger may access the aperture addresses, queue
// scratch base, and queue type.
HSAKMT_STATUS kmt_status;
if (core::Runtime::runtime_singleton_->KfdVersion().supports_exception_debugging) {
queue_rsrc.ErrorReason = &exception_signal_->signal_.value;
kmt_status = hsaKmtCreateQueue(node_id, HSA_QUEUE_COMPUTE_AQL, 100, priority_, ring_buf_,
ring_buf_alloc_bytes_, queue_event_, &queue_rsrc);
} else {
kmt_status = hsaKmtCreateQueue(node_id, HSA_QUEUE_COMPUTE_AQL, 100, priority_, ring_buf_,
ring_buf_alloc_bytes_, NULL, &queue_rsrc);
}
if (kmt_status != HSAKMT_STATUS_SUCCESS)
throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
"Queue create failed at hsaKmtCreateQueue\n");
// Complete populating the doorbell signal structure.
signal_.legacy_hardware_doorbell_ptr = (volatile uint32_t*)queue_rsrc.Queue_DoorBell;
// Bind Id of Queue such that is unique i.e. it is not re-used by another
// queue (AQL, HOST) in the same process during its lifetime.
amd_queue_.hsa_queue.id = this->GetQueueId();
queue_id_ = queue_rsrc.QueueId;
MAKE_NAMED_SCOPE_GUARD(QueueGuard, [&]() { hsaKmtDestroyQueue(queue_id_); });
// Initialize scratch memory related entities
queue_scratch_.queue_retry = amd_queue_.queue_inactive_signal;
InitScratchSRD();
if (AMD::hsa_amd_signal_async_handler(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE,
0, DynamicScratchHandler, this) != HSA_STATUS_SUCCESS)
throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
"Queue event handler failed registration.\n");
if (core::Runtime::runtime_singleton_->KfdVersion().supports_exception_debugging) {
if (AMD::hsa_amd_signal_async_handler(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE,
0, DynamicScratchHandler<false>,
this) != HSA_STATUS_SUCCESS)
throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
"Queue event handler failed registration.\n");
if (AMD::hsa_amd_signal_async_handler(core::Signal::Convert(exception_signal_),
HSA_SIGNAL_CONDITION_NE, 0, ExceptionHandler,
this) != HSA_STATUS_SUCCESS)
throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
"Queue event handler failed registration.\n");
} else {
if (AMD::hsa_amd_signal_async_handler(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE,
0, DynamicScratchHandler<true>,
this) != HSA_STATUS_SUCCESS)
throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES,
"Queue event handler failed registration.\n");
exceptionState = ERROR_HANDLER_DONE;
}
pm4_ib_buf_ = core::Runtime::runtime_singleton_->system_allocator()(
pm4_ib_size_b_, 0x1000, core::MemoryRegion::AllocateExecutable);
@@ -305,9 +332,17 @@ AqlQueue::~AqlQueue() {
HSA::hsa_signal_store_relaxed(amd_queue_.queue_inactive_signal, 0x8000000000000000ull);
}
// Remove kfd exception handler
exceptionState |= ERROR_HANDLER_TERMINATE;
while ((exceptionState & ERROR_HANDLER_DONE) != ERROR_HANDLER_DONE) {
exception_signal_->StoreRelease(-1ull);
exception_signal_->WaitRelaxed(HSA_SIGNAL_CONDITION_NE, -1ull, -1ull, HSA_WAIT_STATE_BLOCKED);
}
Inactivate();
agent_->ReleaseQueueScratch(queue_scratch_);
FreeRegisteredRingBuffer();
exception_signal_->DestroySignal();
HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal);
if (core::g_use_interrupt_wait) {
ScopedAcquire<KernelMutex> lock(&queue_lock_);
@@ -734,6 +769,7 @@ hsa_status_t AqlQueue::SetPriority(HSA_QUEUE_PRIORITY priority) {
return (err == HSAKMT_STATUS_SUCCESS ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR_OUT_OF_RESOURCES);
}
template <bool HandleExceptions>
bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) {
AqlQueue* queue = (AqlQueue*)arg;
hsa_status_t errorCode = HSA_STATUS_SUCCESS;
@@ -849,46 +885,48 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) {
}
}
} else if ((error_code & 2) == 2) { // Invalid dim
errorCode = HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS;
} else if (HandleExceptions) {
if ((error_code & 2) == 2) { // Invalid dim
errorCode = HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS;
} else if ((error_code & 4) == 4) { // Invalid group memory
errorCode = HSA_STATUS_ERROR_INVALID_ALLOCATION;
} else if ((error_code & 4) == 4) { // Invalid group memory
errorCode = HSA_STATUS_ERROR_INVALID_ALLOCATION;
} else if ((error_code & 8) == 8) { // Invalid (or NULL) code
errorCode = HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
} else if ((error_code & 8) == 8) { // Invalid (or NULL) code
errorCode = HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
} else if (((error_code & 32) == 32) || // Invalid format: 32 is generic,
((error_code & 256) == 256)) { // 256 is vendor specific packets
errorCode = HSA_STATUS_ERROR_INVALID_PACKET_FORMAT;
} else if (((error_code & 32) == 32) || // Invalid format: 32 is generic,
((error_code & 256) == 256)) { // 256 is vendor specific packets
errorCode = HSA_STATUS_ERROR_INVALID_PACKET_FORMAT;
} else if ((error_code & 64) == 64) { // Group is too large
errorCode = HSA_STATUS_ERROR_INVALID_ARGUMENT;
} else if ((error_code & 64) == 64) { // Group is too large
errorCode = HSA_STATUS_ERROR_INVALID_ARGUMENT;
} else if ((error_code & 128) == 128) { // Out of VGPRs
errorCode = HSA_STATUS_ERROR_INVALID_ISA;
} else if ((error_code & 128) == 128) { // Out of VGPRs
errorCode = HSA_STATUS_ERROR_INVALID_ISA;
} else if ((error_code & 0x20000000) == 0x20000000) { // Memory violation (>48-bit)
errorCode = hsa_status_t(HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION);
} else if ((error_code & 0x20000000) == 0x20000000) { // Memory violation (>48-bit)
errorCode = hsa_status_t(HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION);
} else if ((error_code & 0x40000000) == 0x40000000) { // Illegal instruction
errorCode = hsa_status_t(HSA_STATUS_ERROR_ILLEGAL_INSTRUCTION);
} else if ((error_code & 0x40000000) == 0x40000000) { // Illegal instruction
errorCode = hsa_status_t(HSA_STATUS_ERROR_ILLEGAL_INSTRUCTION);
} else if ((error_code & 0x80000000) == 0x80000000) { // Debug trap
errorCode = HSA_STATUS_ERROR_EXCEPTION;
fatal = true;
} else if ((error_code & 0x80000000) == 0x80000000) { // Debug trap
errorCode = HSA_STATUS_ERROR_EXCEPTION;
fatal = true;
} else { // Undefined code
assert(false && "Undefined queue error code");
errorCode = HSA_STATUS_ERROR;
fatal = true;
} else { // Undefined code
assert(false && "Undefined queue error code");
errorCode = HSA_STATUS_ERROR;
fatal = true;
}
}
if (errorCode == HSA_STATUS_SUCCESS) {
if (changeWait) {
core::Runtime::runtime_singleton_->SetAsyncSignalHandler(
queue->amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE, waitVal,
DynamicScratchHandler, queue);
DynamicScratchHandler<HandleExceptions>, queue);
return false;
}
return true;
@@ -913,6 +951,74 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) {
return false;
}
bool AqlQueue::ExceptionHandler(hsa_signal_value_t error_code, void* arg) {
struct queue_error_t {
uint32_t code;
hsa_status_t status;
};
static const queue_error_t QueueErrors[] = {
// EC_QUEUE_TRAP
2, HSA_STATUS_ERROR_EXCEPTION,
// EC_QUEUE_ILLEGAL_INSTRUCTION
3, (hsa_status_t)HSA_STATUS_ERROR_ILLEGAL_INSTRUCTION,
// EC_QUEUE_MEMORY_VIOLATION
4, HSA_STATUS_ERROR,
// EC_QUEUE_APERTURE_VIOLATION
5, (hsa_status_t)HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION,
// EC_QUEUE_PACKET_DISPATCH_DIM_INVALID
16, HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS,
// EC_QUEUE_PACKET_DISPATCH_GROUP_SEGMENT_SIZE_INVALID
17, HSA_STATUS_ERROR_INVALID_ALLOCATION,
// EC_QUEUE_PACKET_DISPATCH_CODE_INVALID
18, HSA_STATUS_ERROR_INVALID_CODE_OBJECT,
// EC_QUEUE_PACKET_UNSUPPORTED
20, HSA_STATUS_ERROR_INVALID_PACKET_FORMAT,
// EC_QUEUE_PACKET_DISPATCH_WORK_GROUP_SIZE_INVALID
21, HSA_STATUS_ERROR_INVALID_ARGUMENT,
// EC_QUEUE_PACKET_DISPATCH_REGISTER_SIZE_INVALID
22, HSA_STATUS_ERROR_INVALID_ISA,
// EC_QUEUE_PACKET_VENDOR_UNSUPPORTED
23, HSA_STATUS_ERROR_INVALID_PACKET_FORMAT,
// EC_QUEUE_PREEMPTION_ERROR
31, HSA_STATUS_ERROR,
// EC_DEVICE_MEMORY_VIOLATION
33, (hsa_status_t)HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION,
// EC_DEVICE_RAS_ERROR
34, HSA_STATUS_ERROR,
// EC_DEVICE_FATAL_HALT
35, HSA_STATUS_ERROR,
// EC_DEVICE_NEW
36, HSA_STATUS_ERROR,
// EC_PROCESS_DEVICE_REMOVE
50, HSA_STATUS_ERROR};
AqlQueue* queue = (AqlQueue*)arg;
hsa_status_t errorCode = HSA_STATUS_ERROR;
if (queue->exceptionState == ERROR_HANDLER_TERMINATE) {
Signal* signal = queue->exception_signal_;
queue->exceptionState = ERROR_HANDLER_DONE;
signal->StoreRelease(0);
return false;
}
for (auto& error : QueueErrors) {
if (error_code & (1 << (error.code - 1))) {
errorCode = error.status;
break;
}
}
// Undefined or unexpected code
assert((errorCode != HSA_STATUS_ERROR) && "Undefined or unexpected queue error code");
queue->Suspend();
if (queue->errors_callback_ != nullptr) {
queue->errors_callback_(errorCode, queue->public_handle(), queue->errors_data_);
}
return false;
}
hsa_status_t AqlQueue::SetCUMasking(const uint32_t num_cu_mask_count,
const uint32_t* cu_mask) {
HSAKMT_STATUS ret = hsaKmtSetQueueCUMask(
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+14 -1
View File
@@ -215,6 +215,15 @@ void GpuAgent::AssembleShader(const char* func_name, AssembleTarget assemble_tar
{kCodeTrapHandler1010, sizeof(kCodeTrapHandler1010), 2, 4},
{kCodeTrapHandler10, sizeof(kCodeTrapHandler10), 2, 4},
}},
{"TrapHandlerKfdExceptions",
{
{NULL, 0, 0, 0},
{kCodeTrapHandler8, sizeof(kCodeTrapHandler8), 2, 4},
{kCodeTrapHandler9, sizeof(kCodeTrapHandler9), 2, 4},
{kCodeTrapHandler90a, sizeof(kCodeTrapHandler90a), 2, 4},
{kCodeTrapHandler1010, sizeof(kCodeTrapHandler1010), 2, 4},
{kCodeTrapHandler10, sizeof(kCodeTrapHandler10), 2, 4},
}},
{"CopyAligned",
{
{kCodeCopyAligned7, sizeof(kCodeCopyAligned7), 32, 12},
@@ -1397,7 +1406,11 @@ void GpuAgent::BindTrapHandler() {
}
// Assemble the trap handler source code.
AssembleShader("TrapHandler", AssembleTarget::ISA, trap_code_buf_, trap_code_buf_size_);
if (core::Runtime::runtime_singleton_->KfdVersion().supports_exception_debugging)
AssembleShader("TrapHandlerKfdExceptions", AssembleTarget::ISA, trap_code_buf_,
trap_code_buf_size_);
else
AssembleShader("TrapHandler", AssembleTarget::ISA, trap_code_buf_, trap_code_buf_size_);
// Bind the trap handler to this node.
HSAKMT_STATUS err = hsaKmtSetTrapHandler(node_id(), trap_code_buf_, trap_code_buf_size_,
projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_topology.cpp
+12 -1
View File
@@ -49,6 +49,7 @@
#include <map>
#include <string>
#include <sstream>
#include <link.h>
#ifndef NDBEUG
#include <iostream>
@@ -62,6 +63,8 @@
#include "core/inc/amd_memory_region.h"
#include "core/util/utils.h"
extern r_debug _amdgpu_r_debug;
namespace rocr {
namespace AMD {
// Minimum acceptable KFD version numbers
@@ -120,7 +123,7 @@ GpuAgent* DiscoverGpu(HSAuint32 node_id, HsaNodeProperties& node_prop, bool xnac
try {
gpu = new GpuAgent(node_id, node_prop, xnack_mode);
const HsaVersionInfo& kfd_version = core::Runtime::runtime_singleton_->KfdVersion();
const HsaVersionInfo& kfd_version = core::Runtime::runtime_singleton_->KfdVersion().version;
// Check for sramecc incompatibility due to sramecc not being reported correctly in kfd before
// 1.4.
@@ -358,6 +361,12 @@ bool Load() {
return false;
}
// Register runtime and optionally enable the debugger
HSAKMT_STATUS err =
hsaKmtRuntimeEnable(&_amdgpu_r_debug, core::Runtime::runtime_singleton_->flag().debug());
if (err != HSAKMT_STATUS_SUCCESS) return false;
core::Runtime::runtime_singleton_->KfdVersion(err != HSAKMT_STATUS_NOT_SUPPORTED);
// Build topology table.
BuildTopology();
@@ -365,6 +374,8 @@ bool Load() {
}
bool Unload() {
hsaKmtRuntimeDisable();
hsaKmtReleaseSystemProperties();
// Close connection to kernel driver.
projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h
+6
View File
@@ -146,6 +146,9 @@ class Flag {
// be interpreted as not defining the env variable.
var = os::GetEnvVar("HSA_XNACK");
xnack_ = (var == "0") ? XNACK_DISABLE : ((var == "1") ? XNACK_ENABLE : XNACK_UNCHANGED);
var = os::GetEnvVar("HSA_ENABLE_DEBUG");
debug_ = (var == "1") ? true : false;
}
bool check_flat_scratch() const { return check_flat_scratch_; }
@@ -201,6 +204,8 @@ class Flag {
XNACK_REQUEST xnack() const { return xnack_; }
bool debug() const { return debug_; }
private:
bool check_flat_scratch_;
bool enable_vm_fault_message_;
@@ -220,6 +225,7 @@ class Flag {
bool check_sramecc_validity_;
bool patch_xgmi_link_weight_;
bool patch_link_override_;
bool debug_;
SDMA_OVERRIDE enable_sdma_;