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


[ROCm/hip commit: 3c4af7c371]

projects/hip/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);