AI/rocm-systems
2
0
Derivar 0
Código Questões Pedidos de integração Operações Pacotes Planeamentos Lançamentos Wiki Trabalho

Support app(hcc compiled/Hip-Vdi runtime linked)

Ver a proveniência 
The issues of the following functions have been fixed.
hipModuleLoad: Make Hip-Vdi runtime able to read code object module
generated by Hcc compiler.
hipLaunchKernel: Use introspect method to find function if it cannot
be found from platform state instance.

Change-Id: Id740e5a96614ec6a0b6c704f8f74600bfdc4983e


[ROCm/clr commit: 62ef029288]
Este cometimento está contido em:
Tao Sang
2020-02-06 13:56:41 -05:00
cometido por Tao Sang
ascendente 6982f99f6e
cometimento 94a70b3447
1 ficheiros modificados com 42 adições e 30 eliminações
Descarregar ficheiro patch Descarregar ficheiro diff Expandir todos os ficheiros Colapsar todos os ficheiros
projects/clr/hipamd/vdi/hip_platform.cpp
+42 -30
Ver ficheiro
@@ -96,12 +96,17 @@ bool __hipExtractCodeObjectFromFatBinary(const void* data,
desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
reinterpret_cast<uintptr_t>(&desc->triple[0]) + desc->tripleSize)) {
std::string triple(desc->triple, sizeof(HIP_AMDGCN_AMDHSA_TRIPLE) - 1);
if (triple.compare(HIP_AMDGCN_AMDHSA_TRIPLE))
std:size_t offset = 0;
if (!std::strncmp(desc->triple, HIP_AMDGCN_AMDHSA_TRIPLE,
sizeof(HIP_AMDGCN_AMDHSA_TRIPLE) - 1)) {
offset = sizeof(HIP_AMDGCN_AMDHSA_TRIPLE); //For code objects created by CLang
} else if (!std::strncmp(desc->triple, HCC_AMDGCN_AMDHSA_TRIPLE,
sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1)) {
offset = sizeof(HCC_AMDGCN_AMDHSA_TRIPLE); //For code objects created by Hcc
} else {
continue;
std::string target(desc->triple + sizeof(HIP_AMDGCN_AMDHSA_TRIPLE),
desc->tripleSize - sizeof(HIP_AMDGCN_AMDHSA_TRIPLE));
}
std::string target(desc->triple + offset, desc->tripleSize - offset);
const void *image = reinterpret_cast<const void*>(
reinterpret_cast<uintptr_t>(obheader) + desc->offset);
@@ -556,31 +561,6 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
exec.sharedMem_, exec.hStream_, nullptr, extra));
}
extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMemBytes,
hipStream_t stream)
{
HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes,
stream);
int deviceId = ihipGetDevice();
if (deviceId == -1) {
HIP_RETURN(hipErrorNoDevice);
}
hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId);
if (func == nullptr) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
HIP_RETURN(hipModuleLaunchKernel(func, gridDim.x, gridDim.y, gridDim.z,
blockDim.x, blockDim.y, blockDim.z,
sharedMemBytes, stream, args, nullptr));
}
hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
size_t size = 0;
if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr,
@@ -950,6 +930,38 @@ void hipLaunchCooperativeKernelGGLImpl(
#endif // defined(ATI_OS_LINUX)
extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMemBytes,
hipStream_t stream)
{
HIP_INIT_API(NONE, hostFunction, gridDim, blockDim, args, sharedMemBytes,
stream);
int deviceId = ihipGetDevice();
if (deviceId == -1) {
HIP_RETURN(hipErrorNoDevice);
}
hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId);
if (func == nullptr) {
#ifdef ATI_OS_LINUX
const auto it = hip_impl::functions().find(reinterpret_cast<uintptr_t>(hostFunction));
if (it == hip_impl::functions().cend()) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
func = it->second;
#else
HIP_RETURN(hipErrorInvalidDeviceFunction);
#endif
}
HIP_RETURN(hipModuleLaunchKernel(func, gridDim.x, gridDim.y, gridDim.z,
blockDim.x, blockDim.y, blockDim.z,
sharedMemBytes, stream, args, nullptr));
}
// conversion routines between float and half precision
static inline std::uint32_t f32_as_u32(float f) { union { float f; std::uint32_t u; } v; v.f = f; return v.u; }
static inline float u32_as_f32(std::uint32_t u) { union { float f; std::uint32_t u; } v; v.u = u; return v.f; }