Added hipHostAlloc feature for CUDA

[ROCm/clr commit: 2212d35e2d]

projects/clr/hipamd/bin/hipify

@@ -286,7 +286,12 @@ while (@ARGV) {
         $ft{'mem'} += s/\bcudaMallocHost\b/hipMallocHost/g;
         $ft{'mem'} += s/\bcudaFree\b/hipFree/g;
         $ft{'mem'} += s/\bcudaFreeHost\b/hipFreeHost/g;
-
+		$ft{'mem'} += s/\bcudaHostAlloc\b/hipHostAlloc/g;
+		$ft{'mem'} += s/\bcudaHostGetDevicePointer\b/hipHostGetDevicePointer/g;
+		$ft{'mem'} += s/\bcudaHostAllocDefault\b/hipHostAllocDefault/g;
+		$ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostAllocPortable/g;
+		$ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostAllocMapped/g;
+		$ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostAllocWriteCombined/g;
 
 
         #--------

projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h

@@ -50,6 +50,10 @@ hipMemcpyHostToHost
 } hipTextureFilterMode;*/
 #define hipFilterModePoint cudaFilterModePoint
 
+#define hipHostAllocDefault cudaHostAllocDefault
+#define hipHostAllocPortable cudaHostAllocPortable
+#define hipHostAllocMapped cudaHostAllocMapped
+#define hipHostAllocWriteCombined cudaHostAllocWriteCombined
 
 typedef cudaEvent_t hipEvent_t;
 typedef cudaStream_t hipStream_t;
@@ -115,6 +119,15 @@ inline static hipError_t hipFree(void* ptr) {
 inline static hipError_t hipMallocHost(void** ptr, size_t size) {
     return hipCUDAErrorTohipError(cudaMallocHost(ptr, size));
 }
+
+inline static hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags){
+	return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
+}
+
+inline static hipError_t hipHostGetDevicePointer(void** devPtr, void* hostPtr, unsigned int flags){
+	return hipCUDAErrorTohipError(cudaHostGetDevicePointer(devPtr, hostPtr, flags));
+}
+
 inline static hipError_t hipFreeHost(void* ptr) {
     return hipCUDAErrorTohipError(cudaFreeHost(ptr));
 }

projects/clr/hipamd/tests/src/CMakeLists.txt

@@ -126,7 +126,7 @@ make_hip_executable (hipIntrinsics hipMathFunctions.cpp hipSinglePrecisionIntrin
 make_hip_executable (hipPointerAttrib hipPointerAttrib.cpp) 
 make_hip_executable (hipMultiThreadStreams1 hipMultiThreadStreams1.cpp)
 make_hip_executable (hipMultiThreadStreams2 hipMultiThreadStreams2.cpp) 
-
+make_hip_executable (hipHostAlloc hipHostAlloc.cpp)
 make_hip_executable (hipStreamL5 hipStreamL5.cpp)
 target_link_libraries(hipMathFunctionsHost m)
 
@@ -146,7 +146,7 @@ make_test(hipEnvVarDriver " " )
 make_test(hipPointerAttrib " " )
 make_test(hipMultiThreadStreams1 " " )
 make_test(hipMultiThreadStreams2 " " )
-
+make_test(hipHostAlloc " ")
 make_test(hipMemcpy  " " )
 make_test(hipMemcpyAsync  " " )
 

projects/clr/hipamd/tests/src/hipHostAlloc.cpp

@@ -0,0 +1,42 @@
+#include"test_common.h"
+
+#define LEN 1024*1024
+#define SIZE LEN*sizeof(float)
+
+__global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd){
+int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+Cd[tx] = Ad[tx] + Bd[tx];
+}
+
+int main(){
+float *A, *B, *C;
+float *Ad, *Bd, *Cd;
+
+hipDeviceProp_t prop;
+int device;
+HIPCHECK(hipGetDevice(&device));
+HIPCHECK(hipDeviceGetProperties(&prop, device));
+if(prop.canMapHostMemory != 1){
+std::cout<<"Exiting..."<<std::endl;
+}
+HIPCHECK(hipHostAlloc((void**)&A, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
+HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
+HIPCHECK(hipHostAlloc((void**)&C, SIZE, hipHostAllocMapped));
+
+HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
+HIPCHECK(hipHostGetDevicePointer((void**)&Bd, B, 0));
+HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
+
+for(int i=0;i<LEN;i++){
+A[i] = 1.0f;
+B[i] = 2.0f;
+}
+
+dim3 dimGrid(LEN/512,1,1);
+dim3 dimBlock(512,1,1);
+
+hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);
+
+passed();
+
+}