From 5e5f4e23d52505d4f8f0667cbf7cef662348c6ec Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Wed, 4 Mar 2020 16:05:37 -0500 Subject: [PATCH] [vdi] Fix hipGetSymbol{Address|Size} - Use symbol value as the qeury key. Compared to the symbol name, the symbol value is more robust as developers may use unqualified or qualified identifiers. It also removes the mangling and/or demangling requirement for the runtime API. Change-Id: I9d4259f3842612c7cc98551269fc2092d8b5c19e [ROCm/hip commit: b72196613a60ceacb0c167c204739194260a9893] --- .../hip/include/hip/hcc_detail/hip_runtime.h | 2 +- .../include/hip/hcc_detail/hip_runtime_api.h | 52 ++++++++++++++++--- projects/hip/vdi/hip_internal.hpp | 8 ++- projects/hip/vdi/hip_memory.cpp | 40 +++++++++----- projects/hip/vdi/hip_platform.cpp | 39 +++++++++++--- 5 files changed, 114 insertions(+), 27 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index ef591b8c51..a3a0963ba9 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -379,7 +379,7 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri #elif defined(__clang__) && defined(__HIP__) #define HIP_KERNEL_NAME(...) __VA_ARGS__ -#define HIP_SYMBOL(X) #X +#define HIP_SYMBOL(X) X typedef int hipLaunchParm; diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index d35d426a5e..2b5649e721 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1488,18 +1488,18 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, const char* name); -hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName); -hipError_t hipGetSymbolSize(size_t* size, const void* symbolName); -hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, +hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol); +hipError_t hipGetSymbolSize(size_t* size, const void* symbol); +hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeBytes, size_t offset __dparm(0), hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)); -hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, +hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream __dparm(0)); -hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, +hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol, size_t sizeBytes, size_t offset __dparm(0), hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)); -hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, +hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream __dparm(0)); @@ -3517,6 +3517,46 @@ static hipError_t __host__ inline hipOccupancyMaxActiveBlocksPerMultiprocessorWi } #endif // defined(__cplusplus) && !defined(__HCC__) && defined(__clang__) && defined(__HIP__) +#if defined(__cplusplus) && !defined(__HCC__) + +template +hipError_t hipGetSymbolAddress(void** devPtr, const T &symbol) { + return ::hipGetSymbolAddress(devPtr, (const void *)&symbol); +} + +template +hipError_t hipGetSymbolSize(size_t* size, const T &symbol) { + return ::hipGetSymbolSize(size, (const void *)&symbol); +} + +template +hipError_t hipMemcpyToSymbol(const T& symbol, const void* src, size_t sizeBytes, + size_t offset __dparm(0), + hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)) { + return ::hipMemcpyToSymbol((const void*)&symbol, src, sizeBytes, offset, kind); +} + +template +hipError_t hipMemcpyToSymbolAsync(const T& symbol, const void* src, size_t sizeBytes, size_t offset, + hipMemcpyKind kind, hipStream_t stream __dparm(0)) { + return ::hipMemcpyToSymbolAsync((const void*)&symbol, src, sizeBytes, offset, kind, stream); +} + +template +hipError_t hipMemcpyFromSymbol(void* dst, const T &symbol, + size_t sizeBytes, size_t offset __dparm(0), + hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)) { + return ::hipMemcpyFromSymbol(dst, (const void*)&symbol, sizeBytes, offset, kind); +} + +template +hipError_t hipMemcpyFromSymbolAsync(void* dst, const T& symbol, size_t sizeBytes, size_t offset, + hipMemcpyKind kind, hipStream_t stream __dparm(0)) { + return ::hipMemcpyFromSymbolAsync(dst, (const void*)&symbol, sizeBytes, offset, kind, stream); +} + +#endif + #if USE_PROF_API #include #endif diff --git a/projects/hip/vdi/hip_internal.hpp b/projects/hip/vdi/hip_internal.hpp index d5d35fcf1d..92397d1255 100755 --- a/projects/hip/vdi/hip_internal.hpp +++ b/projects/hip/vdi/hip_internal.hpp @@ -237,6 +237,8 @@ private: std::unordered_map functions_; std::unordered_multimap vars_; + // Map from the host shadow symbol to its device name. + std::unordered_map symbols_; static PlatformState* platform_; @@ -251,8 +253,10 @@ public: std::vector< std::pair >* unregisterVar(hipModule_t hmod); + bool findSymbol(const void *hostVar, std::string &devName); PlatformState::DeviceVar* findVar(std::string hostVar, int deviceId, hipModule_t hmod); - void registerVar(const void* hostvar, const DeviceVar& var); + void registerVarSym(const void *hostVar, const char *symbolName); + void registerVar(const char* symbolName, const DeviceVar& var); void registerFunction(const void* hostFunction, const DeviceFunction& func); bool registerModFuncs(std::vector& func_names, hipModule_t* module); @@ -260,7 +264,7 @@ public: bool createFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name); hipFunction_t getFunc(const void* hostFunction, int deviceId); bool getFuncAttr(const void* hostFunction, hipFuncAttributes* func_attr); - bool getGlobalVar(const void* hostVar, int deviceId, hipModule_t hmod, + bool getGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod, hipDeviceptr_t* dev_ptr, size_t* size_ptr); bool getTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef); diff --git a/projects/hip/vdi/hip_memory.cpp b/projects/hip/vdi/hip_memory.cpp index 083b855fa9..78288a6a78 100644 --- a/projects/hip/vdi/hip_memory.cpp +++ b/projects/hip/vdi/hip_memory.cpp @@ -717,15 +717,19 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { }; -hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count, +hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t count, size_t offset, hipMemcpyKind kind) { - HIP_INIT_API(hipMemcpyToSymbol, symbolName, src, count, offset, kind); + HIP_INIT_API(hipMemcpyToSymbol, symbol, src, count, offset, kind); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; + std::string symbolName; + if (!PlatformState::instance().findSymbol(symbol, symbolName)) { + HIP_RETURN(hipErrorInvalidSymbol); + } /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr, &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } @@ -741,15 +745,19 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou HIP_RETURN(hipMemcpy(device_ptr, src, count, kind)); } -hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, +hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind) { - HIP_INIT_API(hipMemcpyFromSymbol, symbolName, dst, count, offset, kind); + HIP_INIT_API(hipMemcpyFromSymbol, symbol, dst, count, offset, kind); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; + std::string symbolName; + if (!PlatformState::instance().findSymbol(symbol, symbolName)) { + HIP_RETURN(hipErrorInvalidSymbol); + } /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr, &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } @@ -765,15 +773,19 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count, HIP_RETURN(hipMemcpy(dst, device_ptr, count, kind)); } -hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_t count, +hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(hipMemcpyToSymbolAsync, symbolName, src, count, offset, kind, stream); + HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, count, offset, kind, stream); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; + std::string symbolName; + if (!PlatformState::instance().findSymbol(symbol, symbolName)) { + HIP_RETURN(hipErrorInvalidSymbol); + } /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr, &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } @@ -789,15 +801,19 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_ HIP_RETURN(hipMemcpyAsync(device_ptr, src, count, kind, stream)); } -hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count, +hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream) { - HIP_INIT_API(hipMemcpyFromSymbolAsync, symbolName, dst, count, offset, kind, stream); + HIP_INIT_API(hipMemcpyFromSymbolAsync, symbol, dst, count, offset, kind, stream); size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; + std::string symbolName; + if (!PlatformState::instance().findSymbol(symbol, symbolName)) { + HIP_RETURN(hipErrorInvalidSymbol); + } /* Get address and size for the global symbol */ - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr, &device_ptr, &sym_size)) { HIP_RETURN(hipErrorInvalidSymbol); } diff --git a/projects/hip/vdi/hip_platform.cpp b/projects/hip/vdi/hip_platform.cpp index 0ddc2fe8af..ca79c5300d 100755 --- a/projects/hip/vdi/hip_platform.cpp +++ b/projects/hip/vdi/hip_platform.cpp @@ -266,7 +266,21 @@ PlatformState::DeviceVar* PlatformState::findVar(std::string hostVar, int device return dvar; } -void PlatformState::registerVar(const void* hostvar, +bool PlatformState::findSymbol(const void *hostVar, std::string &symbolName) { + auto it = symbols_.find(hostVar); + if (it != symbols_.end()) { + symbolName = it->second; + return true; + } + return false; +} + +void PlatformState::registerVarSym(const void *hostVar, const char *symbolName) { + amd::ScopedLock lock(lock_); + symbols_.insert(std::make_pair(hostVar, std::string(symbolName))); +} + +void PlatformState::registerVar(const char* hostvar, const DeviceVar& rvar) { amd::ScopedLock lock(lock_); vars_.insert(std::make_pair(std::string(reinterpret_cast(hostvar)), rvar)); @@ -447,7 +461,7 @@ bool PlatformState::getTexRef(const char* hostVar, hipModule_t hmod, textureRefe return true; } -bool PlatformState::getGlobalVar(const void* hostVar, int deviceId, hipModule_t hmod, +bool PlatformState::getGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod, hipDeviceptr_t* dev_ptr, size_t* size_ptr) { amd::ScopedLock lock(lock_); DeviceVar* dvar = findVar(std::string(reinterpret_cast(hostVar)), deviceId, hmod); @@ -542,6 +556,7 @@ extern "C" void __hipRegisterVar( std::vector{g_devices.size()}, false }; PlatformState::instance().registerVar(hostVar, dvar); + PlatformState::instance().registerVarSym(var, deviceVar); } extern "C" void __hipUnregisterFatBinary(std::vector< std::pair >* modules) @@ -643,18 +658,30 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction) exec.sharedMem_, exec.hStream_, nullptr, extra)); } -hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { +hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol) { + HIP_INIT_API(hipGetSymbolAddress, devPtr, symbol); + + std::string symbolName; + if (!PlatformState::instance().findSymbol(symbol, symbolName)) { + HIP_RETURN(hipErrorInvalidSymbol); + } size_t size = 0; - if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + if(!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr, devPtr, &size)) { HIP_RETURN(hipErrorInvalidSymbol); } HIP_RETURN(hipSuccess); } -hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbolName) { +hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbol) { + HIP_INIT_API(hipGetSymbolSize, sizePtr, symbol); + + std::string symbolName; + if (!PlatformState::instance().findSymbol(symbol, symbolName)) { + HIP_RETURN(hipErrorInvalidSymbol); + } hipDeviceptr_t devPtr = nullptr; - if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr, + if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr, &devPtr, sizePtr)) { HIP_RETURN(hipErrorInvalidSymbol); }