[AMDGPU][Libomptarget][NFC] Remove atmi_mem_place_t

[pdhaliwal]

Summary

This struct was used to specify the device on which memory was
being allocated/free in atmi_malloc/free. It has now been replaced
with int DeviceId.

[pdhaliwal]

Remove whitespace changes.

[JonChesterfield]

Nice, like that a lot. I hadn't noticed that we are passing device id == 0 around to mean a CPU device, that seems a path to considerable confusion. Probably better to split the allocators into gpu (which takes a device id) and cpu (which doesn't) in some later patch.

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

typedef struct atmi_place_s {
   atmi_devtype_t type;
   /**
    * The device ordinal number ordered by runtime; -1 for any
    */
   int device_id;
 } atmi_place_t;
 
 /**
- * @brief ATMI Memory Place
- */
-typedef struct atmi_mem_place_s {
-  /**
-   * The node in a cluster where computation should occur.
-   * Default is node_id = 0 for local computations.
-   */
-  unsigned int node_id;
-  /**
-   * Device type: CPU, GPU or DSP
-   */
-  atmi_devtype_t dev_type;
-  /**
-   * The device ordinal number ordered by runtime; -1 for any
-   */
-  int dev_id;
-  // atmi_memtype_t mem_type;        // Fine grained or Coarse grained
-  /**
-   * The memory space/region ordinal number ordered by runtime; -1 for any
-   */
-  int mem_id;
-} atmi_mem_place_t;
-
-/**
  * @brief ATMI Memory Space/region Structure
  */
 typedef struct atmi_memory_s {
   /**
    * Memory capacity
    */
   unsigned long int capacity;
   /**

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

hsa_status_t atmi_memcpy_h2d(hsa_signal_t signal, void *deviceDest,
   // 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 HSA_STATUS_SUCCESS;
   }
 
   void *tempHostPtr;
-  atmi_mem_place_t CPU = ATMI_MEM_PLACE_CPU_MEM(0, 0, 0);
-  hsa_status_t ret = atmi_malloc(&tempHostPtr, size, CPU);
+  hsa_status_t ret =
+      atmi_malloc(&tempHostPtr, size, 0 /* DeviceId */, ATMI_DEVTYPE_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);
 
hsa_status_t atmi_memcpy_d2h(hsa_signal_t signal, void *dest,
   // 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 HSA_STATUS_SUCCESS;
   }
 
   void *tempHostPtr;
-  atmi_mem_place_t CPU = ATMI_MEM_PLACE_CPU_MEM(0, 0, 0);
-  hsa_status_t ret = atmi_malloc(&tempHostPtr, size, CPU);
+
+  hsa_status_t ret =
+      atmi_malloc(&tempHostPtr, size, 0 /* DeviceId */, ATMI_DEVTYPE_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 HSA_STATUS_ERROR;
   }
 
   memcpy(dest, tempHostPtr, size);
   return HSA_STATUS_SUCCESS;
 }
 
 hsa_status_t atmi_free(void *ptr) { return core::Runtime::Memfree(ptr); }
 
-hsa_status_t atmi_malloc(void **ptr, size_t size, atmi_mem_place_t place) {
-  return core::Runtime::Malloc(ptr, size, place);
+hsa_status_t atmi_malloc(void **ptr, size_t size, int DeviceId,
+                         atmi_devtype_t DeviceType) {
+  return core::Runtime::Malloc(ptr, size, DeviceId, DeviceType);
 }

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

  * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
  * @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.
  */
 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);
+    int DeviceId, 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.
  *

  * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
  * @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.
  */
 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);
+    int DeviceId, 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"
 
 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) {
+    int DeviceId, 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 HSA_STATUS_ERROR;
-  if (place.dev_id < 0 ||
-      place.dev_id >= machine->device_count_by_type[place.dev_type])
+  if (DeviceId < 0 ||
+      DeviceId >= machine->device_count_by_type[ATMI_DEVTYPE_GPU])
     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 HSA_STATUS_SUCCESS;
   } else {
     *var_addr = NULL;
     *var_size = 0;
     return HSA_STATUS_ERROR;
   }
 }
 
 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,
+    int DeviceId, 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 HSA_STATUS_ERROR;
-  if (place.dev_id < 0 ||
-      place.dev_id >= machine->device_count_by_type[place.dev_type])
+  if (DeviceId < 0 ||
+      DeviceId >= machine->device_count_by_type[ATMI_DEVTYPE_GPU])
     return HSA_STATUS_ERROR;
 
   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;

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

 /*===--------------------------------------------------------------------------
  *              ATMI (Asynchronous Task and Memory Interface)
  *
  * This file is distributed under the MIT License. See LICENSE.txt for details.
  *===------------------------------------------------------------------------*/
 #ifndef INCLUDE_ATMI_RUNTIME_H_
 #define INCLUDE_ATMI_RUNTIME_H_
 
 #include "atmi.h"
 #include "hsa.h"
 #include <inttypes.h>
 #include <stdlib.h>
 #ifndef __cplusplus
 #include <stdbool.h>
 #endif
 
 #ifdef __cplusplus
 extern "C" {

  *
  * @param[in] place The memory place in the system to perform the allocation.
  *
  * @retval ::HSA_STATUS_SUCCESS The function has executed successfully.
  *
  * @retval ::HSA_STATUS_ERROR The function encountered errors.
  *
  */
-hsa_status_t atmi_malloc(void **ptr, size_t size, atmi_mem_place_t place);
-
+hsa_status_t atmi_malloc(void **ptr, size_t size, int DeviceId,
+                         atmi_devtype_t DeviceType);
 /**
  * @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.

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

 #include <vector>
 
 using core::TaskImpl;
 extern ATLMachine g_atl_machine;
 
 namespace core {
 
 namespace {
-ATLProcessor &get_processor_by_mem_place(atmi_mem_place_t place) {
-  int dev_id = place.dev_id;
-  switch (place.dev_type) {
+ATLProcessor &get_processor_by_mem_place(int DeviceId,
+                                         atmi_devtype_t DeviceType) {
+  switch (DeviceType) {
   case ATMI_DEVTYPE_CPU:
-    return g_atl_machine.processors<ATLCPUProcessor>()[dev_id];
+    return g_atl_machine.processors<ATLCPUProcessor>()[DeviceId];
   case ATMI_DEVTYPE_GPU:
-    return g_atl_machine.processors<ATLGPUProcessor>()[dev_id];
+    return g_atl_machine.processors<ATLGPUProcessor>()[DeviceId];
   }
 }
 
-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);
+hsa_amd_memory_pool_t get_memory_pool_by_mem_place(int DeviceId,
+                                                   atmi_devtype_t DeviceType) {
+  ATLProcessor &proc = get_processor_by_mem_place(DeviceId, DeviceType);
+  return get_memory_pool(proc, 0 /*Memory Type (always zero) */);
 }
 } // namespace
 
 hsa_status_t register_allocation(void *ptr, size_t size,
-                                 atmi_mem_place_t place) {
-  if (place.dev_type == ATMI_DEVTYPE_CPU)
+                                 atmi_devtype_t DeviceType) {
+  if (DeviceType == ATMI_DEVTYPE_CPU)
     return allow_access_to_all_gpu_agents(ptr);
   else
     return HSA_STATUS_SUCCESS;
 }
 
-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 Runtime::Malloc(void **ptr, size_t size, int DeviceId,
+                             atmi_devtype_t DeviceType) {
+  hsa_amd_memory_pool_t pool =
+      get_memory_pool_by_mem_place(DeviceId, DeviceType);
   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);
+              DeviceType == ATMI_DEVTYPE_CPU ? "CPU" : "GPU", DeviceId, *ptr);
 
   if (err == HSA_STATUS_SUCCESS) {
-    err = register_allocation(*ptr, size, place);
+    err = register_allocation(*ptr, size, DeviceType);
   }
 
   return (err == HSA_STATUS_SUCCESS) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
 }
 
 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) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
 }
 
 } // namespace core

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

 #include <cstring>
 #include <deque>
 #include <map>
 #include <queue>
 #include <string>
 #include <utility>
 #include <vector>
 
 #include "hsa.h"
 #include "hsa_ext_amd.h"
 #include "hsa_ext_finalize.h"
 
 #include "atmi.h"
 #include "atmi_runtime.h"
 #include "rt.h"
 
 #define MAX_NUM_KERNELS (1024 * 16)

 }
 
 template <typename T> inline T *alignUp(T *value, size_t alignment) {
   return reinterpret_cast<T *>(
       alignDown((intptr_t)(value + alignment - 1), alignment));
 }
 
 hsa_status_t register_allocation(void *addr, size_t size,
-                                 atmi_mem_place_t place);
+                                 atmi_devtype_t DeviceType);
 
 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(hsa_status_t err);
 
 #endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_

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

static hsa_status_t RegisterModuleFromMemory(
       void *, size_t, atmi_place_t,
       hsa_status_t (*on_deserialized_data)(void *data, size_t size,
                                            void *cb_state),
       void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
 
   // data
   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);
+  static hsa_status_t Malloc(void **ptr, size_t size, int DeviceId,
+                             atmi_devtype_t DeviceType);
 
   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

 /*===--------------------------------------------------------------------------
  *              ATMI (Asynchronous Task and Memory Interface)
  *
  * This file is distributed under the MIT License. See LICENSE.txt for details.
  *===------------------------------------------------------------------------*/
 #include <gelf.h>
 #include <libelf.h>
 
 #include <cassert>
 #include <cstdarg>
 #include <fstream>
 #include <iomanip>
 #include <iostream>
 #include <set>
if (type == HSA_SYMBOL_KIND_KERNEL) {
     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));
       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);
     err = register_allocation(reinterpret_cast<void *>(info.addr),
-                              (size_t)info.size, place);
+                              (size_t)info.size, ATMI_DEVTYPE_GPU);
     if (err != HSA_STATUS_SUCCESS) {
       return err;
     }
     SymbolInfoTable[std::string(name)] = info;
     free(name);
   } else {
     DEBUG_PRINT("Symbol is an indirect function\n");
   }

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

 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <dlfcn.h>
 #include <elf.h>
 #include <ffi.h>
 #include <fstream>
 #include <iostream>
 #include <libelf.h>
 #include <list>
 #include <memory>
 #include <mutex>
 #include <shared_mutex>
 #include <thread>
 #include <unordered_map>
 #include <vector>
 

 /// List that contains all the kernels.
 /// FIXME: we may need this to be per device and per library.
 std::list<KernelTy> KernelsList;
 
 // ATMI API to get gpu and gpu memory place
 static atmi_place_t get_gpu_place(int device_id) {
   return ATMI_PLACE_GPU(0, device_id);
 }
-static atmi_mem_place_t get_gpu_mem_place(int device_id) {
-  return ATMI_MEM_PLACE_GPU_MEM(0, device_id, 0);
-}
 
 static std::vector<hsa_agent_t> find_gpu_agents() {
   std::vector<hsa_agent_t> res;
 
   hsa_status_t err = hsa_iterate_agents(
       [](hsa_agent_t agent, void *data) -> hsa_status_t {
         std::vector<hsa_agent_t> *res =
             static_cast<std::vector<hsa_agent_t> *>(data);
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;
         hsa_status_t err = atmi_interop_hsa_get_symbol_info(
-            SymbolInfo, get_gpu_mem_place(device_id), sym(), &state_ptr,
-            &state_ptr_size);
+            SymbolInfo, device_id, sym(), &state_ptr, &state_ptr_size);
         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 HSA_STATUS_ERROR;
         }
 
         return DeviceInfo.freesignalpool_memcpy_h2d(state_ptr, &host_device_env,
                                                     state_ptr_size, device_id);
       }
     }
     return HSA_STATUS_SUCCESS;
   }
 };
 
-static hsa_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, int DeviceId) {
   uint64_t rounded = 4 * ((size + 3) / 4);
   void *ptr;
-  hsa_status_t err = atmi_malloc(&ptr, rounded, place);
+  hsa_status_t err = atmi_malloc(&ptr, rounded, DeviceId, ATMI_DEVTYPE_GPU);
   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);
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];
     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);
+        SymbolInfoMap, device_id, "omptarget_nvptx_device_State", &state_ptr,
+        &state_ptr_size);
 
     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 (err != HSA_STATUS_SUCCESS) {
           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;
-          hsa_status_t err = atmi_calloc(&ptr, device_State_bytes,
-                                         get_gpu_mem_place(device_id));
+          hsa_status_t err = atmi_calloc(&ptr, device_State_bytes, device_id);
           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,
           };
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];
       hsa_status_t err = atmi_interop_hsa_get_symbol_info(
-          SymbolInfoMap, get_gpu_mem_place(device_id), e->name, &varptr,
-          &varsize);
+          SymbolInfoMap, device_id, e->name, &varptr, &varsize);
 
       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) {
if (e->size) {
            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];
     hsa_status_t err = atmi_interop_hsa_get_kernel_info(
-        KernelInfoMap, place, e->name,
+        KernelInfoMap, device_id, 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();
          it != arg_sizes.end(); it++) {
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;
   }
 
-  hsa_status_t err = atmi_malloc(&ptr, size, get_gpu_mem_place(device_id));
+  hsa_status_t err = atmi_malloc(&ptr, size, device_id, ATMI_DEVTYPE_GPU);
   DP("Tgt alloc data %ld bytes, (tgt:%016llx).\n", size,
      (long long unsigned)(Elf64_Addr)ptr);
   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) {