Fix for hipMemcpyFromSymbol (sync)

Change-Id: I66afec5443ce904a63ced1fafece5144ca59393e


[ROCm/clr commit: e22044de36]

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);
   }
 }
 

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);
 
     //---

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);