fixed symbol memcpy issue

Change-Id: I89d7401be51d194bcbf771020ba66e3d3b6a18f8


[ROCm/clr commit: 0d4e6ae60a]

projects/clr/hipamd/src/hip_hcc.cpp

@@ -1780,6 +1780,28 @@ 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)
+{
+  if(kind == hipMemcpyHostToHost){
+    acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyHostToHost);
+  }
+  if(kind == hipMemcpyHostToDevice){
+    acc.memcpy_symbol(dst, (void*)src, sizeBytes);
+  }
+  if(kind == hipMemcpyDeviceToDevice){
+    acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToDevice);
+  }
+  if(kind == hipMemcpyDeviceToHost){
+    acc.memcpy_symbol(dst, (void*)src, sizeBytes, Kalmar::hcMemcpyDeviceToHost);
+  }
+}
+
+void ihipStream_t::lockedSymbolCopyAsync(hc::accelerator &acc, void* dst, const 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);
+}
 
 void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes, unsigned kind)
 {
@@ -1881,6 +1903,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

@@ -273,7 +273,7 @@ extern void recordApiTrace(std::string *fullStr, const std::string &apiStr);
     API_TRACE(0, __VA_ARGS__);
 
 
-// Like above, but will trace with DB_CMD.  
+// Like above, but will trace with DB_CMD.
 // Replace HIP_INIT_API with this call inside important APIs that launch work on the GPU:
 // kernel launches, copy commands, memory sets, etc.
 #define HIP_INIT_CMD_API(...) \
@@ -479,8 +479,6 @@ private:
 typedef ihipStreamCriticalBase_t<StreamMutex> ihipStreamCritical_t;
 typedef LockedAccessor<ihipStreamCritical_t> LockedAccessor_StreamCrit_t;
 
-
-
 //---
 // Internal stream structure.
 class ihipStream_t {
@@ -494,10 +492,10 @@ public:
 
     // kind is hipMemcpyKind
     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);
 
     //---
     // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.

projects/clr/hipamd/src/hip_memory.cpp

@@ -449,17 +449,23 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
 
     hc::accelerator acc = ctx->getDevice()->_acc;
 
-    void *ptr = acc.get_symbol_address(symbolName);
-    tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr);
+    void *dst = acc.get_symbol_address(symbolName);
+    tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
 
-    if(ptr == nullptr)
+    if(dst == nullptr)
     {
         return ihipLogStatus(hipErrorInvalidSymbol);
     }
 
     hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
 
-    stream->locked_copySync(ptr, src, count + offset, kind);
+    if(kind == hipMemcpyHostToDevice || kind == hipMemcpyDeviceToHost || kind == hipMemcpyDeviceToDevice || kind == hipMemcpyHostToHost)
+    {
+      stream->lockedSymbolCopySync(acc, dst, src, count + offset, kind);
+    //  acc.memcpy_symbol(dst, (void*)src, count+offset);
+    } else {
+      return ihipLogStatus(hipErrorInvalidValue);
+    }
 
     return ihipLogStatus(hipSuccess);
 }
@@ -479,17 +485,17 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
 
     hc::accelerator acc = ctx->getDevice()->_acc;
 
-    void *ptr = acc.get_symbol_address(symbolName);
-    tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr);
+    void *dst = acc.get_symbol_address(symbolName);
+    tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, dst);
 
-    if(ptr == nullptr)
+    if(dst == nullptr)
     {
         return ihipLogStatus(hipErrorInvalidSymbol);
     }
 
     if (stream) {
         try {
-            stream->locked_copyAsync(ptr, src, count + offset, kind);
+          stream->lockedSymbolCopyAsync(acc, dst, src, count + offset, kind);
         }
         catch (ihipException ex) {
             e = ex._code;

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

@@ -47,7 +47,7 @@ __global__ void Assign(hipLaunchParm lp, int* Out)
 
 int main()
 {
-    int *A, *B, *Ad;
+    int *A, *Am, *B, *Ad;
     A = new int[NUM];
     B = new int[NUM];
     for(unsigned i=0;i<NUM;i++) {
@@ -56,10 +56,14 @@ int main()
     }
 
     hipMalloc((void**)&Ad, SIZE);
+    hipHostMalloc((void**)&Am, SIZE);
+    for(unsigned i=0;i<NUM;i++) {
+        Am[i] = -1*i;
+    }
 
     hipStream_t stream;
     hipStreamCreate(&stream);
-    hipMemcpyToSymbolAsync(HIP_SYMBOL(global), A, SIZE, 0, hipMemcpyHostToDevice, stream);
+    hipMemcpyToSymbolAsync(HIP_SYMBOL(global), 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);