From 130545d322986192105f09c4e07a385474578948 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Tue, 21 Mar 2017 23:48:04 +0530
Subject: [PATCH]  Fix for hipMemcpyFromSymbol (sync)

Change-Id: I66afec5443ce904a63ced1fafece5144ca59393e


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

diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp
index 4d922e65e7..17cffbc013 100644
--- a/projects/clr/hipamd/src/hip_hcc.cpp
+++ b/projects/clr/hipamd/src/hip_hcc.cpp
@@ -1872,19 +1872,19 @@ void ihipStream_t::addSymbolPtrToTracker(hc::accelerator& acc, void* ptr, size_t
   hc::am_memtracker_add(ptr, ptrInfo);
 }
 
-void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind)
+void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, size_t offset, unsigned kind)
 {
   if(kind == hipMemcpyHostToHost){
-    acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyHostToHost);
+    acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset, Kalmar::hcMemcpyHostToHost);
   }
   if(kind == hipMemcpyHostToDevice){
-    acc.memcpy_symbol(dst, (void*)src, sizeBytes);
+    acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset);
   }
   if(kind == hipMemcpyDeviceToDevice){
-    acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToDevice);
+    acc.memcpy_symbol(dst, (void*)src, sizeBytes, offset, Kalmar::hcMemcpyDeviceToDevice);
   }
   if(kind == hipMemcpyDeviceToHost){
-    acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToHost);
+    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 b23aead072..1c287bfc44 100644
--- a/projects/clr/hipamd/src/hip_hcc.h
+++ b/projects/clr/hipamd/src/hip_hcc.h
@@ -503,7 +503,7 @@ public:
     void locked_copySync (void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn = true);
     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, 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);
 
     //---
diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp
index a92d11b847..b888c5054c 100644
--- a/projects/clr/hipamd/src/hip_memory.cpp
+++ b/projects/clr/hipamd/src/hip_memory.cpp
@@ -497,7 +497,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
 
     if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost)
     {
-      stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind);
+      stream->lockedSymbolCopySync(acc, dst, (void*)src, count, offset, kind);
     //  acc.memcpy_symbol(dst, (void*)src, count+offset);
     } else {
       return ihipLogStatus(hipErrorInvalidValue);
@@ -532,7 +532,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const char* symbolName, size_t count,
 
     if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost)
     {
-      stream->lockedSymbolCopySync(acc, dst, (void*)src, count + offset, kind);
+      stream->lockedSymbolCopySync(acc, dst, (void*)src, count,  offset, kind);
     }
     else {
       return ihipLogStatus(hipErrorInvalidValue);