From cfe930f9d67ba4e11daa7fb751e3f30ab9922c9a Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Thu, 28 Mar 2019 22:45:04 -0500
Subject: [PATCH] Make hipModuleGetGlobal be in HIP runtime so it can be
 discovered at runtime (#981)

* Make hipModuleGetGlobal be in HIP runtime so it can be discovered at runtime

In HIP PR #929, quite a few HIP public APIs were made as inline functions with
hidden visibility. It was necessary to support applications with shared
libraries with GPU kernels launched via hipLaunchKernelGGL(), after HIP runtime
is initialized.

In empirical tests, the implementation has been proved to be a bit too
excessive, especially for hipModuleGetGlobal(). The function is used by another
type of client applications which relies on the existence of this function
within HIP runtime so global symbols from HSA code objects loaded dynamically
at runtime can be retrieved programmtically.

This commit moves hipModuleGetGlobal() back to src/hip_module.cpp, and makes it
visible and not inline, to fulfill requirements for applications
aforementioned. It does not change the behavior of applications depending on
hipLaunchKernelGGL().

* Add HIP_INIT_API into the implementation of hipModuleGetGlobal

Address review comments.

* Fix failing HIP unit tests


[ROCm/clr commit: 04915cea2f3ee33a37ac545c340bb98fdb395408]
---
 .../include/hip/hcc_detail/hip_runtime_api.h  | 25 +++++++------------
 projects/clr/hipamd/src/hip_module.cpp        | 10 ++++++++
 2 files changed, 19 insertions(+), 16 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
index 7d9f08ac69..cee4e81054 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
@@ -1419,9 +1419,14 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
                                     hipMemcpyKind kind,
                                     hipStream_t stream __dparm(0));
 #else
-__attribute__((visibility("hidden")))
 hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*);
 
+namespace hip_impl {
+inline
+__attribute__((visibility("hidden")))
+hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                                          const char* name);
+} // Namespace hip_impl.
 
 /**
  *  @brief Copies the memory address of symbol @p symbolName to @p devPtr
@@ -1439,7 +1444,7 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
     //HIP_INIT_API(hipGetSymbolAddress, devPtr, symbolName);
     hip_impl::hip_init();
     size_t size = 0;
-    return hipModuleGetGlobal(devPtr, &size, 0, (const char*)symbolName);
+    return hip_impl::read_agent_global_from_process(devPtr, &size, (const char*)symbolName);
 }
 
 
@@ -1459,7 +1464,7 @@ hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) {
     // HIP_INIT_API(hipGetSymbolSize, size, symbolName);
     hip_impl::hip_init();
     void* devPtr = nullptr;
-    return hipModuleGetGlobal(&devPtr, size, 0, (const char*)symbolName);
+    return hip_impl::read_agent_global_from_process(&devPtr, size, (const char*)symbolName);
 }
 
 #if defined(__cplusplus)
@@ -2710,20 +2715,8 @@ extern "C" {
  *
  * @returns hipSuccess, hipErrorInvalidValue, hipErrorNotInitialized
  */
-inline
-__attribute__((visibility("hidden")))
 hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
-                              hipModule_t hmod, const char* name) {
-    if (!dptr || !bytes) return hipErrorInvalidValue;
-
-    if (!name) return hipErrorNotInitialized;
-
-    const auto r = hmod ?
-        hip_impl::read_agent_global_from_module(dptr, bytes, hmod, name) :
-        hip_impl::read_agent_global_from_process(dptr, bytes, name);
-
-    return r;
-}
+                              hipModule_t hmod, const char* name);
 #endif // __HIP_VDI__
 
 hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp
index 6401e13727..e81204d86b 100644
--- a/projects/clr/hipamd/src/hip_module.cpp
+++ b/projects/clr/hipamd/src/hip_module.cpp
@@ -282,6 +282,16 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
         localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0));
 }
 
+hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
+                              hipModule_t hmod, const char* name) {
+    HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
+    if (!dptr || !bytes || !hmod) return hipErrorInvalidValue;
+
+    if (!name) return hipErrorNotInitialized;
+
+    return hip_impl::read_agent_global_from_module(dptr, bytes, hmod, name);
+}
+
 namespace hip_impl {
     hsa_executable_t executable_for(hipModule_t hmod) {
         return hmod->executable;