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 {