[libomptarget][amdgpu] Drop env variables

[JonChesterfield]

Summary

Use the same debug print as the rest of libomptarget plugins with
the same environment control. Also drop the max queue size debugging hook as
I don't believe it is still in use, can bring it back near the rest of the env
handling in rtl.cpp if someone objects.

That makes most of rt.h and all of utils.cpp unused. Clean that up and simplify
control flow in a couple of places.

Behaviour change is that debug prints that used to use the old environment
variable now use the new one and print in slightly different format, and the
removal of the max queue size variable.

[JonChesterfield]

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

24

This was previously forward declared in rtl.cpp, moving it to a header to catch divergence in prototype vs implementation

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

31

Potentially contentious point is here. Previously we let stoi(ATMI_MAX_HSA_QUEUE_SIZE) override the maximum size HSA queue created. This patch leaves the 4096 limit in place (which seems low to me) but removes the environment control.

We might want a whole suite of environment variables for overriding things like queue size, but I don't see anything particularly special about this one. Suggest we drop it and if someone objects reinstate it (in rtl.cpp, near the other environment uses)

[JonChesterfield]

• ensure debug.h has the DEBUG_PREFIX macro defined before including

[jdoerfert]

I like that the plugin shrinks. No objections here.

[pdhaliwal]

LGTM

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

 endif()
 
 add_library(omptarget.rtl.amdgpu SHARED
       impl/impl.cpp
       impl/interop_hsa.cpp
       impl/data.cpp
       impl/get_elf_mach_gfx_name.cpp
       impl/system.cpp
-      impl/utils.cpp
       impl/msgpack.cpp
       src/rtl.cpp
       ${LIBOMPTARGET_EXTRA_SOURCE}
       )
 
 # Install plugin under the lib destination folder.
 # When we build for debug, OPENMP_LIBDIR_SUFFIX get set to -debug
 install(TARGETS omptarget.rtl.amdgpu LIBRARY DESTINATION "lib${OPENMP_LIBDIR_SUFFIX}")

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

 #include <cassert>
 #include <stdio.h>
 #include <string.h>
 #include <vector>
 
 using core::TaskImpl;
 
 namespace core {
-
-hsa_status_t Runtime::HostMalloc(void **ptr, size_t size,
-                                 hsa_amd_memory_pool_t MemoryPool) {
+namespace Runtime {
+hsa_status_t HostMalloc(void **ptr, size_t size,
   hsa_status_t err = hsa_amd_memory_pool_allocate(MemoryPool, size, 0, ptr);
-  DEBUG_PRINT("Malloced %p\n", *ptr);
-
+  DP("Malloced %p\n", *ptr);
   if (err == HSA_STATUS_SUCCESS) {
     err = core::allow_access_to_all_gpu_agents(*ptr);
   }
-  return (err == HSA_STATUS_SUCCESS) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
+  return err;
 }
 
-hsa_status_t Runtime::Memfree(void *ptr) {
+hsa_status_t 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;
+  DP("Freed %p\n", ptr);
+  return err;
 }
-
+} // namespace Runtime
 } // namespace core

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

 //===--- amdgpu/impl/impl.cpp ------------------------------------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-#include "impl_runtime.h"
 #include "hsa_api.h"
+#include "impl_runtime.h"
 #include "internal.h"
 #include "rt.h"
 #include <memory>
 
 /*
  * Data
  */
 
static hsa_status_t invoke_hsa_copy(hsa_signal_t sig, void *dest,
   if (err != HSA_STATUS_SUCCESS) {
     return err;
   }
 
   // async_copy reports success by decrementing and failure by setting to < 0
   hsa_signal_value_t got = init;
   while (got == init) {
     got = hsa_signal_wait_scacquire(sig, HSA_SIGNAL_CONDITION_NE, init,
-                                    UINT64_MAX, ATMI_WAIT_STATE);
+                                    UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
   }
 
   if (got != success) {
     return HSA_STATUS_ERROR;
   }
 
   return err;
 }
hsa_status_t impl_memcpy_h2d(hsa_signal_t signal, void *deviceDest,
   // locking part of a read only segment. Fall back for now.
   if (rc == HSA_STATUS_SUCCESS) {
     return HSA_STATUS_SUCCESS;
   }
 
   void *tempHostPtr;
   hsa_status_t ret = core::Runtime::HostMalloc(&tempHostPtr, size, MemoryPool);
   if (ret != HSA_STATUS_SUCCESS) {
-    DEBUG_PRINT("HostMalloc: Unable to alloc %zu bytes for temp scratch\n",
-                size);
+    DP("HostMalloc: Unable to alloc %zu bytes for temp scratch\n", size);
     return ret;
   }
   std::unique_ptr<void, implFreePtrDeletor> del(tempHostPtr);
   memcpy(tempHostPtr, hostSrc, size);
 
   if (invoke_hsa_copy(signal, deviceDest, tempHostPtr, size, agent) !=
       HSA_STATUS_SUCCESS) {
     return HSA_STATUS_ERROR;
hsa_status_t impl_memcpy_d2h(hsa_signal_t signal, void *dest,
   // locking part of a read only segment. Fall back for now.
   if (rc == HSA_STATUS_SUCCESS) {
     return HSA_STATUS_SUCCESS;
   }
 
   void *tempHostPtr;
   hsa_status_t ret = core::Runtime::HostMalloc(&tempHostPtr, size, MemoryPool);
   if (ret != HSA_STATUS_SUCCESS) {
-    DEBUG_PRINT("HostMalloc: Unable to alloc %zu bytes for temp scratch\n",
-                size);
+    DP("HostMalloc: Unable to alloc %zu bytes for temp scratch\n", size);
     return ret;
   }
   std::unique_ptr<void, implFreePtrDeletor> 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;
 }

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

 #include <queue>
 #include <string>
 #include <utility>
 #include <vector>
 
 #include "hsa_api.h"
 
 #include "impl_runtime.h"
-#include "rt.h"
+
+#ifndef TARGET_NAME
+#error "Missing TARGET_NAME macro"
+#endif
+#define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
+#include "Debug.h"
 
 #define MAX_NUM_KERNELS (1024 * 16)
 
 typedef struct impl_implicit_args_s {
   unsigned long offset_x;
   unsigned long offset_y;
   unsigned long offset_z;
   unsigned long hostcall_ptr;
   char num_gpu_queues;
   unsigned long gpu_queue_ptr;
   char num_cpu_queues;
   unsigned long cpu_worker_signals;
   unsigned long cpu_queue_ptr;
   unsigned long kernarg_template_ptr;
 } impl_implicit_args_t;
 
-extern "C" {
-
-#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 {                                                                         \
-  } while (false)
-#endif
-
-#ifndef HSA_RUNTIME_INC_HSA_H_
-typedef struct hsa_signal_s {
-  uint64_t handle;
-} hsa_signal_t;
-#endif
-
-}
-
-/* ---------------------------------------------------------------------------------
- * Simulated CPU Data Structures and API
- * ---------------------------------------------------------------------------------
- */
-
-#define ATMI_WAIT_STATE HSA_WAIT_STATE_BLOCKED
-
 // ---------------------- Kernel Start -------------
 typedef struct atl_kernel_info_s {
   uint64_t kernel_object;
   uint32_t group_segment_size;
   uint32_t private_segment_size;
   uint32_t sgpr_count;
   uint32_t vgpr_count;
   uint32_t sgpr_spill_count;
struct SignalPoolT {
   SignalPoolT(SignalPoolT &&) = delete;
   ~SignalPoolT() {
     size_t N = state.size();
     for (size_t i = 0; i < N; i++) {
       hsa_signal_t signal = state.front();
       state.pop();
       hsa_status_t rc = hsa_signal_destroy(signal);
       if (rc != HSA_STATUS_SUCCESS) {
-        DEBUG_PRINT("Signal pool destruction failed\n");
+        DP("Signal pool destruction failed\n");
       }
     }
   }
   size_t size() {
     lock l(&mutex);
     return state.size();
   }
   void push(hsa_signal_t s) {

 
 extern bool atl_is_impl_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);
+inline const char *get_error_string(hsa_status_t err) {
+  const char *res;
+  hsa_status_t rc = hsa_status_string(err, &res);
+  return (rc == HSA_STATUS_SUCCESS) ? res : "HSA_STATUS UNKNOWN.";
+}
 
 #endif // SRC_RUNTIME_INCLUDE_INTERNAL_H_

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

 //===--- amdgpu/impl/rt.h ----------------------------------------- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 #ifndef SRC_RUNTIME_INCLUDE_RT_H_
 #define SRC_RUNTIME_INCLUDE_RT_H_
 
-#include "impl_runtime.h"
 #include "hsa_api.h"
+#include "impl_runtime.h"
+#include "internal.h"
+
 #include <string>
 
 namespace core {
+namespace Runtime {
+hsa_status_t Memfree(void *);
+hsa_status_t HostMalloc(void **ptr, size_t size,
+                        hsa_amd_memory_pool_t MemoryPool);
 
-#define DEFAULT_MAX_QUEUE_SIZE 4096
-#define DEFAULT_DEBUG_MODE 0
-class Environment {
-public:
-  Environment()
-      : max_queue_size_(DEFAULT_MAX_QUEUE_SIZE),
-        debug_mode_(DEFAULT_DEBUG_MODE) {
-    GetEnvAll();
-  }
-
-  void GetEnvAll();
-
-  int getMaxQueueSize() const { return max_queue_size_; }
-  int getDebugMode() const { return debug_mode_; }
-
-private:
-  std::string GetEnv(const char *name) {
-    char *env = getenv(name);
-    std::string ret;
-    if (env) {
-      ret = env;
-    }
-    return ret;
-  }
-
-  int max_queue_size_;
-  int debug_mode_;
-};
-
-class Runtime final {
-public:
-  static Runtime &getInstance() {
-    static Runtime instance;
-    return instance;
-  }
-
-  // modules
-  static hsa_status_t RegisterModuleFromMemory(
-      void *, size_t, hsa_agent_t agent,
-      hsa_status_t (*on_deserialized_data)(void *data, size_t size,
-                                           void *cb_state),
-      void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
+} // namespace Runtime
+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, hsa_agent_t agent,
 
-  // data
-  static hsa_status_t Memcpy(hsa_signal_t, void *, const void *, size_t);
-  static hsa_status_t Memfree(void *);
-  static hsa_status_t HostMalloc(void **ptr, size_t size,
-                                 hsa_amd_memory_pool_t MemoryPool);
-
-  int getMaxQueueSize() const { return env_.getMaxQueueSize(); }
-  int getDebugMode() const { return env_.getDebugMode(); }
-
-protected:
-  Runtime() = default;
-  ~Runtime() = default;
-  Runtime(const Runtime &) = delete;
-  Runtime &operator=(const Runtime &) = delete;
-
-protected:
-  // variable to track environment variables
-  Environment env_;
-};
-
 } // namespace core
 
 #endif // SRC_RUNTIME_INCLUDE_RT_H_

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

if (kernel_segment_size > 0) {
         }
         // populate info with sizes and offsets
         info.arg_sizes.push_back(lcArg.size_);
         // v3 has offset field and not align field
         size_t new_offset = lcArg.offset_;
         size_t padding = new_offset - offset;
         offset = new_offset;
         info.arg_offsets.push_back(lcArg.offset_);
-        DEBUG_PRINT("Arg[%lu] \"%s\" (%u, %u)\n", i, lcArg.name_.c_str(),
-                    lcArg.size_, lcArg.offset_);
+        DP("Arg[%lu] \"%s\" (%u, %u)\n", i, lcArg.name_.c_str(), lcArg.size_,
+           lcArg.offset_);
         offset += lcArg.size_;
 
         // check if the arg is a hidden/implicit arg
         // this logic assumes that all hidden args are 8-byte aligned
         if (!isImplicit(lcArg.valueKind_)) {
           kernel_explicit_args_size += lcArg.size_;
         } else {
           hasHiddenArgs = true;
         }
         kernel_explicit_args_size += padding;
       }
     }
 
     // add size of implicit args, e.g.: offset x, y and z and pipe pointer, but
-    // in ATMI, do not count the compiler set implicit args, but set your own
-    // implicit args by discounting the compiler set implicit args
+    // do not count the compiler set implicit args, but set your own implicit
+    // args by discounting the compiler set implicit args
     info.kernel_segment_size =
         (hasHiddenArgs ? kernel_explicit_args_size : kernel_segment_size) +
         sizeof(impl_implicit_args_t);
-    DEBUG_PRINT("[%s: kernarg seg size] (%lu --> %u)\n", kernelName.c_str(),
-                kernel_segment_size, info.kernel_segment_size);
+    DP("[%s: kernarg seg size] (%lu --> %u)\n", kernelName.c_str(),
 
     // kernel received, now add it to the kernel info table
     KernelInfoTable[kernelName] = info;
   }
 
   return HSA_STATUS_SUCCESS;
 }
 
 static hsa_status_t
 populate_InfoTables(hsa_executable_symbol_t symbol,
                     std::map<std::string, atl_kernel_info_t> &KernelInfoTable,
                     std::map<std::string, atl_symbol_info_t> &SymbolInfoTable) {
   hsa_symbol_kind_t type;
 
   uint32_t name_length;
   hsa_status_t err;
   err = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE,
                                        &type);
   if (err != HSA_STATUS_SUCCESS) {
     printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
            "Symbol info extraction", get_error_string(err));
     return err;
   }
-  DEBUG_PRINT("Exec Symbol type: %d\n", type);
+  DP("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));
       return err;
     }
err = hsa_executable_symbol_get_info(
         &(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));
       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);
+    DP("Kernel %s --> %lx symbol %u group segsize %u pvt segsize %u bytes "
 
     // assign it back to the kernel info table
     KernelInfoTable[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) {
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;
     }
 
-    DEBUG_PRINT("Symbol %s = %p (%u bytes)\n", name, (void *)info.addr,
-                info.size);
+    DP("Symbol %s = %p (%u bytes)\n", name, (void *)info.addr, info.size);
     SymbolInfoTable[std::string(name)] = info;
     free(name);
   } else {
-    DEBUG_PRINT("Symbol is an indirect function\n");
+    DP("Symbol is an indirect function\n");
   }
   return HSA_STATUS_SUCCESS;
 }
 
 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, hsa_agent_t agent,
hsa_status_t RegisterModuleFromMemory(
   {
     {
       // Some metadata info is not available through ROCr API, so use custom
       // code object metadata parsing to collect such metadata info
 
       err = get_code_object_custom_metadata(module_bytes, module_size,
                                             KernelInfoTable);
       if (err != HSA_STATUS_SUCCESS) {
-        DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
-                    "Getting custom code object metadata",
-                    get_error_string(err));
+        DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
+           "Getting custom code object metadata", get_error_string(err));
         continue;
       }
 
       // Deserialize code object.
       hsa_code_object_t code_object = {0};
       err = hsa_code_object_deserialize(module_bytes, module_size, NULL,
                                         &code_object);
       if (err != HSA_STATUS_SUCCESS) {
-        DEBUG_PRINT("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
-                    "Code Object Deserialization", get_error_string(err));
+        DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
         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);
       hsa_status_t impl_err =
           on_deserialized_data(code_object_alloc_data, module_size, cb_state);
       if (impl_err != HSA_STATUS_SUCCESS) {
         printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
                "Error in deserialized_data callback",
                get_error_string(impl_err));
         return impl_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));
+        DP("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
         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);
+  DP("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 HSA_STATUS_ERROR;
     }

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

-//===--- amdgpu/impl/utils.cpp ------------------------------------ C++ -*-===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-#include "internal.h"
-#include "rt.h"
-
-#include <stdio.h>
-#include <string>
-
-const char *get_error_string(hsa_status_t err) {
-  const char *res;
-  hsa_status_t rc = hsa_status_string(err, &res);
-  return (rc == HSA_STATUS_SUCCESS) ? res : "HSA_STATUS UNKNOWN.";
-}
-
-namespace core {
-/*
- * Environment variables
- */
-void Environment::GetEnvAll() {
-  std::string var = GetEnv("ATMI_HELP");
-  if (!var.empty()) {
-    printf("ATMI_MAX_HSA_QUEUE_SIZE : positive integer\n"
-           "ATMI_DEBUG : 1 for printing out trace/debug info\n");
-  }
-
-  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

 #include <shared_mutex>
 #include <unordered_map>
 #include <vector>
 
 #include "interop_hsa.h"
 #include "impl_runtime.h"
 
 #include "internal.h"
+#include "rt.h"
 
-#include "Debug.h"
 #include "get_elf_mach_gfx_name.h"
 #include "omptargetplugin.h"
 #include "print_tracing.h"
 
 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
-#ifndef TARGET_NAME
-#error "Missing TARGET_NAME macro"
-#endif
-#define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
 
 // hostrpc interface, FIXME: consider moving to its own include these are
 // statically linked into amdgpu/plugin if present from hostrpc_services.a,
 // linked as --whole-archive to override the weak symbols that are used to
 // implement a fallback for toolchains that do not yet have a hostrpc library.
 extern "C" {
 unsigned long hostrpc_assign_buffer(hsa_agent_t agent, hsa_queue_t *this_Q,
                                     uint32_t device_id);
#define check(msg, status)                                                     \
   }
 #else
 #define check(msg, status)                                                     \
   {}
 #endif
 
 #include "elf_common.h"
 
-namespace core {
-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,
-    hsa_agent_t agent,
-    hsa_status_t (*on_deserialized_data)(void *data, size_t size,
-                                         void *cb_state),
-    void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
-}
-
 namespace hsa {
 template <typename C> hsa_status_t iterate_agents(C cb) {
   auto L = [](hsa_agent_t agent, void *data) -> hsa_status_t {
     C *unwrapped = static_cast<C *>(data);
     return (*unwrapped)(agent);
   };
   return hsa_iterate_agents(L, static_cast<void *>(&cb));
 }
RTLDeviceInfoTy() {
     hsa_status_t err = core::atl_init_gpu_context();
     if (err == HSA_STATUS_SUCCESS) {
       HSAInitializeSucceeded = true;
     } else {
       DP("Error when initializing " GETNAME(TARGET_NAME) "\n");
       return;
     }
 
-    // Init hostcall soon after initializing ATMI
+    // Init hostcall soon after initializing hsa
     hostrpc_init();
 
     err = FindAgents([&](hsa_device_type_t DeviceType, hsa_agent_t Agent) {
       if (DeviceType == HSA_DEVICE_TYPE_CPU) {
         CPUAgents.push_back(Agent);
       } else {
         HSAAgents.push_back(Agent);
       }
for (int i = 0; i < NumberOfDevices; i++) {
       uint32_t queue_size = 0;
       {
         hsa_status_t err = hsa_agent_get_info(
             HSAAgents[i], HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_size);
         if (err != HSA_STATUS_SUCCESS) {
           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();
+        enum { MaxQueueSize = 4096 };
+        if (queue_size > MaxQueueSize) {
+          queue_size = MaxQueueSize;
         }
       }
 
       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 queue %d\n", i);
~RTLDeviceInfoTy() {
     if (!HSAInitializeSucceeded) {
       // Then none of these can have been set up and they can't be torn down
       return;
     }
     // Run destructors on types that use HSA before
     // impl_finalize removes access to it
     deviceStateStore.clear();
     KernelArgPoolMap.clear();
-    // Terminate hostrpc before finalizing ATMI
+    // Terminate hostrpc before finalizing hsa
     hostrpc_terminate();
 
     hsa_status_t Err;
     for (uint32_t I = 0; I < HSAExecutables.size(); I++) {
       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));
if (!elf_machine_id_is_amdgcn(image)) {
     }
 
     err = env.after_loading();
     if (err != HSA_STATUS_SUCCESS) {
       return NULL;
     }
   }
 
-  DP("ATMI module successfully loaded!\n");
+  DP("AMDGPU 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;
if (!ArgPool) {
         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
-      // uninitialized
+      // Initialize implicit arguments. TODO: Which of these can be dropped
       impl_implicit_args_t *impl_args =
           reinterpret_cast<impl_implicit_args_t *>(
               static_cast<char *>(kernarg) + ArgPool->kernarg_segment_size);
       memset(impl_args, 0,
              sizeof(impl_implicit_args_t)); // may not be necessary
       impl_args->offset_x = 0;
       impl_args->offset_y = 0;
       impl_args->offset_z = 0;