[libomptarget][nfc] Move hostcall required test to rtl

[JonChesterfield]

Summary

[libomptarget][nfc] Move hostcall required test to rtl

Remove a global, fix minor race. First of N patches to bring up hostcall.

[JonChesterfield]

Hostcall exists as a libraries in rocm, aomp and my github. Trunk has hooks - mostly weak symbols - to allow the host plugin source to work with or without it present.

The variable that is checked for (before and after this change) is introduced by a common variable in the GPU entry point for hostcalls. The idea is to only spin up the thread(s) associated with hostcall if the GPU will use it.

__asm__("; hostcall_invoke: record need for hostcall support\n\t"
        ".type needs_hostcall_buffer,@object\n\t"
        ".global needs_hostcall_buffer\n\t"
        ".comm needs_hostcall_buffer,4":::);

This removes one of the remaining uses of the symbol info table.

[JonChesterfield]

Side point on races - I think it is possible for one thread to start loading a binary, get immediately suspended, and a second thread to then launch a kernel that is supposed to come from that binary. The locking scheme in this plugin probably doesn't handle that. Hopefully libomptarget does have locking that makes that work, otherwise the cuda plugin is probably also racy.

[JonChesterfield]

Was reviewed on aomp, will wait a little while before landing this to match.

Edit: though this is NFC on trunk, after landing here it'll reach rocm, where it will be tested by the usual pipeline

[ldionne]

Your commit message looks like

commit df005fa364ae3914231f98f80acc120fcf782c26
Author: Jon Chesterfield <[e-mail]>
Date:   Tue May 25 22:43:16 2021 +0100

    [libomptarget][nfc] Move hostcall required test to rtl

    [libomptarget][nfc] Move hostcall required test to rtl

    Remove a global, fix minor race. First of N patches to bring up hostcall.

    Reviewed By: JonChesterfield

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

Notice the duplicate subject line? This appears to be a recurring issue with your commits, so it might be something in your workflow?

[JonChesterfield]

Thanks for pointing it out! I was under the impression that the contents of 'Summary' as shown in Phabricator turns into the commit message.

Looks like the commit message I'm creating is the title from phabricator, then a newline, then the summary, so the workflow change is probably to write less in the Summary field.

I'm committing a patch by:

git checkout main && git reset --hard origin/main && git pull # clean slate
arc patch D1000
git checkout main
git cherry-pick D1000
git log -n2 # eyeball it
git push

combined with creating patches with arc diff, I think the failure mode must be in my use of arcanist.

Edited the WIP diffs to, I hope, avoid continuing in this pattern.

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

 ATLMachine g_atl_machine;
 
 std::vector<hsa_amd_memory_pool_t> atl_gpu_kernarg_pools;
 
 std::vector<std::map<std::string, atl_kernel_info_t>> KernelInfoTable;
 std::vector<std::map<std::string, atl_symbol_info_t>> SymbolInfoTable;
 
 bool g_atmi_initialized = false;
-bool g_atmi_hostcall_required = false;
 
 /*
    atlc is all internal global values.
    The structure atl_context_t is defined in atl_internal.h
    Most references will use the global structure prefix atlc.
 */
 atl_context_t atlc = {.struct_initialized = false};
 
if (type == HSA_SYMBOL_KIND_KERNEL) {
     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);
     if (err != HSA_STATUS_SUCCESS) {
       return err;
     }
     SymbolInfoTable[gpu][std::string(name)] = info;
-    if (strcmp(name, "needs_hostcall_buffer") == 0)
-      g_atmi_hostcall_required = true;
     free(name);
   } else {
     DEBUG_PRINT("Symbol is an indirect function\n");
   }
   return HSA_STATUS_SUCCESS;
 }
 
 atmi_status_t Runtime::RegisterModuleFromMemory(

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

public:
   int EnvMaxTeamsDefault;
 
   // OpenMP Requires Flags
   int64_t RequiresFlags;
 
   // Resource pools
   SignalPoolT FreeSignalPool;
 
+  bool hostcall_required = false;
+
   std::vector<hsa_executable_t> HSAExecutables;
 
   struct atmiFreePtrDeletor {
     void operator()(void *p) {
       atmi_free(p); // ignore failure to free
     }
   };
 
if (rc != HSA_STATUS_SUCCESS) {
     atmi_free(ptr);
     return ATMI_STATUS_ERROR;
   }
 
   *ret_ptr = ptr;
   return ATMI_STATUS_SUCCESS;
 }
 
+static bool image_contains_symbol(void *data, size_t size, const char *sym) {
+  symbol_info si;
+  int rc = get_symbol_info_without_loading((char *)data, size, sym, &si);
+  return (rc == 0) && (si.addr != nullptr);
+}
+
 __tgt_target_table *__tgt_rtl_load_binary_locked(int32_t device_id,
                                                  __tgt_device_image *image) {
   // This function loads the device image onto gpu[device_id] and does other
   // per-image initialization work. Specifically:
   //
   // - Initialize an omptarget_device_environmentTy instance embedded in the
   //   image at the symbol "omptarget_device_environment"
   //   Fields debug_level, device_num, num_devices. Used by the deviceRTL.
__tgt_target_table *__tgt_rtl_load_binary_locked(int32_t device_id,
   }
 
   {
     auto env = device_environment(device_id, DeviceInfo.NumberOfDevices, image,
                                   img_size);
 
     atmi_status_t err = module_register_from_memory_to_place(
         (void *)image->ImageStart, img_size, get_gpu_place(device_id),
-        [&](void *data, size_t size) { return env.before_loading(data, size); },
+        [&](void *data, size_t size) {
+          if (image_contains_symbol(data, size, "needs_hostcall_buffer")) {
+            __atomic_store_n(&DeviceInfo.hostcall_required, true,
+                             __ATOMIC_RELEASE);
+          }
+          return env.before_loading(data, size);
+        },
         DeviceInfo.HSAExecutables);
 
     check("Module registering", err);
     if (err != ATMI_STATUS_SUCCESS) {
       fprintf(stderr,
               "Possible gpu arch mismatch: device:%s, image:%s please check"
               " compiler flag: -march=<gpu>\n",
               DeviceInfo.GPUName[device_id].c_str(),
static uint64_t acquire_available_packet_id(hsa_queue_t *queue) {
   bool full = true;
   while (full) {
     full =
         packet_id >= (queue->size + hsa_queue_load_read_index_scacquire(queue));
   }
   return packet_id;
 }
 
-extern bool g_atmi_hostcall_required; // declared without header by atmi
-
 static int32_t __tgt_rtl_run_target_team_region_locked(
     int32_t device_id, void *tgt_entry_ptr, void **tgt_args,
     ptrdiff_t *tgt_offsets, int32_t arg_num, int32_t num_teams,
     int32_t thread_limit, uint64_t loop_tripcount);
 
 int32_t __tgt_rtl_run_target_team_region(int32_t device_id, void *tgt_entry_ptr,
                                          void **tgt_args,
                                          ptrdiff_t *tgt_offsets,
if (!ArgPool) {
               static_cast<char *>(kernarg) + ArgPool->kernarg_segment_size);
       memset(impl_args, 0,
              sizeof(atmi_implicit_args_t)); // may not be necessary
       impl_args->offset_x = 0;
       impl_args->offset_y = 0;
       impl_args->offset_z = 0;
 
       // assign a hostcall buffer for the selected Q
-      if (g_atmi_hostcall_required) {
+      if (__atomic_load_n(&DeviceInfo.hostcall_required, __ATOMIC_ACQUIRE)) {
         // hostrpc_assign_buffer is not thread safe, and this function is
         // under a multiple reader lock, not a writer lock.
         static pthread_mutex_t hostcall_init_lock = PTHREAD_MUTEX_INITIALIZER;
         pthread_mutex_lock(&hostcall_init_lock);
         impl_args->hostcall_ptr = hostrpc_assign_buffer(
             DeviceInfo.HSAAgents[device_id], queue, device_id);
         pthread_mutex_unlock(&hostcall_init_lock);
         if (!impl_args->hostcall_ptr) {