added runtime api hipMemcpyFromSymbolAsync

Change-Id: Ibaf925faf0ba464dd0ed6c5ea74c224c2ce38889


[ROCm/hip commit: 639fd4dd5e]

projects/hip/src/hip_hcc.cpp

@@ -1780,7 +1780,7 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
     }
 }
 
-void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, const void* src, size_t sizeBytes, unsigned kind)
+void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind)
 {
   if(kind == hipMemcpyHostToHost){
     acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyHostToHost);
@@ -1796,11 +1796,18 @@ void ihipStream_t::lockedSymbolCopySync(hc::accelerator &acc, void* dst, const v
   }
 }
 
-void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, const void* src, size_t sizeBytes, unsigned kind)
+void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, void* src, size_t sizeBytes, unsigned kind)
 {
-  hc::AmPointerInfo dstPtrInfo(NULL, dst, sizeBytes, acc, true, false);
-  hc::am_memtracker_add(dst, dstPtrInfo);
-  locked_getAv()->copy_async((void*)src, dst, sizeBytes);
+  if(kind == hipMemcpyHostToDevice) {
+    hc::AmPointerInfo dstPtrInfo(NULL, dst, sizeBytes, acc, true, false);
+    hc::am_memtracker_add(dst, dstPtrInfo);
+    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);
+    locked_getAv()->copy_async((void*)src, dst, sizeBytes);
+  }
 }
 
 void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
@@ -1903,7 +1910,7 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
             LockedAccessor_StreamCrit_t crit(_criticalData);
 
             this->ensureHaveQueue(crit);
-            
+
 #if USE_COPY_EXT_V2
             crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? &copyDevice->getDevice()->_acc : nullptr, forceUnpinnedCopy);
 #else

projects/hip/src/hip_hcc.h

@@ -494,8 +494,8 @@ 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, const void* src, size_t sizeBytes, unsigned kind);
-    void lockedSymbolCopyAsync(hc::accelerator &acc, 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 lockedSymbolCopyAsync(hc::accelerator &acc, void *dst, void* src, size_t sizeBytes, unsigned kind);
 
     //---
     // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.

projects/hip/src/hip_memory.cpp

@@ -461,7 +461,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, 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);
@@ -493,6 +493,44 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
         return ihipLogStatus(hipErrorInvalidSymbol);
     }
 
+    if (stream) {
+        try {
+          stream->lockedSymbolCopyAsync(acc, dst, (void*)src, count + offset, kind);
+        }
+        catch (ihipException ex) {
+            e = ex._code;
+        }
+    } else {
+        e = hipErrorInvalidValue;
+    }
+
+    return ihipLogStatus(e);
+}
+
+
+hipError_t hipMemcpyFromSymbolAsync(void* dst, const char* symbolName, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream)
+{
+    HIP_INIT_CMD_API(symbolName, dst, count, offset, kind, stream);
+
+    if(symbolName == nullptr)
+    {
+        return ihipLogStatus(hipErrorInvalidSymbol);
+    }
+
+    hipError_t e = hipSuccess;
+
+    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, src);
+
+    if(src == nullptr || dst == nullptr)
+    {
+        return ihipLogStatus(hipErrorInvalidSymbol);
+    }
+
     if (stream) {
         try {
           stream->lockedSymbolCopyAsync(acc, dst, src, count + offset, kind);
@@ -505,8 +543,6 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
     }
 
     return ihipLogStatus(e);
-
-
 }
 
 //---