From f506183dd148d97e3378eb994b2ac9c948ef0ada Mon Sep 17 00:00:00 2001 From: Craig Citro Date: Thu, 17 May 2018 16:06:21 -0700 Subject: [PATCH] Don't add branch prediction hints when compiling under nvcc. As seen in #19203, the `__builtin_expect` compiler builtin isn't recognized as a builtin in nvcc8, leading to compilation failures of the form ./tensorflow/core/kernels/gather_functor_gpu.cu.h(57): error: calling a __host__ function("__builtin_ expect") from a __global__ function("tensorflow::GatherOpKernel< ::Eigen::half, int, (bool)1> ") is n ot allowed when attempting to build TensorFlow. This change fixes things by adding an additional check for `__NVCC__`, and avoiding any branch prediction hints in that case. PiperOrigin-RevId: 197067418 --- tensorflow/core/platform/macros.h | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/platform/macros.h b/tensorflow/core/platform/macros.h index 3723968..b65eb43 100644 --- a/tensorflow/core/platform/macros.h +++ b/tensorflow/core/platform/macros.h @@ -79,7 +79,11 @@ limitations under the License. // analysis. Giving it this information can help it optimize for the // common case in the absence of better information (ie. // -fprofile-arcs). -#if TF_HAS_BUILTIN(__builtin_expect) || (defined(__GNUC__) && __GNUC__ >= 3) +// +// We need to disable this for GPU builds, though, since nvcc8 and older +// don't recognize `__builtin_expect` as a builtin, and fail compilation. +#if (!defined(__NVCC__)) && \ + (TF_HAS_BUILTIN(__builtin_expect) || (defined(__GNUC__) && __GNUC__ >= 3)) #define TF_PREDICT_FALSE(x) (__builtin_expect(x, 0)) #define TF_PREDICT_TRUE(x) (__builtin_expect(!!(x), 1)) #else -- 2.7.4