From 9c2a33b5c9ac314077f00b89116fd8911082c5d9 Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
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/hip commit: b8a0c14171a1fd6744ea4d82e065c62524237fea]
---
 projects/hip/include/hip/hcc_detail/math_functions.h    | 5 +++++
 projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp | 8 ++++++++
 2 files changed, 13 insertions(+)

diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h
index d9125e1654..10967605d4 100644
--- a/projects/hip/include/hip/hcc_detail/math_functions.h
+++ b/projects/hip/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 <hip/hcc_detail/hip_runtime.h>
diff --git a/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp b/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp
index fd9bff1ad9..6063eee76c 100644
--- a/projects/hip/tests/src/deviceLib/hipTestIncludeMath.cpp
+++ b/projects/hip/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 <hip/hip_runtime.h>
 #include "test_common.h"