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,12 @@ 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); + } else { + }; DEBUG_PRINT("Malloced [%s %d] %p\n", place.dev_type == ATMI_DEVTYPE_CPU ? "CPU" : "GPU", place.dev_id, *ptr); @@ -64,7 +69,12 @@ 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); + } else { + }; 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,12 @@ 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); + } else { + }; } atmi_status_t Runtime::Initialize() { @@ -202,7 +200,12 @@ 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); + } else { + }; } atl_set_atmi_initialized(); @@ -214,7 +217,12 @@ 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); + } else { + }; } for (uint32_t i = 0; i < SymbolInfoTable.size(); i++) { @@ -228,7 +236,12 @@ 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); + } else { + }; return ATMI_STATUS_SUCCESS; } @@ -252,12 +265,22 @@ 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); + } else { + }; 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); + } else { + }; 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 +301,48 @@ 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); + } else { + }; 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); + } else { + }; 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); + } else { + }; 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); + } else { + }; g_atl_machine.addProcessor(new_proc); } break; case HSA_DEVICE_TYPE_DSP: { @@ -353,7 +396,12 @@ 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); + } else { + }; if (err != HSA_STATUS_SUCCESS) return err; @@ -489,7 +537,13 @@ } 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); + } else { + }; } /* Find a memory region that supports kernel arguments. */ atl_gpu_kernarg_region.handle = (uint64_t)-1; @@ -498,7 +552,12 @@ &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); + } else { + }; } if (num_procs > 0) return HSA_STATUS_SUCCESS; @@ -510,14 +569,24 @@ 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); + } else { + }; 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); + } else { + }; int gpu_count = g_atl_machine.processorCount(); KernelInfoTable.resize(gpu_count); @@ -602,7 +671,12 @@ } 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); + } else { + }; init_tasks(); atlc.g_gpu_initialized = true; @@ -817,7 +891,12 @@ 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; + } else { + }; for (size_t i = 0; i < kernelsSize; i++) { assert(msgpack_errors == 0); @@ -826,43 +905,77 @@ 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; + } else { + }; 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; + } else { + }; 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; + } else { + }; + 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; + } else { + }; + 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; + } else { + }; + 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; + } else { + }; + 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; + } else { + }; // create a map from symbol to name DEBUG_PRINT("Kernel symbol %s; Name: %s; Size: %lu\n", symbolName.c_str(), @@ -877,8 +990,12 @@ 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; + } else { + }; info.num_args = argsSize; @@ -887,12 +1004,20 @@ 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; + } else { + }; 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; + } else { + }; // populate info with sizes and offsets info.arg_sizes.push_back(lcArg.size_); @@ -942,23 +1067,44 @@ 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); + } else { + }; 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); + } else { + }; char *name = reinterpret_cast(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); + } else { + }; 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); + } else { + }; } atl_kernel_info_t info; std::string kernelName = KernelNameMap[std::string(name)]; @@ -966,8 +1112,13 @@ // 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); + } else { + }; } // found, so assign and update info = KernelInfoTable[gpu][kernelName]; @@ -976,15 +1127,33 @@ 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); + } else { + }; 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); + } else { + }; 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); + } else { + }; DEBUG_PRINT( "Kernel %s --> %lx symbol %u group segsize %u pvt segsize %u bytes " @@ -998,22 +1167,42 @@ } 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); + } else { + }; char *name = reinterpret_cast(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); + } else { + }; 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); + } else { + }; 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); + } else { + }; 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 +1235,24 @@ 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); + } else { + }; // 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); + } else { + }; bool module_load_success = false; do // Existing control flow used continue, preserve that for this patch @@ -1063,13 +1262,24 @@ // 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; + } else { + }; // 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; + } else { + }; assert(0 != code_object.handle); // Mutating the device image here avoids another allocation & memcpy @@ -1077,12 +1287,23 @@ reinterpret_cast(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); + } else { + }; /* 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; + } else { + }; // cannot iterate over symbols until executable is frozen } @@ -1092,11 +1313,21 @@ 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); + } else { + }; err = hsa_executable_iterate_symbols(executable, populate_InfoTables, static_cast(&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); + } else { + }; // 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,12 @@ 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); + } else { + }; } } @@ -140,7 +141,13 @@ 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); + } else { + }; core::allow_access_to_all_gpu_agents(kernarg_region); for (int i = 0; i < MAX_NUM_KERNELS; i++) { @@ -473,7 +480,13 @@ 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); + } else { + }; if (queue_size > core::Runtime::getInstance().getMaxQueueSize()) { queue_size = core::Runtime::getInstance().getMaxQueueSize(); }