[libomptarget][amdgpu] Remove majority of fatal errors

[JonChesterfield]

Summary

[libomptarget][amdgpu] Remove majority of fatal errors

Replaces most calls to exit() with returning an error to the library entry
point. Minor changes to error handling for clear bugs, remove some dead code.

Each exit() call site replaced is either in a library entry point or a
function that already returns error codes on some paths. The existing handling
is not well tested but replacing exit() with a fallback path should be a strict
improvement.

Remaining two early exit points are an abort() from a callback and exit() from
within msgpack. Fixes for those are less obvious and left for a later patch.

[JonChesterfield]

• typo

[JonChesterfield]

This replaces some debug printing with DP() macro where it is near exit calls, leaving other (f)printf calls for a later patch.

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

23

in internal.h

62

register_allocation doesn't look like it handles invalid pointers gracefully, keeping it guarded by allocate succeeding

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

51

was unused

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

368

redundant checks

1046

branch was always true

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

102

this exit was reasonable in atmi, not so in libomptarget

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

1833

previously, this is a nullptr dereference on partial construction of the global

[pdhaliwal]

Looks good with one minor NIT.

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

137

should check for return status here.

[pdhaliwal]

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

1036

char *name is leaking on these exits.

[JonChesterfield]

• Check error code from allow_access, fold destroy into destructor

[JonChesterfield]

• fix copy error

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

 #include <string.h>
 #include <thread>
 #include <vector>
 
 using core::TaskImpl;
 extern ATLMachine g_atl_machine;
 
 namespace core {
-void allow_access_to_all_gpu_agents(void *ptr);
 
 namespace {
 ATLProcessor &get_processor_by_mem_place(atmi_mem_place_t place) {
   int dev_id = place.dev_id;
   switch (place.dev_type) {
   case ATMI_DEVTYPE_CPU:
     return g_atl_machine.processors<ATLCPUProcessor>()[dev_id];
   case ATMI_DEVTYPE_GPU:
     return g_atl_machine.processors<ATLGPUProcessor>()[dev_id];
   }
 }
 
 hsa_amd_memory_pool_t get_memory_pool_by_mem_place(atmi_mem_place_t place) {
   ATLProcessor &proc = get_processor_by_mem_place(place);
   return get_memory_pool(proc, place.mem_id);
 }
 } // namespace
 
-void register_allocation(void *ptr, size_t size, atmi_mem_place_t place) {
+hsa_status_t 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);
+    return allow_access_to_all_gpu_agents(ptr);
+  else
+    return HSA_STATUS_SUCCESS;
 }
 
 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);
-  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);
+  if (err == HSA_STATUS_SUCCESS) {
+    err = register_allocation(*ptr, size, place);
+  }
 
-  return ret;
+  return (err == HSA_STATUS_SUCCESS) ? ATMI_STATUS_SUCCESS : ATMI_STATUS_ERROR;
 }
 
 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);
-  if (err != HSA_STATUS_SUCCESS) {
-    printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "atmi_free",
-           get_error_string(err));
-    exit(1);
-  }
+  hsa_status_t err = hsa_amd_memory_pool_free(ptr);
   DEBUG_PRINT("Freed %p\n", ptr);
 
-  if (err != HSA_STATUS_SUCCESS)
-    ret = ATMI_STATUS_ERROR;
-  return ret;
+  return (err == HSA_STATUS_SUCCESS) ? ATMI_STATUS_SUCCESS : ATMI_STATUS_ERROR;
 }
 
 } // namespace core

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

typedef struct atmi_implicit_args_s {
   unsigned long cpu_queue_ptr;
   unsigned long kernarg_template_ptr;
 } atmi_implicit_args_t;
 
 #ifdef __cplusplus
 extern "C" {
 #endif
 
-#define check(msg, status)                                                     \
-  if (status != HSA_STATUS_SUCCESS) {                                          \
-    printf("%s failed.\n", #msg);                                              \
-    exit(1);                                                                   \
-  }
-
 #ifdef DEBUG
 #define DEBUG_PRINT(fmt, ...)                                                  \
   if (core::Runtime::getInstance().getDebugMode()) {                           \
     fprintf(stderr, "[%s:%d] " fmt, __FILE__, __LINE__, ##__VA_ARGS__);        \
   }
 #else
 #define DEBUG_PRINT(...)                                                       \
   do {                                                                         \
template <typename T> inline T alignUp(T value, size_t alignment) {
   return alignDown((T)(value + alignment - 1), alignment);
 }
 
 template <typename T> inline T *alignUp(T *value, size_t alignment) {
   return reinterpret_cast<T *>(
       alignDown((intptr_t)(value + alignment - 1), alignment));
 }
 
-extern void register_allocation(void *addr, size_t size,
+hsa_status_t register_allocation(void *addr, size_t size,
                                  atmi_mem_place_t place);
 
 extern bool atl_is_atmi_initialized();
 
 bool handle_group_signal(hsa_signal_value_t value, void *arg);
 
-
-void allow_access_to_all_gpu_agents(void *ptr);
+hsa_status_t 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);
 
 #endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_

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

 
 static void atl_reset_atmi_initialized() {
   // FIXME: thread safe? locks?
   g_atmi_initialized = false;
 }
 
 bool atl_is_atmi_initialized() { return g_atmi_initialized; }
 
-void allow_access_to_all_gpu_agents(void *ptr) {
-  hsa_status_t err;
+hsa_status_t allow_access_to_all_gpu_agents(void *ptr) {
   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);
-  if (err != HSA_STATUS_SUCCESS) {
-    printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
-           "Allow agents ptr access", get_error_string(err));
-    exit(1);
-  }
+  return hsa_amd_agents_allow_access(agents.size(), &agents[0], NULL, ptr);
 }
 
 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) {
-    if (atl_init_gpu_context() != ATMI_STATUS_SUCCESS) {
+    atmi_status_t rc = atl_init_gpu_context();
+    if (rc != 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);
+      return rc;
     }
   }
 
   atl_set_atmi_initialized();
   return ATMI_STATUS_SUCCESS;
 }
 
 atmi_status_t Runtime::Finalize() {
-  hsa_status_t err;
-
+  atmi_status_t rc = ATMI_STATUS_SUCCESS;
   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();
+  hsa_status_t err = hsa_shut_down();
   if (err != HSA_STATUS_SUCCESS) {
     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__, "Shutting down HSA",
            get_error_string(err));
-    exit(1);
+    rc = ATMI_STATUS_ERROR;
   }
 
-  return ATMI_STATUS_SUCCESS;
+  return rc;
 }
 
 static void atmi_init_context_structs() {
   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);
   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);
+    return 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);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Get memory pool info", get_error_string(err));
-      exit(1);
+      return 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);
   if (err != HSA_STATUS_SUCCESS) {
     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
            "Get device type info", get_error_string(err));
-    exit(1);
+    return 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);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Iterate all memory pools", get_error_string(err));
-      exit(1);
+      return 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);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Query the agent profile", get_error_string(err));
-      exit(1);
+      return 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);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Iterate all memory pools", get_error_string(err));
-      exit(1);
+      return err;
     }
     g_atl_machine.addProcessor(new_proc);
   } break;
   case HSA_DEVICE_TYPE_DSP: {
     err = HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
   } break;
   }
 

 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;
   }
   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>();
   /* For CPU memory pools, add other devices that can access them directly
    * or indirectly */
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;
     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);
+      return err;
     }
   }
   hsa_region_t atl_gpu_kernarg_region;
   /* 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;
     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);
+      return 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();
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Initializing the hsa runtime", get_error_string(err));
-      exit(1);
+      return err;
     }
     if (err != HSA_STATUS_SUCCESS)
       return err;
 
     err = init_compute_and_memory();
     if (err != HSA_STATUS_SUCCESS)
       return 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);
+      return 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++)
atmi_status_t atl_init_gpu_context() {
   err = init_hsa();
   if (err != HSA_STATUS_SUCCESS)
     return ATMI_STATUS_ERROR;
 
   err = hsa_amd_register_system_event_handler(callbackEvent, NULL);
   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);
+    return ATMI_STATUS_ERROR;
   }
 
   init_tasks();
   atlc.g_gpu_initialized = true;
   return ATMI_STATUS_SUCCESS;
 }
 
 static bool isImplicit(KernelArgMD::ValueKind value_kind) {
static hsa_status_t populate_InfoTables(hsa_executable_t executable,
 
   uint32_t name_length;
   hsa_status_t err;
   err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE,
                                        &type);
   if (err != HSA_STATUS_SUCCESS) {
     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
            "Symbol info extraction", get_error_string(err));
-    exit(1);
+    return 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);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Symbol info extraction", get_error_string(err));
-      exit(1);
+      return err;
     }
     char *name = reinterpret_cast<char *>(malloc(name_length + 1));
     err = hsa_executable_symbol_get_info(symbol,
                                          HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Symbol info extraction", get_error_string(err));
-      exit(1);
+      return 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
-      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);
-      }
+      return HSA_STATUS_ERROR;
     }
     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()) {
-      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);
-      }
+      return HSA_STATUS_ERROR;
     }
     // 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));
     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);
+      return err;
     }
     err = hsa_executable_symbol_get_info(
         symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
         &(info.group_segment_size));
     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);
+      return err;
     }
     err = hsa_executable_symbol_get_info(
         symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
         &(info.private_segment_size));
     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);
+      return 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);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Symbol info extraction", get_error_string(err));
-      exit(1);
+      return err;
     }
     char *name = reinterpret_cast<char *>(malloc(name_length + 1));
     err = hsa_executable_symbol_get_info(symbol,
                                          HSA_EXECUTABLE_SYMBOL_INFO_NAME, name);
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Symbol info extraction", get_error_string(err));
-      exit(1);
+      return 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));
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Symbol info address extraction", get_error_string(err));
-      exit(1);
+      return err;
     }
 
     err = hsa_executable_symbol_get_info(
         symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, &(info.size));
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Symbol info size extraction", get_error_string(err));
-      exit(1);
+      return 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);
+    err = register_allocation(reinterpret_cast<void *>(info.addr),
+                              (size_t)info.size, place);
+    if (err != HSA_STATUS_SUCCESS) {
+      return err;
+    }
     SymbolInfoTable[gpu][std::string(name)] = info;
     if (strcmp(name, "needs_hostcall_buffer") == 0)
       g_atmi_hostcall_required = true;
     free(name);
   } else {
     DEBUG_PRINT("Symbol is an indirect function\n");
   }
   return HSA_STATUS_SUCCESS;
atmi_status_t Runtime::RegisterModuleFromMemory(
   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);
   if (err != HSA_STATUS_SUCCESS) {
     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
            "Query the agent profile", get_error_string(err));
-    exit(1);
+    return ATMI_STATUS_ERROR;
   }
   // 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);
   if (err != HSA_STATUS_SUCCESS) {
     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
            "Create the executable", get_error_string(err));
-    exit(1);
+    return ATMI_STATUS_ERROR;
   }
 
   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
do // Existing control flow used continue, preserve that for this patch
       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);
       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);
+        return atmi_err;
       }
 
       /* Load the code object.  */
       err =
           hsa_executable_load_code_object(executable, agent, code_object, NULL);
       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, "");
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Freeze the executable", get_error_string(err));
-      exit(1);
+      return ATMI_STATUS_ERROR;
     }
 
     err = hsa_executable_iterate_symbols(executable, populate_InfoTables,
                                          static_cast<void *>(&gpu));
     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);
+      return ATMI_STATUS_ERROR;
     }
 
     // save the executable and destroy during finalize
     HSAExecutables.push_back(executable);
     return ATMI_STATUS_SUCCESS;
   } else {
     return ATMI_STATUS_ERROR;
   }
 }
 
 } // namespace core

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

  * Environment variables
  */
 void Environment::GetEnvAll() {
   std::string var = GetEnv("ATMI_HELP");
   if (!var.empty()) {
     std::cout << "ATMI_MAX_HSA_QUEUE_SIZE : positive integer" << std::endl
               << "ATMI_DEBUG : 1 for printing out trace/debug info"
               << std::endl;
-    exit(0);
   }
 
   var = GetEnv("ATMI_MAX_HSA_QUEUE_SIZE");
   if (!var.empty())
     max_queue_size_ = std::stoi(var);
 
   var = GetEnv("ATMI_DEBUG");
   if (!var.empty())
     debug_mode_ = std::stoi(var);
 }
 } // namespace core

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

DP("Warning: Attempting to assign hostrpc to device %u, but hostrpc library "
      "missing\n",
      device_id);
   return 0;
 }
 }
 
 int print_kernel_trace;
 
-#undef check // Drop definition from internal.h
 #ifdef OMPTARGET_DEBUG
 #define check(msg, status)                                                     \
   if (status != ATMI_STATUS_SUCCESS) {                                         \
     DP(#msg " failed\n");                                                      \
   } else {                                                                     \
     DP(#msg " succeeded\n");                                                   \
   }
 #else
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);
       if (r != HSA_STATUS_SUCCESS) {
-        printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
-               "Memory pool free", get_error_string(r));
-        exit(1);
+        DP("hsa_amd_memory_pool_free failed: %s\n", get_error_string(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);
-      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);
+
+
+    if (err != HSA_STATUS_SUCCESS) {
+      DP("hsa_amd_memory_pool_allocate failed: %s\n", get_error_string(err));
+      kernarg_region = nullptr; // paranoid
+      return;
+    }
+
+    err = core::allow_access_to_all_gpu_agents(kernarg_region);
+      DP("hsa allow_access_to_all_gpu_agents failed: %s\n",
+      auto r = hsa_amd_memory_pool_free(kernarg_region);
+      if (r != HSA_STATUS_SUCCESS) {
+        // if free failed, can't do anything more to resolve it
+        DP("hsa memory poll free failed: %s\n", get_error_string(err));
+      }
+      kernarg_region = nullptr;
+      return;
 
-      for (int i = 0; i < MAX_NUM_KERNELS; i++) {
-        free_kernarg_segments.push(i);
-      }
-    }
-  }
 
   void *allocate(uint64_t arg_num) {
     assert((arg_num * sizeof(void *)) == kernarg_segment_size);
     lock l(&mutex);
     void *res = nullptr;
     if (!free_kernarg_segments.empty()) {
 
       int free_idx = free_kernarg_segments.front();
RTLDeviceInfoTy() {
     GPUName.resize(NumberOfDevices);
     GroupsPerDevice.resize(NumberOfDevices);
     WarpSize.resize(NumberOfDevices);
     NumTeams.resize(NumberOfDevices);
     NumThreads.resize(NumberOfDevices);
     deviceStateStore.resize(NumberOfDevices);
 
     for (int i = 0; i < NumberOfDevices; i++) {
+      HSAQueues[i] = nullptr;
+    }
+
+    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);
+        hsa_status_t err = hsa_agent_get_info(
+            HSAAgents[i], HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size);
         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);
+          DP("HSA query QUEUE_MAX_SIZE failed for agent %d\n", i);
+          return;
         }
         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]);
       if (rc != HSA_STATUS_SUCCESS) {
-        DP("Failed to create HSA queues\n");
+        DP("Failed to create HSA queue %d\n", i);
         return;
       }
 
       deviceStateStore[i] = {nullptr, 0};
     }
 
     for (int i = 0; i < NumberOfDevices; i++) {
       ThreadsPerGroup[i] = RTLDeviceInfoTy::Default_WG_Size;
RTLDeviceInfoTy() {
     } else {
       EnvMaxTeamsDefault = -1;
     }
 
     // Default state.
     RequiresFlags = OMP_REQ_UNDEFINED;
   }
 
-  void DestroyHSAExecutables() {
-    hsa_status_t Err;
-    for (uint32_t I = 0; I < HSAExecutables.size(); I++) {
-      Err = hsa_executable_destroy(HSAExecutables[I]);
-      if (Err != HSA_STATUS_SUCCESS) {
-        printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
-               "Destroying executable", get_error_string(Err));
-        return;
-      }
-    }
-  }
-
   ~RTLDeviceInfoTy() {
     DP("Finalizing the HSA-ATMI DeviceInfo.\n");
     // Run destructors on types that use HSA before
     // atmi_finalize removes access to it
     deviceStateStore.clear();
     KernelArgPoolMap.clear();
     // Terminate hostrpc before finalizing ATMI
     hostrpc_terminate();
 
-    DestroyHSAExecutables();
+    for (uint32_t I = 0; I < HSAExecutables.size(); I++) {
+      hsa_status_t Err = hsa_executable_destroy(HSAExecutables[I]);
+      if (Err != HSA_STATUS_SUCCESS) {
+        DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+           "Destroying executable", get_error_string(Err));
+      }
+    }
+
     atmi_finalize();
   }
 };
 
 pthread_mutex_t SignalPoolT::mutex = PTHREAD_MUTEX_INITIALIZER;
 
 // TODO: May need to drop the trailing to fields until deviceRTL is updated
 struct omptarget_device_environmentTy {
static uint64_t get_device_State_bytes(char *ImageStart, size_t img_size) {
   {
     // If this is the deviceRTL, get the state variable size
     symbol_info size_si;
     int rc = get_symbol_info_without_loading(
         ImageStart, img_size, "omptarget_nvptx_device_State_size", &size_si);
 
     if (rc == 0) {
       if (size_si.size != sizeof(uint64_t)) {
-        fprintf(stderr,
-                "Found device_State_size variable with wrong size, aborting\n");
-        exit(1);
+        DP("Found device_State_size variable with wrong size\n");
+        return 0;
       }
 
       // Read number of bytes directly from the elf
       memcpy(&device_State_bytes, size_si.addr, sizeof(uint64_t));
     }
   }
   return device_State_bytes;
 }
fprintf(traceToStdout ? stdout : stderr,
             arg_num, num_groups, threadsPerGroup, num_teams, thread_limit,
             sgpr_count, vgpr_count, sgpr_spill_count, vgpr_spill_count,
             loop_tripcount, KernelInfo->Name);
   }
 
   // Run on the device.
   {
     hsa_queue_t *queue = DeviceInfo.HSAQueues[device_id];
+    if (!queue) {
+      return OFFLOAD_FAIL;
+    }
     uint64_t packet_id = acquire_available_packet_id(queue);
 
     const uint32_t mask = queue->size - 1; // size is a power of 2
     hsa_kernel_dispatch_packet_t *packet =
         (hsa_kernel_dispatch_packet_t *)queue->base_address +
         (packet_id & mask);
 
     // packet->header is written last
// Run on the device.
     KernelArgPool *ArgPool = nullptr;
     {
       auto it = KernelArgPoolMap.find(std::string(KernelInfo->Name));
       if (it != KernelArgPoolMap.end()) {
         ArgPool = (it->second).get();
       }
     }
     if (!ArgPool) {
-      fprintf(stderr, "Warning: No ArgPool for %s on device %d\n",
-              KernelInfo->Name, device_id);
+      DP("Warning: No ArgPool for %s on device %d\n", KernelInfo->Name,
+         device_id);
     }
     {
       void *kernarg = nullptr;
       if (ArgPool) {
         assert(ArgPool->kernarg_segment_size == (arg_num * sizeof(void *)));
         kernarg = ArgPool->allocate(arg_num);
       }
       if (!kernarg) {
-        printf("Allocate kernarg failed\n");
-        exit(1);
+        DP("Allocate kernarg failed\n");
+        return OFFLOAD_FAIL;
       }
 
       // Copy explicit arguments
       for (int i = 0; i < arg_num; i++) {
         memcpy((char *)kernarg + sizeof(void *) * i, args[i], sizeof(void *));
       }
 
       // Initialize implicit arguments. ATMI seems to leave most fields
if (!ArgPool) {
       }
 
       packet->kernarg_address = kernarg;
     }
 
     {
       hsa_signal_t s = DeviceInfo.FreeSignalPool.pop();
       if (s.handle == 0) {
-        printf("Failed to get signal instance\n");
-        exit(1);
+        DP("Failed to get signal instance\n");
+        return OFFLOAD_FAIL;
       }
       packet->completion_signal = s;
       hsa_signal_store_relaxed(packet->completion_signal, 1);
     }
 
     core::packet_store_release(reinterpret_cast<uint32_t *>(packet),
                                core::create_header(), packet->setup);