Merge in the rocclr based hip runtime (#2032)

* Merge master-next changes in master (include vdi development in master branch)

hipamd/src/h2f.cpp

@@ -57,12 +57,14 @@ static inline std::uint32_t __convert_float_to_half(float a) noexcept {
 
 // On machines without fp16 instructions, clang lowers llvm.convert.from.fp16
 // to call of this function.
-extern "C" float __gnu_h2f_ieee(unsigned short h){
+extern "C" __attribute__((visibility("default")))
+float __gnu_h2f_ieee(unsigned short h){
   return __convert_half_to_float((std::uint32_t) h);
 }
 
 // On machines without fp16 instructions, clang lowers llvm.convert.to.fp16
 // to call of this function.
-extern "C" unsigned short __gnu_f2h_ieee(float f){
+extern "C" __attribute__((visibility("default")))
+unsigned short __gnu_f2h_ieee(float f){
   return (unsigned short)__convert_float_to_half(f);
 }

hipamd/src/hip_clang.cpp

@@ -28,6 +28,7 @@ THE SOFTWARE.
 #include "hip_hcc_internal.h"
 #include "hip_fatbin.h"
 #include "trace_helper.h"
+#include "program_state.inl"
 
 #ifdef __GNUC__
 #pragma GCC visibility push (default)
@@ -94,8 +95,10 @@ __hipRegisterFatBinary(const void* data)
         module->executable, agent);
 
       if (module->executable.handle) {
-        modules->at(deviceId) = module;
-        tprintf(DB_FB, "Loaded code object for %s\n", name);
+         hip_impl::program_state_impl::read_kernarg_metadata(image, module->kernargs);
+         modules->at(deviceId) = module;
+
+         tprintf(DB_FB, "Loaded code object for %s, args size=%ld\n", name, module->kernargs.size());
       } else {
         fprintf(stderr, "Failed to load code object for %s\n", name);
         abort();
@@ -157,16 +160,215 @@ extern "C" void __hipRegisterFunction(
   g_functions.insert(std::make_pair(hostFunction, std::move(functions)));
 }
 
+static inline const char* hsa_strerror(hsa_status_t status) {
+  const char* str = nullptr;
+  if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) {
+    return str;
+  }
+  return "Unknown error";
+}
+
+struct RegisteredVar {
+public:
+  RegisteredVar(): size_(0), devicePtr_(nullptr) {}
+  ~RegisteredVar() {}
+
+  static inline const char* hsa_strerror(hsa_status_t status) {
+  const char* str = nullptr;
+  if (hsa_status_string(status, &str) == HSA_STATUS_SUCCESS) {
+    return str;
+  }
+  return "Unknown error";
+}
+
+hipDeviceptr_t getdeviceptr() const { return devicePtr_; };
+  size_t getvarsize() const { return size_; };
+
+  size_t size_;               // Size of the variable
+  hipDeviceptr_t devicePtr_;  //Device Memory Address of the variable.
+};
+
+struct DeviceVar {
+  void* shadowVptr;
+  std::string hostVar;
+  size_t size;
+  std::vector<hipModule_t>* modules;
+  std::vector<RegisteredVar> rvars;
+  bool dyn_undef;
+};
+
+std::unordered_multimap<std::string, DeviceVar > g_vars;
+
+//The logic follows PlatformState::getGlobalVar in VDI RT
+static DeviceVar* findVar(std::string hostVar, int deviceId, hipModule_t hmod) {
+  DeviceVar* dvar = nullptr;
+  if (hmod != nullptr) {
+    // If module is provided, then get the var only from that module
+    auto var_range = g_vars.equal_range(hostVar);
+    for (auto it = var_range.first; it != var_range.second; ++it) {
+      if ((*it->second.modules)[deviceId] == hmod) {
+        dvar = &(it->second);
+        break;
+      }
+    }
+  } else {
+    // If var count is < 2, return the var
+    if (g_vars.count(hostVar) < 2) {
+      auto it = g_vars.find(hostVar);
+      dvar = ((it == g_vars.end()) ? nullptr : &(it->second));
+    } else {
+      // If var count is > 2, return the original var,
+      // if original var count != 1, return g_vars.end()/Invalid
+      size_t orig_global_count = 0;
+      auto var_range = g_vars.equal_range(hostVar);
+      for (auto it = var_range.first; it != var_range.second; ++it) {
+        // when dyn_undef is set, it is a shadow var
+        if (it->second.dyn_undef == false) {
+          ++orig_global_count;
+          dvar = &(it->second);
+        }
+      }
+      dvar = ((orig_global_count == 1) ? dvar : nullptr);
+    }
+  }
+  return dvar;
+}
+
+hipError_t ihipGetGlobalVar(hipDeviceptr_t* dev_ptr, size_t* size_ptr,
+                             const char* hostVar, hipModule_t hmod) {
+  GET_TLS();
+  auto ctx = ihipGetTlsDefaultCtx();
+
+  if (!ctx) return hipErrorInvalidValue;
+
+  auto device = ctx->getDevice();
+
+  if (!device) return hipErrorInvalidValue;
+
+  ihipDevice_t* currentDevice = ihipGetDevice(device->_deviceId);
+
+  if (!currentDevice) return hipErrorInvalidValue;
+
+  int deviceId = device->_deviceId;
+
+  DeviceVar* dvar = findVar(std::string(hostVar), deviceId, hmod);
+  if (dvar == nullptr) return hipErrorInvalidValue;
+
+  if (dvar->rvars[deviceId].getdeviceptr() == nullptr) return hipErrorInvalidValue;
+
+  *size_ptr = dvar->rvars[deviceId].getvarsize();
+  *dev_ptr = dvar->rvars[deviceId].getdeviceptr();
+  return hipSuccess;
+}
+
+static bool createGlobalVarObj(const hsa_executable_t& hsaExecutable, const hsa_agent_t& hasAgent,
+                               const char* global_name, void** device_pptr, size_t* bytes) {
+  hsa_status_t status = HSA_STATUS_SUCCESS;
+  hsa_symbol_kind_t sym_type;
+  hsa_executable_symbol_t global_symbol;
+  std::string buildLog;
+
+  /* Find HSA Symbol by name */
+  status = hsa_executable_get_symbol_by_name(hsaExecutable, global_name, &hasAgent,
+                                             &global_symbol);
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog += "Error: Failed to find the Symbol by Name: ";
+    buildLog += hsa_strerror(status);
+    tprintf(DB_FB, "createGlobalVarObj: %s\n", buildLog.c_str());
+    return false;
+  }
+
+  /* Find HSA Symbol Type */
+  status = hsa_executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE,
+                                          &sym_type);
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog += "Error: Failed to find the Symbol Type : ";
+    buildLog += hsa_strerror(status);
+    tprintf(DB_FB, "createGlobalVarObj: %s\n", buildLog.c_str());
+    return false;
+  }
+
+  /* Make sure symbol type is VARIABLE */
+  if (sym_type != HSA_SYMBOL_KIND_VARIABLE) {
+    buildLog += "Error: Symbol is not of type VARIABLE : ";
+    buildLog += hsa_strerror(status);
+    tprintf(DB_FB, "createGlobalVarObj: %s\n", buildLog.c_str());
+    return false;
+  }
+
+  /* Retrieve the size of the variable */
+  status = hsa_executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE, bytes);
+
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog += "Error: Failed to retrieve the Symbol Size : ";
+    buildLog += hsa_strerror(status);
+    tprintf(DB_FB, "createGlobalVarObj: %s\n", buildLog.c_str());
+    return false;
+  }
+
+  /* Find HSA Symbol Address */
+  status = hsa_executable_symbol_get_info(global_symbol,
+                                          HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, device_pptr);
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog += "Error: Failed to find the Symbol Address : ";
+    buildLog += hsa_strerror(status);
+    tprintf(DB_FB, "createGlobalVarObj: %s\n", buildLog.c_str());
+    return false;
+  } else {
+    tprintf(DB_FB, "createGlobalVarObj: var %s : device=%p, size=%zu\n", global_name, *device_pptr, *bytes);
+  }
+
+  return true;
+}
+
+// Registers a device-side global variable.
+// For each global variable in device code, there is a corresponding shadow
+// global variable in host code. The shadow host variable is used to keep
+// track of the value of the device side global variable between kernel
+// executions.
+// The basic logic is taken from VDI RT, but there is much difference.
 extern "C" void __hipRegisterVar(
-  std::vector<hipModule_t>* modules,
-  char*       hostVar,
-  char*       deviceVar,
-  const char* deviceName,
-  int         ext,
-  int         size,
-  int         constant,
-  int         global)
+  std::vector<hipModule_t>* modules,   // The device modules containing code object
+  char*       var,       // The shadow variable in host code
+  char*       hostVar,   // Variable name in host code
+  const char* deviceVar, // Variable name in device code
+  int         ext,       // Whether this variable is external
+  int         size,      // Size of the variable
+  int         constant,  // Whether this variable is constant
+  int         global)    // Unknown, always 0
 {
+    HIP_INIT_API(__hipRegisterVar, modules, var, hostVar, deviceVar, ext, size, constant, global);
+
+    DeviceVar dvar{var, std::string{ hostVar }, static_cast<size_t>(size), modules,
+           std::vector<RegisteredVar>{ g_deviceCnt }, false };
+
+    for (int deviceId = 0; deviceId < g_deviceCnt; deviceId++) {
+        auto device = ihipGetDevice(deviceId);
+        if(!device) {
+           continue;
+        }
+        hsa_executable_t& executable = (*modules)[deviceId]->executable;
+        hsa_agent_t& agent = g_allAgents[deviceId + 1];
+        size_t bytes = 0;
+        hipDeviceptr_t devicePtr = nullptr;
+
+        bool success = createGlobalVarObj(executable, agent, hostVar, &devicePtr, &bytes);
+        if(!success) {
+           return;
+        }
+        dvar.rvars[deviceId].devicePtr_ = devicePtr;
+        dvar.rvars[deviceId].size_ = bytes;
+
+        hc::AmPointerInfo ptrInfo(nullptr, devicePtr, devicePtr, bytes, device->_acc, true, false);
+        hc::am_memtracker_add(devicePtr, ptrInfo);
+
+    #if USE_APP_PTR_FOR_CTX
+        hc::am_memtracker_update(devicePtr, device->_deviceId, 0u, ihipGetTlsDefaultCtx());
+    #else
+        hc::am_memtracker_update(devicePtr, device->_deviceId, 0u);
+    #endif
+    }
+    g_vars.insert(std::make_pair(std::string(hostVar), dvar));
 }
 
 extern "C" void __hipUnregisterFatBinary(std::vector<hipModule_t>* modules)
@@ -226,6 +428,41 @@ extern "C" hipError_t __hipPopCallConfiguration(
   return hipSuccess;
 }
 
+int getCurrentDeviceId()
+{
+  GET_TLS();
+
+  int deviceId = 0;
+  auto ctx = ihipGetTlsDefaultCtx();
+
+  if(!ctx) return deviceId;
+
+  LockedAccessor_CtxCrit_t crit(ctx->criticalData());
+
+  if(crit->_execStack.size() != 0)
+  {
+    auto &exec = crit->_execStack.top();
+
+    if (exec._hStream) {
+      deviceId = exec._hStream->getDevice()->_deviceId;
+    } else if (ctx->getDevice()) {
+      deviceId = ctx->getDevice()->_deviceId;
+    }
+  } else if (ctx->getDevice()) {
+    deviceId = ctx->getDevice()->_deviceId;
+  }
+  return deviceId;
+}
+
+hipFunction_t ihipGetDeviceFunction(const void *hostFunction)
+{
+  int deviceId = getCurrentDeviceId();
+  auto it = g_functions.find(hostFunction);
+  if (it == g_functions.end() || !it->second[deviceId]) {
+    return nullptr;
+  }
+  return it->second[deviceId];
+}
 
 hipError_t hipSetupArgument(
   const void *arg,

hipamd/src/hip_hcc_internal.h

@@ -33,7 +33,7 @@ THE SOFTWARE.
 #include "hip_prof_api.h"
 #include "hip_util.h"
 #include "env.h"
-
+#include <unordered_map>
 
 #if (__hcc_workweek__ < 16354)
 #error("This version of HIP requires a newer version of HCC.");
@@ -1009,6 +1009,18 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod,
 hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0);
 hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream);
 
+/**
+ * @brief Copies the memory address and size of symbol @p symbolName
+ *
+ * @param[in]  symbolName - Symbol on device
+ * @param[out] devPtr - Pointer to a pointer to the memory referred to by the symbol
+ * @param[out] size - Pointer to the size of the symbol
+ * @return #hipSuccess, #hipErrorNotInitialized, #hipErrorNotFound, #hipErrorInvalidValue
+ *
+ */
+hipError_t ihipGetGlobalVar(hipDeviceptr_t* dev_ptr, size_t* size_ptr, const char* hostVar,
+                            hipModule_t hmod = nullptr);
+
 // Stream printf functions:
 inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) {
     os << "stream:";
@@ -1080,4 +1092,14 @@ static inline ihipCtx_t* iihipGetTlsDefaultCtx(TlsData* tls) {
     return tls->defaultCtx;
 }
 
+/**
+ *  @brief Get device function from host kernel function pointer
+ *  Needed only for clang + HIP-HCC RT
+ *
+ *  @param [in] hostFunction host kernel function pointer
+ *
+ *  @returns hipFuntion_t, nullptr
+ */
+hipFunction_t ihipGetDeviceFunction(const void *hostFunction);
+
 #endif

hipamd/src/hip_module.cpp

@@ -344,6 +344,8 @@ hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList
         if (kds[i] == nullptr) {
             return hipErrorInvalidValue;
         }
+        if (!kds[i]->_kernarg_layout.empty()) continue;
+
         hip_impl::kernargs_size_align kargs = ps.get_kernargs_size_align(
                 reinterpret_cast<std::uintptr_t>(lp.func));
         kds[i]->_kernarg_layout = *reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(
@@ -397,6 +399,14 @@ hipError_t ihipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList
     return result;
 }
 
+__attribute__((visibility("default")))
+hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
+                                              int  numDevices, unsigned int  flags) {
+    HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList, numDevices, flags);
+    auto& ps = hip_impl::get_program_state();
+    return ihipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags, ps);
+}
+
 void getGprsLdsUsage(hipFunction_t f, size_t* usedVGPRS, size_t* usedSGPRS, size_t* usedLDS)
 {
     if (f->_is_code_object_v3) {
@@ -736,7 +746,6 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
     mg_sync *mg_sync_ptr = 0;
     vector<mg_info *> mg_info_ptr;
 
-
     result = hip_internal::ihipHostMalloc(tls, (void **)&mg_sync_ptr, sizeof(mg_sync), hipHostMallocDefault, true);
     if (result != hipSuccess) {
         return hipErrorInvalidValue;
@@ -1091,7 +1100,12 @@ namespace hip_impl {
 
     hipError_t agent_globals::read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
                                               const char* name) {
-        return impl->read_agent_global_from_process(dptr, bytes, name);
+        hipError_t result = impl->read_agent_global_from_process(dptr, bytes, name);
+        if(result != hipSuccess) {
+           // For Clang Compiler + Hcc Rt
+           result = ihipGetGlobalVar(dptr, bytes, name);
+        }
+        return result;
     }
 
 } // Namespace hip_impl.
@@ -1259,19 +1273,34 @@ hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t
     if (!*func) return hipErrorInvalidValue;
 
     std::string name_str(name);
+    std::string namekd_str(name_str + ".kd");
+    bool kernel_by_namekd = false;
+
     auto kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent);
 
     if (kernel.handle == 0u) {
-        name_str.append(".kd");
-        kernel = find_kernel_by_name(hmod->executable, name_str.c_str(), agent);
+        kernel_by_namekd = true; //Find kernel by namekd_str
+        kernel = find_kernel_by_name(hmod->executable, namekd_str.c_str(), agent);
     }
 
     if (kernel.handle == 0u) return hipErrorNotFound;
 
+    //For hipModuleLoad(), hmod->kernargs must contain an args with key
+    //name_str or namekd_str.
+    //For hipLaunchKernelGGL(), hmod->kernargs is empty, thus we need
+    //insert hmod->kernargs[name_str]
+    auto it = hmod->kernargs.find(name_str); //Look up args from the original name
+    if (it == hmod->kernargs.end()) {
+        it = hmod->kernargs.find(namekd_str); //Look up args from .kd name
+    }
+
     // TODO: refactor the whole ihipThisThat, which is a mess and yields the
     //       below, due to hipFunction_t being a pointer to ihipModuleSymbol_t.
+
     func[0][0] = *static_cast<hipFunction_t>(
-        Kernel_descriptor{kernel_object(kernel), name_str, hmod->kernargs[name_str]});
+        Kernel_descriptor{kernel_object(kernel),
+        kernel_by_namekd ? namekd_str : name_str,
+        it != hmod->kernargs.end() ? it->second : hmod->kernargs[name_str]});
 
     return hipSuccess;
 }

hipamd/src/hiprtc.cpp

@@ -250,7 +250,7 @@ struct _hiprtcProgram {
 
         const auto it{find_if(reader.sections.begin(), reader.sections.end(),
                               [](const section* x) {
-            return x->get_name() == ".kernel";
+            return (x->get_name() == ".hip_fatbin") || (x->get_name() == ".kernel");
         })};
 
         if (it == reader.sections.end()) return false;
@@ -513,7 +513,7 @@ extern "C" hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char*
 
     const auto src{p->writeTemporaryFiles(tmp.path())};
 
-    vector<string> args{hipcc, "-shared"};
+    vector<string> args{hipcc, "-fPIC -shared"};
     if (n) args.insert(args.cend(), o, o + n);
 
     handleTarget(args);

hipamd/src/program_state.inl

@@ -19,6 +19,8 @@
 #include <hsa/hsa_ven_amd_loader.h>
 #include <amd_comgr.h>
 #include "hc.hpp"
+#include "hip_hcc_internal.h"
+#include "trace_helper.h"
 
 #include <link.h>
 
@@ -734,6 +736,27 @@ public:
                 != AMD_COMGR_STATUS_SUCCESS)
                 return;
 
+            //Look up “.value_kind” to decide whether to ignore it
+            //See http://llvm.org/docs/AMDGPUUsage.html#code-object-v3-metadata-mattr-code-object-v3
+            amd_comgr_metadata_node_t arg_value_kind_md;
+            if (amd_comgr_metadata_lookup(arg_md, ".value_kind", &arg_value_kind_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
+
+            std::string arg_value_kind{ metadata_to_string(arg_value_kind_md) };
+
+            if (amd_comgr_destroy_metadata(arg_value_kind_md)
+                != AMD_COMGR_STATUS_SUCCESS)
+                return;
+
+            if (arg_value_kind.find("hidden_") == 0) {
+                if (amd_comgr_destroy_metadata(arg_md)
+                    != AMD_COMGR_STATUS_SUCCESS)
+                    return;
+
+                continue; //Ignore hidden arg
+            }
+
             amd_comgr_metadata_node_t arg_size_md;
             if (amd_comgr_metadata_lookup(arg_md, ".size", &arg_size_md)
                 != AMD_COMGR_STATUS_SUCCESS)
@@ -937,14 +960,16 @@ public:
 
         auto it0 = get_functions(agent).find(function_address);
 
-        if (it0 == get_functions(agent).cend()) {
-            hip_throw(std::runtime_error{
+        if (it0 != get_functions(agent).cend()) return it0->second;
+
+        // For hip-clang compiler + Hcc RT
+        hipFunction_t f = ihipGetDeviceFunction((const void*)function_address);
+        if (f) return reinterpret_cast<Kernel_descriptor&>(*f);
+
+        hip_throw(std::runtime_error{
                     "No device code available for function: " +
                     std::string(name(function_address)) +
                     ", for agent: " + name(agent)});
-        }
-
-        return it0->second;
     }
 
     const std::vector<std::pair<std::size_t, std::size_t>>&