From 94a70b3447d38bbb672d554e8123c07df0a01eb8 Mon Sep 17 00:00:00 2001
From: Tao Sang <tao.sang@amd.com>
Date: Thu, 6 Feb 2020 13:56:41 -0500
Subject: [PATCH] Support app(hcc compiled/Hip-Vdi runtime linked)

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: 62ef02928804cb92aecc19a18c0d50f87a5a3f81]
---
 projects/clr/hipamd/vdi/hip_platform.cpp | 72 ++++++++++++++----------
 1 file changed, 42 insertions(+), 30 deletions(-)

diff --git a/projects/clr/hipamd/vdi/hip_platform.cpp b/projects/clr/hipamd/vdi/hip_platform.cpp
index b71a2befad..a33390091c 100644
--- a/projects/clr/hipamd/vdi/hip_platform.cpp
+++ b/projects/clr/hipamd/vdi/hip_platform.cpp
@@ -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; }