AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[vdi] Fix hipGetSymbol{Address|Size}

Browse Source 
- 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: b72196613a]
This commit is contained in:
Michael LIAO
2020-03-04 16:05:37 -05:00
parent 24cd888b3a
commit 5e5f4e23d5
5 changed files with 114 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/hip_runtime.h
+1 -1
View File
@@ -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;
projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+46 -6
View File
@@ -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 <typename T>
hipError_t hipGetSymbolAddress(void** devPtr, const T &symbol) {
return ::hipGetSymbolAddress(devPtr, (const void *)&symbol);
}
template <typename T>
hipError_t hipGetSymbolSize(size_t* size, const T &symbol) {
return ::hipGetSymbolSize(size, (const void *)&symbol);
}
template <typename T>
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 <typename T>
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 <typename T>
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 <typename T>
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 <hip/hcc_detail/hip_prof_str.h>
#endif
projects/hip/vdi/hip_internal.hpp
+6 -2
View File
@@ -237,6 +237,8 @@ private:
std::unordered_map<const void*, DeviceFunction > functions_;
std::unordered_multimap<std::string, DeviceVar > vars_;
// Map from the host shadow symbol to its device name.
std::unordered_map<const void*, std::string> symbols_;
static PlatformState* platform_;
@@ -251,8 +253,10 @@ public:
std::vector< std::pair<hipModule_t, bool> >* 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<std::string>& 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);
projects/hip/vdi/hip_memory.cpp
+28 -12
View File
@@ -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);
}
projects/hip/vdi/hip_platform.cpp
+33 -6
View File
@@ -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<const char*>(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<const char*>(hostVar)), deviceId, hmod);
@@ -542,6 +556,7 @@ extern "C" void __hipRegisterVar(
std::vector<PlatformState::RegisteredVar>{g_devices.size()}, false };
PlatformState::instance().registerVar(hostVar, dvar);
PlatformState::instance().registerVarSym(var, deviceVar);
}
extern "C" void __hipUnregisterFatBinary(std::vector< std::pair<hipModule_t, bool> >* 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);
}