diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h index 17f68b2b59..2ac6b3503d 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -222,8 +222,12 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo void Suspend(); /// @brief Handler for hardware queue events. + template 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 dynamicScratchState; + std::atomic 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_; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h index d232e6f955..3bb560a52f 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/runtime.h @@ -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 diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 8263741278..f86a5dd304 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -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 _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, + 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, + 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 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 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, 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( diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index 0c35286f6c..a43d91f845 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -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_, diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_topology.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_topology.cpp index 71bc738215..664048904a 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_topology.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_topology.cpp @@ -49,6 +49,7 @@ #include #include #include +#include #ifndef NDBEUG #include @@ -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. diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h b/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h index 88fbe93be2..c7149363ad 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h @@ -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_;