[libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin

[JonChesterfield]

Summary

[libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin

Drops an enum that was identical to a HSA one, localises some functions where
they were only called from one TU. Covers everything internalize + adce can
identify as dead, except for msgpack::dump which is useful when debugging.

[JonChesterfield]

• drop semicolon

[jdoerfert]

LG

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

  * @brief Memory Access Type.
  */
 typedef enum atmi_memtype_s {
   ATMI_MEMTYPE_FINE_GRAINED = 0,
   ATMI_MEMTYPE_COARSE_GRAINED = 1,
   ATMI_MEMTYPE_ANY
 } atmi_memtype_t;
 
-/**
- * @brief ATMI Memory Fences for Tasks.
- */
-typedef enum atmi_task_fence_scope_s {
-  /**
-   * No memory fence applied; external fences have to be applied around the task
-   * launch/completion.
-   */
-  ATMI_FENCE_SCOPE_NONE = 0,
-  /**
-   * The fence is applied to the device.
-   */
-  ATMI_FENCE_SCOPE_DEVICE = 1,
-  /**
-   * The fence is applied to the entire system.
-   */
-  ATMI_FENCE_SCOPE_SYSTEM = 2
-} atmi_task_fence_scope_t;
-
 /** @} */
 
 /** \defgroup common Common ATMI Structures
  *  @{
  */
 
 /**
  * @brief ATMI Compute Place
typedef struct atmi_machine_s {
   /**
    * The device structures categorized by the device type
    */
   atmi_device_t *devices_by_type[ATMI_DEVTYPE_ALL];
 } atmi_machine_t;
 
 // Below are some helper macros that can be used to setup
 // some of the ATMI data structures.
-#define ATMI_PLACE_CPU(node, cpu_id)                                           \
-  { .node_id = node, .type = ATMI_DEVTYPE_CPU, .device_id = cpu_id }
 #define ATMI_PLACE_GPU(node, gpu_id)                                           \
   { .node_id = node, .type = ATMI_DEVTYPE_GPU, .device_id = gpu_id }
-#define ATMI_MEM_PLACE_CPU(node, cpu_id)                                       \
-  {                                                                            \
-    .node_id = node, .dev_type = ATMI_DEVTYPE_CPU, .dev_id = cpu_id,           \
-    .mem_id = -1                                                               \
-  }
-#define ATMI_MEM_PLACE_GPU(node, gpu_id)                                       \
-  {                                                                            \
-    .node_id = node, .dev_type = ATMI_DEVTYPE_GPU, .dev_id = gpu_id,           \
-    .mem_id = -1                                                               \
-  }
+
 #define ATMI_MEM_PLACE_CPU_MEM(node, cpu_id, cpu_mem_id)                       \
   {                                                                            \
     .node_id = node, .dev_type = ATMI_DEVTYPE_CPU, .dev_id = cpu_id,           \
     .mem_id = cpu_mem_id                                                       \
   }
 #define ATMI_MEM_PLACE_GPU_MEM(node, gpu_id, gpu_mem_id)                       \
   {                                                                            \
     .node_id = node, .dev_type = ATMI_DEVTYPE_GPU, .dev_id = gpu_id,           \
     .mem_id = gpu_mem_id                                                       \
   }
 #define ATMI_MEM_PLACE(d_type, d_id, m_id)                                     \
   { .node_id = 0, .dev_type = d_type, .dev_id = d_id, .mem_id = m_id }
 
 #endif // INCLUDE_ATMI_H_

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

 #include <vector>
 
 using core::TaskImpl;
 extern ATLMachine g_atl_machine;
 
 namespace core {
 void allow_access_to_all_gpu_agents(void *ptr);
 
-const char *getPlaceStr(atmi_devtype_t type) {
-  switch (type) {
-  case ATMI_DEVTYPE_CPU:
-    return "CPU";
-  case ATMI_DEVTYPE_GPU:
-    return "GPU";
-  default:
-    return NULL;
-  }
-}
-
+namespace {
 ATLProcessor &get_processor_by_mem_place(atmi_mem_place_t place) {
   int dev_id = place.dev_id;
   switch (place.dev_type) {
   case ATMI_DEVTYPE_CPU:
     return g_atl_machine.processors<ATLCPUProcessor>()[dev_id];
   case ATMI_DEVTYPE_GPU:
     return g_atl_machine.processors<ATLGPUProcessor>()[dev_id];
   }
 }
 
 hsa_amd_memory_pool_t get_memory_pool_by_mem_place(atmi_mem_place_t place) {
   ATLProcessor &proc = get_processor_by_mem_place(place);
   return get_memory_pool(proc, place.mem_id);
 }
+} // namespace
 
 void register_allocation(void *ptr, size_t size, atmi_mem_place_t place) {
   if (place.dev_type == ATMI_DEVTYPE_CPU)
     allow_access_to_all_gpu_agents(ptr);
 }
 
 atmi_status_t Runtime::Malloc(void **ptr, size_t size, atmi_mem_place_t place) {
   atmi_status_t ret = ATMI_STATUS_SUCCESS;

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));
 }
 
 extern void register_allocation(void *addr, size_t size,
                                 atmi_mem_place_t place);
-extern hsa_amd_memory_pool_t
-get_memory_pool_by_mem_place(atmi_mem_place_t place);
+
 extern bool atl_is_atmi_initialized();
 
 bool handle_group_signal(hsa_signal_value_t value, void *arg);
 
-void packet_store_release(uint32_t *packet, uint16_t header, uint16_t rest);
-uint16_t
-create_header(hsa_packet_type_t type, int barrier,
-              atmi_task_fence_scope_t acq_fence = ATMI_FENCE_SCOPE_SYSTEM,
-              atmi_task_fence_scope_t rel_fence = ATMI_FENCE_SCOPE_SYSTEM);
 
 void 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);
 
 #define ATMIErrorCheck(msg, status)                                            \

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>

 namespace core {
 /* Machine Info */
 atmi_machine_t *Runtime::GetMachineInfo() {
   if (!atlc.g_hsa_initialized)
     return NULL;
   return &g_atmi_machine;
 }
 
-void atl_set_atmi_initialized() {
+static void atl_set_atmi_initialized() {
   // FIXME: thread safe? locks?
   g_atmi_initialized = true;
 }
 
-void atl_reset_atmi_initialized() {
+static void atl_reset_atmi_initialized() {
   // FIXME: thread safe? locks?
   g_atmi_initialized = false;
 }
 
 bool atl_is_atmi_initialized() { return g_atmi_initialized; }
 
 void allow_access_to_all_gpu_agents(void *ptr) {
   hsa_status_t err;
atmi_status_t Runtime::Finalize() {
 
   atl_reset_atmi_initialized();
   err = hsa_shut_down();
   ErrorCheck(Shutting down HSA, err);
 
   return ATMI_STATUS_SUCCESS;
 }
 
-void atmi_init_context_structs() {
+static void atmi_init_context_structs() {
   atlc_p = &atlc;
   atlc.struct_initialized = true; /* This only gets called one time */
   atlc.g_hsa_initialized = false;
   atlc.g_gpu_initialized = false;
   atlc.g_tasks_initialized = false;
 }
 
 // Implement memory_pool iteration function
atmi_status_t atl_init_gpu_context() {
   err = hsa_amd_register_system_event_handler(callbackEvent, NULL);
     ErrorCheck(Registering the system for memory faults, err);
 
     init_tasks();
     atlc.g_gpu_initialized = true;
     return ATMI_STATUS_SUCCESS;
 }
 
-bool isImplicit(KernelArgMD::ValueKind value_kind) {
+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:
   case KernelArgMD::ValueKind::HiddenPrintfBuffer:
   case KernelArgMD::ValueKind::HiddenDefaultQueue:
   case KernelArgMD::ValueKind::HiddenCompletionAction:

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>
 
DP("Warning: Attempting to assign hostrpc to device %u, but hostrpc library "
      "missing\n",
      device_id);
   return 0;
 }
 }
 
 int print_kernel_trace;
 
-// Size of the target call stack struture
-uint32_t TgtStackItemSize = 0;
-
 #undef check // Drop definition from internal.h
 #ifdef OMPTARGET_DEBUG
 #define check(msg, status)                                                     \
   if (status != ATMI_STATUS_SUCCESS) {                                         \
     /* fprintf(stderr, "[%s:%d] %s failed.\n", __FILE__, __LINE__, #msg);*/    \
     DP(#msg " failed\n");                                                      \
     /*assert(0);*/                                                             \
   } else {                                                                     \
if (status != HSA_STATUS_SUCCESS) {
     }
     fprintf(stderr, "[%s:%d] GPU error in queue %p %d (%s)\n", __FILE__,
             __LINE__, source, status, status_string);
     abort();
   }
 }
 
 namespace core {
+namespace {
 void packet_store_release(uint32_t *packet, uint16_t header, uint16_t rest) {
   __atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE);
 }
 
-uint16_t create_header(hsa_packet_type_t type, int barrier,
-                       atmi_task_fence_scope_t acq_fence,
-                       atmi_task_fence_scope_t rel_fence) {
-  uint16_t header = type << HSA_PACKET_HEADER_TYPE;
-  header |= barrier << HSA_PACKET_HEADER_BARRIER;
-  header |= (hsa_fence_scope_t) static_cast<int>(
-      acq_fence << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE);
-  header |= (hsa_fence_scope_t) static_cast<int>(
-      rel_fence << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
+uint16_t create_header() {
+  uint16_t header = HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE;
+  header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
+  header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
   return header;
 }
+} // namespace
 } // namespace core
 
 /// Class containing all the device information
 class RTLDeviceInfoTy {
   std::vector<std::list<FuncOrGblEntryTy>> FuncGblEntries;
 
 public:
   // load binary populates symbol tables and mutates various global state
if (!ArgPool) {
       if (s.handle == 0) {
         printf("Failed to get signal instance\n");
         exit(1);
       }
       packet->completion_signal = s;
       hsa_signal_store_relaxed(packet->completion_signal, 1);
     }
 
-    core::packet_store_release(
-        reinterpret_cast<uint32_t *>(packet),
-        core::create_header(HSA_PACKET_TYPE_KERNEL_DISPATCH, 0,
-                            ATMI_FENCE_SCOPE_SYSTEM, ATMI_FENCE_SCOPE_SYSTEM),
-        packet->setup);
+    core::packet_store_release(reinterpret_cast<uint32_t *>(packet),
+                               core::create_header(), packet->setup);
 
     hsa_signal_store_relaxed(queue->doorbell_signal, packet_id);
 
     while (hsa_signal_wait_scacquire(packet->completion_signal,
                                      HSA_SIGNAL_CONDITION_EQ, 0, UINT64_MAX,
                                      HSA_WAIT_STATE_BLOCKED) != 0)
       ;