diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h
index 1243934d01..4cbaf9ea4e 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h
@@ -1062,6 +1062,27 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
 hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind);
 
 
+/**
+ *  @brief Copies @p sizeBytes bytes from the memory area pointed to by @p src to the memory area pointed to by @p offset bytes from the start of symbol @p symbol 
+ *
+ *  The memory areas may not overlap. Symbol can either be a variable that resides in global or constant memory space, or it can be a character string,
+ *  naming a variable that resides in global or constant memory space. Kind can be either hipMemcpyHostToDevice or hipMemcpyDeviceToDevice
+ *  hipMemcpyToSymbolAsync() is asynchronous with respect to the host, so the call may return before copy is complete.
+ *  TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now.
+ *
+ *  @param[in]  symbolName - Symbol destination on device
+ *  @param[in]  src - Data being copy from
+ *  @param[in]  sizeBytes - Data size in bytes
+ *  @param[in]  offset - Offset from start of symbol in bytes
+ *  @param[in]  kind - Type of transfer
+ *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree, #hipErrorUnknown
+ *
+ *  @see hipMemcpy, hipMemcpy2D, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpy2DFromArray, hipMemcpyArrayToArray, hipMemcpy2DArrayToArray, hipMemcpyFromSymbol, hipMemcpyAsync, hipMemcpy2DAsync, hipMemcpyToArrayAsync, hipMemcpy2DToArrayAsync, hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync
+ */
+hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream);
+
+
+
 /**
  *  @brief Copy data from src to dst asynchronously.
  *
diff --git a/hipamd/include/hip/hip_runtime_api.h b/hipamd/include/hip/hip_runtime_api.h
index 3406bcbbc9..f08f03e6f5 100644
--- a/hipamd/include/hip/hip_runtime_api.h
+++ b/hipamd/include/hip/hip_runtime_api.h
@@ -183,7 +183,7 @@ typedef enum hipError_t {
     hipErrorInvalidHandle           = 400,
     hipErrorNotFound                = 500,
     hipErrorIllegalAddress          = 700,
-
+    hipErrorInvalidSymbol           = 701,
 // Runtime Error Codes start here.
     hipErrorMissingConfiguration    = 1001,
     hipErrorMemoryAllocation        = 1002,    ///< Memory allocation error.
diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index 278502973a..a7c19e8949 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -385,24 +385,67 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
 {
     HIP_INIT_API(symbolName, src, count, offset, kind);
 
-#ifdef USE_MEMCPYTOSYMBOL
-    if(kind != hipMemcpyHostToDevice)
+    if(symbolName == nullptr)
     {
-        return ihipLogStatus(hipErrorInvalidValue);
+        return ihipLogStatus(hipErrorInvalidSymbol);
     }
-    auto ctx = ihipGetTlsDefaultCtx();
 
-    //hsa_signal_t depSignal;
-    //int depSignalCnt = ctx._default_stream->preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D);
-    assert(0);  // Need to properly synchronize the copy - do something with depSignal if != NULL.
+    auto ctx = ihipGetTlsDefaultCtx();
 
     hc::accelerator acc = ctx->getDevice()->_acc;
 
-    acc.memcpy_symbol(symbolName, (void*) src,count, offset);
-#endif
+    void *ptr = acc.get_symbol_address(symbolName);
+
+    if(ptr == nullptr)
+    {
+        return ihipLogStatus(hipErrorInvalidSymbol);
+    }
+
+    hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull);
+
+    stream->locked_copySync(ptr, src, count + offset, kind);
+
     return ihipLogStatus(hipSuccess);
 }
 
+hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_t count, size_t offset, hipMemcpyKind kind, hipStream_t stream)
+{
+    HIP_INIT_API(symbolName, src, count, offset, kind, stream);
+
+    if(symbolName == nullptr)
+    {
+        return ihipLogStatus(hipErrorInvalidSymbol);
+    }
+
+    hipError_t e = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+
+    hc::accelerator acc = ctx->getDevice()->_acc;
+
+    void *ptr = acc.get_symbol_address(symbolName);
+
+    if(ptr == nullptr)
+    {
+        return ihipLogStatus(hipErrorInvalidSymbol);
+    }
+
+    if (stream) {
+        try {
+            stream->locked_copyAsync(ptr, src, count + offset, kind);
+        }
+        catch (ihipException ex) {
+            e = ex._code;
+        }
+    } else {
+        e = hipErrorInvalidValue;
+    }
+
+    return ihipLogStatus(e);
+
+
+}
+
 //---
 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind)
 {
diff --git a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
index 359bc5d6db..6a71295bd2 100644
--- a/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
@@ -38,13 +38,29 @@ int main()
     int *A, *B, *Ad;
     A = new int[NUM];
     B = new int[NUM];
-    for(unsigned i=0;i<NUM;i++)
-    {
+    for(unsigned i=0;i<NUM;i++) {
         A[i] = -1*i;
         B[i] = 0;
     }
 
     hipMalloc((void**)&Ad, SIZE);
+
+    hipStream_t stream;
+    hipStreamCreate(&stream);
+    hipMemcpyToSymbolAsync("global", A, SIZE, 0, hipMemcpyHostToDevice, stream);
+    hipStreamSynchronize(stream);
+    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]);
+    }
+
+    for(unsigned i=0;i<NUM;i++) {
+        A[i] = -2*i;
+        B[i] = 0;
+    }
+
     hipMemcpyToSymbol("global", A, SIZE, 0, hipMemcpyHostToDevice);
     hipLaunchKernel(Assign, dim3(1,1,1), dim3(NUM,1,1), 0, 0, Ad);
     hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost);