added more changes to memcpytosymbol

1. Refactored code to use HCC internal APIs rather than HCC copy APIs
2. Added hipMemcpyToSymbolAsync
3. Added test for hipMemcpyToSymbolAsync
4. Added new error hipErrorInvalidSymbol

Change-Id: I0e359b2d0ff5d682bbccdf9c2923e16b35e39497

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.
  *

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.

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)
 {

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