From 62c3bd71416ac94e7761eca1d20bcce6ca7bb08b Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Fri, 24 Mar 2017 10:30:33 +0530
Subject: [PATCH] Fix for hipMemcpyFromSymbolAsync

Change-Id: I449c669c8f0ef041deaf0a1bc812a71b2f0cc5a6


[ROCm/clr commit: dfa516f80432dada86167ce8c00626d364dc208b]
---
 projects/clr/hipamd/src/hip_hcc.cpp    | 26 +++++++++++++++++++++-----
 projects/clr/hipamd/src/hip_hcc.h      |  2 +-
 projects/clr/hipamd/src/hip_memory.cpp |  5 +++--
 3 files changed, 25 insertions(+), 8 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp
index 17cffbc013..e422a6d4db 100644
--- a/projects/clr/hipamd/src/hip_hcc.cpp
+++ b/projects/clr/hipamd/src/hip_hcc.cpp
@@ -1888,15 +1888,31 @@ 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)
+void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind)
 {
   if(kind == hipMemcpyHostToDevice) {
-    addSymbolPtrToTracker(acc, dst, sizeBytes);
-    locked_getAv()->copy_async((void*)src, dst, sizeBytes);
+    hc::AmPointerInfo srcPtrInfo(NULL, NULL, 0, acc, 0, 0);
+    bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);
+    if(srcTracked) {
+        addSymbolPtrToTracker(acc, dst, sizeBytes);
+        locked_getAv()->copy_async((void*)src, dst, sizeBytes);
+    } else {
+        LockedAccessor_StreamCrit_t crit(_criticalData);
+        this->wait(crit);
+        acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset);
+    }
   }
   if(kind == hipMemcpyDeviceToHost) {
-    addSymbolPtrToTracker(acc, src, sizeBytes);
-    locked_getAv()->copy_async((void*)src, dst, sizeBytes);
+    hc::AmPointerInfo dstPtrInfo(NULL, NULL, 0, acc, 0, 0);
+    bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
+    if(dstTracked) {
+        addSymbolPtrToTracker(acc, src, sizeBytes);
+        locked_getAv()->copy_async((void*)src, dst, sizeBytes);
+    } else {
+        LockedAccessor_StreamCrit_t crit(_criticalData);
+        this->wait(crit);
+        acc.memcpy_symbol((void*)src, (void*)dst, sizeBytes, offset, Kalmar::hcMemcpyDeviceToHost);
+    }
   }
 }
 
diff --git a/projects/clr/hipamd/src/hip_hcc.h b/projects/clr/hipamd/src/hip_hcc.h
index 1c287bfc44..245f154305 100644
--- a/projects/clr/hipamd/src/hip_hcc.h
+++ b/projects/clr/hipamd/src/hip_hcc.h
@@ -504,7 +504,7 @@ public:
     void locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind);
 
     void lockedSymbolCopySync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, size_t offset, unsigned kind);
-    void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, unsigned kind);
+    void lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, size_t offset, unsigned kind);
 
     //---
     // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.
diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index b888c5054c..94121838fd 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -567,7 +567,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
 
     if (stream) {
         try {
-          stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count + offset, kind);
+          stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count, offset, kind);
         }
         catch (ihipException ex) {
             e = ex._code;
@@ -603,9 +603,10 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t co
         return ihipLogStatus(hipErrorInvalidSymbol);
     }
 
+    stream = ihipSyncAndResolveStream(stream);
     if (stream) {
         try {
-          stream->lockedSymbolCopyAsync(acc, dst, src, count + offset, kind);
+          stream->lockedSymbolCopyAsync(acc, dst, src, count, offset, kind);
         }
         catch (ihipException ex) {
             e = ex._code;