From 2e245ae58cd412fd72363326a789f1365fee4f39 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Thu, 23 Feb 2017 11:29:06 -0600
Subject: [PATCH] Added initial support for hipMemcpyFromSymbol. But not
 working!

Change-Id: I48d8c7de4ec9f85c6c942be995fb488a3931f5d7
---
 include/hip/hcc_detail/hip_runtime_api.h    |  2 ++
 src/hip_hcc.cpp                             | 12 ++++---
 src/hip_hcc.h                               |  2 ++
 src/hip_memory.cpp                          | 36 +++++++++++++++++++++
 tests/src/deviceLib/hipTestDeviceSymbol.cpp |  2 ++
 5 files changed, 49 insertions(+), 5 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 2c75b584c4..f156d3fdbd 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -1159,6 +1159,8 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz
  */
 hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
 
+hipError_t hipMemcpyFromSymbol(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind);
+
 hipError_t hipMemcpyFromSymbolAsync(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
 
 /**
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index a119ea1c54..efc0265ba8 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -1750,7 +1750,6 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
     bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
     bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
 
-
     hc::hcCommandKind hcCopyDir;
     ihipCtx_t *copyDevice;
     bool forceUnpinnedCopy;
@@ -1780,6 +1779,11 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
     }
 }
 
+void ihipStream_t::addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes) {
+  hc::AmPointerInfo ptrInfo(NULL, ptr, sizeBytes, acc, true, false);
+  hc::am_memtracker_add(ptr, ptrInfo);
+}
+
 void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind)
 {
   if(kind == hipMemcpyHostToHost){
@@ -1799,13 +1803,11 @@ void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* s
 void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind)
 {
   if(kind == hipMemcpyHostToDevice) {
-    hc::AmPointerInfo dstPtrInfo(NULL, dst, sizeBytes, acc, true, false);
-    hc::am_memtracker_add(dst, dstPtrInfo);
+    addSymbolPtrToTracker(acc, dst, sizeBytes);
     locked_getAv()->copy_async((void*)src, dst, sizeBytes);
   }
   if(kind == hipMemcpyDeviceToHost) {
-    hc::AmPointerInfo srcPtrInfo(NULL, src, sizeBytes, acc, true, false);
-    hc::am_memtracker_add(src, srcPtrInfo);
+    addSymbolPtrToTracker(acc, src, sizeBytes);
     locked_getAv()->copy_async((void*)src, dst, sizeBytes);
   }
 }
diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index d7d92a221c..9ebac73cad 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -551,6 +551,8 @@ private:
 
     bool canSeeMemory(const ihipCtx_t *thisCtx, const hc::AmPointerInfo *dstInfo, const hc::AmPointerInfo *srcInfo);
 
+    void addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t sizeBytes);
+
 public: // TODO - move private
     // Critical Data - MUST be accessed through LockedAccessor_StreamCrit_t
     ihipStreamCritical_t        _criticalData;
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index cb265159ba..479040c099 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -470,6 +470,42 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
     return ihipLogStatus(hipSuccess);
 }
 
+
+hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count, size_t offset, hipMemcpyKind kind)
+{
+    HIP_INIT_CMD_API(symbolName, dst, count, offset, kind);
+
+    if(symbolName == nullptr)
+    {
+        return ihipLogStatus(hipErrorInvalidSymbol);
+    }
+
+    auto ctx = ihipGetTlsDefaultCtx();
+
+    hc::accelerator acc = ctx->getDevice()->_acc;
+
+    void *src = acc.get_symbol_address(symbolName);
+    tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
+
+    if(dst == nullptr)
+    {
+        return ihipLogStatus(hipErrorInvalidSymbol);
+    }
+
+    hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
+
+    if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost)
+    {
+      stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind);
+    }
+    else {
+      return ihipLogStatus(hipErrorInvalidValue);
+    }
+
+    return ihipLogStatus(hipSuccess);
+}
+
+
 hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream)
 {
     HIP_INIT_CMD_API(symbolName, src, count, offset, kind, stream);
diff --git a/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/tests/src/deviceLib/hipTestDeviceSymbol.cpp
index 476c5e0997..00a1c52565 100644
--- a/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+++ b/tests/src/deviceLib/hipTestDeviceSymbol.cpp
@@ -89,8 +89,10 @@ int main()
     hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice);
     hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
     hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
+    hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost);
     for(unsigned i=0;i<NUM;i++) {
         assert(A[i] == B[i]);
+//        assert(A[i] == C[i]);
     }
 
     passed();