From 655ce09f679a90ecd561538227c703b42d0fc5fa Mon Sep 17 00:00:00 2001
From: Christian Sigg <csigg@google.com>
Date: Sat, 27 Jun 2020 01:35:38 -0700
Subject: [PATCH] Provide ldexp float overload for HIP, it's missing in their
 headers.

PiperOrigin-RevId: 318607361
Change-Id: Id4d3058d3960484e9f2136de7dc72b055c1a3fad
---
 tensorflow/core/kernels/cwise_ops_gpu_common.cu.h | 6 ++++++
 tensorflow/core/kernels/rnn/blas_gemm.h           | 5 +++++
 2 files changed, 11 insertions(+)

diff --git a/tensorflow/core/kernels/cwise_ops_gpu_common.cu.h b/tensorflow/core/kernels/cwise_ops_gpu_common.cu.h
index ecc58da315f..8849c3f4edd 100644
--- a/tensorflow/core/kernels/cwise_ops_gpu_common.cu.h
+++ b/tensorflow/core/kernels/cwise_ops_gpu_common.cu.h
@@ -30,6 +30,12 @@ limitations under the License.
 #include "tensorflow/core/platform/types.h"
 
 #include "tensorflow/core/platform/logging.h"
+
+#ifdef __HIP_DEVICE_COMPILE__
+// Provide ldexp float overload for HIP, it's missing in their headers.
+__device__ inline float ldexp(float x, int exp) { return ldexpf(x, exp); }
+#endif
+
 namespace tensorflow {
 namespace functor {
 
diff --git a/tensorflow/core/kernels/rnn/blas_gemm.h b/tensorflow/core/kernels/rnn/blas_gemm.h
index 126e1edef17..74f4cd2bb39 100644
--- a/tensorflow/core/kernels/rnn/blas_gemm.h
+++ b/tensorflow/core/kernels/rnn/blas_gemm.h
@@ -25,6 +25,11 @@ limitations under the License.
 #include "tensorflow/core/kernels/eigen_contraction_kernel.h"
 #endif
 
+#ifdef __HIP_DEVICE_COMPILE__
+// Provide ldexp float overload for HIP, it's missing in their headers.
+__device__ inline float ldexp(float x, int exp) { return ldexpf(x, exp); }
+#endif
+
 namespace tensorflow {
 class OpKernelContext;
 namespace functor {