Fix for hipMemcpyFromSymbolAsync

Change-Id: I449c669c8f0ef041deaf0a1bc812a71b2f0cc5a6


[ROCm/clr commit: dfa516f804]

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

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.

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;