diff --git a/openmp/libomptarget/plugins/amdgpu/impl/data.cpp b/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
--- a/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
@@ -48,7 +48,11 @@
atmi_status_t ret = ATMI_STATUS_SUCCESS;
hsa_amd_memory_pool_t pool = get_memory_pool_by_mem_place(place);
hsa_status_t err = hsa_amd_memory_pool_allocate(pool, size, 0, ptr);
- ErrorCheck("atmi_malloc", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "atmi_malloc",
+ get_error_string(err));
+ exit(1);
+ }
DEBUG_PRINT("Malloced [%s %d] %p\n",
place.dev_type == ATMI_DEVTYPE_CPU ? "CPU" : "GPU", place.dev_id,
*ptr);
@@ -64,7 +68,11 @@
atmi_status_t ret = ATMI_STATUS_SUCCESS;
hsa_status_t err;
err = hsa_amd_memory_pool_free(ptr);
- ErrorCheck("atmi_free", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "atmi_free",
+ get_error_string(err));
+ exit(1);
+ }
DEBUG_PRINT("Freed %p\n", ptr);
if (err != HSA_STATUS_SUCCESS)
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/internal.h b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
--- a/openmp/libomptarget/plugins/amdgpu/impl/internal.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
@@ -234,31 +234,4 @@
const char *get_error_string(hsa_status_t err);
const char *get_atmi_error_string(atmi_status_t err);
-#define ATMIErrorCheck(msg, status) \
- if (status != ATMI_STATUS_SUCCESS) { \
- printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, msg, \
- get_atmi_error_string(status)); \
- exit(1); \
- } else { \
- /* printf("%s succeeded.\n", #msg);*/ \
- }
-
-#define ErrorCheck(msg, status) \
- if (status != HSA_STATUS_SUCCESS) { \
- printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, msg, \
- get_error_string(status)); \
- exit(1); \
- } else { \
- /* printf("%s succeeded.\n", #msg);*/ \
- }
-
-#define ErrorCheckAndContinue(msg, status) \
- if (status != HSA_STATUS_SUCCESS) { \
- DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, msg, \
- get_error_string(status)); \
- continue; \
- } else { \
- /* printf("%s succeeded.\n", #msg);*/ \
- }
-
#endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
--- a/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/system.cpp
@@ -20,13 +20,6 @@
#include "msgpack.h"
-#define msgpackErrorCheck(msg, status) \
- if (status != 0) { \
- printf("[%s:%d] %s failed\n", __FILE__, __LINE__, msg); \
- return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; \
- } else { \
- }
-
typedef unsigned char *address;
/*
* Note descriptors.
@@ -193,7 +186,11 @@
agents.push_back(gpu_procs[i].agent());
}
err = hsa_amd_agents_allow_access(agents.size(), &agents[0], NULL, ptr);
- ErrorCheck("Allow agents ptr access", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Allow agents ptr access", get_error_string(err));
+ exit(1);
+ }
}
atmi_status_t Runtime::Initialize() {
@@ -202,7 +199,11 @@
return ATMI_STATUS_SUCCESS;
if (devtype == ATMI_DEVTYPE_ALL || devtype & ATMI_DEVTYPE_GPU) {
- ATMIErrorCheck("GPU context init", atl_init_gpu_context());
+ if (atl_init_gpu_context() != ATMI_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "GPU context init",
+ get_atmi_error_string(atl_init_gpu_context()));
+ exit(1);
+ }
}
atl_set_atmi_initialized();
@@ -214,7 +215,11 @@
for (uint32_t i = 0; i < g_executables.size(); i++) {
err = hsa_executable_destroy(g_executables[i]);
- ErrorCheck("Destroying executable", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Destroying executable", get_error_string(err));
+ exit(1);
+ }
}
for (uint32_t i = 0; i < SymbolInfoTable.size(); i++) {
@@ -228,7 +233,11 @@
atl_reset_atmi_initialized();
err = hsa_shut_down();
- ErrorCheck("Shutting down HSA", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Shutting down HSA",
+ get_error_string(err));
+ exit(1);
+ }
return ATMI_STATUS_SUCCESS;
}
@@ -252,12 +261,20 @@
err = hsa_amd_memory_pool_get_info(
memory_pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED,
&alloc_allowed);
- ErrorCheck("Alloc allowed in memory pool check", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Alloc allowed in memory pool check", get_error_string(err));
+ exit(1);
+ }
if (alloc_allowed) {
uint32_t global_flag = 0;
err = hsa_amd_memory_pool_get_info(
memory_pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &global_flag);
- ErrorCheck("Get memory pool info", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Get memory pool info", get_error_string(err));
+ exit(1);
+ }
if (HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED & global_flag) {
ATLMemory new_mem(memory_pool, *proc, ATMI_MEMTYPE_FINE_GRAINED);
proc->addMemory(new_mem);
@@ -278,28 +295,44 @@
hsa_status_t err = HSA_STATUS_SUCCESS;
hsa_device_type_t device_type;
err = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
- ErrorCheck("Get device type info", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Get device type info", get_error_string(err));
+ exit(1);
+ }
switch (device_type) {
case HSA_DEVICE_TYPE_CPU: {
;
ATLCPUProcessor new_proc(agent);
err = hsa_amd_agent_iterate_memory_pools(agent, get_memory_pool_info,
&new_proc);
- ErrorCheck("Iterate all memory pools", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Iterate all memory pools", get_error_string(err));
+ exit(1);
+ }
g_atl_machine.addProcessor(new_proc);
} break;
case HSA_DEVICE_TYPE_GPU: {
;
hsa_profile_t profile;
err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &profile);
- ErrorCheck("Query the agent profile", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Query the agent profile", get_error_string(err));
+ exit(1);
+ }
atmi_devtype_t gpu_type;
gpu_type =
(profile == HSA_PROFILE_FULL) ? ATMI_DEVTYPE_iGPU : ATMI_DEVTYPE_dGPU;
ATLGPUProcessor new_proc(agent, gpu_type);
err = hsa_amd_agent_iterate_memory_pools(agent, get_memory_pool_info,
&new_proc);
- ErrorCheck("Iterate all memory pools", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Iterate all memory pools", get_error_string(err));
+ exit(1);
+ }
g_atl_machine.addProcessor(new_proc);
} break;
case HSA_DEVICE_TYPE_DSP: {
@@ -353,7 +386,11 @@
if (err == HSA_STATUS_INFO_BREAK) {
err = HSA_STATUS_SUCCESS;
}
- ErrorCheck("Getting a gpu agent", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Getting a gpu agent",
+ get_error_string(err));
+ exit(1);
+ }
if (err != HSA_STATUS_SUCCESS)
return err;
@@ -489,7 +526,12 @@
}
err = (atl_cpu_kernarg_region.handle == (uint64_t)-1) ? HSA_STATUS_ERROR
: HSA_STATUS_SUCCESS;
- ErrorCheck("Finding a CPU kernarg memory region handle", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Finding a CPU kernarg memory region handle",
+ get_error_string(err));
+ exit(1);
+ }
}
/* Find a memory region that supports kernel arguments. */
atl_gpu_kernarg_region.handle = (uint64_t)-1;
@@ -498,7 +540,11 @@
&atl_gpu_kernarg_region);
err = (atl_gpu_kernarg_region.handle == (uint64_t)-1) ? HSA_STATUS_ERROR
: HSA_STATUS_SUCCESS;
- ErrorCheck("Finding a kernarg memory region", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Finding a kernarg memory region", get_error_string(err));
+ exit(1);
+ }
}
if (num_procs > 0)
return HSA_STATUS_SUCCESS;
@@ -510,14 +556,22 @@
if (atlc.g_hsa_initialized == false) {
DEBUG_PRINT("Initializing HSA...");
hsa_status_t err = hsa_init();
- ErrorCheck("Initializing the hsa runtime", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Initializing the hsa runtime", get_error_string(err));
+ exit(1);
+ }
if (err != HSA_STATUS_SUCCESS)
return err;
err = init_compute_and_memory();
if (err != HSA_STATUS_SUCCESS)
return err;
- ErrorCheck("After initializing compute and memory", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "After initializing compute and memory", get_error_string(err));
+ exit(1);
+ }
int gpu_count = g_atl_machine.processorCount<ATLGPUProcessor>();
KernelInfoTable.resize(gpu_count);
@@ -602,7 +656,11 @@
}
err = hsa_amd_register_system_event_handler(callbackEvent, NULL);
- ErrorCheck("Registering the system for memory faults", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Registering the system for memory faults", get_error_string(err));
+ exit(1);
+ }
init_tasks();
atlc.g_gpu_initialized = true;
@@ -817,7 +875,11 @@
msgpack_errors =
map_lookup_array({metadata.first, metadata.second}, "amdhsa.kernels",
&kernel_array, &kernelsSize);
- msgpackErrorCheck("kernels lookup in program metadata", msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "kernels lookup in program metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
for (size_t i = 0; i < kernelsSize; i++) {
assert(msgpack_errors == 0);
@@ -826,43 +888,70 @@
msgpack::byte_range element;
msgpack_errors += array_lookup_element(kernel_array, i, &element);
- msgpackErrorCheck("element lookup in kernel metadata", msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "element lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
msgpack_errors += map_lookup_string(element, ".name", &kernelName);
msgpack_errors += map_lookup_string(element, ".symbol", &symbolName);
- msgpackErrorCheck("strings lookup in kernel metadata", msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "strings lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
atl_kernel_info_t info = {0, 0, 0, 0, 0, 0, 0, 0, 0, {}, {}, {}};
uint64_t sgpr_count, vgpr_count, sgpr_spill_count, vgpr_spill_count;
msgpack_errors += map_lookup_uint64_t(element, ".sgpr_count", &sgpr_count);
- msgpackErrorCheck("sgpr count metadata lookup in kernel metadata",
- msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "sgpr count metadata lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
+
info.sgpr_count = sgpr_count;
msgpack_errors += map_lookup_uint64_t(element, ".vgpr_count", &vgpr_count);
- msgpackErrorCheck("vgpr count metadata lookup in kernel metadata",
- msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "vgpr count metadata lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
+
info.vgpr_count = vgpr_count;
msgpack_errors +=
map_lookup_uint64_t(element, ".sgpr_spill_count", &sgpr_spill_count);
- msgpackErrorCheck("sgpr spill count metadata lookup in kernel metadata",
- msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "sgpr spill count metadata lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
+
info.sgpr_spill_count = sgpr_spill_count;
msgpack_errors +=
map_lookup_uint64_t(element, ".vgpr_spill_count", &vgpr_spill_count);
- msgpackErrorCheck("vgpr spill count metadata lookup in kernel metadata",
- msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "vgpr spill count metadata lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
+
info.vgpr_spill_count = vgpr_spill_count;
size_t kernel_explicit_args_size = 0;
uint64_t kernel_segment_size;
msgpack_errors += map_lookup_uint64_t(element, ".kernarg_segment_size",
&kernel_segment_size);
- msgpackErrorCheck("kernarg segment size metadata lookup in kernel metadata",
- msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "kernarg segment size metadata lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
// create a map from symbol to name
DEBUG_PRINT("Kernel symbol %s; Name: %s; Size: %lu\n", symbolName.c_str(),
@@ -877,8 +966,11 @@
msgpack::byte_range args_array;
msgpack_errors +=
map_lookup_array(element, ".args", &args_array, &argsSize);
- msgpackErrorCheck("kernel args metadata lookup in kernel metadata",
- msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "kernel args metadata lookup in kernel metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
info.num_args = argsSize;
@@ -887,13 +979,18 @@
msgpack::byte_range args_element;
msgpack_errors += array_lookup_element(args_array, i, &args_element);
- msgpackErrorCheck("iterate args map in kernel args metadata",
- msgpack_errors);
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "iterate args map in kernel args metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
msgpack_errors += populate_kernelArgMD(args_element, &lcArg);
- msgpackErrorCheck("iterate args map in kernel args metadata",
- msgpack_errors);
-
+ if (msgpack_errors != 0) {
+ printf("[%s:%d] %s failed\n", __FILE__, __LINE__,
+ "iterate args map in kernel args metadata");
+ return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+ }
// populate info with sizes and offsets
info.arg_sizes.push_back(lcArg.size_);
// v3 has offset field and not align field
@@ -942,23 +1039,40 @@
hsa_status_t err;
err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE,
&type);
- ErrorCheck("Symbol info extraction", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Symbol info extraction", get_error_string(err));
+ exit(1);
+ }
DEBUG_PRINT("Exec Symbol type: %d\n", type);
if (type == HSA_SYMBOL_KIND_KERNEL) {
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length);
- ErrorCheck("Symbol info extraction", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Symbol info extraction", get_error_string(err));
+ exit(1);
+ }
char *name = reinterpret_cast<char *>(malloc(name_length + 1));
err = hsa_executable_symbol_get_info(symbol,
HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
- ErrorCheck("Symbol info extraction", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Symbol info extraction", get_error_string(err));
+ exit(1);
+ }
name[name_length] = 0;
if (KernelNameMap.find(std::string(name)) == KernelNameMap.end()) {
// did not find kernel name in the kernel map; this can happen only
// if the ROCr API for getting symbol info (name) is different from
// the comgr method of getting symbol info
- ErrorCheck("Invalid kernel name", HSA_STATUS_ERROR_INVALID_CODE_OBJECT);
+ if (HSA_STATUS_ERROR_INVALID_CODE_OBJECT != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Invalid kernel name",
+ get_error_string(HSA_STATUS_ERROR_INVALID_CODE_OBJECT));
+ exit(1);
+ }
}
atl_kernel_info_t info;
std::string kernelName = KernelNameMap[std::string(name)];
@@ -966,8 +1080,12 @@
// because the non-ROCr custom code object parsing is called before
// iterating over the code object symbols using ROCr
if (KernelInfoTable[gpu].find(kernelName) == KernelInfoTable[gpu].end()) {
- ErrorCheck("Finding the entry kernel info table",
- HSA_STATUS_ERROR_INVALID_CODE_OBJECT);
+ if (HSA_STATUS_ERROR_INVALID_CODE_OBJECT != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Finding the entry kernel info table",
+ get_error_string(HSA_STATUS_ERROR_INVALID_CODE_OBJECT));
+ exit(1);
+ }
}
// found, so assign and update
info = KernelInfoTable[gpu][kernelName];
@@ -976,15 +1094,30 @@
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
&(info.kernel_object));
- ErrorCheck("Extracting the symbol from the executable", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Extracting the symbol from the executable",
+ get_error_string(err));
+ exit(1);
+ }
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
&(info.group_segment_size));
- ErrorCheck("Extracting the group segment size from the executable", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Extracting the group segment size from the executable",
+ get_error_string(err));
+ exit(1);
+ }
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
&(info.private_segment_size));
- ErrorCheck("Extracting the private segment from the executable", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Extracting the private segment from the executable",
+ get_error_string(err));
+ exit(1);
+ }
DEBUG_PRINT(
"Kernel %s --> %lx symbol %u group segsize %u pvt segsize %u bytes "
@@ -998,22 +1131,38 @@
} else if (type == HSA_SYMBOL_KIND_VARIABLE) {
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &name_length);
- ErrorCheck("Symbol info extraction", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Symbol info extraction", get_error_string(err));
+ exit(1);
+ }
char *name = reinterpret_cast<char *>(malloc(name_length + 1));
err = hsa_executable_symbol_get_info(symbol,
HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
- ErrorCheck("Symbol info extraction", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Symbol info extraction", get_error_string(err));
+ exit(1);
+ }
name[name_length] = 0;
atl_symbol_info_t info;
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &(info.addr));
- ErrorCheck("Symbol info address extraction", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Symbol info address extraction", get_error_string(err));
+ exit(1);
+ }
err = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &(info.size));
- ErrorCheck("Symbol info size extraction", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Symbol info size extraction", get_error_string(err));
+ exit(1);
+ }
atmi_mem_place_t place = ATMI_MEM_PLACE(ATMI_DEVTYPE_GPU, gpu, 0);
DEBUG_PRINT("Symbol %s = %p (%u bytes)\n", name, (void *)info.addr,
@@ -1046,14 +1195,22 @@
hsa_profile_t agent_profile;
err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_profile);
- ErrorCheck("Query the agent profile", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Query the agent profile", get_error_string(err));
+ exit(1);
+ }
// FIXME: Assume that every profile is FULL until we understand how to build
// GCN with base profile
agent_profile = HSA_PROFILE_FULL;
/* Create the empty executable. */
err = hsa_executable_create(agent_profile, HSA_EXECUTABLE_STATE_UNFROZEN, "",
&executable);
- ErrorCheck("Create the executable", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Create the executable", get_error_string(err));
+ exit(1);
+ }
bool module_load_success = false;
do // Existing control flow used continue, preserve that for this patch
@@ -1063,13 +1220,22 @@
// code object metadata parsing to collect such metadata info
err = get_code_object_custom_metadata(module_bytes, module_size, gpu);
- ErrorCheckAndContinue("Getting custom code object metadata", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Getting custom code object metadata",
+ get_error_string(err));
+ continue;
+ }
// Deserialize code object.
hsa_code_object_t code_object = {0};
err = hsa_code_object_deserialize(module_bytes, module_size, NULL,
&code_object);
- ErrorCheckAndContinue("Code Object Deserialization", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Code Object Deserialization", get_error_string(err));
+ continue;
+ }
assert(0 != code_object.handle);
// Mutating the device image here avoids another allocation & memcpy
@@ -1077,12 +1243,21 @@
reinterpret_cast<void *>(code_object.handle);
atmi_status_t atmi_err =
on_deserialized_data(code_object_alloc_data, module_size, cb_state);
- ATMIErrorCheck("Error in deserialized_data callback", atmi_err);
+ if (atmi_err != ATMI_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Error in deserialized_data callback",
+ get_atmi_error_string(atmi_err));
+ exit(1);
+ }
/* Load the code object. */
err =
hsa_executable_load_code_object(executable, agent, code_object, NULL);
- ErrorCheckAndContinue("Loading the code object", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Loading the code object", get_error_string(err));
+ continue;
+ }
// cannot iterate over symbols until executable is frozen
}
@@ -1092,11 +1267,19 @@
if (module_load_success) {
/* Freeze the executable; it can now be queried for symbols. */
err = hsa_executable_freeze(executable, "");
- ErrorCheck("Freeze the executable", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Freeze the executable", get_error_string(err));
+ exit(1);
+ }
err = hsa_executable_iterate_symbols(executable, populate_InfoTables,
static_cast<void *>(&gpu));
- ErrorCheck("Iterating over symbols for execuatable", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Iterating over symbols for execuatable", get_error_string(err));
+ exit(1);
+ }
// save the executable and destroy during finalize
g_executables.push_back(executable);
diff --git a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -76,12 +76,8 @@
#ifdef OMPTARGET_DEBUG
#define check(msg, status) \
if (status != ATMI_STATUS_SUCCESS) { \
- /* fprintf(stderr, "[%s:%d] %s failed.\n", __FILE__, __LINE__, #msg);*/ \
DP(#msg " failed\n"); \
- /*assert(0);*/ \
} else { \
- /* fprintf(stderr, "[%s:%d] %s succeeded.\n", __FILE__, __LINE__, #msg); \
- */ \
DP(#msg " succeeded\n"); \
}
#else
@@ -121,7 +117,11 @@
if (kernarg_region) {
auto r = hsa_amd_memory_pool_free(kernarg_region);
assert(r == HSA_STATUS_SUCCESS);
- ErrorCheck("Memory pool free", r);
+ if (r != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Memory pool free", get_error_string(r));
+ exit(1);
+ }
}
}
@@ -140,7 +140,12 @@
atl_gpu_kernarg_pools[0],
kernarg_size_including_implicit() * MAX_NUM_KERNELS, 0,
&kernarg_region);
- ErrorCheck("Allocating memory for the executable-kernel", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Allocating memory for the executable-kernel",
+ get_error_string(err));
+ exit(1);
+ }
core::allow_access_to_all_gpu_agents(kernarg_region);
for (int i = 0; i < MAX_NUM_KERNELS; i++) {
@@ -473,7 +478,12 @@
hsa_status_t err;
err = hsa_agent_get_info(HSAAgents[i], HSA_AGENT_INFO_QUEUE_MAX_SIZE,
&queue_size);
- ErrorCheck("Querying the agent maximum queue size", err);
+ if (err != HSA_STATUS_SUCCESS) {
+ printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+ "Querying the agent maximum queue size",
+ get_error_string(err));
+ exit(1);
+ }
if (queue_size > core::Runtime::getInstance().getMaxQueueSize()) {
queue_size = core::Runtime::getInstance().getMaxQueueSize();
}