[libomptarget][amdgpu] Drop env variables

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.

Reviewed By: pdhaliwal

Differential Revision: https://reviews.llvm.org/D108784

Author: JonChesterfield

openmp/libomptarget/plugins/amdgpu/CMakeLists.txt

@@ -67,7 +67,6 @@ add_library(omptarget.rtl.amdgpu SHARED
       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}

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

@@ -17,23 +17,21 @@
 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_amd_memory_pool_t MemoryPool) {
   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

@@ -5,8 +5,8 @@
 // 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>
@@ -32,7 +32,7 @@ static hsa_status_t invoke_hsa_copy(hsa_signal_t sig, void *dest,
   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) {
@@ -64,8 +64,7 @@ hsa_status_t impl_memcpy_h2d(hsa_signal_t signal, void *deviceDest,
   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);
@@ -94,8 +93,7 @@ hsa_status_t impl_memcpy_d2h(hsa_signal_t signal, void *dest,
   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);

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

@@ -24,7 +24,12 @@
 #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)
 
@@ -41,34 +46,6 @@ typedef struct impl_implicit_args_s {
   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;
@@ -110,7 +87,7 @@ struct SignalPoolT {
       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");
       }
     }
   }
@@ -183,6 +160,10 @@ 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

@@ -8,74 +8,26 @@
 #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 {
-
-#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);
-
-  // 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 Runtime {
+hsa_status_t Memfree(void *);
+hsa_status_t HostMalloc(void **ptr, size_t size,
+                        hsa_amd_memory_pool_t MemoryPool);
+
+} // 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,
+    hsa_status_t (*on_deserialized_data)(void *data, size_t size,
+                                         void *cb_state),
+    void *cb_state, std::vector<hsa_executable_t> &HSAExecutables);
 
 } // namespace core
 

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

@@ -525,8 +525,8 @@ static hsa_status_t get_code_object_custom_metadata(
         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
@@ -541,13 +541,13 @@ static hsa_status_t get_code_object_custom_metadata(
     }
 
     // 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_segment_size, info.kernel_segment_size);
 
     // kernel received, now add it to the kernel info table
     KernelInfoTable[kernelName] = info;
@@ -571,7 +571,7 @@ populate_InfoTables(hsa_executable_symbol_t symbol,
            "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);
@@ -636,11 +636,10 @@ populate_InfoTables(hsa_executable_symbol_t symbol,
       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 "
+       "kernarg\n",
+       kernelName.c_str(), info.kernel_object, info.group_segment_size,
+       info.private_segment_size, info.kernel_segment_size);
 
     // assign it back to the kernel info table
     KernelInfoTable[kernelName] = info;
@@ -681,12 +680,11 @@ populate_InfoTables(hsa_executable_symbol_t symbol,
       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;
 }
@@ -730,9 +728,8 @@ hsa_status_t RegisterModuleFromMemory(
       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;
       }
 
@@ -741,8 +738,8 @@ hsa_status_t RegisterModuleFromMemory(
       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__,
+           "Code Object Deserialization", get_error_string(err));
         continue;
       }
       assert(0 != code_object.handle);
@@ -763,16 +760,16 @@ hsa_status_t RegisterModuleFromMemory(
       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__,
+           "Loading the code object", get_error_string(err));
         continue;
       }
 
       // cannot iterate over symbols until executable is frozen
     }
     module_load_success = true;
   } while (0);
-  DEBUG_PRINT("Modules loaded successful? %d\n", module_load_success);
+  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, "");