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

Support async. queue errors and dynamic scratch without KFD events.

Browse Source 
Change-Id: I4e9e7a37aa7b9c96b28ce79f562760283e02b1e0
This commit is contained in:
Sean Keely
2017-03-28 03:12:59 -05:00
parent c4544906b9
commit 7dfeee5074
5 changed files with 38 additions and 54 deletions
Download Patch File Download Diff File Expand all files Collapse all files
runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+18 -26
View File
@@ -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<KernelMutex> _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<KernelMutex> _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<KernelMutex> lock(&queue_lock_);
queue_count_--;
@@ -307,8 +301,6 @@ AqlQueue::~AqlQueue() {
queue_event_ = NULL;
}
}
#endif
core::Runtime::runtime_singleton_->system_deallocator()(pm4_ib_buf_);
}
runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp
+14 -26
View File
@@ -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);
runtime/hsa-runtime/core/runtime/hsa.cpp
+4
View File
@@ -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 {
//---------------------------------------------------------------------------//
runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp
+1 -1
View File
@@ -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);
runtime/hsa-runtime/core/runtime/runtime.cpp
+1 -1
View File
@@ -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();