From 8724273f28c18cf246c1dcaa7ea0a8431c556638 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 10 Nov 2016 10:53:10 -0600
Subject: [PATCH] Doc change only - add comments to test.

Change-Id: Ie42087cf3c78e49337b18bb71f3f0e1e7950ee1b
---
 tests/README.md                        | 12 ++++++------
 tests/src/hipFuncDeviceSynchronize.cpp |  7 ++++++-
 tests/src/hipHostGetFlags.cpp          |  2 +-
 3 files changed, 13 insertions(+), 8 deletions(-)

diff --git a/tests/README.md b/tests/README.md
index 56bb4e7edd..223bd149dc 100644
--- a/tests/README.md
+++ b/tests/README.md
@@ -53,11 +53,11 @@ ctest -R Memcpy
 
 ### If a test fails - how to debug a test
 
-Extract the commandline from the testing log:
+Find the test and commandline that fail:
 
 (From the test build directory, perhaps hip/tests/build)
-$ grep -A3 -m2  hipMemcpy-size  Testing/Temporary/LastTest.log
-36/47 Testing: hipMemcpy-size
-36/47 Test: hipMemcpy-size
-Command: "/home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory/hipMemcpy" "--tests" "0x6"
-Directory: /home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory
+grep -IR hipMemcpy-modes -IR ../tests/
+../tests/src/runtimeApi/memory/hipMemcpy.cpp: * RUN_NAMED: %t hipMemcpy-modes --tests 0x1
+
+
+
diff --git a/tests/src/hipFuncDeviceSynchronize.cpp b/tests/src/hipFuncDeviceSynchronize.cpp
index 6d05253536..930bc37b8b 100644
--- a/tests/src/hipFuncDeviceSynchronize.cpp
+++ b/tests/src/hipFuncDeviceSynchronize.cpp
@@ -35,6 +35,7 @@ THE SOFTWARE.
 
 __global__ void Iter(hipLaunchParm lp, int *Ad, int num){
     int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    // Kernel loop designed to execute very slowly... ... ...   so we can test timing-related behavior below
     if(tx == 0){
         for(int i = 0; i<num;i++){
             Ad[tx] += 1;
@@ -61,7 +62,11 @@ int main(){
     for(int i=0;i<NUM_STREAMS;i++){
         HIPCHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i]));
     }
-    
+
+   
+   // This first check but relies on the kernel running for so long that the D2H async memcopy has not started yet.   
+   // This will be true in an optimal asynchronous implementation.  
+   // Conservative implementations which synchronize the hipMemcpyAsync will fail, ie if HIP_LAUNCH_BLOCKING=true
     HIPASSERT(1<<30 != A[NUM_STREAMS-1][0]-1);
     HIPCHECK(hipDeviceSynchronize());
     HIPASSERT(1<<30 == A[NUM_STREAMS-1][0]-1);
diff --git a/tests/src/hipHostGetFlags.cpp b/tests/src/hipHostGetFlags.cpp
index 8664095422..a58ba7aa43 100644
--- a/tests/src/hipHostGetFlags.cpp
+++ b/tests/src/hipHostGetFlags.cpp
@@ -76,7 +76,7 @@ dim3 dimBlock(512,1,1);
 
 hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);
 
-HIPCHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
+HIPCHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));  // Note this really HostToHost not DeviceToHost, since memory is mapped...
 HIPASSERT(C[10] == 3.0f);
 HIPASSERT(flagA == FlagA);
 HIPASSERT(flagB == FlagB);