From 3301deb07cde3f6c228e5c022edaa1d607bd3f5d Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sun, 17 Apr 2016 07:54:39 -0500
Subject: [PATCH] test update

---
 hipamd/CMakeLists.txt                     |  1 +
 hipamd/tests/src/hipPeerToPeer_simple.cpp | 71 ++++++++++++++---------
 2 files changed, 46 insertions(+), 26 deletions(-)

diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt
index d498038c24..25459766ba 100644
--- a/hipamd/CMakeLists.txt
+++ b/hipamd/CMakeLists.txt
@@ -92,6 +92,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc")
     set(CMAKE_C_COMPILER   "${HCC_HOME}/bin/hcc")
 
+    # Set HIP_HCC so we know this is HIP compile, some files are shared with HCC (staging_buffer).
     set(CMAKE_CXX_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -stdlib=libc++ -DHIP_HCC")
     set(CMAKE_C_FLAGS   " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -stdlib=libc++ -DHIP_HCC")
 
diff --git a/hipamd/tests/src/hipPeerToPeer_simple.cpp b/hipamd/tests/src/hipPeerToPeer_simple.cpp
index d756ee3e18..e5b61714b2 100644
--- a/hipamd/tests/src/hipPeerToPeer_simple.cpp
+++ b/hipamd/tests/src/hipPeerToPeer_simple.cpp
@@ -99,44 +99,61 @@ void enablePeerFirst()
         assert(canAccessPeer);
 
         HIPCHECK(hipSetDevice(g_peerDevice));
-            HIPCHECK(hipDeviceEnablePeerAccess(g_currentDevice, 0));
-        }
+        HIPCHECK(hipDeviceEnablePeerAccess(g_currentDevice, 0));
+    }
 
-        size_t Nbytes = N*sizeof(char);
+    size_t Nbytes = N*sizeof(char);
 
-        char *A_d0, *A_d1;
-        char *A_h;
+    char *A_d0, *A_d1;
+    char *A_h;
 
-        A_h = (char*)malloc(Nbytes);
+    A_h = (char*)malloc(Nbytes);
 
-        // allocate and initialize memory on device0
+    // allocate and initialize memory on device0
+    HIPCHECK (hipSetDevice(g_currentDevice));
+    HIPCHECK (hipMalloc(&A_d0, Nbytes) );
+    HIPCHECK (hipMemset(A_d0, memsetval, Nbytes) ); 
+    // TODO - remove me:
+    HIPCHECK (hipDeviceSynchronize());
+
+    // allocate and initialize memory on peer device
+    HIPCHECK (hipSetDevice(g_peerDevice));
+    HIPCHECK (hipMalloc(&A_d1, Nbytes) );
+    HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) ); 
+
+    // TODO - remove me:
+    HIPCHECK (hipDeviceSynchronize());
+
+
+    // Device0 push to device1, using P2P:
+    HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
+    HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // This is P2P copy.
+
+    // TODO - remove me:
+    if (1) {
         HIPCHECK (hipSetDevice(g_currentDevice));
-        HIPCHECK (hipMalloc(&A_d0, Nbytes) );
-        HIPCHECK (hipMemset(A_d0, memsetval, Nbytes) ); 
-
-        // allocate and initialize memory on peer device
+        HIPCHECK (hipDeviceSynchronize());
         HIPCHECK (hipSetDevice(g_peerDevice));
-        HIPCHECK (hipMalloc(&A_d1, Nbytes) );
-        HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) ); 
+        HIPCHECK (hipDeviceSynchronize());
+    }
 
+    // Copy data back to host:
+    HIPCHECK (hipSetDevice(g_peerDevice));
+    HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));
 
+    // TODO - remove me:
+    HIPCHECK (hipDeviceSynchronize());
 
-        // Device0 push to device1, using P2P:
-        HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
-        HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault));
-
-        // Copy data back to host:
-        HIPCHECK (hipSetDevice(g_peerDevice));
-        HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));
-
-        // Check host data:
-        for (int i=0; i<N; i++) {
-            if (A_h[i] != memsetval) {
-                failed("mismatch at index:%d computed:0x%02x, golden memsetval:0x%02x\n", i, (int)A_h[i], (int)memsetval);
-            }
+    // Check host data:
+    for (int i=0; i<N; i++) {
+        if (A_h[i] != memsetval) {
+            failed("mismatch at index:%d computed:0x%02x, golden memsetval:0x%02x\n", i, (int)A_h[i], (int)memsetval);
         }
     }
 
+    printf ("==done: %s\n\n", __func__);
+}
+
 
 //---
 // Test which allocated memory first, then enables peer2peer.
@@ -199,6 +216,7 @@ void allocMemoryFirst()
             failed("mismatch at index:%d computed:0x%02x, golden memsetval:0x%02x\n", i, (int)A_h[i], (int)memsetval);
         }
     }
+    printf ("==done: %s\n\n", __func__);
 }
 
 void simpleNegative()
@@ -233,6 +251,7 @@ void simpleNegative()
 
     
     // More tests here:
+    printf ("==done: %s\n\n", __func__);
 }