[libomptarget][nfc] Drop stringify in macro

[JonChesterfield]

Summary

[libomptarget][nfc] Drop stringify in macro
A step towards deleting the macros entirely.

[pdhaliwal]

LGTM, thanks!

[JonChesterfield]

• add missing file

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);
+  ErrorCheck("atmi_malloc", err);
   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);
+  ErrorCheck("atmi_free", err);
   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,                \
+    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,                \
+    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,           \
+    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 "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);                   \
+    printf("[%s:%d] %s failed\n", __FILE__, __LINE__, msg);                    \
     return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;                               \
   } else {                                                                     \
   }
 
 typedef unsigned char *address;
 /*
  * Note descriptors.
  */
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);
+  ErrorCheck("Allow agents ptr access", err);
 }
 
 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());
+    ATMIErrorCheck("GPU context init", atl_init_gpu_context());
   }
 
   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);
+    ErrorCheck("Destroying executable", err);
   }
 
   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);
+  ErrorCheck("Shutting down HSA", err);
 
   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);
+  ErrorCheck("Alloc allowed in memory pool check", err);
   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);
+    ErrorCheck("Get memory pool info", err);
     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);
+  ErrorCheck("Get device type info", err);
   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);
+    ErrorCheck("Iterate all memory pools", err);
     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);
+    ErrorCheck("Query the agent profile", err);
     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);
+    ErrorCheck("Iterate all memory pools", err);
     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);
+  ErrorCheck("Getting a gpu agent", err);
   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);
+    ErrorCheck("Finding a CPU kernarg memory region handle", err);
   }
   /* 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);
+    ErrorCheck("Finding a kernarg memory region", err);
   }
   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);
+    ErrorCheck("Initializing the hsa runtime", err);
     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);
+    ErrorCheck("After initializing compute and memory", err);
 
     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);
+  ErrorCheck("Registering the system for memory faults", err);
 
-    init_tasks();
-    atlc.g_gpu_initialized = true;
-    return ATMI_STATUS_SUCCESS;
 }
 
 static bool isImplicit(KernelArgMD::ValueKind value_kind) {
   switch (value_kind) {
   case KernelArgMD::ValueKind::HiddenGlobalOffsetX:
   case KernelArgMD::ValueKind::HiddenGlobalOffsetY:
   case KernelArgMD::ValueKind::HiddenGlobalOffsetZ:
   case KernelArgMD::ValueKind::HiddenNone:
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);
+  msgpackErrorCheck("kernels lookup in program metadata", msgpack_errors);
 
   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);
+    msgpackErrorCheck("element lookup in kernel metadata", msgpack_errors);
 
     msgpack_errors += map_lookup_string(element, ".name", &kernelName);
     msgpack_errors += map_lookup_string(element, ".symbol", &symbolName);
-    msgpackErrorCheck(strings lookup in kernel metadata, msgpack_errors);
+    msgpackErrorCheck("strings lookup in kernel metadata", msgpack_errors);
 
     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,
+    msgpackErrorCheck("sgpr count metadata lookup in kernel metadata",
                       msgpack_errors);
     info.sgpr_count = sgpr_count;
 
     msgpack_errors += map_lookup_uint64_t(element, ".vgpr_count", &vgpr_count);
-    msgpackErrorCheck(vgpr count metadata lookup in kernel metadata,
+    msgpackErrorCheck("vgpr count metadata lookup in kernel metadata",
                       msgpack_errors);
     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,
+    msgpackErrorCheck("sgpr spill count metadata lookup in kernel metadata",
                       msgpack_errors);
     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,
+    msgpackErrorCheck("vgpr spill count metadata lookup in kernel metadata",
                       msgpack_errors);
     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,
+    msgpackErrorCheck("kernarg segment size metadata lookup in kernel metadata",
                       msgpack_errors);
 
     // 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,
+      msgpackErrorCheck("kernel args metadata lookup in kernel metadata",
                         msgpack_errors);
 
       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,
+        msgpackErrorCheck("iterate args map in kernel args metadata",
                           msgpack_errors);
 
         msgpack_errors += populate_kernelArgMD(args_element, &lcArg);
-        msgpackErrorCheck(iterate args map in kernel args metadata,
+        msgpackErrorCheck("iterate args map in kernel args metadata",
                           msgpack_errors);
 
         // 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;
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);
+  ErrorCheck("Symbol info extraction", err);
   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);
+    ErrorCheck("Symbol info extraction", err);
     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);
+    ErrorCheck("Symbol info extraction", err);
     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);
+      ErrorCheck("Invalid kernel name", HSA_STATUS_ERROR_INVALID_CODE_OBJECT);
     }
     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,
+      ErrorCheck("Finding the entry kernel info table",
                  HSA_STATUS_ERROR_INVALID_CODE_OBJECT);
     }
     // 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);
+    ErrorCheck("Extracting the symbol from the executable", err);
     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);
+    ErrorCheck("Extracting the group segment size from the executable", err);
     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);
+    ErrorCheck("Extracting the private segment from the executable", err);
 
     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);
+    ErrorCheck("Symbol info extraction", err);
     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);
+    ErrorCheck("Symbol info extraction", err);
     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);
+    ErrorCheck("Symbol info address extraction", err);
 
     err = hsa_executable_symbol_get_info(
         symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &(info.size));
-    ErrorCheck(Symbol info size extraction, err);
+    ErrorCheck("Symbol info size extraction", err);
 
     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);
+  ErrorCheck("Query the agent profile", err);
   // 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);
+  ErrorCheck("Create the executable", err);
 
   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);
+      ErrorCheckAndContinue("Getting custom code object metadata", err);
 
       // 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);
+      ErrorCheckAndContinue("Code Object Deserialization", err);
       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);
+      ATMIErrorCheck("Error in deserialized_data callback", atmi_err);
 
       /* Load the code object.  */
       err =
           hsa_executable_load_code_object(executable, agent, code_object, NULL);
-      ErrorCheckAndContinue(Loading the code object, err);
+      ErrorCheckAndContinue("Loading the code object", err);
 
       // 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);
+    ErrorCheck("Freeze the executable", err);
 
     err = hsa_executable_iterate_symbols(executable, populate_InfoTables,
                                          static_cast<void *>(&gpu));
-    ErrorCheck(Iterating over symbols for execuatable, err);
+    ErrorCheck("Iterating over symbols for execuatable", err);
 
     // 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

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);
+      ErrorCheck("Memory pool free", r);
     }
   }
 
   // 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);
+      ErrorCheck("Allocating memory for the executable-kernel", err);
       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);
+        ErrorCheck("Querying the agent maximum queue size", err);
         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]);