From b6f0248f536ebd6021b6bd6a3faa7d37b881394d Mon Sep 17 00:00:00 2001 From: Sean Keely Date: Wed, 14 Mar 2018 03:52:44 -0500 Subject: [PATCH] Respect new memory model requirements at queue destroy. Spec requires GPU release fences and CPU acquire fences at queue destroy. Also update the recognized status codes. Change-Id: If9166f5149f65417c7057ff7c0f69f6ac094d6ab --- runtime/hsa-runtime/core/inc/amd_aql_queue.h | 3 + runtime/hsa-runtime/core/inc/queue.h | 3 +- .../core/runtime/amd_aql_queue.cpp | 29 +- runtime/hsa-runtime/core/runtime/hsa.cpp | 14 +- runtime/hsa-runtime/inc/hsa.h | 287 +++++++++--------- 5 files changed, 181 insertions(+), 155 deletions(-) diff --git a/runtime/hsa-runtime/core/inc/amd_aql_queue.h b/runtime/hsa-runtime/core/inc/amd_aql_queue.h index 04eaa3d904..a46471c820 100644 --- a/runtime/hsa-runtime/core/inc/amd_aql_queue.h +++ b/runtime/hsa-runtime/core/inc/amd_aql_queue.h @@ -208,6 +208,9 @@ class AqlQueue : public core::Queue, private core::LocalSignal, public core::Doo /// that enable kernel access scratch memory void InitScratchSRD(); + /// @brief Halt the queue without destroying it or fencing memory. + void Suspend(); + /// @brief Handler for hardware queue events. static bool DynamicScratchHandler(hsa_signal_value_t error_code, void* arg); diff --git a/runtime/hsa-runtime/core/inc/queue.h b/runtime/hsa-runtime/core/inc/queue.h index d5fc232095..bf148cc241 100644 --- a/runtime/hsa-runtime/core/inc/queue.h +++ b/runtime/hsa-runtime/core/inc/queue.h @@ -318,10 +318,11 @@ class Queue : public Checked<0xFA3906A679F9DB49>, private LocalQueue { virtual void do_set_public_handle(hsa_queue_t* handle) { public_handle_ = handle; } - hsa_queue_t* public_handle_; virtual bool _IsA(rtti_t id) const = 0; + hsa_queue_t* public_handle_; + private: DISALLOW_COPY_AND_ASSIGN(Queue); }; diff --git a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp index 5ad9140db2..9b72bfc3f9 100644 --- a/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp @@ -290,8 +290,7 @@ AqlQueue::~AqlQueue() { HSA::hsa_signal_store_relaxed(amd_queue_.queue_inactive_signal, 0x8000000000000000ull); } - auto err = hsaKmtDestroyQueue(queue_id_); - assert(err == HSAKMT_STATUS_SUCCESS && "hsaKmtDestroyQueue failed."); + Inactivate(); FreeRegisteredRingBuffer(); agent_->ReleaseQueueScratch(queue_scratch_.queue_base); HSA::hsa_signal_destroy(amd_queue_.queue_inactive_signal); @@ -686,11 +685,17 @@ int AqlQueue::CreateRingBufferFD(const char* ring_buf_shm_path, #endif } +void AqlQueue::Suspend() { + auto err = hsaKmtUpdateQueue(queue_id_, 0, HSA_QUEUE_PRIORITY_NORMAL, NULL, 0, NULL); + assert(err == HSAKMT_STATUS_SUCCESS && "hsaKmtUpdateQueue failed."); +} + hsa_status_t AqlQueue::Inactivate() { bool active = active_.exchange(false, std::memory_order_relaxed); if (active) { - auto err = hsaKmtUpdateQueue(queue_id_, 0, HSA_QUEUE_PRIORITY_NORMAL, NULL, 0, NULL); - assert(err == HSAKMT_STATUS_SUCCESS && "hsaKmtUpdateQueue failed."); + auto err = hsaKmtDestroyQueue(queue_id_); + assert(err == HSAKMT_STATUS_SUCCESS && "hsaKmtDestroyQueue failed."); + atomic::Fence(std::memory_order_acquire); } return HSA_STATUS_SUCCESS; } @@ -698,6 +703,7 @@ hsa_status_t AqlQueue::Inactivate() { bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { AqlQueue* queue = (AqlQueue*)arg; hsa_status_t errorCode = HSA_STATUS_SUCCESS; + bool fatal = false; // Process errors only if queue is not terminating. if ((queue->dynamicScratchState & ERROR_HANDLER_TERMINATE) != ERROR_HANDLER_TERMINATE) { @@ -750,10 +756,12 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { } 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; } if (errorCode == HSA_STATUS_SUCCESS) { @@ -761,12 +769,19 @@ bool AqlQueue::DynamicScratchHandler(hsa_signal_value_t error_code, void* arg) { return true; } - queue->Inactivate(); - if (queue->errors_callback_ != nullptr) + queue->Suspend(); + if (queue->errors_callback_ != nullptr) { queue->errors_callback_(errorCode, queue->public_handle(), queue->errors_data_); + } + if (fatal) { + //Temporarilly removed until there is clarity on exactly what debugtrap's semantics are. + //assert(false && "Fatal queue error"); + //std::abort(); + } } // Copy here is to protect against queue being released between setting the scratch state and - // updating the signal value. + // updating the signal value. The signal itself is safe to use because it is ref counted rather + // than being released with the queue. hsa_signal_t signal = queue->amd_queue_.queue_inactive_signal; queue->dynamicScratchState = ERROR_HANDLER_DONE; HSA::hsa_signal_store_screlease(signal, -1ull); diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index fdbdfb1538..5513ca2366 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -2626,22 +2626,26 @@ hsa_status_t hsa_status_string( break; case HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: *status_string = - "HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: *The code object reader is invalid."; + "HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER: The code object reader is invalid."; break; case HSA_STATUS_ERROR_INVALID_CACHE: - *status_string = "HSA_STATUS_ERROR_INVALID_CACHE: *The cache is invalid."; + *status_string = "HSA_STATUS_ERROR_INVALID_CACHE: The cache is invalid."; break; case HSA_STATUS_ERROR_INVALID_WAVEFRONT: - *status_string = "HSA_STATUS_ERROR_INVALID_WAVEFRONT: *The wavefront is invalid."; + *status_string = "HSA_STATUS_ERROR_INVALID_WAVEFRONT: The wavefront is invalid."; break; case HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP: - *status_string = "HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP: *The signal group is invalid."; + *status_string = "HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP: The signal group is invalid."; break; case HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: *status_string = - "HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: *The HSA runtime is not in the configuration " + "HSA_STATUS_ERROR_INVALID_RUNTIME_STATE: The HSA runtime is not in the configuration " "state."; break; + case HSA_STATUS_ERROR_FATAL: + *status_string = + "HSA_STATUS_ERROR_FATAL: The queue received an error that may require process " + "termination."; case HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED: *status_string = "HSA_EXT_STATUS_ERROR_IMAGE_FORMAT_UNSUPPORTED: Image " diff --git a/runtime/hsa-runtime/inc/hsa.h b/runtime/hsa-runtime/inc/hsa.h index b8ec90a667..0ed2b689e2 100644 --- a/runtime/hsa-runtime/inc/hsa.h +++ b/runtime/hsa-runtime/inc/hsa.h @@ -116,148 +116,151 @@ extern "C" { * @brief Status codes. */ typedef enum { - /** - * The function has been executed successfully. - */ - HSA_STATUS_SUCCESS = 0x0, - /** - * A traversal over a list of elements has been interrupted by the - * application before completing. - */ - HSA_STATUS_INFO_BREAK = 0x1, - /** - * A generic error has occurred. - */ - HSA_STATUS_ERROR = 0x1000, - /** - * One of the actual arguments does not meet a precondition stated in the - * documentation of the corresponding formal argument. - */ - HSA_STATUS_ERROR_INVALID_ARGUMENT = 0x1001, - /** - * The requested queue creation is not valid. - */ - HSA_STATUS_ERROR_INVALID_QUEUE_CREATION = 0x1002, - /** - * The requested allocation is not valid. - */ - HSA_STATUS_ERROR_INVALID_ALLOCATION = 0x1003, - /** - * The agent is invalid. - */ - HSA_STATUS_ERROR_INVALID_AGENT = 0x1004, - /** - * The memory region is invalid. - */ - HSA_STATUS_ERROR_INVALID_REGION = 0x1005, - /** - * The signal is invalid. - */ - HSA_STATUS_ERROR_INVALID_SIGNAL = 0x1006, - /** - * The queue is invalid. - */ - HSA_STATUS_ERROR_INVALID_QUEUE = 0x1007, - /** - * The HSA runtime failed to allocate the necessary resources. This error - * may also occur when the HSA runtime needs to spawn threads or create - * internal OS-specific events. - */ - HSA_STATUS_ERROR_OUT_OF_RESOURCES = 0x1008, - /** - * The AQL packet is malformed. - */ - HSA_STATUS_ERROR_INVALID_PACKET_FORMAT = 0x1009, - /** - * An error has been detected while releasing a resource. - */ - HSA_STATUS_ERROR_RESOURCE_FREE = 0x100A, - /** - * An API other than ::hsa_init has been invoked while the reference count - * of the HSA runtime is 0. - */ - HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B, - /** - * The maximum reference count for the object has been reached. - */ - HSA_STATUS_ERROR_REFCOUNT_OVERFLOW = 0x100C, - /** - * The arguments passed to a functions are not compatible. - */ - HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS = 0x100D, - /** - * The index is invalid. - */ - HSA_STATUS_ERROR_INVALID_INDEX = 0x100E, - /** - * The instruction set architecture is invalid. - */ - HSA_STATUS_ERROR_INVALID_ISA = 0x100F, - /** - * The instruction set architecture name is invalid. - */ - HSA_STATUS_ERROR_INVALID_ISA_NAME = 0x1017, - /** - * The code object is invalid. - */ - HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010, - /** - * The executable is invalid. - */ - HSA_STATUS_ERROR_INVALID_EXECUTABLE = 0x1011, - /** - * The executable is frozen. - */ - HSA_STATUS_ERROR_FROZEN_EXECUTABLE = 0x1012, - /** - * There is no symbol with the given name. - */ - HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 0x1013, - /** - * The variable is already defined. - */ - HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED = 0x1014, - /** - * The variable is undefined. - */ - HSA_STATUS_ERROR_VARIABLE_UNDEFINED = 0x1015, - /** - * An HSAIL operation resulted in a hardware exception. - */ - HSA_STATUS_ERROR_EXCEPTION = 0x1016, - /** - * The code object symbol is invalid. - */ - HSA_STATUS_ERROR_INVALID_CODE_SYMBOL = 0x1018, - /** - * The executable symbol is invalid. - */ - HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL = 0x1019, - /** - * The file descriptor is invalid. - */ - HSA_STATUS_ERROR_INVALID_FILE = 0x1020, - /** - * The code object reader is invalid. - */ - HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER = 0x1021, - /** - * The cache is invalid. - */ - HSA_STATUS_ERROR_INVALID_CACHE = 0x1022, - /** - * The wavefront is invalid. - */ - HSA_STATUS_ERROR_INVALID_WAVEFRONT = 0x1023, - /** - * The signal group is invalid. - */ - HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP = 0x1024, - /** - * The HSA runtime is not in the configuration state. - */ - HSA_STATUS_ERROR_INVALID_RUNTIME_STATE = 0x1025 - + /** + * The function has been executed successfully. + */ + HSA_STATUS_SUCCESS = 0x0, + /** + * A traversal over a list of elements has been interrupted by the + * application before completing. + */ + HSA_STATUS_INFO_BREAK = 0x1, + /** + * A generic error has occurred. + */ + HSA_STATUS_ERROR = 0x1000, + /** + * One of the actual arguments does not meet a precondition stated in the + * documentation of the corresponding formal argument. + */ + HSA_STATUS_ERROR_INVALID_ARGUMENT = 0x1001, + /** + * The requested queue creation is not valid. + */ + HSA_STATUS_ERROR_INVALID_QUEUE_CREATION = 0x1002, + /** + * The requested allocation is not valid. + */ + HSA_STATUS_ERROR_INVALID_ALLOCATION = 0x1003, + /** + * The agent is invalid. + */ + HSA_STATUS_ERROR_INVALID_AGENT = 0x1004, + /** + * The memory region is invalid. + */ + HSA_STATUS_ERROR_INVALID_REGION = 0x1005, + /** + * The signal is invalid. + */ + HSA_STATUS_ERROR_INVALID_SIGNAL = 0x1006, + /** + * The queue is invalid. + */ + HSA_STATUS_ERROR_INVALID_QUEUE = 0x1007, + /** + * The HSA runtime failed to allocate the necessary resources. This error + * may also occur when the HSA runtime needs to spawn threads or create + * internal OS-specific events. + */ + HSA_STATUS_ERROR_OUT_OF_RESOURCES = 0x1008, + /** + * The AQL packet is malformed. + */ + HSA_STATUS_ERROR_INVALID_PACKET_FORMAT = 0x1009, + /** + * An error has been detected while releasing a resource. + */ + HSA_STATUS_ERROR_RESOURCE_FREE = 0x100A, + /** + * An API other than ::hsa_init has been invoked while the reference count + * of the HSA runtime is 0. + */ + HSA_STATUS_ERROR_NOT_INITIALIZED = 0x100B, + /** + * The maximum reference count for the object has been reached. + */ + HSA_STATUS_ERROR_REFCOUNT_OVERFLOW = 0x100C, + /** + * The arguments passed to a functions are not compatible. + */ + HSA_STATUS_ERROR_INCOMPATIBLE_ARGUMENTS = 0x100D, + /** + * The index is invalid. + */ + HSA_STATUS_ERROR_INVALID_INDEX = 0x100E, + /** + * The instruction set architecture is invalid. + */ + HSA_STATUS_ERROR_INVALID_ISA = 0x100F, + /** + * The instruction set architecture name is invalid. + */ + HSA_STATUS_ERROR_INVALID_ISA_NAME = 0x1017, + /** + * The code object is invalid. + */ + HSA_STATUS_ERROR_INVALID_CODE_OBJECT = 0x1010, + /** + * The executable is invalid. + */ + HSA_STATUS_ERROR_INVALID_EXECUTABLE = 0x1011, + /** + * The executable is frozen. + */ + HSA_STATUS_ERROR_FROZEN_EXECUTABLE = 0x1012, + /** + * There is no symbol with the given name. + */ + HSA_STATUS_ERROR_INVALID_SYMBOL_NAME = 0x1013, + /** + * The variable is already defined. + */ + HSA_STATUS_ERROR_VARIABLE_ALREADY_DEFINED = 0x1014, + /** + * The variable is undefined. + */ + HSA_STATUS_ERROR_VARIABLE_UNDEFINED = 0x1015, + /** + * An HSAIL operation resulted in a hardware exception. + */ + HSA_STATUS_ERROR_EXCEPTION = 0x1016, + /** + * The code object symbol is invalid. + */ + HSA_STATUS_ERROR_INVALID_CODE_SYMBOL = 0x1018, + /** + * The executable symbol is invalid. + */ + HSA_STATUS_ERROR_INVALID_EXECUTABLE_SYMBOL = 0x1019, + /** + * The file descriptor is invalid. + */ + HSA_STATUS_ERROR_INVALID_FILE = 0x1020, + /** + * The code object reader is invalid. + */ + HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER = 0x1021, + /** + * The cache is invalid. + */ + HSA_STATUS_ERROR_INVALID_CACHE = 0x1022, + /** + * The wavefront is invalid. + */ + HSA_STATUS_ERROR_INVALID_WAVEFRONT = 0x1023, + /** + * The signal group is invalid. + */ + HSA_STATUS_ERROR_INVALID_SIGNAL_GROUP = 0x1024, + /** + * The HSA runtime is not in the configuration state. + */ + HSA_STATUS_ERROR_INVALID_RUNTIME_STATE = 0x1025, + /** + * The queue received an error that may require process termination. + */ + HSA_STATUS_ERROR_FATAL = 0x1026 } hsa_status_t; /**