From 02760925a991a2fa7da72a31a215ffbea803e230 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <adityaavinash1@gmail.com>
Date: Mon, 7 Mar 2016 03:42:50 -0600
Subject: [PATCH] Added hipHostRegister feature for CUDA backend and its tests

[ROCm/hip commit: 4ed0b1cb1a905f7fd190eefce6cbf0f27453d3d7]
---
 .../hip/include/nvcc_detail/hip_runtime_api.h |  8 +++++
 projects/hip/tests/src/CMakeLists.txt         |  3 +-
 projects/hip/tests/src/hipHostRegister.cpp    | 30 +++++++++++++++++++
 3 files changed, 40 insertions(+), 1 deletion(-)
 create mode 100644 projects/hip/tests/src/hipHostRegister.cpp

diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h
index 6da81c742b..9a46a773e0 100644
--- a/projects/hip/include/nvcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h
@@ -132,6 +132,14 @@ inline static hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr){
 	return hipCUDAErrorTohipError(cudaHostGetFlags(flagsPtr, hostPtr));
 }
 
+inline static hipError_t hipHostRegister(void* ptr, size_t size, unsigned int flags){
+	return hipCUDAErrorTohipError(cudaHostRegister(ptr, size, flags));
+}
+
+inline static hipError_t hipHostUnregister(void* ptr){
+	return hipCUDAErrorTohipError(cudaHostUnregister(ptr));
+}
+
 inline static hipError_t hipFreeHost(void* ptr) {
     return hipCUDAErrorTohipError(cudaFreeHost(ptr));
 }
diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt
index 0b69839e72..b4627c9ed4 100644
--- a/projects/hip/tests/src/CMakeLists.txt
+++ b/projects/hip/tests/src/CMakeLists.txt
@@ -129,6 +129,7 @@ make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp)
 make_hip_executable (hipHostAlloc hipHostAlloc.cpp)
 make_hip_executable (hipStreamL5 hipStreamL5.cpp)
 make_hip_executable (hipHostGetFlags hipHostGetFlags.cpp)
+make_hip_executable (hipHostRegister hipHostRegister.cpp)
 target_link_libraries(hipMathFunctionsHost m)
 
 make_test(hip_ballot " " )
@@ -152,7 +153,7 @@ make_test(hipMemcpy  " " )
 make_test(hipMemcpyAsync  " " )
 make_test(hipHostGetFlags " ")
 make_test(hipHcc  " " )
-
+make_test(hipHostRegister " ")
 make_test(hipStreamL5 " ")
 
 make_hipify_test(specialFunc.cu )
diff --git a/projects/hip/tests/src/hipHostRegister.cpp b/projects/hip/tests/src/hipHostRegister.cpp
new file mode 100644
index 0000000000..241d37a802
--- /dev/null
+++ b/projects/hip/tests/src/hipHostRegister.cpp
@@ -0,0 +1,30 @@
+#include"test_common.h"
+
+__global__ void Inc(hipLaunchParm lp, float *Ad){
+int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+Ad[tx] = Ad[tx] + float(1);
+}
+
+int main(){
+	float *A, *Ad;
+	const size_t size = N * sizeof(float);
+	A = (float*)malloc(size);
+	HIPCHECK(hipHostRegister(A, size, 0));
+
+	for(int i=0;i<N;i++){
+		A[i] = float(1);
+	}
+
+	HIPCHECK(hipMalloc(&Ad, size));
+
+	HIPCHECK(hipMemcpy(Ad, A, size, hipMemcpyHostToDevice));
+
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, 0, Ad);
+	HIPCHECK(hipDeviceSynchronize());
+
+	HIPCHECK(hipMemcpy(A, Ad, size, hipMemcpyDeviceToHost));
+
+	HIPASSERT(A[10] == 2.0f);
+	HIPCHECK(hipHostUnregister(A));
+	passed();
+}