diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt
--- a/clang/lib/Headers/CMakeLists.txt
+++ b/clang/lib/Headers/CMakeLists.txt
@@ -142,6 +142,7 @@
   cuda_wrappers/algorithm
   cuda_wrappers/complex
   cuda_wrappers/new
+  cuda_wrappers/ymath.h
 )
 
 set(ppc_wrapper_files
diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
--- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -61,6 +61,28 @@
 #include <algorithm>
 #include <complex>
 #include <new>
+
+// Define math functions from <ymath.h> on MSVC for the device compilation
+// only to avoid conflicts with MSVC runtime in the host compilation.
+#if defined(_MSC_VER) && defined(__HIP_DEVICE_COMPILE__)
+#if defined(__cplusplus)
+extern "C" {
+#endif // defined(__cplusplus)
+#pragma push_macro("__HOST_DEVICE__")
+#define __HOST_DEVICE__                                                        \
+  static __host__ __device__ inline __attribute__((always_inline))
+__HOST_DEVICE__ double _Cosh(double x, double y) { return cosh(x) * y; }
+__HOST_DEVICE__ float _FCosh(float x, float y) { return coshf(x) * y; }
+__HOST_DEVICE__ short _Dtest(double *p) { return fpclassify(*p); }
+__HOST_DEVICE__ short _FDtest(float *p) { return fpclassify(*p); }
+__HOST_DEVICE__ double _Sinh(double x, double y) { return sinh(x) * y; }
+__HOST_DEVICE__ float _FSinh(float x, float y) { return sinhf(x) * y; }
+#pragma pop_macro("__HOST_DEVICE__")
+#if defined(__cplusplus)
+}
+#endif // defined(__cplusplus)
+#endif // defined(_MSC_VER) && defined(__HIP_DEVICE_COMPILE__)
+
 #endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
 
 #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1
diff --git a/clang/lib/Headers/cuda_wrappers/ymath.h b/clang/lib/Headers/cuda_wrappers/ymath.h
new file mode 100644
--- /dev/null
+++ b/clang/lib/Headers/cuda_wrappers/ymath.h
@@ -0,0 +1,22 @@
+/*===---- ymath.h - HIP wrapper for <ymath.h> ------------------------------===
+ *
+ * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+ * See https://llvm.org/LICENSE.txt for license information.
+ * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#ifndef __CLANG_CUDA_WRAPPERS_YMATH_H
+#define __CLANG_CUDA_WRAPPERS_YMATH_H
+
+// Wrapper around <ymath.h> that forces its functions to be __host__
+// __device__.
+
+#pragma clang force_cuda_host_device begin
+
+#include_next <ymath.h>
+
+#pragma clang force_cuda_host_device end
+
+#endif // include guard