[libomptarget][amdgpu][nfc] Expand errorcheck macros

[JonChesterfield]

Summary

[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.

[JonChesterfield]

• remove else {}; noise

[JonChesterfield]

• remove else {}; noise

[pdhaliwal]

openmp/libomptarget/plugins/amdgpu/impl/data.cpp

52–54

Will it be better to use assert's here?

[JonChesterfield]

This patch doesn't fix any of the control flow, just expands macros and deletes else{}; clauses

openmp/libomptarget/plugins/amdgpu/impl/data.cpp

52–54

No, malloc can fail, and its failure can be handled. Assert might be better than exit, but what we actually need is return some-error-code

[pdhaliwal]

LGTM.

[protze.joachim]

openmp/libomptarget/plugins/amdgpu/impl/data.cpp

52–54

Isn't that was the lines below do?

60

This line cannot be reached with your change.

[JonChesterfield]

openmp/libomptarget/plugins/amdgpu/impl/data.cpp

52–54

Yep. Nice example of something that is obvious once the macro is expanded, and was missed by repeated reading with the macro in use.

60

Then it was unreachable before. Not great, kind of why I don't like control flow escaping macros.

All this change does is macroexpand ErrorCheck et al and s_else {};__g then clang format. Quite literally, I write this with regex replace / run preprocessor / regex replace.

Functional changes / stuff that requires thought intentionally left for the next round. See also D102228 for the previous incremental change which made that scripting simple.

[JonChesterfield]

openmp/libomptarget/plugins/amdgpu/impl/data.cpp

60

I.e. before we had

#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);  }

ErrorCheck("atmi_malloc", err);

which expands to
if (err != HSA_STATUS_SUCCESS) {printf(); exit(1); }

[JonChesterfield]

I'm going to land this so I can follow up with fixing the above points and simplify system.cpp. In D102346

openmp/libomptarget/plugins/amdgpu/impl/data.cpp

void register_allocation(void *ptr, size_t size, atmi_mem_place_t place) {
   if (place.dev_type == ATMI_DEVTYPE_CPU)
     allow_access_to_all_gpu_agents(ptr);
 }
 
 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);
   if (err != HSA_STATUS_SUCCESS)
     ret = ATMI_STATUS_ERROR;
 
   register_allocation(*ptr, size, place);
 
   return ret;
 }
 
 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)
     ret = ATMI_STATUS_ERROR;
   return ret;
 }
 
 } // namespace core

openmp/libomptarget/plugins/amdgpu/impl/internal.h

 
 
 void allow_access_to_all_gpu_agents(void *ptr);
 } // namespace core
 
 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_

openmp/libomptarget/plugins/amdgpu/impl/system.cpp

 #include <string>
 
 #include "internal.h"
 #include "machine.h"
 #include "rt.h"
 
 #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.
  */
 typedef struct {
   uint32_t n_namesz; /* Length of note's name. */
   uint32_t n_descsz; /* Length of note's value. */
   uint32_t n_type;   /* Type of note. */
void allow_access_to_all_gpu_agents(void *ptr) {
   hsa_status_t err;
   std::vector<ATLGPUProcessor> &gpu_procs =
       g_atl_machine.processors<ATLGPUProcessor>();
   std::vector<hsa_agent_t> agents;
   for (uint32_t i = 0; i < gpu_procs.size(); i++) {
     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() {
   atmi_devtype_t devtype = ATMI_DEVTYPE_GPU;
   if (atl_is_atmi_initialized())
     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();
   return ATMI_STATUS_SUCCESS;
 }
 
 atmi_status_t Runtime::Finalize() {
   hsa_status_t err;
 
   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++) {
     SymbolInfoTable[i].clear();
   }
   SymbolInfoTable.clear();
   for (uint32_t i = 0; i < KernelInfoTable.size(); i++) {
     KernelInfoTable[i].clear();
   }
   KernelInfoTable.clear();
 
   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;
 }
 
 static void atmi_init_context_structs() {
   atlc_p = &atlc;
   atlc.struct_initialized = true; /* This only gets called one time */
   atlc.g_hsa_initialized = false;
   atlc.g_gpu_initialized = false;
   atlc.g_tasks_initialized = false;
 }
 
 // Implement memory_pool iteration function
 static hsa_status_t get_memory_pool_info(hsa_amd_memory_pool_t memory_pool,
                                          void *data) {
   ATLProcessor *proc = reinterpret_cast<ATLProcessor *>(data);
   hsa_status_t err = HSA_STATUS_SUCCESS;
   // Check if the memory_pool is allowed to allocate, i.e. do not return group
   // memory
   bool alloc_allowed = false;
   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);
       if (HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT & global_flag) {
         DEBUG_PRINT("GPU kernel args pool handle: %lu\n", memory_pool.handle);
         atl_gpu_kernarg_pools.push_back(memory_pool);
       }
     } else {
       ATLMemory new_mem(memory_pool, *proc, ATMI_MEMTYPE_COARSE_GRAINED);
       proc->addMemory(new_mem);
     }
   }
 
   return err;
 }
 
 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: {
     err = HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
   } break;
   }
 
   return err;

 static hsa_status_t init_compute_and_memory() {
   hsa_status_t err;
 
   /* Iterate over the agents and pick the gpu agent */
   err = hsa_iterate_agents(get_agent_info, NULL);
   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;
 
   /* Init all devices or individual device types? */
   std::vector<ATLCPUProcessor> &cpu_procs =
       g_atl_machine.processors<ATLCPUProcessor>();
   std::vector<ATLGPUProcessor> &gpu_procs =
       g_atl_machine.processors<ATLGPUProcessor>();
static hsa_status_t init_compute_and_memory() {
   if (cpu_procs.size() > 0) {
     err = hsa_agent_iterate_regions(
         cpu_procs[0].agent(), get_fine_grained_region, &atl_cpu_kernarg_region);
     if (err == HSA_STATUS_INFO_BREAK) {
       err = HSA_STATUS_SUCCESS;
     }
     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;
   if (gpu_procs.size() > 0) {
     hsa_agent_iterate_regions(gpu_procs[0].agent(), get_kernarg_memory_region,
                               &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;
   else
     return HSA_STATUS_ERROR_NOT_INITIALIZED;
 }
 
 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);
     SymbolInfoTable.resize(gpu_count);
     for (uint32_t i = 0; i < SymbolInfoTable.size(); i++)
       SymbolInfoTable[i].clear();
     for (uint32_t i = 0; i < KernelInfoTable.size(); i++)
       KernelInfoTable[i].clear();
if (err != HSA_STATUS_SUCCESS)
     return ATMI_STATUS_ERROR;
 
   if (context_init_time_init == 0) {
     clock_gettime(CLOCK_MONOTONIC_RAW, &context_init_time);
     context_init_time_init = 1;
   }
 
   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;
   return ATMI_STATUS_SUCCESS;
 }
 
 static bool isImplicit(KernelArgMD::ValueKind value_kind) {
   switch (value_kind) {
static hsa_status_t get_code_object_custom_metadata(void *binary,
   }
 
   uint64_t kernelsSize = 0;
   int msgpack_errors = 0;
   msgpack::byte_range kernel_array;
   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);
     std::string kernelName;
     std::string symbolName;
 
     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(),
                 kernelName.c_str(), kernel_segment_size);
     KernelNameMap[symbolName] = kernelName;
 
     bool hasHiddenArgs = false;
     if (kernel_segment_size > 0) {
       uint64_t argsSize;
       size_t offset = 0;
 
       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;
 
       for (size_t i = 0; i < argsSize; ++i) {
         KernelArgMD lcArg;
 
         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
         size_t new_offset = lcArg.offset_;
         size_t padding = new_offset - offset;
         offset = new_offset;
         info.arg_offsets.push_back(lcArg.offset_);
         DEBUG_PRINT("Arg[%lu] \"%s\" (%u, %u)\n", i, lcArg.name_.c_str(),
static hsa_status_t populate_InfoTables(hsa_executable_t executable,
                                         void *data) {
   int gpu = *static_cast<int *>(data);
   hsa_symbol_kind_t type;
 
   uint32_t name_length;
   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)];
     // by now, the kernel info table should already have an entry
     // 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];
 
     /* Extract dispatch information from the symbol */
     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 "
         "kernarg\n",
         kernelName.c_str(), info.kernel_object, info.group_segment_size,
         info.private_segment_size, info.kernel_segment_size);
 
     // assign it back to the kernel info table
     KernelInfoTable[gpu][kernelName] = info;
     free(name);
   } 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,
                 info.size);
     register_allocation(reinterpret_cast<void *>(info.addr), (size_t)info.size,
                         place);
     SymbolInfoTable[gpu][std::string(name)] = info;
     if (strcmp(name, "needs_hostcall_buffer") == 0)
atmi_status_t Runtime::RegisterModuleFromMemory(
 
   DEBUG_PRINT("Trying to load module to GPU-%d\n", gpu);
   ATLGPUProcessor &proc = get_processor<ATLGPUProcessor>(place);
   hsa_agent_t agent = proc.agent();
   hsa_executable_t executable = {0};
   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
   {
     {
       // Some metadata info is not available through ROCr API, so use custom
       // 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
       void *code_object_alloc_data =
           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
     }
     module_load_success = true;
   } while (0);
   DEBUG_PRINT("Modules loaded successful? %d\n", module_load_success);
   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);
     return ATMI_STATUS_SUCCESS;
   } else {
     return ATMI_STATUS_ERROR;
   }
 }
 
 } // namespace core

openmp/libomptarget/plugins/amdgpu/src/rtl.cpp

 }
 
 int print_kernel_trace;
 
 #undef check // Drop definition from internal.h
 #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
 #define check(msg, status)                                                     \
   {}
 #endif
 
 #include "elf_common.h"
public:
   uint32_t kernarg_size_including_implicit() {
     return kernarg_segment_size + sizeof(atmi_implicit_args_t);
   }
 
   ~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);
+      }
     }
   }
 
   // Can't really copy or move a mutex
   KernelArgPool() = default;
   KernelArgPool(const KernelArgPool &) = delete;
   KernelArgPool(KernelArgPool &&) = delete;
 
   KernelArgPool(uint32_t kernarg_segment_size)
       : kernarg_segment_size(kernarg_segment_size) {
 
     // atmi uses one pool per kernel for all gpus, with a fixed upper size
     // preserving that exact scheme here, including the queue<int>
     {
       hsa_status_t err = hsa_amd_memory_pool_allocate(
           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++) {
         free_kernarg_segments.push(i);
       }
     }
   }
 
RTLDeviceInfoTy() {
     deviceStateStore.resize(NumberOfDevices);
 
     for (int i = 0; i < NumberOfDevices; i++) {
       uint32_t queue_size = 0;
       {
         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();
         }
       }
 
       hsa_status_t rc = hsa_queue_create(
           HSAAgents[i], queue_size, HSA_QUEUE_TYPE_MULTI, callbackQueue, NULL,
           UINT32_MAX, UINT32_MAX, &HSAQueues[i]);