From b75a8d9ed0e2c8d96643e47121e6b0391d62351b Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Tue, 20 Dec 2016 05:50:15 -0600 Subject: [PATCH] Allow reducing max occupancy (max scratch waves) when applications request large amounts of scratch. Also emit error messages to stderr if no async queue error callback was registered and queue fault messages are enabled (on by default). Queue fault messages are controlled with env key HSA_ENABLE_QUEUE_FAULT_MESSAGE. Change-Id: I496487b8d048b83aa95b9784e92928211f167b17 [ROCm/ROCR-Runtime commit: 0e17cc28879b2baf5186d232d3053a9dc1651058] --- .../runtime/hsa-runtime/CMakeLists.txt | 1 + .../runtime/hsa-runtime/core/inc/queue.h | 3 + .../core/runtime/amd_aql_queue.cpp | 5 +- .../core/runtime/amd_gpu_agent.cpp | 57 ++++++++++++++++--- .../core/runtime/amd_memory_region.cpp | 1 + .../runtime/hsa-runtime/core/runtime/hsa.cpp | 2 + .../hsa-runtime/core/runtime/queue.cpp | 56 ++++++++++++++++++ .../runtime/hsa-runtime/core/util/flag.h | 3 +- 8 files changed, 114 insertions(+), 14 deletions(-) create mode 100644 projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp diff --git a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt b/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt index 5561d0cf2b..00ac394805 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt +++ b/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt @@ -149,6 +149,7 @@ set ( SRCS "core/util/lnx/os_linux.cpp" "core/runtime/isa.cpp" "core/runtime/runtime.cpp" "core/runtime/signal.cpp" + "core/runtime/queue.cpp" "core/runtime/cache.cpp" "core/common/shared.cpp" "core/common/hsa_table_interface.cpp" diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h index b73294faf8..5ec43bd7b9 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/queue.h @@ -300,6 +300,9 @@ class Queue : public Checked<0xFA3906A679F9DB49>, // @brief Submits a block of PM4 and waits until it has been executed. virtual void ExecutePM4(uint32_t* cmd_data, size_t cmd_size_b) = 0; + /// @ brief Reports async queue errors to stderr if no other error handler was registered. + static void DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data); + // Handle of AMD Queue struct amd_queue_t& amd_queue_; 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 3b01d2bc97..1d1633f66e 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 @@ -683,10 +683,7 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { // Align whole waves to 1KB. scratch.size_per_thread = AlignUp(scratch.size_per_thread, 16); scratch.size = scratch.size_per_thread * (queue->amd_queue_.max_cu_id + 1) * - 32 * 64; // TODO: replace constants. - - // printf("Growing scratch to %u - %u\n", uint32_t(scratch.size_per_thread), - // uint32_t(scratch.size)); + queue->agent_->properties().MaxSlotsScratchCU * queue->agent_->properties().WaveFrontSize; queue->agent_->AcquireQueueScratch(scratch); if (scratch.queue_base == NULL) { 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 e698099d18..9fcf50bf73 100755 --- 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 @@ -921,18 +921,57 @@ void GpuAgent::AcquireQueueScratch(ScratchInfo& scratch) { ScopedAcquire lock(&scratch_lock_); scratch.queue_base = scratch_pool_.alloc(scratch.size); - scratch.queue_process_offset = - uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base()); + scratch.queue_process_offset = uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base()); - if ((scratch.queue_base != NULL) && (profile_ == HSA_PROFILE_BASE)) { - HSAuint64 alternate_va; - if (HSAKMT_STATUS_SUCCESS != - hsaKmtMapMemoryToGPU(scratch.queue_base, scratch.size, &alternate_va)) { - assert(false && "Map scratch subrange failed!"); - scratch_pool_.free(scratch.queue_base); - scratch.queue_base = NULL; + if (scratch.queue_base != NULL) { + if (profile_ == HSA_PROFILE_FULL) return; + if (profile_ == HSA_PROFILE_BASE) { + HSAuint64 alternate_va; + if (HSAKMT_STATUS_SUCCESS == + hsaKmtMapMemoryToGPU(scratch.queue_base, scratch.size, &alternate_va)) + return; } } + + // Scratch request failed allocation or mapping. + scratch_pool_.free(scratch.queue_base); + scratch.queue_base = NULL; + +// Attempt to trim the maximum number of concurrent waves to allow scratch to fit. +// This is somewhat dangerous as it limits the number of concurrent waves from future dispatches +// on the queue if those waves use even small amounts of scratch. +#ifndef NDEBUG + if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) + fprintf(stderr, "Failed to map requested scratch - reducing queue occupancy.\n"); +#endif + uint64_t num_cus = properties_.NumFComputeCores / properties_.NumSIMDPerCU; + uint64_t size_per_wave = AlignUp(scratch.size_per_thread * properties_.WaveFrontSize, 1024); + uint64_t total_waves = scratch.size / size_per_wave; + uint64_t waves_per_cu = total_waves / num_cus; + while (waves_per_cu != 0) { + size_t size = waves_per_cu * num_cus * size_per_wave; + void* base = scratch_pool_.alloc(size); + HSAuint64 alternate_va; + if ((base != NULL) && + ((profile_ == HSA_PROFILE_FULL) || + (hsaKmtMapMemoryToGPU(base, size, &alternate_va) == HSAKMT_STATUS_SUCCESS))) { + // Scratch allocated and either full profile or map succeeded. + scratch.queue_base = base; + scratch.size = size; + scratch.queue_process_offset = + uintptr_t(scratch.queue_base) - uintptr_t(scratch_pool_.base()); + return; + } + scratch_pool_.free(base); + waves_per_cu--; + } + + // Failed to allocate minimal scratch + assert(scratch.queue_base == NULL && "bad scratch data"); +#ifndef NDEBUG + if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) + fprintf(stderr, "Could not allocate scratch for one wave per CU.\n"); +#endif } void GpuAgent::ReleaseQueueScratch(void* base) { diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp index 95b4f259ab..295be378d9 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_memory_region.cpp @@ -281,6 +281,7 @@ hsa_status_t MemoryRegion::GetInfo(hsa_region_info_t attribute, case HSA_HEAPTYPE_FRAME_BUFFER_PRIVATE: case HSA_HEAPTYPE_FRAME_BUFFER_PUBLIC: case HSA_HEAPTYPE_SYSTEM: + case HSA_HEAPTYPE_GPU_SCRATCH: *((size_t*)value) = max_single_alloc_size_; break; default: diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp index 66927b5fe3..08852b9b40 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -564,6 +564,8 @@ hsa_status_t hsa_queue_create( return HSA_STATUS_ERROR_INVALID_QUEUE_CREATION; } + if (callback == NULL) callback = core::Queue::DefaultErrorHandler; + core::Queue* cmd_queue = NULL; status = agent->QueueCreate(size, type, callback, data, private_segment_size, group_segment_size, &cmd_queue); diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp new file mode 100644 index 0000000000..9246d4a04d --- /dev/null +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/queue.cpp @@ -0,0 +1,56 @@ +//////////////////////////////////////////////////////////////////////////////// +// +// The University of Illinois/NCSA +// Open Source License (NCSA) +// +// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved. +// +// Developed by: +// +// AMD Research and AMD HSA Software Development +// +// Advanced Micro Devices, Inc. +// +// www.amd.com +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to +// deal with the Software without restriction, including without limitation +// the rights to use, copy, modify, merge, publish, distribute, sublicense, +// and/or sell copies of the Software, and to permit persons to whom the +// Software is furnished to do so, subject to the following conditions: +// +// - Redistributions of source code must retain the above copyright notice, +// this list of conditions and the following disclaimers. +// - Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimers in +// the documentation and/or other materials provided with the distribution. +// - Neither the names of Advanced Micro Devices, Inc, +// nor the names of its contributors may be used to endorse or promote +// products derived from this Software without specific prior written +// permission. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL +// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR +// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, +// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER +// DEALINGS WITH THE SOFTWARE. +// +//////////////////////////////////////////////////////////////////////////////// + +#include "core/inc/queue.h" +#include "core/inc/runtime.h" + +namespace core { + +void Queue::DefaultErrorHandler(hsa_status_t status, hsa_queue_t* source, void* data) { + if (core::Runtime::runtime_singleton_->flag().enable_queue_fault_message()) { + const char* msg = "UNKNOWN ERROR"; + HSA::hsa_status_string(status, &msg); + fprintf(stderr, "Queue at %p inactivated due to async error:\n\t%s\n", source, msg); + } +} + +} 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 1e1fe0f833..d7add470c5 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/util/flag.h @@ -64,7 +64,7 @@ class Flag { enable_vm_fault_message_ = (var == "0") ? false : true; var = os::GetEnvVar("HSA_ENABLE_QUEUE_FAULT_MESSAGE"); - enable_queue_fault_message_ = (var == "1") ? true : false; + enable_queue_fault_message_ = (var == "0") ? false : true; var = os::GetEnvVar("HSA_ENABLE_INTERRUPT"); enable_interrupt_ = (var == "0") ? false : true; @@ -105,6 +105,7 @@ class Flag { bool enable_interrupt() const { return enable_interrupt_; } bool enable_thread_trace() const { return enable_thread_trace_; } + bool thread_trace_buff_size() const { return thread_trace_buff_size_; } bool enable_sdma() const { return enable_sdma_; }