From fca3c4b260eaabab34682db02b29b75af579b066 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Fri, 25 Mar 2016 06:41:49 -0500
Subject: [PATCH] updated hipdevicesync test

[ROCm/hip commit: b371dd58ea2cf37ed8ef5179036827c241401a02]
---
 projects/hip/tests/src/CMakeLists.txt         |  3 +-
 .../tests/src/hipFuncDeviceSynchronize.cpp    | 29 ++++++++++++++-----
 2 files changed, 23 insertions(+), 9 deletions(-)

diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt
index e5adf4e4d3..acd2060647 100644
--- a/projects/hip/tests/src/CMakeLists.txt
+++ b/projects/hip/tests/src/CMakeLists.txt
@@ -154,6 +154,7 @@ make_hip_executable (hipMemoryAllocate hipMemoryAllocate.cpp)
 make_hip_executable (hipFuncSetDeviceFlags hipFuncSetDeviceFlags.cpp)
 make_hip_executable (hipFuncGetDevice hipFuncGetDevice.cpp)
 make_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp)
+make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp)
 
 make_test(hip_ballot " " )
 make_test(hip_anyall " " )
@@ -194,6 +195,6 @@ make_test(hipMemoryAllocate " ")
 make_test(hipFuncSetDeviceFlags " ")
 make_test(hipFuncGetDevice " ")
 make_test(hipFuncSetDevice " ")
-
+make_test(hipFuncDeviceSynchronize " ")
 
 make_hipify_test(specialFunc.cu )
diff --git a/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp b/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp
index 306136d1a9..bb274f47bc 100644
--- a/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp
+++ b/projects/hip/tests/src/hipFuncDeviceSynchronize.cpp
@@ -17,18 +17,23 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
+/*
+ * Test for checking the functionality of
+ * hipError_t hipDeviceSynchronize();
+*/
+
 #include"test_common.h"
 
 #define _SIZE sizeof(int)*1024*1024
-#define NUM_STREAMS 10
+#define NUM_STREAMS 2
 
 __global__ void Iter(hipLaunchParm lp, int *Ad, int num){
-int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
-if(tx == 0){
-for(int i = 0; i<num;i++){
-Ad[tx] += 1;
-}
-}
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    if(tx == 0){
+        for(int i = 0; i<num;i++){
+            Ad[tx] += 1;
+        }
+    }
 }
 
 int main(){
@@ -36,16 +41,24 @@ int main(){
     int *Ad[NUM_STREAMS];
     hipStream_t stream[NUM_STREAMS];
     for(int i=0;i<NUM_STREAMS;i++){
-        A[i] = (int*)malloc(_SIZE);
+        HIPCHECK(hipHostMalloc((void**)&A[i], _SIZE, hipHostMallocDefault));
         A[i][0] = 1;
         HIPCHECK(hipMalloc((void**)&Ad[i], _SIZE));
         HIPCHECK(hipStreamCreate(&stream[i]));
     }
     for(int i=0;i<NUM_STREAMS;i++){
         HIPCHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i]));
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
         hipLaunchKernel(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], 1<<30);
+    }
+    for(int i=0;i<NUM_STREAMS;i++){
         HIPCHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i]));
     }
+    
     HIPASSERT(1<<30 != A[NUM_STREAMS-1][0]-1);
+    HIPCHECK(hipDeviceSynchronize());
+    HIPASSERT(1<<30 == A[NUM_STREAMS-1][0]-1);
+    passed();
 }