[Openmp-commits] [openmp] 9934571 - [libomptarget][amdgpu][nfc] Expand errorcheck macros
Jon Chesterfield via Openmp-commits
openmp-commits at lists.llvm.org
Wed May 12 09:30:57 PDT 2021
Author: Jon Chesterfield
Date: 2021-05-12T17:30:41+01:00
New Revision: 9934571eab9c6b3be22c4c8857bd3f4280b77843
URL: https://github.com/llvm/llvm-project/commit/9934571eab9c6b3be22c4c8857bd3f4280b77843
DIFF: https://github.com/llvm/llvm-project/commit/9934571eab9c6b3be22c4c8857bd3f4280b77843.diff
LOG: [libomptarget][amdgpu][nfc] Expand errorcheck macros
[libomptarget][amdgpu][nfc] Expand errorcheck macros
These macros expand to continue, which is confusing, or exit,
which is incompatible with continuing execution on offloading fail.
Expanding the macros in place makes the code look untidy but the
control flow obvious and amenable to improving. In particular, exit
becomes easier to eliminate.
Reviewed By: pdhaliwal
Differential Revision: https://reviews.llvm.org/D102230
Added:
Modified:
openmp/libomptarget/plugins/amdgpu/impl/data.cpp
openmp/libomptarget/plugins/amdgpu/impl/internal.h
openmp/libomptarget/plugins/amdgpu/impl/system.cpp
openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
Removed:
################################################################################
diff --git a/openmp/libomptarget/plugins/amdgpu/impl/data.cpp b/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
index b169e60e9a33f..345e7647bb29f 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/impl/data.cpp
@@ -48,7 +48,11 @@ atmi_status_t Runtime::Malloc(void **ptr, size_t size, atmi_mem_place_t place) {
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 Runtime::Memfree(void *ptr) {
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
index 4e26167b3b1ff..174a94086b2db 100644
--- a/openmp/libomptarget/plugins/amdgpu/impl/internal.h
+++ b/openmp/libomptarget/plugins/amdgpu/impl/internal.h
@@ -234,31 +234,4 @@ void allow_access_to_all_gpu_agents(void *ptr);
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
index be96cca2bec17..90de4bb171c32 100644
--- 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 @@ void allow_access_to_all_gpu_agents(void *ptr) {
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 @@ atmi_status_t Runtime::Initialize() {
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 @@ atmi_status_t Runtime::Finalize() {
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 @@ atmi_status_t Runtime::Finalize() {
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 @@ static hsa_status_t get_memory_pool_info(hsa_amd_memory_pool_t memory_pool,
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 @@ static hsa_status_t get_agent_info(hsa_agent_t agent, void *data) {
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 @@ static hsa_status_t init_compute_and_memory() {
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 @@ static hsa_status_t init_compute_and_memory() {
}
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 @@ static hsa_status_t init_compute_and_memory() {
&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 @@ hsa_status_t init_hsa() {
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 @@ atmi_status_t atl_init_gpu_context() {
}
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 @@ static hsa_status_t get_code_object_custom_metadata(void *binary,
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 @@ static hsa_status_t get_code_object_custom_metadata(void *binary,
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 @@ static hsa_status_t get_code_object_custom_metadata(void *binary,
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 @@ static hsa_status_t get_code_object_custom_metadata(void *binary,
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 @@ static hsa_status_t populate_InfoTables(hsa_executable_t executable,
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
diff erent 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 @@ static hsa_status_t populate_InfoTables(hsa_executable_t executable,
// 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 @@ static hsa_status_t populate_InfoTables(hsa_executable_t executable,
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 @@ static hsa_status_t populate_InfoTables(hsa_executable_t executable,
} 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 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
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 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
// 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 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
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 @@ atmi_status_t Runtime::RegisterModuleFromMemory(
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
index 98bf3da2a3ea3..d0c6c2c1ee82c 100644
--- a/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
+++ b/openmp/libomptarget/plugins/amdgpu/src/rtl.cpp
@@ -76,12 +76,8 @@ int print_kernel_trace;
#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 @@ struct KernelArgPool {
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 @@ struct KernelArgPool {
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 @@ class RTLDeviceInfoTy {
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();
}
More information about the Openmp-commits
mailing list