From f312a7386e5f87acff4cfd94c9d7e59fefe33c40 Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Fri, 20 Oct 2017 07:12:04 -0500 Subject: [PATCH] Exception support for Queue. Remove "zombie" queue state and report queue creation failure via exceptions. Make Shared object a final container and support array objects with Shared. Add message printing to hsa_exception in debug builds. Change-Id: I459f38c80846018acbf45538874e95f91dd6b195 --- runtime/hsa-runtime/core/common/shared.h | 100 +++++++++++++++--- runtime/hsa-runtime/core/inc/amd_aql_queue.h | 14 +-- runtime/hsa-runtime/core/inc/host_queue.h | 3 - .../hsa-runtime/core/inc/intercept_queue.h | 7 +- runtime/hsa-runtime/core/inc/queue.h | 36 +++---- runtime/hsa-runtime/core/inc/signal.h | 1 - .../core/runtime/amd_aql_queue.cpp | 60 ++++++----- .../core/runtime/amd_gpu_agent.cpp | 21 ++-- .../hsa-runtime/core/runtime/host_queue.cpp | 25 ++--- runtime/hsa-runtime/core/runtime/hsa.cpp | 29 ++--- .../hsa-runtime/core/runtime/hsa_ext_amd.cpp | 12 +-- .../core/runtime/intercept_queue.cpp | 24 +---- runtime/hsa-runtime/core/util/utils.h | 11 +- 13 files changed, 181 insertions(+), 162 deletions(-) diff --git a/runtime/hsa-runtime/core/common/shared.h b/runtime/hsa-runtime/core/common/shared.h index 5ea97932d6..900b414fd9 100644 --- a/runtime/hsa-runtime/core/common/shared.h +++ b/runtime/hsa-runtime/core/common/shared.h @@ -52,7 +52,8 @@ namespace core { /// @brief Base class encapsulating the allocator and deallocator for -/// shared shared object. +/// shared shared object. As used this will allocate GPU visible host +/// memory mapped to all GPUs. class BaseShared { public: static void SetAllocateAndFree( @@ -67,10 +68,9 @@ class BaseShared { static std::function free_; }; -/// @brief Base class for classes that encapsulates object shared between -/// host and agents. Alignment defaults to __alignof(T) but may be increased. -template -class Shared : public BaseShared { +/// @brief Container for object located in GPU visible host memory. +/// Alignment defaults to __alignof(T) but may be increased. +template class Shared final : private BaseShared { public: Shared() { assert(allocate_ != nullptr && free_ != nullptr && @@ -80,29 +80,105 @@ class Shared : public BaseShared { shared_object_ = reinterpret_cast(allocate_(sizeof(T), Max(__alignof(T), Align), 0)); + if (shared_object_ == nullptr) throw std::bad_alloc(); - if (shared_object_ != NULL) new (shared_object_) T; + MAKE_NAMED_SCOPE_GUARD(throwGuard, [&]() { free_(shared_object_); }); + + new (shared_object_) T; + + throwGuard.Dismiss(); } - virtual ~Shared() { + ~Shared() { assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set"); - if (IsSharedObjectAllocationValid()) { + if (shared_object_ != nullptr) { shared_object_->~T(); free_(shared_object_); } } - T* shared_object() const { return shared_object_; } - - bool IsSharedObjectAllocationValid() const { - return (shared_object_ != NULL); + Shared(Shared&& rhs) { + this->~Shared(); + shared_object_ = rhs.shared_object_; + rhs.shared_object_ = nullptr; } + Shared& operator=(Shared&& rhs) { + this->~Shared(); + shared_object_ = rhs.shared_object_; + rhs.shared_object_ = nullptr; + return *this; + } + + T* shared_object() const { return shared_object_; } private: T* shared_object_; }; +/// @brief Container for array located in GPU visible host memory. +/// Alignment defaults to __alignof(T) but may be increased. +template class SharedArray final : private BaseShared { + public: + SharedArray() : shared_object_(nullptr) {} + + explicit SharedArray(size_t length) : shared_object_(nullptr), len(length) { + assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set"); + static_assert((__alignof(T) <= Align) || (Align == 0), "Align is less than alignof(T)"); + + shared_object_ = + reinterpret_cast(allocate_(sizeof(T) * length, Max(__alignof(T), Align), 0)); + if (shared_object_ == nullptr) throw std::bad_alloc(); + + size_t i = 0; + + MAKE_NAMED_SCOPE_GUARD(loopGuard, [&]() { + for (size_t t = 0; t < i - 1; t++) shared_object_[t].~T(); + free_(shared_object_); + }); + + for (; i < length; i++) new (&shared_object_[i]) T; + + loopGuard.Dismiss(); + } + + ~SharedArray() { + assert(allocate_ != nullptr && free_ != nullptr && "Shared object allocator is not set"); + + if (shared_object_ != nullptr) { + for (size_t i = 0; i < len; i++) shared_object_[i].~T(); + free_(shared_object_); + } + } + + SharedArray(SharedArray&& rhs) { + this->~SharedArray(); + shared_object_ = rhs.shared_object_; + rhs.shared_object_ = nullptr; + len = rhs.len; + } + SharedArray& operator=(SharedArray&& rhs) { + this->~SharedArray(); + shared_object_ = rhs.shared_object_; + rhs.shared_object_ = nullptr; + len = rhs.len; + return *this; + } + + T& operator[](size_t index) { + assert(index < len && "Index out of bounds."); + return shared_object_[index]; + } + const T& operator[](size_t index) const { + assert(index < len && "Index out of bounds."); + return shared_object_[index]; + } + + private: + T* shared_object_; + size_t len; +}; + } // namespace core #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/runtime/hsa-runtime/core/inc/amd_aql_queue.h index 4033ba61cd..5d551d4da8 100644 --- a/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -68,9 +68,6 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig ~AqlQueue(); - /// @brief Indicates if queue is valid or not - bool IsValid() const { return valid_; } - /// @brief Queue interfaces hsa_status_t Inactivate() override; @@ -368,11 +365,8 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig // Id of the Queue used in communication with thunk HSA_QUEUEID queue_id_; - // Indicates is queue is valid - bool valid_; - - // Indicates if queue is inactive - int32_t active_; + // Indicates if queue is active + std::atomic active_; // Cached value of HsaNodeProperties.HSA_CAPABILITY.DoorbellType int doorbell_type_; @@ -401,7 +395,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig static HsaEvent* queue_event_; // Queue count - used to ref count queue_event_ - static volatile uint32_t queue_count_; + static std::atomic queue_count_; // Mutex for queue_event_ manipulation static KernelMutex queue_lock_; @@ -411,5 +405,7 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Sig // Forbid copying and moving of this object DISALLOW_COPY_AND_ASSIGN(AqlQueue); }; + } // namespace amd + #endif // header guard diff --git a/runtime/hsa-runtime/core/inc/host_queue.h b/runtime/hsa-runtime/core/inc/host_queue.h index 5a883ed327..20ecd0d683 100644 --- a/runtime/hsa-runtime/core/inc/host_queue.h +++ b/runtime/hsa-runtime/core/inc/host_queue.h @@ -148,8 +148,6 @@ class HostQueue : public Queue { assert(false && "HostQueue::ExecutePM4 is unimplemented"); } - bool active() const { return active_; } - void* operator new(size_t size) { return _aligned_malloc(size, HSA_QUEUE_ALIGN_BYTES); } @@ -167,7 +165,6 @@ class HostQueue : public Queue { static int rtti_id_; static const size_t kRingAlignment = 256; const uint32_t size_; - bool active_; void* ring_; // Host queue id counter, starting from 0x80000000 to avoid overlaping diff --git a/runtime/hsa-runtime/core/inc/intercept_queue.h b/runtime/hsa-runtime/core/inc/intercept_queue.h index 2ad2abc23b..bc51842434 100644 --- a/runtime/hsa-runtime/core/inc/intercept_queue.h +++ b/runtime/hsa-runtime/core/inc/intercept_queue.h @@ -63,10 +63,6 @@ class QueueWrapper : public Queue { std::unique_ptr wrapped; explicit QueueWrapper(std::unique_ptr queue) : Queue(), wrapped(std::move(queue)) { - if (Queue::Shared::shared_object() == NULL) { - return; - } - memcpy(&amd_queue_, &wrapped->amd_queue_, sizeof(amd_queue_t)); wrapped->set_public_handle(wrapped.get(), public_handle_); } @@ -226,6 +222,9 @@ class InterceptQueue : public QueueProxy, private LocalSignal, public Signal { // Indicates queue active/inactive state. std::atomic active_; + // Proxy packet buffer + SharedArray buffer_; + static const hsa_signal_value_t DOORBELL_MAX = 0xFFFFFFFFFFFFFFFFull; static bool HandleAsyncDoorbell(hsa_signal_value_t value, void* arg); diff --git a/runtime/hsa-runtime/core/inc/queue.h b/runtime/hsa-runtime/core/inc/queue.h index 1394fd50af..98824ed5d4 100644 --- a/runtime/hsa-runtime/core/inc/queue.h +++ b/runtime/hsa-runtime/core/inc/queue.h @@ -124,6 +124,14 @@ struct SharedQueue { Queue* core_queue; }; +class LocalQueue { + public: + SharedQueue* queue() const { return local_queue_.shared_object(); } + + private: + Shared local_queue_; +}; + /// @brief Class Queue which encapsulate user mode queues and /// provides Api to access its Read, Write indices using Acquire, /// Release and Relaxed semantics. @@ -132,16 +140,10 @@ Queue is intended to be an pure interface class and may be wrapped or replaced by tools. All funtions other than Convert and public_handle must be virtual. */ -class Queue : public Checked<0xFA3906A679F9DB49>, - public Shared { +class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue { public: - Queue() : Shared(), amd_queue_(shared_object()->amd_queue) { - if (!Shared::IsSharedObjectAllocationValid()) { - return; - } - - shared_object()->core_queue = this; - + Queue() : LocalQueue(), amd_queue_(queue()->amd_queue) { + queue()->core_queue = this; public_handle_ = Convert(this); } @@ -153,9 +155,7 @@ class Queue : public Checked<0xFA3906A679F9DB49>, /// /// @return hsa_queue_t * Pointer to the public data type of a queue static __forceinline hsa_queue_t* Convert(Queue* queue) { - return ((queue != NULL) && (queue->IsSharedObjectAllocationValid())) - ? &queue->amd_queue_.hsa_queue - : NULL; + return (queue != nullptr) ? &queue->amd_queue_.hsa_queue : nullptr; } /// @brief Transform the public data type of a Queue's data type into an @@ -165,14 +165,10 @@ class Queue : public Checked<0xFA3906A679F9DB49>, /// /// @return Queue * Pointer to the Queue's implementation object static __forceinline Queue* Convert(const hsa_queue_t* queue) { - return (queue != NULL) - ? reinterpret_cast( - reinterpret_cast(queue) - - (reinterpret_cast( - &reinterpret_cast(1234) - ->amd_queue.hsa_queue) - - uintptr_t(1234)))->core_queue - : NULL; + return (queue != nullptr) + ? reinterpret_cast(reinterpret_cast(queue) - + offsetof(SharedQueue, amd_queue.hsa_queue))->core_queue + : nullptr; } /// @brief Inactivate the queue object. Once inactivate a diff --git a/runtime/hsa-runtime/core/inc/signal.h b/runtime/hsa-runtime/core/inc/signal.h index 028138a9a1..ddc31ebd73 100644 --- a/runtime/hsa-runtime/core/inc/signal.h +++ b/runtime/hsa-runtime/core/inc/signal.h @@ -101,7 +101,6 @@ static_assert(std::is_trivially_destructible::value, class LocalSignal { public: explicit LocalSignal(hsa_signal_value_t initial_value) { - if (!local_signal_.IsSharedObjectAllocationValid()) throw std::bad_alloc(); local_signal_.shared_object()->amd_signal.value = initial_value; } diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index c1f295eb33..2bab8840f4 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -73,30 +73,26 @@ namespace amd { const uint32_t kAmdQueueAlignBytes = 0x40; HsaEvent* AqlQueue::queue_event_ = NULL; -volatile uint32_t AqlQueue::queue_count_ = 0; +std::atomic AqlQueue::queue_count_(0); KernelMutex AqlQueue::queue_lock_; -int AqlQueue::rtti_id_; +int AqlQueue::rtti_id_ = 0; AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, ScratchInfo& scratch, core::HsaEventCallback callback, void* err_data, bool is_kv) : Queue(), LocalSignal(0), Signal(signal()), - ring_buf_(NULL), + ring_buf_(nullptr), ring_buf_alloc_bytes_(0), queue_id_(HSA_QUEUEID(-1)), - valid_(false), + active_(false), agent_(agent), queue_scratch_(scratch), errors_callback_(callback), errors_data_(err_data), is_kv_queue_(is_kv), - pm4_ib_buf_(NULL), + pm4_ib_buf_(nullptr), pm4_ib_size_b_(0x1000) { - if (!Queue::Shared::IsSharedObjectAllocationValid()) { - return; - } - // 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. @@ -122,14 +118,16 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr queue_size_pkts = Max(queue_size_pkts, min_pkts); uint32_t queue_size_bytes = queue_size_pkts * sizeof(core::AqlPacket); - if ((queue_size_bytes & (queue_size_bytes - 1)) != 0) return; + if ((queue_size_bytes & (queue_size_bytes - 1)) != 0) + throw AMD::hsa_exception(HSA_STATUS_ERROR_INVALID_QUEUE_CREATION, + "Requested queue with non-power of two packet capacity.\n"); // Allocate the AQL packet ring buffer. AllocRegisteredRingBuffer(queue_size_pkts); - if (ring_buf_ == NULL) return; + if (ring_buf_ == nullptr) throw std::bad_alloc(); MAKE_NAMED_SCOPE_GUARD(RingGuard, [&]() { FreeRegisteredRingBuffer(); }); - // Fill the ring buffer with ALWAYS_RESERVED packet headers. + // Fill the ring buffer with invalid packet headers. // Leave packet content uninitialized to help track errors. for (uint32_t pkt_id = 0; pkt_id < queue_size_pkts; ++pkt_id) { ((uint32_t*)ring_buf_)[16 * pkt_id] = HSA_PACKET_TYPE_INVALID; @@ -154,7 +152,9 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr kmt_status = hsaKmtCreateQueue(node_id, HSA_QUEUE_COMPUTE_AQL, 100, HSA_QUEUE_PRIORITY_NORMAL, ring_buf_, ring_buf_alloc_bytes_, NULL, &queue_rsrc); - if (kmt_status != HSAKMT_STATUS_SUCCESS) return; + if (kmt_status != HSAKMT_STATUS_SUCCESS) + throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, + "Queue create failed at hsaKmtCreateQueue\n"); queue_id_ = queue_rsrc.QueueId; MAKE_NAMED_SCOPE_GUARD(QueueGuard, [&]() { hsaKmtDestroyQueue(queue_id_); }); @@ -228,7 +228,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr queue_count_--; if (queue_count_ == 0) { core::InterruptSignal::DestroyEvent(queue_event_); - queue_event_ = NULL; + queue_event_ = nullptr; } }); @@ -239,35 +239,38 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr if (core::g_use_interrupt_wait) { ScopedAcquire _lock(&queue_lock_); queue_count_++; - if (queue_event_ == NULL) { + if (queue_event_ == nullptr) { 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 (queue_event_ == nullptr) + throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, + "Queue event creation failed.\n"); } auto Signal = new core::InterruptSignal(0, queue_event_); - if (Signal == nullptr) return; + assert(Signal != nullptr && "Should have thrown!\n"); amd_queue_.queue_inactive_signal = core::InterruptSignal::Convert(Signal); } else { EventGuard.Dismiss(); auto Signal = new core::DefaultSignal(0); - if (Signal == nullptr) return; + assert(Signal != nullptr && "Should have thrown!\n"); amd_queue_.queue_inactive_signal = core::DefaultSignal::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; + throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, + "Queue event handler failed registration.\n"); pm4_ib_buf_ = core::Runtime::runtime_singleton_->system_allocator()( pm4_ib_size_b_, 0x1000, core::MemoryRegion::AllocateExecutable); - if (pm4_ib_buf_ == NULL) return; + if (pm4_ib_buf_ == nullptr) + throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "PM4 IB allocation failed.\n"); MAKE_NAMED_SCOPE_GUARD(PM4IBGuard, [&]() { core::Runtime::runtime_singleton_->system_deallocator()(pm4_ib_buf_); }); - valid_ = true; - active_ = 1; + active_ = true; PM4IBGuard.Dismiss(); RingGuard.Dismiss(); @@ -277,11 +280,7 @@ AqlQueue::AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id, Scr } AqlQueue::~AqlQueue() { - if (!IsValid()) { - return; - } - - if (active_ == 1) hsaKmtDestroyQueue(queue_id_); + Inactivate(); FreeRegisteredRingBuffer(); agent_->ReleaseQueueScratch(queue_scratch_.queue_base); @@ -654,8 +653,11 @@ void AqlQueue::FreeRegisteredRingBuffer() { } hsa_status_t AqlQueue::Inactivate() { - int32_t active = atomic::Exchange((volatile int32_t*)&active_, 0); - if (active == 1) hsaKmtDestroyQueue(this->queue_id_); + bool active = active_.exchange(false, std::memory_order_relaxed); + if (active) { + auto err = hsaKmtDestroyQueue(this->queue_id_); + assert(err == HSAKMT_STATUS_SUCCESS && "hsaKmtDestroyQueue failed."); + } return HSA_STATUS_SUCCESS; } diff --git a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp index 00a4b66b53..9379bd8aac 100644 --- a/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp @@ -910,26 +910,21 @@ hsa_status_t GpuAgent::QueueCreate(size_t size, hsa_queue_type32_t queue_type, const uint32_t num_cu = properties_.NumFComputeCores / properties_.NumSIMDPerCU; scratch.size = scratch.size_per_thread * 32 * 64 * num_cu; - scratch.queue_base = NULL; + scratch.queue_base = nullptr; + + MAKE_NAMED_SCOPE_GUARD(scratchGuard, [&]() { ReleaseQueueScratch(scratch.queue_base); }); + if (scratch.size != 0) { AcquireQueueScratch(scratch); - if (scratch.queue_base == NULL) { + if (scratch.queue_base == nullptr) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } } // Create an HW AQL queue - AqlQueue* hw_queue = new AqlQueue(this, size, node_id(), scratch, - event_callback, data, is_kv_device_); - if (hw_queue && hw_queue->IsValid()) { - // return queue - *queue = hw_queue; - return HSA_STATUS_SUCCESS; - } - // If reached here its always an ERROR. - delete hw_queue; - ReleaseQueueScratch(scratch.queue_base); - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + *queue = new AqlQueue(this, size, node_id(), scratch, event_callback, data, is_kv_device_); + scratchGuard.Dismiss(); + return HSA_STATUS_SUCCESS; } void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { diff --git a/runtime/hsa-runtime/core/runtime/host_queue.cpp b/runtime/hsa-runtime/core/runtime/host_queue.cpp index ca6d2aca3c..ec9df377f9 100644 --- a/runtime/hsa-runtime/core/runtime/host_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/host_queue.cpp @@ -50,23 +50,19 @@ namespace core { int HostQueue::rtti_id_ = 0; std::atomic HostQueue::queue_count_(0x80000000); -HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, - hsa_queue_type32_t type, uint32_t features, - hsa_signal_t doorbell_signal) - : Queue(), - size_(ring_size), - active_(false) { - if (!Shared::IsSharedObjectAllocationValid()) { - return; - } - +HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type32_t type, + uint32_t features, hsa_signal_t doorbell_signal) + : Queue(), size_(ring_size) { HSA::hsa_memory_register(this, sizeof(HostQueue)); + MAKE_NAMED_SCOPE_GUARD(registerGuard, + [&]() { HSA::hsa_memory_deregister(this, sizeof(HostQueue)); }); const size_t queue_buffer_size = size_ * sizeof(AqlPacket); if (HSA_STATUS_SUCCESS != HSA::hsa_memory_allocate(region, queue_buffer_size, &ring_)) { - return; + throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "Host queue buffer alloc failed\n"); } + MAKE_NAMED_SCOPE_GUARD(bufferGuard, [&]() { HSA::hsa_memory_free(&ring_); }); assert(IsMultipleOf(ring_, kRingAlignment)); assert(ring_ != NULL); @@ -88,14 +84,11 @@ HostQueue::HostQueue(hsa_region_t region, uint32_t ring_size, AMD_HSA_BITS_SET( amd_queue_.queue_properties, AMD_QUEUE_PROPERTIES_ENABLE_PROFILING, 0); - active_ = true; + bufferGuard.Dismiss(); + registerGuard.Dismiss(); } HostQueue::~HostQueue() { - if (!Shared::IsSharedObjectAllocationValid()) { - return; - } - HSA::hsa_memory_free(ring_); HSA::hsa_memory_deregister(this, sizeof(HostQueue)); } diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index bf953d10b8..2e2b50841c 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -644,8 +644,8 @@ hsa_status_t hsa_queue_create( TRY; IS_OPEN(); - if ((queue == NULL) || (size == 0) || (!IsPowerOfTwo(size)) || - (type < HSA_QUEUE_TYPE_MULTI) || (type > HSA_QUEUE_TYPE_SINGLE)) { + if ((queue == nullptr) || (size == 0) || (!IsPowerOfTwo(size)) || (type < HSA_QUEUE_TYPE_MULTI) || + (type > HSA_QUEUE_TYPE_SINGLE)) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } @@ -662,22 +662,17 @@ hsa_status_t hsa_queue_create( return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION; } - if (callback == NULL) callback = core::Queue::DefaultErrorHandler; + if (callback == nullptr) callback = core::Queue::DefaultErrorHandler; - core::Queue* cmd_queue = NULL; + core::Queue* cmd_queue = nullptr; status = agent->QueueCreate(size, type, callback, data, private_segment_size, group_segment_size, &cmd_queue); - if (cmd_queue != NULL) { - *queue = core::Queue::Convert(cmd_queue); - if (*queue == NULL) { - delete cmd_queue; - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - } else { - *queue = NULL; - } + if (status != HSA_STATUS_SUCCESS) return status; + assert(cmd_queue != nullptr && "Queue not returned but status was success.\n"); + *queue = core::Queue::Convert(cmd_queue); return status; + CATCH; } @@ -701,13 +696,7 @@ hsa_status_t hsa_soft_queue_create(hsa_region_t region, uint32_t size, const core::Signal* signal = core::Signal::Convert(doorbell_signal); IS_VALID(signal); - core::HostQueue* host_queue = - new core::HostQueue(region, size, type, features, doorbell_signal); - - if (!host_queue->active()) { - delete host_queue; - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } + core::HostQueue* host_queue = new core::HostQueue(region, size, type, features, doorbell_signal); *queue = core::Queue::Convert(host_queue); diff --git a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp index 3819e4cdb9..9869451601 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ext_amd.cpp @@ -130,9 +130,9 @@ hsa_status_t handleException() { } catch (const std::bad_alloc& e) { return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } catch (const hsa_exception& e) { + debug_print("HSA exception: %s\n", e.what()); return e.error_code(); - // } catch (std::nested_exception& e) { - // Rethrow exceptions caught from callbacks. + // } catch (std::nested_exception& e) { // Rethrow exceptions from callbacks // e.rethrow_nested(); // return HSA_STATUS_ERROR; } catch (const std::exception& e) { @@ -803,6 +803,8 @@ hsa_status_t hsa_amd_queue_intercept_create( void (*callback)(hsa_status_t status, hsa_queue_t* source, void* data), void* data, uint32_t private_segment_size, uint32_t group_segment_size, hsa_queue_t** queue) { TRY; + IS_OPEN(); + IS_BAD_PTR(queue); hsa_queue_t* lower_queue; hsa_status_t err = HSA::hsa_queue_create(agent_handle, size, type, callback, data, private_segment_size, group_segment_size, &lower_queue); @@ -811,11 +813,7 @@ hsa_status_t hsa_amd_queue_intercept_create( std::unique_ptr upperQueue(new core::InterceptQueue(std::move(lowerQueue))); - if (!upperQueue->IsSharedObjectAllocationValid()) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - - *queue = core::Queue::Convert(upperQueue.get()); - upperQueue.release(); - + *queue = core::Queue::Convert(upperQueue.release()); return HSA_STATUS_SUCCESS; CATCH; } diff --git a/runtime/hsa-runtime/core/runtime/intercept_queue.cpp b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp index 387d352578..707effdad4 100644 --- a/runtime/hsa-runtime/core/runtime/intercept_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/intercept_queue.cpp @@ -75,20 +75,8 @@ InterceptQueue::InterceptQueue(std::unique_ptr queue) retry_index_(0), quit_(false), active_(true) { - if (Queue::Shared::shared_object() == NULL) { - return; - // TODO skeely: cleanup queue constructor failure. - // throw AMD::hsa_exception(HSA_STATUS_ERROR_OUT_OF_RESOURCES, "Failed to allocate Queue ABI - // block.\n"); - } - - void* buffer = Runtime::runtime_singleton_->system_allocator()( - wrapped->amd_queue_.hsa_queue.size * sizeof(hsa_agent_dispatch_packet_t), 4096, - MemoryRegion::AllocateNoFlags); - if (buffer == nullptr) throw std::bad_alloc(); - MAKE_NAMED_SCOPE_GUARD(buffGuard, - [&]() { Runtime::runtime_singleton_->system_deallocator()(buffer); }); - amd_queue_.hsa_queue.base_address = buffer; + buffer_ = SharedArray(wrapped->amd_queue_.hsa_queue.size); + amd_queue_.hsa_queue.base_address = reinterpret_cast(&buffer_[0]); // Match the queue's signal ABI block to async_doorbell_'s // This allows us to use the queue's signal ABI block from devices to trigger async_doorbell while @@ -109,14 +97,9 @@ InterceptQueue::InterceptQueue(std::unique_ptr queue) AddInterceptor(Submit, this); sigGuard.Dismiss(); - buffGuard.Dismiss(); } InterceptQueue::~InterceptQueue() { - if (Queue::Shared::shared_object() == NULL) { - return; - } - active_ = false; // Kill the async doorbell handler @@ -128,9 +111,6 @@ InterceptQueue::~InterceptQueue() { if (val != 0) async_doorbell_->WaitRelaxed(HSA_SIGNAL_CONDITION_EQ, 0, -1, HSA_WAIT_STATE_BLOCKED); async_doorbell_->DestroySignal(); - - // Release buffer resources - Runtime::runtime_singleton_->system_deallocator()(amd_queue_.hsa_queue.base_address); } bool InterceptQueue::HandleAsyncDoorbell(hsa_signal_value_t value, void* arg) { diff --git a/runtime/hsa-runtime/core/util/utils.h b/runtime/hsa-runtime/core/util/utils.h index 2da2d09668..ab279d7f2f 100755 --- a/runtime/hsa-runtime/core/util/utils.h +++ b/runtime/hsa-runtime/core/util/utils.h @@ -120,12 +120,11 @@ static __forceinline unsigned long long int strtoull(const char* str, #endif // A macro to disallow the copy and move constructor and operator= functions -// This should be used in the private: declarations for a class -#define DISALLOW_COPY_AND_ASSIGN(TypeName) \ - TypeName(const TypeName&); \ - TypeName(TypeName&&); \ - void operator=(const TypeName&); \ - void operator=(TypeName&&); +#define DISALLOW_COPY_AND_ASSIGN(TypeName) \ + TypeName(const TypeName&) = delete; \ + TypeName(TypeName&&) = delete; \ + void operator=(const TypeName&) = delete; \ + void operator=(TypeName&&) = delete; template class ScopeGuard {