[libomptarget][nfc][amdgpu] Remove atmi_status_t type

[JonChesterfield]

Summary

ATMI_STATUS_UNKNOWN was unused, deleted references to it.
Replaced ATMI_STATUS_{SUCCESS,ERROR} with HSA_STATUS_{SUCCESS,ERROR}
Replaced atmi_status_t with hsa_status_t

Otherwise no change. In particular, conversions between atmi_status_t and
hsa_status_t will now be conversions between hsa_status_t and itself.

[JonChesterfield]

Deleted the references to ATMI_STATUS_UNKNOWN manually, then

for f in `find . -type f`; do sed -i 's$ATMI_STATUS_SUCCESS$HSA_STATUS_SUCCESS$g' $f; sed -i 's$ATMI_STATUS_ERROR$HSA_STATUS_ERROR$g' $f; sed -i 's$atmi_status_t$hsa_status_t$g' $f; done

and clang-format

Landing this after the other refactors is probably easier. It can be reliably rebased/reimplemented on top of other changes and will otherwise create conflicts with anything near an error path.

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

60

This can be return err;, leaving that for a later patch to help with rebasing. Several patches to these files currently live.

[pdhaliwal]

Looks good. Thanks!

[JonChesterfield]

• rebase

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

 /*===--------------------------------------------------------------------------
  *              ATMI (Asynchronous Task and Memory Interface)
  *
  * This file is distributed under the MIT License. See LICENSE.txt for details.
  *===------------------------------------------------------------------------*/
 #ifndef INCLUDE_ATMI_H_
 #define INCLUDE_ATMI_H_
 
 #define ROCM_VERSION_MAJOR 3
 #define ROCM_VERSION_MINOR 2
 
 /** \defgroup enumerations Enumerated Types
  * @{
  */
 
 /**
- * @brief Status codes.
- */
-typedef enum atmi_status_t {
-  /**
-   * The function has been executed successfully.
-   */
-  ATMI_STATUS_SUCCESS = 0,
-  /**
-   * A undocumented error has occurred.
-   */
-  ATMI_STATUS_UNKNOWN = 1,
-  /**
-   * A generic error has occurred.
-   */
-  ATMI_STATUS_ERROR = 2,
-} atmi_status_t;
-
-/**
  * @brief Device Types.
  */
 typedef enum atmi_devtype_s {
   ATMI_DEVTYPE_CPU = 0x0001,
   ATMI_DEVTYPE_iGPU = 0x0010,                               // Integrated GPU
   ATMI_DEVTYPE_dGPU = 0x0100,                               // Discrete GPU
   ATMI_DEVTYPE_GPU = ATMI_DEVTYPE_iGPU | ATMI_DEVTYPE_dGPU, // Any GPU
   ATMI_DEVTYPE_ALL = 0x111 // Union of all device types

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

 }
 
 struct atmiFreePtrDeletor {
   void operator()(void *p) {
     atmi_free(p); // ignore failure to free
   }
 };
 
-atmi_status_t atmi_memcpy_h2d(hsa_signal_t signal, void *deviceDest,
-                              const void *hostSrc, size_t size,
-                              hsa_agent_t agent) {
+hsa_status_t atmi_memcpy_h2d(hsa_signal_t signal, void *deviceDest,
   hsa_status_t rc = hsa_memory_copy(deviceDest, hostSrc, size);
 
   // hsa_memory_copy sometimes fails in situations where
   // allocate + copy succeeds. Looks like it might be related to
   // locking part of a read only segment. Fall back for now.
   if (rc == HSA_STATUS_SUCCESS) {
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
   }
 
   void *tempHostPtr;
   atmi_mem_place_t CPU = ATMI_MEM_PLACE_CPU_MEM(0, 0, 0);
-  atmi_status_t ret = atmi_malloc(&tempHostPtr, size, CPU);
-  if (ret != ATMI_STATUS_SUCCESS) {
+  hsa_status_t ret = atmi_malloc(&tempHostPtr, size, CPU);
+  if (ret != HSA_STATUS_SUCCESS) {
     DEBUG_PRINT("atmi_malloc: Unable to alloc %d bytes for temp scratch\n",
                 size);
     return ret;
   }
   std::unique_ptr<void, atmiFreePtrDeletor> del(tempHostPtr);
   memcpy(tempHostPtr, hostSrc, size);
 
   if (invoke_hsa_copy(signal, deviceDest, tempHostPtr, size, agent) !=
       HSA_STATUS_SUCCESS) {
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   }
-  return ATMI_STATUS_SUCCESS;
+  return HSA_STATUS_SUCCESS;
 }
 
-atmi_status_t atmi_memcpy_d2h(hsa_signal_t signal, void *dest,
-                              const void *deviceSrc, size_t size,
-                              hsa_agent_t agent) {
+hsa_status_t atmi_memcpy_d2h(hsa_signal_t signal, void *dest,
   hsa_status_t rc = hsa_memory_copy(dest, deviceSrc, size);
 
   // hsa_memory_copy sometimes fails in situations where
   // allocate + copy succeeds. Looks like it might be related to
   // locking part of a read only segment. Fall back for now.
   if (rc == HSA_STATUS_SUCCESS) {
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
   }
 
   void *tempHostPtr;
   atmi_mem_place_t CPU = ATMI_MEM_PLACE_CPU_MEM(0, 0, 0);
-  atmi_status_t ret = atmi_malloc(&tempHostPtr, size, CPU);
-  if (ret != ATMI_STATUS_SUCCESS) {
+  hsa_status_t ret = atmi_malloc(&tempHostPtr, size, CPU);
+  if (ret != HSA_STATUS_SUCCESS) {
     DEBUG_PRINT("atmi_malloc: Unable to alloc %d bytes for temp scratch\n",
                 size);
     return ret;
   }
   std::unique_ptr<void, atmiFreePtrDeletor> del(tempHostPtr);
 
   if (invoke_hsa_copy(signal, tempHostPtr, deviceSrc, size, agent) !=
       HSA_STATUS_SUCCESS) {
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   }
 
   memcpy(dest, tempHostPtr, size);
-  return ATMI_STATUS_SUCCESS;
+  return HSA_STATUS_SUCCESS;
 }
 
-atmi_status_t atmi_free(void *ptr) { return core::Runtime::Memfree(ptr); }
+hsa_status_t atmi_free(void *ptr) { return core::Runtime::Memfree(ptr); }
 
-atmi_status_t atmi_malloc(void **ptr, size_t size, atmi_mem_place_t place) {
+hsa_status_t atmi_malloc(void **ptr, size_t size, atmi_mem_place_t place) {
   return core::Runtime::Malloc(ptr, size, place);
 }

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

  * @param[in] symbol Pointer to a non-NULL global symbol name
  *
  * @param[in] var_addr Pointer to a non-NULL @p void* variable that will
  * hold the device address of the global symbol object.
  *
  * @param[in] var_size Pointer to a non-NULL @p uint variable that will
  * hold the size of the global symbol object.
  *
- * @retval ::ATMI_STATUS_SUCCESS The function has executed successfully.
+ * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
- * @retval ::ATMI_STATUS_ERROR If @p symbol, @p var_addr or @p var_size are
+ * @retval ::HSA_STATUS_ERROR If @p symbol, @p var_addr or @p var_size are
  * invalid
  * location in the current node, or if ATMI is not initialized.
- *
- * @retval ::ATMI_STATUS_UNKNOWN The function encountered errors.
  */
-atmi_status_t atmi_interop_hsa_get_symbol_info(
+hsa_status_t atmi_interop_hsa_get_symbol_info(
     const std::map<std::string, atl_symbol_info_t> &SymbolInfoTable,
     atmi_mem_place_t place, const char *symbol, void **var_addr,
     unsigned int *var_size);
 /**
  * @brief Get the HSA-specific kernel info from a kernel name
  *
  * @detail Use this function to query the HSA-specific kernel info from the
  * kernel name.
  * This function is meaningful only after calling one
  * of the @p atmi_module_register functions.
  *
  * @param[in] place The ATMI memory place
  *
  * @param[in] kernel_name Pointer to a char array with the kernel name
  *
  * @param[in] info The different possible kernel properties
  *
  * @param[in] value Pointer to a non-NULL @p uint variable that will
  * hold the return value of the kernel property.
  *
- * @retval ::ATMI_STATUS_SUCCESS The function has executed successfully.
+ * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
- * @retval ::ATMI_STATUS_ERROR If @p symbol, @p var_addr or @p var_size are
+ * @retval ::HSA_STATUS_ERROR If @p symbol, @p var_addr or @p var_size are
  * invalid
  * location in the current node, or if ATMI is not initialized.
- *
- * @retval ::ATMI_STATUS_UNKNOWN The function encountered errors.
  */
-atmi_status_t atmi_interop_hsa_get_kernel_info(
+hsa_status_t atmi_interop_hsa_get_kernel_info(
     const std::map<std::string, atl_kernel_info_t> &KernelInfoTable,
     atmi_mem_place_t place, const char *kernel_name,
     hsa_executable_symbol_info_t info, uint32_t *value);
 
 /** @} */
 
 #ifdef __cplusplus
 }
 #endif
 
 #endif // INCLUDE_ATMI_INTEROP_HSA_H_

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

 /*===--------------------------------------------------------------------------
  *              ATMI (Asynchronous Task and Memory Interface)
  *
  * This file is distributed under the MIT License. See LICENSE.txt for details.
  *===------------------------------------------------------------------------*/
 #include "atmi_interop_hsa.h"
 #include "internal.h"
 
-atmi_status_t atmi_interop_hsa_get_symbol_info(
+hsa_status_t atmi_interop_hsa_get_symbol_info(
     const std::map<std::string, atl_symbol_info_t> &SymbolInfoTable,
     atmi_mem_place_t place, const char *symbol, void **var_addr,
     unsigned int *var_size) {
   /*
      // Typical usage:
      void *var_addr;
      size_t var_size;
      atmi_interop_hsa_get_symbol_addr(gpu_place, "symbol_name", &var_addr,
      &var_size);
      atmi_memcpy(signal, host_add, var_addr, var_size);
   */
 
   atmi_machine_t *machine = atmi_machine_get_info();
   if (!symbol || !var_addr || !var_size || !machine)
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   if (place.dev_id < 0 ||
       place.dev_id >= machine->device_count_by_type[place.dev_type])
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
 
   // get the symbol info
   std::string symbolStr = std::string(symbol);
   auto It = SymbolInfoTable.find(symbolStr);
   if (It != SymbolInfoTable.end()) {
     atl_symbol_info_t info = It->second;
     *var_addr = reinterpret_cast<void *>(info.addr);
     *var_size = info.size;
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
   } else {
     *var_addr = NULL;
     *var_size = 0;
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   }
 }
 
-atmi_status_t atmi_interop_hsa_get_kernel_info(
+hsa_status_t atmi_interop_hsa_get_kernel_info(
     const std::map<std::string, atl_kernel_info_t> &KernelInfoTable,
     atmi_mem_place_t place, const char *kernel_name,
     hsa_executable_symbol_info_t kernel_info, uint32_t *value) {
   /*
      // Typical usage:
      uint32_t value;
      atmi_interop_hsa_get_kernel_addr(gpu_place, "kernel_name",
                                   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE,
                                   &val);
   */
 
   atmi_machine_t *machine = atmi_machine_get_info();
   if (!kernel_name || !value || !machine)
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   if (place.dev_id < 0 ||
       place.dev_id >= machine->device_count_by_type[place.dev_type])
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
 
-  atmi_status_t status = ATMI_STATUS_SUCCESS;
+  hsa_status_t status = HSA_STATUS_SUCCESS;
   // get the kernel info
   std::string kernelStr = std::string(kernel_name);
   auto It = KernelInfoTable.find(kernelStr);
   if (It != KernelInfoTable.end()) {
     atl_kernel_info_t info = It->second;
     switch (kernel_info) {
     case HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE:
       *value = info.group_segment_size;
       break;
     case HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE:
       *value = info.private_segment_size;
       break;
     case HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE:
       // return the size for non-implicit args
       *value = info.kernel_segment_size - sizeof(atmi_implicit_args_t);
       break;
     default:
       *value = 0;
-      status = ATMI_STATUS_ERROR;
+      status = HSA_STATUS_ERROR;
       break;
     }
   } else {
     *value = 0;
-    status = ATMI_STATUS_ERROR;
+    status = HSA_STATUS_ERROR;
   }
 
   return status;
 }

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

  * @param[in] place Denotes the execution place (device) on which the module
  * should be registered and loaded.
  *
  * @param[in] on_deserialized_data Callback run on deserialized code object,
  * before loading it
  *
  * @param[in] cb_state void* passed to on_deserialized_data callback
  *
- * @retval ::ATMI_STATUS_SUCCESS The function has executed successfully.
+ * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
- * @retval ::ATMI_STATUS_ERROR The function encountered errors.
- *
- * @retval ::ATMI_STATUS_UNKNOWN The function encountered errors.
+ * @retval ::HSA_STATUS_ERROR The function encountered errors.
  *
  */
-atmi_status_t atmi_module_register_from_memory_to_place(
+hsa_status_t atmi_module_register_from_memory_to_place(
     void *module_bytes, size_t module_size, atmi_place_t place,
-    atmi_status_t (*on_deserialized_data)(void *data, size_t size,
-                                          void *cb_state),
+    hsa_status_t (*on_deserialized_data)(void *data, size_t size,
     void *cb_state);
 
 /** @} */
 
 /** \defgroup machine ATMI Machine
  * @{
  */
 /**

  * then it cannot be accessed by other devices in the system.
  *
  * @param[in] ptr The pointer to the memory that will be allocated.
  *
  * @param[in] size The size of the allocation in bytes.
  *
  * @param[in] place The memory place in the system to perform the allocation.
  *
- * @retval ::ATMI_STATUS_SUCCESS The function has executed successfully.
- *
- * @retval ::ATMI_STATUS_ERROR The function encountered errors.
+ * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
- * @retval ::ATMI_STATUS_UNKNOWN The function encountered errors.
+ * @retval ::HSA_STATUS_ERROR The function encountered errors.
  *
  */
-atmi_status_t atmi_malloc(void **ptr, size_t size, atmi_mem_place_t place);
+hsa_status_t atmi_malloc(void **ptr, size_t size, atmi_mem_place_t place);
 
 /**
  * @brief Frees memory that was previously allocated.
  *
  * @detail This function frees memory that was previously allocated by calling
  * @p atmi_malloc. It throws an error otherwise. It is illegal to access a
  * pointer after a call to this function.
  *
  * @param[in] ptr The pointer to the memory that has to be freed.
  *
- * @retval ::ATMI_STATUS_SUCCESS The function has executed successfully.
- *
- * @retval ::ATMI_STATUS_ERROR The function encountered errors.
+ * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
- * @retval ::ATMI_STATUS_UNKNOWN The function encountered errors.
+ * @retval ::HSA_STATUS_ERROR The function encountered errors.
  *
  */
-atmi_status_t atmi_free(void *ptr);
+hsa_status_t atmi_free(void *ptr);
 
-atmi_status_t atmi_memcpy_h2d(hsa_signal_t signal, void *deviceDest,
-                              const void *hostSrc, size_t size,
-                              hsa_agent_t agent);
+hsa_status_t atmi_memcpy_h2d(hsa_signal_t signal, void *deviceDest,
 
-atmi_status_t atmi_memcpy_d2h(hsa_signal_t sig, void *hostDest,
-                              const void *deviceSrc, size_t size,
-                              hsa_agent_t agent);
+hsa_status_t atmi_memcpy_d2h(hsa_signal_t sig, void *hostDest,
 
 /** @} */
 
 #ifdef __cplusplus
 }
 #endif
 
 #endif // INCLUDE_ATMI_RUNTIME_H_

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

 hsa_status_t register_allocation(void *ptr, size_t size,
                                  atmi_mem_place_t place) {
   if (place.dev_type == ATMI_DEVTYPE_CPU)
     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) {
+hsa_status_t Runtime::Malloc(void **ptr, size_t size, atmi_mem_place_t place) {
   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);
   DEBUG_PRINT("Malloced [%s %d] %p\n",
               place.dev_type == ATMI_DEVTYPE_CPU ? "CPU" : "GPU", place.dev_id,
               *ptr);
 
   if (err == HSA_STATUS_SUCCESS) {
     err = register_allocation(*ptr, size, place);
   }
 
-  return (err == HSA_STATUS_SUCCESS) ? ATMI_STATUS_SUCCESS : ATMI_STATUS_ERROR;
+  return (err == HSA_STATUS_SUCCESS) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
 }
 
-atmi_status_t Runtime::Memfree(void *ptr) {
+hsa_status_t Runtime::Memfree(void *ptr) {
   hsa_status_t err = hsa_amd_memory_pool_free(ptr);
   DEBUG_PRINT("Freed %p\n", ptr);
 
-  return (err == HSA_STATUS_SUCCESS) ? ATMI_STATUS_SUCCESS : ATMI_STATUS_ERROR;
+  return (err == HSA_STATUS_SUCCESS) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
 }
 
 } // namespace core

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

struct lock {
     ~lock() { pthread_mutex_unlock(m); }
     pthread_mutex_t *m;
   };
 };
 
 extern std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
 
 namespace core {
-atmi_status_t atl_init_gpu_context();
+hsa_status_t atl_init_gpu_context();
 
 hsa_status_t init_hsa();
 hsa_status_t finalize_hsa();
 /*
  * Generic utils
  */
 template <typename T> inline T alignDown(T value, size_t alignment) {
   return (T)(value & ~(alignment - 1));

 extern bool atl_is_atmi_initialized();
 
 bool handle_group_signal(hsa_signal_value_t value, void *arg);
 
 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);
+const char *get_atmi_error_string(hsa_status_t err);
 
 #endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_

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

public:
   static Runtime &getInstance() {
     static Runtime instance;
     return instance;
   }
 
   // machine info
   static atmi_machine_t *GetMachineInfo();
   // modules
-  static atmi_status_t RegisterModuleFromMemory(
+  static hsa_status_t RegisterModuleFromMemory(
       void *, size_t, atmi_place_t,
-      atmi_status_t (*on_deserialized_data)(void *data, size_t size,
-                                            void *cb_state),
+      hsa_status_t (*on_deserialized_data)(void *data, size_t size,
       void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
 
   // data
-  static atmi_status_t Memcpy(hsa_signal_t, void *, const void *, size_t);
-  static atmi_status_t Memfree(void *);
-  static atmi_status_t Malloc(void **, size_t, atmi_mem_place_t);
+  static hsa_status_t Memcpy(hsa_signal_t, void *, const void *, size_t);
+  static hsa_status_t Memfree(void *);
+  static hsa_status_t Malloc(void **, size_t, atmi_mem_place_t);
 
   int getMaxQueueSize() const { return env_.getMaxQueueSize(); }
   int getDebugMode() const { return env_.getDebugMode(); }
 
 protected:
   Runtime() = default;
   ~Runtime() = default;
   Runtime(const Runtime &) = delete;

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

if (!(memory_fault.fault_reason_mask & 0x00111111)) {
         err_string += "Can't determine the exact fault address. ";
     }
     fprintf(stderr, "%s\n", err_string.c_str());
     return HSA_STATUS_ERROR;
   }
   return HSA_STATUS_SUCCESS;
 }
 
-atmi_status_t atl_init_gpu_context() {
+hsa_status_t atl_init_gpu_context() {
   if (atlc.struct_initialized == false)
     atmi_init_context_structs();
   if (atlc.g_gpu_initialized != false)
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
 
   hsa_status_t err;
   err = init_hsa();
   if (err != HSA_STATUS_SUCCESS)
-    return ATMI_STATUS_ERROR;
+    return HSA_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));
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   }
 
   init_tasks();
   atlc.g_gpu_initialized = true;
-  return ATMI_STATUS_SUCCESS;
+  return HSA_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:
if (type == HSA_SYMBOL_KIND_KERNEL) {
     SymbolInfoTable[std::string(name)] = info;
     free(name);
   } else {
     DEBUG_PRINT("Symbol is an indirect function\n");
   }
   return HSA_STATUS_SUCCESS;
 }
 
-atmi_status_t RegisterModuleFromMemory(
+hsa_status_t RegisterModuleFromMemory(
     std::map<std::string, atl_kernel_info_t> &KernelInfoTable,
     std::map<std::string, atl_symbol_info_t> &SymbolInfoTable,
     void *module_bytes, size_t module_size, atmi_place_t place,
-    atmi_status_t (*on_deserialized_data)(void *data, size_t size,
-                                          void *cb_state),
+    hsa_status_t (*on_deserialized_data)(void *data, size_t size,
     void *cb_state, std::vector<hsa_executable_t> &HSAExecutables) {
   hsa_status_t err;
   int gpu = place.device_id;
   assert(gpu >= 0);
 
   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);
   if (err != HSA_STATUS_SUCCESS) {
     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
            "Query the agent profile", get_error_string(err));
-    return ATMI_STATUS_ERROR;
+    return HSA_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));
-    return ATMI_STATUS_ERROR;
+    return HSA_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
                     "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 =
+      hsa_status_t atmi_err =
           on_deserialized_data(code_object_alloc_data, module_size, cb_state);
-      if (atmi_err != ATMI_STATUS_SUCCESS) {
+      if (atmi_err != HSA_STATUS_SUCCESS) {
         printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
                "Error in deserialized_data callback",
                get_atmi_error_string(atmi_err));
         return atmi_err;
       }
 
       /* Load the code object.  */
       err =
hsa_status_t RegisterModuleFromMemory(
   } 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));
-      return ATMI_STATUS_ERROR;
+      return HSA_STATUS_ERROR;
     }
 
     err = hsa::executable_iterate_symbols(
         executable,
         [&](hsa_executable_t, hsa_executable_symbol_t symbol) -> hsa_status_t {
           return populate_InfoTables(symbol, gpu, KernelInfoTable,
                                      SymbolInfoTable);
         });
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Iterating over symbols for execuatable", get_error_string(err));
-      return ATMI_STATUS_ERROR;
+      return HSA_STATUS_ERROR;
     }
 
     // save the executable and destroy during finalize
     HSAExecutables.push_back(executable);
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
   } else {
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   }
 }
 
 } // namespace core

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

 #include <iostream>
 #include <pthread.h>
 #include <sched.h>
 #include <stdio.h>
 
 /*
  * Helper functions
  */
-const char *get_atmi_error_string(atmi_status_t err) {
+const char *get_atmi_error_string(hsa_status_t err) {
   switch (err) {
-  case ATMI_STATUS_SUCCESS:
-    return "ATMI_STATUS_SUCCESS";
-  case ATMI_STATUS_UNKNOWN:
-    return "ATMI_STATUS_UNKNOWN";
-  case ATMI_STATUS_ERROR:
-    return "ATMI_STATUS_ERROR";
+  case HSA_STATUS_SUCCESS:
+    return "HSA_STATUS_SUCCESS";
+  case HSA_STATUS_ERROR:
+    return "HSA_STATUS_ERROR";
   default:
     return "";
   }
 }
 
 const char *get_error_string(hsa_status_t err) {
   switch (err) {
   case HSA_STATUS_SUCCESS:

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

hostrpc_assign_buffer(hsa_agent_t, hsa_queue_t *, uint32_t device_id) {
   return 0;
 }
 }
 
 int print_kernel_trace;
 
 #ifdef OMPTARGET_DEBUG
 #define check(msg, status)                                                     \
-  if (status != ATMI_STATUS_SUCCESS) {                                         \
+  if (status != HSA_STATUS_SUCCESS) {                                          \
     DP(#msg " failed\n");                                                      \
   } else {                                                                     \
     DP(#msg " succeeded\n");                                                   \
   }
 #else
 #define check(msg, status)                                                     \
   {}
 #endif
 
 #include "elf_common.h"
 
 namespace core {
-atmi_status_t RegisterModuleFromMemory(
+hsa_status_t RegisterModuleFromMemory(
     std::map<std::string, atl_kernel_info_t> &KernelInfo,
     std::map<std::string, atl_symbol_info_t> &SymbolInfoTable, void *, size_t,
     atmi_place_t,
-    atmi_status_t (*on_deserialized_data)(void *data, size_t size,
-                                          void *cb_state),
+    hsa_status_t (*on_deserialized_data)(void *data, size_t size,
     void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
 }
 
 /// Keep entries table per device
 struct FuncOrGblEntryTy {
   __tgt_target_table Table;
   std::vector<__tgt_offload_entry> Entries;
 };
static const int Max_Teams =
       llvm::omp::AMDGPUGpuGridValues[llvm::omp::GVIDX::GV_Max_Teams];
   static const int Warp_Size =
       llvm::omp::AMDGPUGpuGridValues[llvm::omp::GVIDX::GV_Warp_Size];
   static const int Max_WG_Size =
       llvm::omp::AMDGPUGpuGridValues[llvm::omp::GVIDX::GV_Max_WG_Size];
   static const int Default_WG_Size =
       llvm::omp::AMDGPUGpuGridValues[llvm::omp::GVIDX::GV_Default_WG_Size];
 
-  using MemcpyFunc = atmi_status_t (*)(hsa_signal_t, void *, const void *,
-                                       size_t size, hsa_agent_t);
-  atmi_status_t freesignalpool_memcpy(void *dest, const void *src, size_t size,
-                                      MemcpyFunc Func, int32_t deviceId) {
+  using MemcpyFunc = hsa_status_t (*)(hsa_signal_t, void *, const void *,
+  hsa_status_t freesignalpool_memcpy(void *dest, const void *src, size_t size,
     hsa_agent_t agent = HSAAgents[deviceId];
     hsa_signal_t s = FreeSignalPool.pop();
     if (s.handle == 0) {
-      return ATMI_STATUS_ERROR;
+      return HSA_STATUS_ERROR;
     }
-    atmi_status_t r = Func(s, dest, src, size, agent);
+    hsa_status_t r = Func(s, dest, src, size, agent);
     FreeSignalPool.push(s);
     return r;
   }
 
-  atmi_status_t freesignalpool_memcpy_d2h(void *dest, const void *src,
-                                          size_t size, int32_t deviceId) {
+  hsa_status_t freesignalpool_memcpy_d2h(void *dest, const void *src,
     return freesignalpool_memcpy(dest, src, size, atmi_memcpy_d2h, deviceId);
   }
 
-  atmi_status_t freesignalpool_memcpy_h2d(void *dest, const void *src,
-                                          size_t size, int32_t deviceId) {
+  hsa_status_t freesignalpool_memcpy_h2d(void *dest, const void *src,
     return freesignalpool_memcpy(dest, src, size, atmi_memcpy_h2d, deviceId);
   }
 
   // Record entry point associated with device
   void addOffloadEntry(int32_t device_id, __tgt_offload_entry entry) {
     assert(device_id < (int32_t)FuncGblEntries.size() &&
            "Unexpected device id!");
     FuncOrGblEntryTy &E = FuncGblEntries[device_id].back();
RTLDeviceInfoTy() {
     //  1 => tracing dispatch only
     // >1 => verbosity increase
     if (char *envStr = getenv("LIBOMPTARGET_KERNEL_TRACE"))
       print_kernel_trace = atoi(envStr);
     else
       print_kernel_trace = 0;
 
     DP("Start initializing HSA-ATMI\n");
-    atmi_status_t err = core::atl_init_gpu_context();
-    if (err != ATMI_STATUS_SUCCESS) {
+    hsa_status_t err = core::atl_init_gpu_context();
+    if (err != HSA_STATUS_SUCCESS) {
       DP("Error when initializing HSA-ATMI\n");
       return;
     }
 
     // Init hostcall soon after initializing ATMI
     hostrpc_init();
 
     HSAAgents = find_gpu_agents();

 
 int32_t dataRetrieve(int32_t DeviceId, void *HstPtr, void *TgtPtr, int64_t Size,
                      __tgt_async_info *AsyncInfo) {
   assert(AsyncInfo && "AsyncInfo is nullptr");
   assert(DeviceId < DeviceInfo.NumberOfDevices && "Device ID too large");
   // Return success if we are not copying back to host from target.
   if (!HstPtr)
     return OFFLOAD_SUCCESS;
-  atmi_status_t err;
+  hsa_status_t err;
   DP("Retrieve data %ld bytes, (tgt:%016llx) -> (hst:%016llx).\n", Size,
      (long long unsigned)(Elf64_Addr)TgtPtr,
      (long long unsigned)(Elf64_Addr)HstPtr);
 
   err = DeviceInfo.freesignalpool_memcpy_d2h(HstPtr, TgtPtr, (size_t)Size,
                                              DeviceId);
 
-  if (err != ATMI_STATUS_SUCCESS) {
+  if (err != HSA_STATUS_SUCCESS) {
     DP("Error when copying data from device to host. Pointers: "
        "host = 0x%016lx, device = 0x%016lx, size = %lld\n",
        (Elf64_Addr)HstPtr, (Elf64_Addr)TgtPtr, (unsigned long long)Size);
     return OFFLOAD_FAIL;
   }
   DP("DONE Retrieve data %ld bytes, (tgt:%016llx) -> (hst:%016llx).\n", Size,
      (long long unsigned)(Elf64_Addr)TgtPtr,
      (long long unsigned)(Elf64_Addr)HstPtr);
   return OFFLOAD_SUCCESS;
 }
 
 int32_t dataSubmit(int32_t DeviceId, void *TgtPtr, void *HstPtr, int64_t Size,
                    __tgt_async_info *AsyncInfo) {
   assert(AsyncInfo && "AsyncInfo is nullptr");
-  atmi_status_t err;
+  hsa_status_t err;
   assert(DeviceId < DeviceInfo.NumberOfDevices && "Device ID too large");
   // Return success if we are not doing host to target.
   if (!HstPtr)
     return OFFLOAD_SUCCESS;
 
   DP("Submit data %ld bytes, (hst:%016llx) -> (tgt:%016llx).\n", Size,
      (long long unsigned)(Elf64_Addr)HstPtr,
      (long long unsigned)(Elf64_Addr)TgtPtr);
   err = DeviceInfo.freesignalpool_memcpy_h2d(TgtPtr, HstPtr, (size_t)Size,
                                              DeviceId);
-  if (err != ATMI_STATUS_SUCCESS) {
+  if (err != HSA_STATUS_SUCCESS) {
     DP("Error when copying data from host to device. Pointers: "
        "host = 0x%016lx, device = 0x%016lx, size = %lld\n",
        (Elf64_Addr)HstPtr, (Elf64_Addr)TgtPtr, (unsigned long long)Size);
     return OFFLOAD_FAIL;
   }
   return OFFLOAD_SUCCESS;
 }
 
int get_symbol_info_without_loading(char *base, size_t img_size,
   if (elf) {
     int rc = get_symbol_info_without_loading(elf, base, symname, res);
     elf_end(elf);
     return rc;
   }
   return 1;
 }
 
-atmi_status_t interop_get_symbol_info(char *base, size_t img_size,
-                                      const char *symname, void **var_addr,
-                                      uint32_t *var_size) {
+hsa_status_t interop_get_symbol_info(char *base, size_t img_size,
   symbol_info si;
   int rc = get_symbol_info_without_loading(base, img_size, symname, &si);
   if (rc == 0) {
     *var_addr = si.addr;
     *var_size = si.size;
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
   } else {
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   }
 }
 
 template <typename C>
-atmi_status_t module_register_from_memory_to_place(
+hsa_status_t module_register_from_memory_to_place(
     std::map<std::string, atl_kernel_info_t> &KernelInfoTable,
     std::map<std::string, atl_symbol_info_t> &SymbolInfoTable,
     void *module_bytes, size_t module_size, atmi_place_t place, C cb,
     std::vector<hsa_executable_t> &HSAExecutables) {
-  auto L = [](void *data, size_t size, void *cb_state) -> atmi_status_t {
+  auto L = [](void *data, size_t size, void *cb_state) -> hsa_status_t {
     C *unwrapped = static_cast<C *>(cb_state);
     return (*unwrapped)(data, size);
   };
   return core::RegisterModuleFromMemory(
       KernelInfoTable, SymbolInfoTable, module_bytes, module_size, place, L,
       static_cast<void *>(&cb), HSAExecutables);
 }
 } // namespace
if (si.size > sizeof(host_device_env)) {
       return;
     }
 
     valid = true;
   }
 
   bool in_image() { return si.sh_type != SHT_NOBITS; }
 
-  atmi_status_t before_loading(void *data, size_t size) {
+  hsa_status_t before_loading(void *data, size_t size) {
     if (valid) {
       if (in_image()) {
         DP("Setting global device environment before load (%u bytes)\n",
            si.size);
         uint64_t offset = (char *)si.addr - (char *)image->ImageStart;
         void *pos = (char *)data + offset;
         memcpy(pos, &host_device_env, si.size);
       }
     }
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
   }
 
-  atmi_status_t after_loading() {
+  hsa_status_t after_loading() {
     if (valid) {
       if (!in_image()) {
         DP("Setting global device environment after load (%u bytes)\n",
            si.size);
         int device_id = host_device_env.device_num;
         auto &SymbolInfo = DeviceInfo.SymbolInfoTable[device_id];
         void *state_ptr;
         uint32_t state_ptr_size;
-        atmi_status_t err = atmi_interop_hsa_get_symbol_info(
+        hsa_status_t err = atmi_interop_hsa_get_symbol_info(
             SymbolInfo, get_gpu_mem_place(device_id), sym(), &state_ptr,
             &state_ptr_size);
-        if (err != ATMI_STATUS_SUCCESS) {
+        if (err != HSA_STATUS_SUCCESS) {
           DP("failed to find %s in loaded image\n", sym());
           return err;
         }
 
         if (state_ptr_size != si.size) {
           DP("Symbol had size %u before loading, %u after\n", state_ptr_size,
              si.size);
-          return ATMI_STATUS_ERROR;
+          return HSA_STATUS_ERROR;
         }
 
         return DeviceInfo.freesignalpool_memcpy_h2d(state_ptr, &host_device_env,
                                                     state_ptr_size, device_id);
       }
     }
-    return ATMI_STATUS_SUCCESS;
+    return HSA_STATUS_SUCCESS;
   }
 };
 
-static atmi_status_t atmi_calloc(void **ret_ptr, size_t size,
-                                 atmi_mem_place_t place) {
+static hsa_status_t atmi_calloc(void **ret_ptr, size_t size,
   uint64_t rounded = 4 * ((size + 3) / 4);
   void *ptr;
-  atmi_status_t err = atmi_malloc(&ptr, rounded, place);
-  if (err != ATMI_STATUS_SUCCESS) {
+  hsa_status_t err = atmi_malloc(&ptr, rounded, place);
+  if (err != HSA_STATUS_SUCCESS) {
     return err;
   }
 
   hsa_status_t rc = hsa_amd_memory_fill(ptr, 0, rounded / 4);
   if (rc != HSA_STATUS_SUCCESS) {
     fprintf(stderr, "zero fill device_state failed with %u\n", rc);
     atmi_free(ptr);
-    return ATMI_STATUS_ERROR;
+    return HSA_STATUS_ERROR;
   }
 
   *ret_ptr = ptr;
-  return ATMI_STATUS_SUCCESS;
+  return HSA_STATUS_SUCCESS;
 }
 
 static bool image_contains_symbol(void *data, size_t size, const char *sym) {
   symbol_info si;
   int rc = get_symbol_info_without_loading((char *)data, size, sym, &si);
   return (rc == 0) && (si.addr != nullptr);
 }
 
__tgt_target_table *__tgt_rtl_load_binary_locked(int32_t device_id,
   }
 
   {
     auto env = device_environment(device_id, DeviceInfo.NumberOfDevices, image,
                                   img_size);
 
     auto &KernelInfo = DeviceInfo.KernelInfoTable[device_id];
     auto &SymbolInfo = DeviceInfo.SymbolInfoTable[device_id];
-    atmi_status_t err = module_register_from_memory_to_place(
+    hsa_status_t err = module_register_from_memory_to_place(
         KernelInfo, SymbolInfo, (void *)image->ImageStart, img_size,
         get_gpu_place(device_id),
         [&](void *data, size_t size) {
           if (image_contains_symbol(data, size, "needs_hostcall_buffer")) {
             __atomic_store_n(&DeviceInfo.hostcall_required, true,
                              __ATOMIC_RELEASE);
           }
           return env.before_loading(data, size);
         },
         DeviceInfo.HSAExecutables);
 
     check("Module registering", err);
-    if (err != ATMI_STATUS_SUCCESS) {
+    if (err != HSA_STATUS_SUCCESS) {
       fprintf(stderr,
               "Possible gpu arch mismatch: device:%s, image:%s please check"
               " compiler flag: -march=<gpu>\n",
               DeviceInfo.GPUName[device_id].c_str(),
               get_elf_mach_gfx_name(elf_e_flags(image)));
       return NULL;
     }
 
     err = env.after_loading();
-    if (err != ATMI_STATUS_SUCCESS) {
+    if (err != HSA_STATUS_SUCCESS) {
       return NULL;
     }
   }
 
   DP("ATMI module successfully loaded!\n");
 
   {
     // the device_State array is either large value in bss or a void* that
     // needs to be assigned to a pointer to an array of size device_state_bytes
     // If absent, it has been deadstripped and needs no setup.
 
     void *state_ptr;
     uint32_t state_ptr_size;
     auto &SymbolInfoMap = DeviceInfo.SymbolInfoTable[device_id];
-    atmi_status_t err = atmi_interop_hsa_get_symbol_info(
+    hsa_status_t err = atmi_interop_hsa_get_symbol_info(
         SymbolInfoMap, get_gpu_mem_place(device_id),
         "omptarget_nvptx_device_State", &state_ptr, &state_ptr_size);
 
-    if (err != ATMI_STATUS_SUCCESS) {
+    if (err != HSA_STATUS_SUCCESS) {
       DP("No device_state symbol found, skipping initialization\n");
     } else {
       if (state_ptr_size < sizeof(void *)) {
         DP("unexpected size of state_ptr %u != %zu\n", state_ptr_size,
            sizeof(void *));
         return NULL;
       }
 
       // if it's larger than a void*, assume it's a bss array and no further
       // initialization is required. Only try to set up a pointer for
       // sizeof(void*)
       if (state_ptr_size == sizeof(void *)) {
         uint64_t device_State_bytes =
             get_device_State_bytes((char *)image->ImageStart, img_size);
         if (device_State_bytes == 0) {
           DP("Can't initialize device_State, missing size information\n");
           return NULL;
         }
 
         auto &dss = DeviceInfo.deviceStateStore[device_id];
         if (dss.first.get() == nullptr) {
           assert(dss.second == 0);
           void *ptr = NULL;
-          atmi_status_t err = atmi_calloc(&ptr, device_State_bytes,
-                                          get_gpu_mem_place(device_id));
-          if (err != ATMI_STATUS_SUCCESS) {
+          hsa_status_t err = atmi_calloc(&ptr, device_State_bytes,
+          if (err != HSA_STATUS_SUCCESS) {
             DP("Failed to allocate device_state array\n");
             return NULL;
           }
           dss = {
               std::unique_ptr<void, RTLDeviceInfoTy::atmiFreePtrDeletor>{ptr},
               device_State_bytes,
           };
         }
 
         void *ptr = dss.first.get();
         if (device_State_bytes != dss.second) {
           DP("Inconsistent sizes of device_State unsupported\n");
           return NULL;
         }
 
         // write ptr to device memory so it can be used by later kernels
         err = DeviceInfo.freesignalpool_memcpy_h2d(state_ptr, &ptr,
                                                    sizeof(void *), device_id);
-        if (err != ATMI_STATUS_SUCCESS) {
+        if (err != HSA_STATUS_SUCCESS) {
           DP("memcpy install of state_ptr failed\n");
           return NULL;
         }
       }
     }
   }
 
   // Here, we take advantage of the data that is appended after img_end to get
for (__tgt_offload_entry *e = HostBegin; e != HostEnd; ++e) {
 
     if (e->size) {
       __tgt_offload_entry entry = *e;
 
       void *varptr;
       uint32_t varsize;
 
       auto &SymbolInfoMap = DeviceInfo.SymbolInfoTable[device_id];
-      atmi_status_t err = atmi_interop_hsa_get_symbol_info(
+      hsa_status_t err = atmi_interop_hsa_get_symbol_info(
           SymbolInfoMap, get_gpu_mem_place(device_id), e->name, &varptr,
           &varsize);
 
-      if (err != ATMI_STATUS_SUCCESS) {
+      if (err != HSA_STATUS_SUCCESS) {
         // Inform the user what symbol prevented offloading
         DP("Loading global '%s' (Failed)\n", e->name);
         return NULL;
       }
 
       if (varsize != e->size) {
         DP("Loading global '%s' - size mismatch (%u != %lu)\n", e->name,
            varsize, e->size);
         return NULL;
       }
 
       DP("Entry point " DPxMOD " maps to global %s (" DPxMOD ")\n",
          DPxPTR(e - HostBegin), e->name, DPxPTR(varptr));
       entry.addr = (void *)varptr;
 
       DeviceInfo.addOffloadEntry(device_id, entry);
 
       if (DeviceInfo.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
           e->flags & OMP_DECLARE_TARGET_LINK) {
         // If unified memory is present any target link variables
         // can access host addresses directly. There is no longer a
         // need for device copies.
         err = DeviceInfo.freesignalpool_memcpy_h2d(varptr, e->addr,
                                                    sizeof(void *), device_id);
-        if (err != ATMI_STATUS_SUCCESS)
+        if (err != HSA_STATUS_SUCCESS)
           DP("Error when copying USM\n");
         DP("Copy linked variable host address (" DPxMOD ")"
            "to device address (" DPxMOD ")\n",
            DPxPTR(*((void **)e->addr)), DPxPTR(varptr));
       }
 
       continue;
     }
 
     DP("to find the kernel name: %s size: %lu\n", e->name, strlen(e->name));
 
     atmi_mem_place_t place = get_gpu_mem_place(device_id);
     uint32_t kernarg_segment_size;
     auto &KernelInfoMap = DeviceInfo.KernelInfoTable[device_id];
-    atmi_status_t err = atmi_interop_hsa_get_kernel_info(
+    hsa_status_t err = atmi_interop_hsa_get_kernel_info(
         KernelInfoMap, place, e->name,
         HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE,
         &kernarg_segment_size);
 
     // each arg is a void * in this openmp implementation
     uint32_t arg_num = kernarg_segment_size / sizeof(void *);
     std::vector<size_t> arg_sizes(arg_num);
     for (std::vector<size_t>::iterator it = arg_sizes.begin();
for (__tgt_offload_entry *e = HostBegin; e != HostEnd; ++e) {
     const char *KernDescName = KernDescNameStr.c_str();
 
     void *KernDescPtr;
     uint32_t KernDescSize;
     void *CallStackAddr = nullptr;
     err = interop_get_symbol_info((char *)image->ImageStart, img_size,
                                   KernDescName, &KernDescPtr, &KernDescSize);
 
-    if (err == ATMI_STATUS_SUCCESS) {
+    if (err == HSA_STATUS_SUCCESS) {
       if ((size_t)KernDescSize != sizeof(KernDescVal))
         DP("Loading global computation properties '%s' - size mismatch (%u != "
            "%lu)\n",
            KernDescName, KernDescSize, sizeof(KernDescVal));
 
       memcpy(&KernDescVal, KernDescPtr, (size_t)KernDescSize);
 
       // Check structure size against recorded size.
if (err == HSA_STATUS_SUCCESS) {
       ExecModeNameStr += "_exec_mode";
       const char *ExecModeName = ExecModeNameStr.c_str();
 
       void *ExecModePtr;
       uint32_t varsize;
       err = interop_get_symbol_info((char *)image->ImageStart, img_size,
                                     ExecModeName, &ExecModePtr, &varsize);
 
-      if (err == ATMI_STATUS_SUCCESS) {
+      if (err == HSA_STATUS_SUCCESS) {
         if ((size_t)varsize != sizeof(int8_t)) {
           DP("Loading global computation properties '%s' - size mismatch(%u != "
              "%lu)\n",
              ExecModeName, varsize, sizeof(int8_t));
           return NULL;
         }
 
         memcpy(&ExecModeVal, ExecModePtr, (size_t)varsize);
if (err == HSA_STATUS_SUCCESS) {
       WGSizeNameStr += "_wg_size";
       const char *WGSizeName = WGSizeNameStr.c_str();
 
       void *WGSizePtr;
       uint32_t WGSize;
       err = interop_get_symbol_info((char *)image->ImageStart, img_size,
                                     WGSizeName, &WGSizePtr, &WGSize);
 
-      if (err == ATMI_STATUS_SUCCESS) {
+      if (err == HSA_STATUS_SUCCESS) {
         if ((size_t)WGSize != sizeof(int16_t)) {
           DP("Loading global computation properties '%s' - size mismatch (%u "
              "!= "
              "%lu)\n",
              WGSizeName, WGSize, sizeof(int16_t));
           return NULL;
         }
 
void *__tgt_rtl_data_alloc(int device_id, int64_t size, void *, int32_t kind) {
   assert(device_id < DeviceInfo.NumberOfDevices && "Device ID too large");
 
   if (kind != TARGET_ALLOC_DEFAULT) {
     REPORT("Invalid target data allocation kind or requested allocator not "
            "implemented yet\n");
     return NULL;
   }
 
-  atmi_status_t err = atmi_malloc(&ptr, size, get_gpu_mem_place(device_id));
+  hsa_status_t err = atmi_malloc(&ptr, size, get_gpu_mem_place(device_id));
   DP("Tgt alloc data %ld bytes, (tgt:%016llx).\n", size,
      (long long unsigned)(Elf64_Addr)ptr);
-  ptr = (err == ATMI_STATUS_SUCCESS) ? ptr : NULL;
+  ptr = (err == HSA_STATUS_SUCCESS) ? ptr : NULL;
   return ptr;
 }
 
 int32_t __tgt_rtl_data_submit(int device_id, void *tgt_ptr, void *hst_ptr,
                               int64_t size) {
   assert(device_id < DeviceInfo.NumberOfDevices && "Device ID too large");
   __tgt_async_info AsyncInfo;
   int32_t rc = dataSubmit(device_id, tgt_ptr, hst_ptr, size, &AsyncInfo);
int32_t __tgt_rtl_data_retrieve_async(int device_id, void *hst_ptr,
   assert(AsyncInfo && "AsyncInfo is nullptr");
   assert(device_id < DeviceInfo.NumberOfDevices && "Device ID too large");
   initAsyncInfo(AsyncInfo);
   return dataRetrieve(device_id, hst_ptr, tgt_ptr, size, AsyncInfo);
 }
 
 int32_t __tgt_rtl_data_delete(int device_id, void *tgt_ptr) {
   assert(device_id < DeviceInfo.NumberOfDevices && "Device ID too large");
-  atmi_status_t err;
+  hsa_status_t err;
   DP("Tgt free data (tgt:%016llx).\n", (long long unsigned)(Elf64_Addr)tgt_ptr);
   err = atmi_free(tgt_ptr);
-  if (err != ATMI_STATUS_SUCCESS) {
+  if (err != HSA_STATUS_SUCCESS) {
     DP("Error when freeing CUDA memory\n");
     return OFFLOAD_FAIL;
   }
   return OFFLOAD_SUCCESS;
 }
 
 // Determine launch values for threadsPerGroup and num_groups.
 // Outputs: treadsPerGroup, num_groups