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
This commit is contained in:
Michael LIAO
2020-03-04 16:05:37 -05:00
parent 47718cbf16
commit a14695d4eb
5 changed files with 114 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/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;
hipamd/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