[AMDGPU][Libomptarget] Drop dead code related to g_atl_machine

[pdhaliwal]

Summary

This patch includes some changes which deletes the code accessing
g_atl_machine global. Some accesses related to memory_pools are
still remaining.

[JonChesterfield]

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

27

int DeviceID is somewhat error prone, I wonder if it's worth changing to something like

struct DeviceIDType
{
 // operations
private:
  uint8_t data;
};

where there is a check at the entry point that it is positive and within the number of available GPUs, and then we pass that pre-validated type instance around

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

45

Are these header-only? They're more likely to cause problems if there is a LLVM library that needs to be linked in to use them

249

Not totally sure about this factoring. What do you think of populating one vector for GPU agents and a second one for CPU agents, instead of tagging each element with a dynamic type?

394

This seems to answer the above comment. Is the llvm::Optional<std::vector<DeviceTypeAgentPair>> type replaceable by passing a stateful lambda to FindAgents(), with a reference to each vector?

[pdhaliwal]

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

27

I don't understand. How would it be different from just using int deviceId and verifying it before passing to the methods?

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

45

I haven't checked that. I just compiled and ran libomptarget tests (all of which are fine). But, there is some code which is present in corresponding .cpps in lib/Support/.

394

Will work on it.

[pdhaliwal]

Use lambdas to accumulate agents.

[pdhaliwal]

Ping!

[JonChesterfield]

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

27

Would move some programming mistake detection from assert time to compile time

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

529

These appear to be unused

572

I'm not sure this needs AllAgents - shouldn't one of the subsets (iirc it would be HSAAgents) suffice?

[pdhaliwal]

Remove unused code.

[pdhaliwal]

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

572

Normally kernarg is found in CPUAgents, so I included both in case HSAAgents also have kernarg pool.

[JonChesterfield]

LG. My preference would be as commented above, but the vector containing the contents of the other two works too.

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

572

I think we should go with passing the list we currently see the kernarg segment in, partly to alert us if a future system changes that association, and partly because it's fractionally faster code that way. I think that also drops the AllAgents vector so there's less redundant state lying around.

[pdhaliwal]

Remove AllAgents vector.

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

/*
      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);
   */
 
   if (!symbol || !var_addr || !var_size)
     return HSA_STATUS_ERROR;
-  if (DeviceId < 0 ||
-      DeviceId >= g_atl_machine.processors<ATLGPUProcessor>().size())
-    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;
/*
      uint32_t value;
      atmi_interop_hsa_get_kernel_addr(gpu_place, "kernel_name",
                                   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE,
                                   &val);
   */
 
   if (!kernel_name || !value)
     return HSA_STATUS_ERROR;
-  if (DeviceId < 0 ||
-      DeviceId >= g_atl_machine.processors<ATLGPUProcessor>().size())
-    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;
     switch (kernel_info) {

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

 class ATLMachine {
 public:
   ATLMachine() {
     cpu_processors_.clear();
     gpu_processors_.clear();
   }
   template <typename T> void addProcessor(const T &p);
   template <typename T> std::vector<T> &processors();
-  template <typename T> size_t processorCount() {
-    return processors<T>().size();
-  }
 
 private:
   std::vector<ATLCPUProcessor> cpu_processors_;
   std::vector<ATLGPUProcessor> gpu_processors_;
 };
 
 hsa_amd_memory_pool_t get_memory_pool(const ATLProcessor &proc,
                                       const int mem_id);
 
-extern ATLMachine g_atl_machine;
-template <typename T> T &get_processor(int dev_id) {
-  if (dev_id == -1) {
-    // user is asking runtime to pick a device
-    // best device of this type? pick 0 for now
-    dev_id = 0;
-  }
-  return g_atl_machine.processors<T>()[dev_id];
-}
-
 #endif // SRC_RUNTIME_INCLUDE_MACHINE_H_

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

 public:
   static Runtime &getInstance() {
     static Runtime instance;
     return instance;
   }
 
   // modules
   static hsa_status_t RegisterModuleFromMemory(
-      void *, size_t, int DeviceId,
+      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);

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 (message_is_string(key, ".name")) {
       });
     }
   });
   return error;
 }
 } // namespace
 
 static hsa_status_t get_code_object_custom_metadata(
-    void *binary, size_t binSize, int gpu,
+    void *binary, size_t binSize,
     std::map<std::string, atl_kernel_info_t> &KernelInfoTable) {
   // parse code object with different keys from v2
   // also, the kernel name is not the same as the symbol name -- so a
   // symbol->name map is needed
 
   std::pair<unsigned char *, unsigned char *> metadata =
       find_metadata(binary, binSize);
   if (!metadata.first) {
for (size_t i = 0; i < kernelsSize; i++) {
     // 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, int gpu,
+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 (type == HSA_SYMBOL_KIND_KERNEL) {
     DEBUG_PRINT("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, int gpu,
+    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) {
   hsa_status_t err;
-  assert(gpu >= 0);
-
-  DEBUG_PRINT("Trying to load module to GPU-%d\n", gpu);
-  ATLGPUProcessor &proc = get_processor<ATLGPUProcessor>(gpu);
-  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 HSA_STATUS_ERROR;
hsa_status_t RegisterModuleFromMemory(
 
   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
 
-      err = get_code_object_custom_metadata(module_bytes, module_size, gpu,
+      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));
         continue;
       }
 
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;
     }
 
     err = hsa::executable_iterate_symbols(
         executable,
         [&](hsa_executable_t, hsa_executable_symbol_t symbol) -> hsa_status_t {
-          return populate_InfoTables(symbol, gpu, KernelInfoTable,
-                                     SymbolInfoTable);
+          return populate_InfoTables(symbol, 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 HSA_STATUS_ERROR;
     }
 
     // save the executable and destroy during finalize
     HSAExecutables.push_back(executable);
     return HSA_STATUS_SUCCESS;
   } else {
     return HSA_STATUS_ERROR;
   }
 }
 
 } // namespace core

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>
 
 // Header from ATMI interface
 #include "atmi_interop_hsa.h"
 #include "atmi_runtime.h"
 
 #include "internal.h"
 
 #include "Debug.h"
 #include "get_elf_mach_gfx_name.h"
 #include "machine.h"
 #include "omptargetplugin.h"
 #include "print_tracing.h"
 
+#include "llvm/ADT/None.h"
+#include "llvm/ADT/Optional.h"
 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
 
 #ifndef TARGET_NAME
 #define TARGET_NAME AMDHSA
 #endif
 #define DEBUG_PREFIX "Target " GETNAME(TARGET_NAME) " RTL"
 
 // hostrpc interface, FIXME: consider moving to its own include these are

 #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,
-    int DeviceId,
+    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);
 }
 
 /// Keep entries table per device
 struct FuncOrGblEntryTy {
   __tgt_target_table Table;
KernelTy(int8_t _ExecutionMode, int16_t _ConstWGSize, int32_t _device_id,
     }
   }
 };
 
 /// List that contains all the kernels.
 /// FIXME: we may need this to be per device and per library.
 std::list<KernelTy> KernelsList;
 
-static std::vector<hsa_agent_t> find_gpu_agents() {
-  std::vector<hsa_agent_t> res;
+typedef std::pair<hsa_device_type_t, hsa_agent_t> DeviceTypeAgentPair;
+
+static llvm::Optional<std::vector<DeviceTypeAgentPair>> FindAgents() {
+  std::vector<DeviceTypeAgentPair> 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);
+        auto *res = static_cast<std::vector<DeviceTypeAgentPair> *>(data);
 
         hsa_device_type_t device_type;
         // get_info fails iff HSA runtime not yet initialized
         hsa_status_t err =
             hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type);
         if (print_kernel_trace > 0 && err != HSA_STATUS_SUCCESS)
           printf("rtl.cpp: err %d\n", err);
         assert(err == HSA_STATUS_SUCCESS);
 
         if (device_type == HSA_DEVICE_TYPE_GPU) {
-          res->push_back(agent);
+          res->push_back({HSA_DEVICE_TYPE_GPU, agent});
+        } else if (device_type == HSA_DEVICE_TYPE_CPU) {
+          res->push_back({HSA_DEVICE_TYPE_CPU, agent});
         }
         return HSA_STATUS_SUCCESS;
       },
       &res);
 
   // iterate_agents fails iff HSA runtime not yet initialized
-  if (print_kernel_trace > 0 && err != HSA_STATUS_SUCCESS)
+  if (print_kernel_trace > 0 && err != HSA_STATUS_SUCCESS) {
     printf("rtl.cpp: err %d\n", err);
-  assert(err == HSA_STATUS_SUCCESS);
+    return llvm::None;
+  }
   return res;
 }
 
 static void callbackQueue(hsa_status_t status, hsa_queue_t *source,
                           void *data) {
   if (status != HSA_STATUS_SUCCESS) {
     const char *status_string;
     if (hsa_status_string(status, &status_string) != HSA_STATUS_SUCCESS) {
hsa_status_t addKernArgPool(hsa_amd_memory_pool_t MemoryPool, void *Data) {
   }
 
   return HSA_STATUS_SUCCESS;
 }
 
 std::pair<hsa_status_t, hsa_amd_memory_pool_t>
 FindKernargPool(const std::vector<hsa_agent_t> &HSAAgents) {
   std::vector<hsa_amd_memory_pool_t> KernArgPools;
-  for (const auto &processor : g_atl_machine.processors<ATLCPUProcessor>()) {
-    hsa_agent_t Agent = processor.agent();
+  for (const auto &Agent : HSAAgents) {
     hsa_status_t err = HSA_STATUS_SUCCESS;
     err = hsa_amd_agent_iterate_memory_pools(
         Agent, addKernArgPool, static_cast<void *>(&KernArgPools));
     if (err != HSA_STATUS_SUCCESS) {
       printf("[%s:%d] %s failed: %s\n", __FILE__, __LINE__,
              "Iterate all memory pools", get_error_string(err));
       return {err, hsa_amd_memory_pool_t{}};
     }
public:
   std::shared_timed_mutex load_run_lock;
 
   int NumberOfDevices;
 
   // GPU devices
   std::vector<hsa_agent_t> HSAAgents;
   std::vector<hsa_queue_t *> HSAQueues; // one per gpu
 
+  // CPUs
+  std::vector<hsa_agent_t> CPUAgents;
+
   // Device properties
   std::vector<int> ComputeUnits;
   std::vector<int> GroupsPerDevice;
   std::vector<int> ThreadsPerGroup;
   std::vector<int> WarpSize;
   std::vector<std::string> GPUName;
 
   // OpenMP properties
assert(device_id < (int32_t)FuncGblEntries.size() &&
            "Unexpected device id!");
     FuncGblEntries[device_id].emplace_back();
     FuncOrGblEntryTy &E = FuncGblEntries[device_id].back();
     // KernelArgPoolMap.clear();
     E.Entries.clear();
     E.Table.EntriesBegin = E.Table.EntriesEnd = 0;
   }
 
   RTLDeviceInfoTy() {
     // LIBOMPTARGET_KERNEL_TRACE provides a kernel launch trace to stderr
     // anytime. You do not need a debug library build.
     //  0 => no tracing
     //  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");
     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();
+    llvm::Optional<std::vector<DeviceTypeAgentPair>> TypeAgentPairs =
+        FindAgents();
+    if (!TypeAgentPairs)
+      return;
+
+    std::vector<hsa_agent_t> AllAgents;
+    for (const auto &TypeAgentPair : *TypeAgentPairs) {
+      if (TypeAgentPair.first == HSA_DEVICE_TYPE_GPU) {
+        HSAAgents.push_back(TypeAgentPair.second);
+      } else {
+        CPUAgents.push_back(TypeAgentPair.second);
+      }
+      AllAgents.push_back(TypeAgentPair.second);
+    }
     NumberOfDevices = (int)HSAAgents.size();
 
     if (NumberOfDevices == 0) {
       DP("There are no devices supporting HSA.\n");
       return;
     } else {
       DP("There are %d devices supporting HSA.\n", NumberOfDevices);
     }
-
-    std::tie(err, KernArgPool) = core::FindKernargPool(HSAAgents);
+    std::tie(err, KernArgPool) = core::FindKernargPool(AllAgents);
     if (err != HSA_STATUS_SUCCESS) {
       DP("Error when reading memory pools\n");
       return;
     }
 
     // Init the device info
     HSAQueues.resize(NumberOfDevices);
     FuncGblEntries.resize(NumberOfDevices);
hsa_status_t module_register_from_memory_to_place(
     std::map<std::string, atl_symbol_info_t> &SymbolInfoTable,
     void *module_bytes, size_t module_size, int DeviceId, C cb,
     std::vector<hsa_executable_t> &HSAExecutables) {
   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, DeviceId, L,
-      static_cast<void *>(&cb), HSAExecutables);
+      KernelInfoTable, SymbolInfoTable, module_bytes, module_size,
+      DeviceInfo.HSAAgents[DeviceId], L, static_cast<void *>(&cb),
+      HSAExecutables);
 }
 } // namespace
 
 static uint64_t get_device_State_bytes(char *ImageStart, size_t img_size) {
   uint64_t device_State_bytes = 0;
   {
     // If this is the deviceRTL, get the state variable size
     symbol_info size_si;