hipStreamWaitEvent update.

Add passed to hipStreamWaitEvent test.
Fix pointerAttrib test to work with new and old HCC versions.
Minor code touchup.

Change-Id: I139ba6ce9f6bf2b4bee89aebdec5981b4346ffc0


[ROCm/hip commit: 0923c2d261]

projects/hip/src/hip_memory.cpp

@@ -148,6 +148,8 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
         hip_status = hipErrorMemoryAllocation;
     }
 
+    //printf ("  hipMalloc allocated %p\n", *ptr);
+
     return ihipLogStatus(hip_status);
 }
 

projects/hip/src/hip_stream.cpp

@@ -74,7 +74,7 @@ hipError_t hipStreamCreate(hipStream_t *stream)
 }
 
 
-#ifndef USE_AV_COPY
+#if USE_AV_COPY==0
 //---
 /**
  * @bug This function conservatively waits for all work in the specified stream to complete.
@@ -93,7 +93,7 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
 
         bool fastWait = false;
 
-#ifdef USE_AV_COPY
+#if USE_AV_COPY
         if (stream != hipStreamNull) {
             printf ("HIP: wait locked stream\n");
             stream->locked_waitEvent(event);

projects/hip/tests/src/hipPointerAttrib.cpp

@@ -28,11 +28,12 @@ THE SOFTWARE.
 #include <vector>
 
 #ifdef __HIP_PLATFORM_HCC__
-//#include "hcc_detail/AM.h"
 #include "hc_am.hpp"
 
 #endif
 
+#define USE_AV_COPY (__hcc_workweek__ >= 16351) 
+
 size_t Nbytes = 0;
 
 //=================================================================================================
@@ -409,11 +410,21 @@ void thread_noise_generator(int iters, size_t numBuffers, Dir addDir, Dir remove
 
         if (addDir == Up) {
             for (char *p = basePtr; p<basePtr + maxSize; p+=bufferSize) {
+#if USE_AV_COPY
+                hc::AmPointerInfo info(p, p, bufferSize, acc, false, false);
+                hc::am_memtracker_add(p, info);
+#else
                 hc::am_memtracker_add(p, bufferSize, acc, false);
+#endif
             }
         } else if (addDir == Down) {
             for (char *p = basePtr+maxSize-bufferSize; p>=0; p-=bufferSize) {
+#if USE_AV_COPY
+                hc::AmPointerInfo info(p, p, bufferSize, acc, false, false);
+                hc::am_memtracker_add(p, info);
+#else
                 hc::am_memtracker_add(p, bufferSize, acc, false);
+#endif
             }
         }
 

projects/hip/tests/src/runtimeApi/stream/hipStreamWaitEvent.cpp

@@ -62,7 +62,7 @@ Streamer<T>::Streamer(size_t numElements) :
 template <typename T>
 void Streamer<T>::runAsync()
 {
-    printf ("testing: %s  numElements=%zu\n", __func__, _numElements);
+    printf ("testing: %s  numElements=%zu size=%6.2fMB\n", __func__, _numElements, _numElements * sizeof(T) / 1024.0/1024.0);
     unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, _numElements);
     hipLaunchKernel(HipTest::vectorADD, dim3(blocks), dim3(threadsPerBlock), 0, _stream, _A_d, _B_d, _C_d, _numElements);
     HIPCHECK(hipEventRecord(_event, _stream));
@@ -117,4 +117,6 @@ int main(int argc, char *argv[])
     }
 
     HIPCHECK(hipDeviceSynchronize());
+
+    passed();
 }