From 7c0b9a005bed4eb9a4bc80841dcc78e1d8d813cc Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Thu, 30 Nov 2017 03:29:04 +0000
Subject: [PATCH] Fix legacy mode detection of the address of an agent
 allocated variable. In this mode, there exist two executables per each code
 object, one created by HCC and one created by HIP. Since we dispatch through
 HCC in legacy mode, we should obtain the address for an agent allocated
 variable from the latter's executable. Also add two omitted validity checks,
 whose absence could lead to segfaults when the current process had no .kernel
 section and / or when an invalid or empty blob was extracted from the latter.

---
 include/hip/hcc_detail/code_object_bundle.hpp |  2 +
 src/hip_memory.cpp                            | 44 ++++++++++++-------
 src/program_state.cpp                         |  2 +-
 3 files changed, 31 insertions(+), 17 deletions(-)

diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp
index 05ba44fcc8..72f9d35c73 100644
--- a/include/hip/hcc_detail/code_object_bundle.hpp
+++ b/include/hip/hcc_detail/code_object_bundle.hpp
@@ -76,6 +76,8 @@ namespace hip_impl
             RandomAccessIterator l,
             Bundled_code_header& x)
         {
+            if (f == l) return false;
+
             std::copy_n(f, sizeof(x.cbuf_), x.cbuf_);
 
             if (valid(x)) {
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index c88a1dabc1..fb25101d7e 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -808,6 +808,26 @@ hipError_t hipHostUnregister(void *hostPtr)
     return ihipLogStatus(hip_status);
 }
 
+namespace
+{
+    inline
+    hipDeviceptr_t agent_address_for_symbol(const char* symbolName)
+    {
+        hipDeviceptr_t r = nullptr;
+
+        #if __hcc_workweek__ >= 17481
+            size_t byte_cnt = 0u;
+            hipModuleGetGlobal(&r, &byte_cnt, 0, symbolName);
+        #else
+            auto ctx = ihipGetTlsDefaultCtx();
+            auto acc = ctx->getDevice()->_acc;
+            r = acc.get_symbol_address(symbolName);
+        #endif
+
+        return r;
+    }
+}
+
 hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind)
 {
     HIP_INIT_SPECIAL_API((TRACE_MCMD), symbolName, src, count, offset, kind);
@@ -821,10 +841,8 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void *src, size_t cou
 
     hc::accelerator acc = ctx->getDevice()->_acc;
 
-    hipDeviceptr_t dst = nullptr;
-    size_t byte_cnt = 0u;
-    auto status = hipModuleGetGlobal(
-        &dst, &byte_cnt, 0, static_cast<const char*>(symbolName));
+    hipDeviceptr_t dst =
+        agent_address_for_symbol(static_cast<const char*>(symbolName));
     tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
 
     if(dst == nullptr)
@@ -859,10 +877,8 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count,
 
     hc::accelerator acc = ctx->getDevice()->_acc;
 
-    hipDeviceptr_t src = nullptr;
-    size_t byte_cnt = 0u;
-    auto status = hipModuleGetGlobal(
-        &src, &byte_cnt, 0, static_cast<const char*>(symbolName));
+    hipDeviceptr_t src =
+        agent_address_for_symbol(static_cast<const char*>(symbolName));
     tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
 
     if(dst == nullptr)
@@ -899,10 +915,8 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void *src, size_
 
     hc::accelerator acc = ctx->getDevice()->_acc;
 
-    hipDeviceptr_t dst = nullptr;
-    size_t byte_cnt = 0u;
-    auto status = hipModuleGetGlobal(
-        &dst, &byte_cnt, 0, static_cast<const char*>(symbolName));
+    hipDeviceptr_t dst =
+        agent_address_for_symbol(static_cast<const char*>(symbolName));
     tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
 
     if(dst == nullptr)
@@ -940,10 +954,8 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co
 
     hc::accelerator acc = ctx->getDevice()->_acc;
 
-    hipDeviceptr_t src = nullptr;
-    size_t byte_cnt = 0u;
-    auto status = hipModuleGetGlobal(
-        &src, &byte_cnt, 0, static_cast<const char*>(symbolName));
+    hipDeviceptr_t src =
+        agent_address_for_symbol(static_cast<const char*>(symbolName));
     tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, src);
 
     if(src == nullptr || dst == nullptr)
diff --git a/src/program_state.cpp b/src/program_state.cpp
index 61c90556be..47071d0236 100644
--- a/src/program_state.cpp
+++ b/src/program_state.cpp
@@ -288,7 +288,7 @@ namespace
                 return x->get_type() == SHT_SYMTAB;
             });
 
-            r = function_names_for(reader, symtab);
+            if (symtab) r = function_names_for(reader, symtab);
         });
 
         return r;