From a9a6b5e8a68b11d5bd4ffc16c17c7f57d6b024ac Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Fri, 24 Mar 2017 10:39:11 +0530
Subject: [PATCH] Update hipTestDeviceSymbol sample

Change-Id: If5ba99c60cd30c4491ca3a4856764224163d3ddf


[ROCm/clr commit: c33c58d6871135e0a324d5123647f78d24554d3b]
---
 .../src/deviceLib/hipTestDeviceSymbol.cpp     | 23 ++++++++++++++++++-
 1 file changed, 22 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
index 6ffaedf659..c2ffb5ce7d 100644
--- a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
+++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp
@@ -93,8 +93,29 @@ int main()
     hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost);
     for(unsigned i=0;i<NUM;i++) {
         assert(A[i] == B[i]);
-//        assert(A[i] == C[i]);
+        assert(A[i] == C[i]);
     }
 
+    for(unsigned i=0;i<NUM;i++) {
+        A[i] = -3*i;
+        B[i] = 0;
+    }
+
+    hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), 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);
+    hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream);
+    hipStreamSynchronize(stream);
+    for(unsigned i=0;i<NUM;i++) {
+        assert(A[i] == B[i]);
+        assert(A[i] == C[i]);
+    }
+    hipHostFree(Am);
+    hipHostFree(Cm);
+    hipFree(Ad);
+    delete[] A;
+    delete[] B;
+    delete[] C;
     passed();
 }