AI/rocm-systems
2
0
Çatalla 0
Kod Konular Değişiklik İstekleri İşlemler Paketler Projeler Sürüm Viki Aktivite

Respect new memory model requirements at queue destroy.

Kaynağa Gözat 
Spec requires GPU release fences and CPU acquire fences at queue destroy.
Also update the recognized status codes.

Change-Id: If9166f5149f65417c7057ff7c0f69f6ac094d6ab
Bu işleme şunda yer alıyor:
Sean Keely
2018-03-14 03:52:44 -05:00
ebeveyn 6df9ba97ce
işleme b6f0248f53
5 değiştirilmiş dosya ile 181 ekleme ve 155 silme
Yama Dosyasını İndir Diff Dosyasını İndir Tüm dosyaları genişlet Tüm dosyaları daralt
runtime/hsa-runtime/core/inc/amd_aql_queue.h
+3
Dosyayı Görüntüle
@@ -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);
runtime/hsa-runtime/core/inc/queue.h
+2 -1
Dosyayı Görüntüle
@@ -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);
};
runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp
+22 -7
Dosyayı Görüntüle
@@ -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);
runtime/hsa-runtime/core/runtime/hsa.cpp
+9 -5
Dosyayı Görüntüle
@@ -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 "
runtime/hsa-runtime/inc/hsa.h
+145 -142
Dosyayı Görüntüle
@@ -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;
/**