From c2fd536c22f0eae437aa3aff9cd4154c66623ff3 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sat, 19 Mar 2016 04:20:10 -0500
Subject: [PATCH] fix nvcc compiler

- MallocHost and FreeHost deprecation.
- Change tests to call new hipHost* equivs.
- Add missing StreamSynchronize.


[ROCm/clr commit: 6984f24d3d2c569501f58a39de197aedc9e50ac8]
---
 .../clr/hipamd/include/nvcc_detail/hip_runtime_api.h  | 11 +++++++++--
 projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp      |  6 +++---
 projects/clr/hipamd/tests/src/hipMemcpy_simple.cpp    |  4 ++--
 projects/clr/hipamd/tests/src/hipPointerAttrib.cpp    |  4 ++--
 projects/clr/hipamd/tests/src/test_common.h           |  6 +++---
 5 files changed, 19 insertions(+), 12 deletions(-)

diff --git a/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h
index f0dee89808..7c37e7f71d 100644
--- a/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h
+++ b/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h
@@ -119,7 +119,9 @@ inline static hipError_t hipFree(void* ptr) {
     return hipCUDAErrorTohipError(cudaFree(ptr));
 }
 
-inline static hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostAlloc instead"))) {
+inline static hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostAlloc instead")));
+
+inline static hipError_t hipMallocHost(void** ptr, size_t size) {
     return hipCUDAErrorTohipError(cudaMallocHost(ptr, size));
 }
 
@@ -143,7 +145,8 @@ inline static hipError_t hipHostUnregister(void* ptr){
 	return hipCUDAErrorTohipError(cudaHostUnregister(ptr));
 }
 
-inline static hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree instead")))  {
+inline static hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree instead")));
+inline static hipError_t hipFreeHost(void* ptr) {
     return hipCUDAErrorTohipError(cudaFreeHost(ptr));
 }
 
@@ -374,6 +377,10 @@ inline static hipError_t hipStreamCreate(hipStream_t *stream)
     return hipCUDAErrorTohipError(cudaStreamCreate(stream));
 }
 
+inline static hipError_t hipStreamSynchronize(hipStream_t stream) 
+{
+    return hipCUDAErrorTohipError(cudaStreamSynchronize(stream));
+}
 
 inline static hipError_t hipStreamDestroy(hipStream_t stream) 
 {
diff --git a/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp b/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp
index bf50f85bbb..74ccb80f20 100644
--- a/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp
+++ b/projects/clr/hipamd/tests/src/hipMemcpyAsync.cpp
@@ -159,7 +159,7 @@ void test_pingpong(hipStream_t stream, size_t numElements, int numInflight, int
     }
 
 
-    HIPCHECK(hipFreeHost(A_h));
+    HIPCHECK(hipHostFree(A_h));
     HIPCHECK(hipFree(A_d));
 }
 
@@ -216,8 +216,8 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
     }
 
 
-    HIPCHECK(hipFreeHost(A_h1));
-    HIPCHECK(hipFreeHost(A_h2));
+    HIPCHECK(hipHostFree(A_h1));
+    HIPCHECK(hipHostFree(A_h2));
     HIPCHECK(hipFree(A_d));
 }
 
diff --git a/projects/clr/hipamd/tests/src/hipMemcpy_simple.cpp b/projects/clr/hipamd/tests/src/hipMemcpy_simple.cpp
index 7d85e2ad96..2dc8399dcb 100644
--- a/projects/clr/hipamd/tests/src/hipMemcpy_simple.cpp
+++ b/projects/clr/hipamd/tests/src/hipMemcpy_simple.cpp
@@ -97,8 +97,8 @@ void simpleTest2(size_t numElements, bool usePinnedHost)
 
     HIPCHECK(hipFree(A_d));
     if (usePinnedHost) {
-        HIPCHECK(hipFreeHost(A_h1));
-        HIPCHECK(hipFreeHost(A_h2));
+        HIPCHECK(hipHostFree(A_h1));
+        HIPCHECK(hipHostFree(A_h2));
     } else {
         free(A_h1);
         free(A_h2);
diff --git a/projects/clr/hipamd/tests/src/hipPointerAttrib.cpp b/projects/clr/hipamd/tests/src/hipPointerAttrib.cpp
index 654b833be9..9fc879b933 100644
--- a/projects/clr/hipamd/tests/src/hipPointerAttrib.cpp
+++ b/projects/clr/hipamd/tests/src/hipPointerAttrib.cpp
@@ -178,7 +178,7 @@ void testSimple()
     HIPASSERT((char*)attribs.hostPointer+Nbytes/2 == (char*)attribs2.hostPointer);
 
 
-    hipFreeHost(A_Pinned_h);
+    hipHostFree(A_Pinned_h);
     e = hipPointerGetAttributes(&attribs, A_Pinned_h);
     HIPASSERT(e == hipErrorUnknown); // Just freed the pointer, this should return an error.
     printf("getAttr:%-20s err=%d (%s), neg-test expected\n", "A_d+NBytes", e, hipGetErrorString(e));
@@ -322,7 +322,7 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
         if (ref._attrib.memoryType == hipMemoryTypeDevice) {
             hipFree(ref._pointer);
         } else {
-            hipFreeHost(ref._pointer);
+            hipHostFree(ref._pointer);
         }
 
     }
diff --git a/projects/clr/hipamd/tests/src/test_common.h b/projects/clr/hipamd/tests/src/test_common.h
index 4afc8fa4bd..0698efe1c7 100644
--- a/projects/clr/hipamd/tests/src/test_common.h
+++ b/projects/clr/hipamd/tests/src/test_common.h
@@ -193,13 +193,13 @@ void freeArrays(T *A_d, T *B_d, T *C_d,
 
     if (usePinnedHost) {
         if (A_h) {
-            HIPCHECK (hipFreeHost(A_h));
+            HIPCHECK (hipHostFree(A_h));
         }
         if (B_h) {
-            HIPCHECK (hipFreeHost(B_h));
+            HIPCHECK (hipHostFree(B_h));
         }
         if (C_h) {
-            HIPCHECK (hipFreeHost(C_h));
+            HIPCHECK (hipHostFree(C_h));
         }
     } else {
         if (A_h) {