diff --git a/projects/clr/hipamd/bin/hipify b/projects/clr/hipamd/bin/hipify index c49598f4ef..811f6ac036 100755 --- a/projects/clr/hipamd/bin/hipify +++ b/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; #-------- 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 1413d735f6..1d99ba28f2 100644 --- a/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h +++ b/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)); } diff --git a/projects/clr/hipamd/tests/src/CMakeLists.txt b/projects/clr/hipamd/tests/src/CMakeLists.txt index 44becb0ed3..4454ac73fb 100644 --- a/projects/clr/hipamd/tests/src/CMakeLists.txt +++ b/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 " " ) diff --git a/projects/clr/hipamd/tests/src/hipHostAlloc.cpp b/projects/clr/hipamd/tests/src/hipHostAlloc.cpp new file mode 100644 index 0000000000..7ad6f4b718 --- /dev/null +++ b/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..."<