added runtime api hipMemcpyFromSymbolAsync

Change-Id: Ibaf925faf0ba464dd0ed6c5ea74c224c2ce38889


[ROCm/clr commit: 3d348b2d81]

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -1159,7 +1159,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz
  */
 hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
 
-
+hipError_t hipMemcpyFromSymbolAsync(void *dst, const char* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
 
 /**
  *  @brief Copy data from src to dst asynchronously.

projects/clr/hipamd/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/clr/hipamd/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/clr/hipamd/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);
-
-
 }
 
 //---

projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp

@@ -18,7 +18,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s 
+ * BUILD: %t %s
  * RUN: %t
  * HIT_END
  */
@@ -32,44 +32,53 @@ THE SOFTWARE.
 #define SIZE 1024*4
 
 #ifdef __HIP_PLATFORM_HCC__
-__attribute__((address_space(1))) int global[NUM];
+__attribute__((address_space(1))) int globalIn[NUM];
+__attribute__((address_space(1))) int globalOut[NUM];
 #endif
 
 #ifdef __HIP_PLATFORM_NVCC__
-__device__ int global[NUM];
+__device__ int globalIn[NUM];
+__device__ int globalOut[NUM];
 #endif
 
 __global__ void Assign(hipLaunchParm lp, int* Out)
 {
     int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
-    Out[tid] = global[tid];
+    Out[tid] = globalIn[tid];
+    globalOut[tid] = globalIn[tid];
 }
 
 int main()
 {
-    int *A, *Am, *B, *Ad;
+    int *A, *Am, *B, *Ad, *C, *Cm;
     A = new int[NUM];
     B = new int[NUM];
+    C = new int[NUM];
     for(unsigned i=0;i<NUM;i++) {
         A[i] = -1*i;
         B[i] = 0;
+        C[i] = 0;
     }
 
     hipMalloc((void**)&Ad, SIZE);
     hipHostMalloc((void**)&Am, SIZE);
+    hipHostMalloc((void**)&Cm, SIZE);
     for(unsigned i=0;i<NUM;i++) {
         Am[i] = -1*i;
+        Cm[i] = 0;
     }
 
     hipStream_t stream;
     hipStreamCreate(&stream);
-    hipMemcpyToSymbolAsync(HIP_SYMBOL(global), Am, SIZE, 0, hipMemcpyHostToDevice, stream);
+    hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream);
     hipStreamSynchronize(stream);
     hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
     hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
-
+    hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
+    hipStreamSynchronize(stream);
     for(unsigned i=0;i<NUM;i++) {
-        assert(A[i] == B[i]);
+        assert(Am[i] == B[i]);
+        assert(Am[i] == Cm[i]);
     }
 
     for(unsigned i=0;i<NUM;i++) {
@@ -77,11 +86,12 @@ int main()
         B[i] = 0;
     }
 
-    hipMemcpyToSymbol(HIP_SYMBOL(global), A, SIZE, 0, hipMemcpyHostToDevice);
+    hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice);
     hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
     hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);
     for(unsigned i=0;i<NUM;i++) {
         assert(A[i] == B[i]);
     }
+
     passed();
 }