From d432dbfe20933775e8bfd578748adb38aa9dedf4 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Fri, 13 Mar 2020 23:05:43 -0400 Subject: [PATCH] Fix failure to get global variables Implement _ihipGetGlobalVar() and ihipGetGlobalVar() to get global variables. Change-Id: I442ab6712e12306c3316f114f5dc42f6daefaad9 --- src/hip_clang.cpp | 215 +++++++++++++++++++++++++++++++++++++++-- src/hip_hcc_internal.h | 12 +++ src/hip_module.cpp | 7 +- 3 files changed, 225 insertions(+), 9 deletions(-) diff --git a/src/hip_clang.cpp b/src/hip_clang.cpp index f17d2673c7..635aa8391f 100644 --- a/src/hip_clang.cpp +++ b/src/hip_clang.cpp @@ -157,16 +157,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* modules; + std::vector rvars; + bool dyn_undef; +}; + +std::unordered_multimap 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* modules, - char* hostVar, - char* deviceVar, - const char* deviceName, - int ext, - int size, - int constant, - int global) + std::vector* 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), modules, + std::vector{ 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* modules) diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index ac63f49dba..8eec7292c2 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -982,6 +982,18 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0); hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream); void ihipStreamCallbackHandler(ihipStreamCallback_t* cb); +/** + * @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:"; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index b527d6ccb0..b6fbc45866 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -879,7 +879,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.