From ff24c82f34db1a405679cb6845ed51eeb977b760 Mon Sep 17 00:00:00 2001
From: Aaron En Ye Shi <enye.shi@gmail.com>
Date: Mon, 19 Oct 2020 17:59:34 +0000
Subject: [PATCH] Fix header for OpenMP AMDGCN support

Add ifndef for __OPENMP_AMDGCN__ target, which shares
these header files with HIP.

Change-Id: I720693cfbcfb7836cd2f361d48530fb55ee3557a


[ROCm/clr commit: 01bd16550726229f737334318ec3c13ba8b0090c]
---
 projects/clr/hipamd/include/hip/hcc_detail/device_functions.h | 3 ++-
 projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h      | 3 +++
 2 files changed, 5 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
index be5ef9d159..82a4c446a3 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h
@@ -1410,6 +1410,7 @@ static inline __device__ void* __hip_hc_memset(void* dst, unsigned char val, siz
 
     return dst;
 }
+#ifndef __OPENMP_AMDGCN__
 static inline __device__ void* memcpy(void* dst, const void* src, size_t size) {
     return __hip_hc_memcpy(dst, src, size);
 }
@@ -1418,5 +1419,5 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) {
     unsigned char val8 = static_cast<unsigned char>(val);
     return __hip_hc_memset(ptr, val8, size);
 }
-
+#endif // !__OPENMP_AMDGCN__
 #endif
diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
index a2d1e4c671..5411bb34b6 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
@@ -312,6 +312,8 @@ static constexpr Coordinates<hip_impl::WorkitemId> threadIdx{};
 #define hipGridDim_z (hc_get_num_groups(2))
 
 #endif // defined __HCC__
+
+#ifndef __OPENMP_AMDGCN__
 #if __HCC_OR_HIP_CLANG__
 #if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
 #if __HIP_ENABLE_DEVICE_MALLOC__
@@ -325,6 +327,7 @@ static inline __device__ void* free(void* ptr) { __builtin_trap(); return nullpt
 #endif
 #endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
 #endif //__HCC_OR_HIP_CLANG__
+#endif // !__OPENMP_AMDGCN__
 
 #ifdef __HCC__