From 7dfeee5074c8bc0883c6474b54b08766c3a1a5db Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Tue, 28 Mar 2017 03:12:59 -0500 Subject: [PATCH] Support async. queue errors and dynamic scratch without KFD events. Change-Id: I4e9e7a37aa7b9c96b28ce79f562760283e02b1e0 --- .../core/runtime/amd_aql_queue.cpp | 44 ++++++++----------- .../core/runtime/amd_gpu_agent.cpp | 40 ++++++----------- runtime/hsa-runtime/core/runtime/hsa.cpp | 4 ++ .../hsa-runtime/core/runtime/hsa_ext_amd.cpp | 2 +- runtime/hsa-runtime/core/runtime/runtime.cpp | 2 +- 5 files changed, 38 insertions(+), 54 deletions(-) diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 152e86c031..9b316fd530 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -64,6 +64,7 @@ #include "core/util/utils.h" #include "core/inc/registers.h" #include "core/inc/interrupt_signal.h" +#include "core/inc/default_signal.h" #include "core/inc/hsa_ext_amd_impl.h" #include "core/inc/amd_gpu_pm4.h" @@ -241,34 +242,28 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, MAKE_NAMED_SCOPE_GUARD(SignalGuard, [&]() { HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal); }); -#if defined(HSA_LARGE_MODEL) && defined(__linux__) - if (core::g_use_interrupt_wait) { - { - ScopedAcquire _lock(&queue_lock_); - queue_count_++; - if (queue_event_ == NULL) { - assert(queue_count_ == 1 && - "Inconsistency in queue event reference counting found.\n"); - queue_event_ = - core::InterruptSignal::CreateEvent(HSA_EVENTTYPE_SIGNAL, false); - if (queue_event_ == NULL) return; - } + if (core::g_use_interrupt_wait) { + ScopedAcquire _lock(&queue_lock_); + queue_count_++; + if (queue_event_ == NULL) { + assert(queue_count_ == 1 && "Inconsistency in queue event reference counting found.\n"); + + queue_event_ = core::InterruptSignal::CreateEvent(HSA_EVENTTYPE_SIGNAL, false); + if (queue_event_ == NULL) return; } - auto signal = new core::InterruptSignal(0, queue_event_); - amd_queue_.queue_inactive_signal = core::InterruptSignal::Convert(signal); - if (AMD::hsa_amd_signal_async_handler( - amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE, 0, - DynamicScratchHandler, this) != HSA_STATUS_SUCCESS) - return; + auto Signal = new core::InterruptSignal(0, queue_event_); + if (Signal == nullptr) return; + amd_queue_.queue_inactive_signal = core::InterruptSignal::Convert(Signal); } else { EventGuard.Dismiss(); - SignalGuard.Dismiss(); + auto Signal = new core::DefaultSignal(0); + if (Signal == nullptr) return; + amd_queue_.queue_inactive_signal = core::DefaultSignal::Convert(Signal); } -#else - EventGuard.Dismiss(); - SignalGuard.Dismiss(); -#endif + if (AMD::hsa_amd_signal_async_handler(amd_queue_.queue_inactive_signal, HSA_SIGNAL_CONDITION_NE, + 0, DynamicScratchHandler, this) != HSA_STATUS_SUCCESS) + return; pm4_ib_buf_ = core::Runtime::runtime_singleton_->system_allocator()( pm4_ib_size_b_, 0x1000, core::MemoryRegion::AllocateExecutable); @@ -298,7 +293,6 @@ AqlQueue::~AqlQueue() { FreeRegisteredRingBuffer(); agent_->ReleaseQueueScratch(queue_scratch_.queue_base); HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal); -#if defined(HSA_LARGE_MODEL) && defined(__linux__) if (core::g_use_interrupt_wait) { ScopedAcquire lock(&queue_lock_); queue_count_--; @@ -307,8 +301,6 @@ AqlQueue::~AqlQueue() { queue_event_ = NULL; } } -#endif - core::Runtime::runtime_singleton_->system_deallocator()(pm4_ib_buf_); } diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index 775b32ff33..82175997d3 100755 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -882,33 +882,21 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, // Allocate scratch memory ScratchInfo scratch; -#if defined(HSA_LARGE_MODEL) && defined(__linux__) - if (core::g_use_interrupt_wait) { - if (private_segment_size == UINT_MAX) { - private_segment_size = - (profile_ == HSA_PROFILE_BASE) ? 0 : scratch_per_thread_; - } - - if (private_segment_size > 262128) { - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - - scratch.size_per_thread = AlignUp(private_segment_size, 16); - if (scratch.size_per_thread > 262128) { - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - - const uint32_t num_cu = - properties_.NumFComputeCores / properties_.NumSIMDPerCU; - scratch.size = scratch.size_per_thread * 32 * 64 * num_cu; - } else { - scratch.size = queue_scratch_len_; - scratch.size_per_thread = scratch_per_thread_; + if (private_segment_size == UINT_MAX) { + private_segment_size = (profile_ == HSA_PROFILE_BASE) ? 0 : scratch_per_thread_; } -#else - scratch.size = queue_scratch_len_; - scratch.size_per_thread = scratch_per_thread_; -#endif + + if (private_segment_size > 262128) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } + + scratch.size_per_thread = AlignUp(private_segment_size, 16); + if (scratch.size_per_thread > 262128) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } + + const uint32_t num_cu = properties_.NumFComputeCores / properties_.NumSIMDPerCU; + scratch.size = scratch.size_per_thread * 32 * 64 * num_cu; scratch.queue_base = NULL; if (scratch.size != 0) { AcquireQueueScratch(scratch); diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index 62060fd825..64129ac03f 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -182,6 +182,10 @@ static_assert(sizeof(void*) == 8, "HSA_LARGE_MODEL is set incorrectly!"); static_assert(sizeof(void*) == 4, "HSA_LARGE_MODEL is set incorrectly!"); #endif +#if !defined(HSA_LARGE_MODEL) || !defined(__linux__) +static_assert(false, "Only HSA_LARGE_MODEL (64bit mode) and Linux supported."); +#endif + namespace HSA { //---------------------------------------------------------------------------// diff --git a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp index 70daffd6af..901866eaf0 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp @@ -348,7 +348,7 @@ hsa_status_t core::Signal* signal = core::Signal::Convert(hsa_signal); IS_VALID(signal); IS_BAD_PTR(handler); - if (!core::InterruptSignal::IsType(signal)) + if (core::g_use_interrupt_wait && (!core::InterruptSignal::IsType(signal))) return HSA_STATUS_ERROR_INVALID_SIGNAL; return core::Runtime::runtime_singleton_->SetAsyncSignalHandler( hsa_signal, cond, value, handler, arg); diff --git a/runtime/hsa-runtime/core/runtime/runtime.cpp b/runtime/hsa-runtime/core/runtime/runtime.cpp index f0def7f6a6..dfdab08614 100755 --- a/runtime/hsa-runtime/core/runtime/runtime.cpp +++ b/runtime/hsa-runtime/core/runtime/runtime.cpp @@ -552,7 +552,7 @@ hsa_status_t Runtime::SetAsyncSignalHandler(hsa_signal_t signal, hsa_amd_signal_handler handler, void* arg) { // Asyncronous signal handler is only supported when KFD events are on. - if (!core::g_use_interrupt_wait) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + // if (!core::g_use_interrupt_wait) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; // Indicate that this signal is in use. if (signal.handle != 0) hsa_signal_handle(signal)->Retain();