From d3295e61e231831028c0250377fff6b831136cd3 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Wed, 8 Aug 2018 08:55:28 -0400 Subject: [PATCH] Fix __HIP_ARCH_* not defined after including math_functions.h hcc_detail/math_functions.h used to include hcc_detail/hip_runtime.h. Removing it has caused regression in TensorFlow 1.8. Put it back for backward compatibiliity. [ROCm/clr commit: 87de95975a79d8156ad2b1a9ebe14a1ff406c725] --- .../clr/hipamd/include/hip/hcc_detail/math_functions.h | 5 +++++ .../clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp | 8 ++++++++ 2 files changed, 13 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h index d9125e1654..10967605d4 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/math_functions.h @@ -1338,3 +1338,8 @@ __HIP_OVERLOAD2(double, min) #pragma pop_macro("__HIP_OVERLOAD2") #pragma pop_macro("__DEVICE__") #pragma pop_macro("__RETURN_TYPE") + +// For backward compatibility. +// There are HIP applications e.g. TensorFlow, expecting __HIP_ARCH_* macros +// defined after including math_functions.h. +#include diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp index fd9bff1ad9..6063eee76c 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestIncludeMath.cpp @@ -45,6 +45,14 @@ __device__ __host__ inline void throw_std_bad_alloc() #endif } +// Test __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ and __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ +// is defined. Eigen HIP/hcc/Half.h __ldg depends on this. +#if !defined(__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__) || \ + !defined(__HIP_ARCH_HAS_DYNAMIC_PARALLEL__) +#error \ + "__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ or __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ not defined" +#endif + #include #include "test_common.h"