Index: clang/docs/HIPSupport.rst
===================================================================
--- /dev/null
+++ clang/docs/HIPSupport.rst
@@ -0,0 +1,23 @@
+============
+HIP Support
+============
+
+.. contents::
+   :local:
+
+Introduction
+============
+
+This document describes HIP support in clang. More details are provided in
+`external document <https://github.com/ROCm-Developer-Tools/HIP>`_\ ,
+which are going to be added to clang documentation in the future.
+
+Standard Library Support
+========================
+
+<std::functional>
+-----------------
+
+Clang supports calling std::functioinal functors in HIP device code. However
+this is limited to `-stdlib=libc++`.
+
Index: clang/lib/Headers/CMakeLists.txt
===================================================================
--- clang/lib/Headers/CMakeLists.txt
+++ clang/lib/Headers/CMakeLists.txt
@@ -140,8 +140,14 @@
 
 set(cuda_wrapper_files
   cuda_wrappers/algorithm
+  cuda_wrappers/array
   cuda_wrappers/complex
+  cuda_wrappers/functional
   cuda_wrappers/new
+  cuda_wrappers/type_traits
+  cuda_wrappers/tuple
+  cuda_wrappers/utility
+  cuda_wrappers/__tuple
 )
 
 set(ppc_wrapper_files
Index: clang/lib/Headers/__clang_hip_runtime_wrapper.h
===================================================================
--- clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -18,6 +18,10 @@
 
 #if __HIP__
 
+#if __has_include(<__libcpp_version>)
+#define __HIP_USE_LIBCPP 1
+#endif // __has_include(<__libcpp_version>)
+
 #if !defined(__HIPCC_RTC__)
 #include <cmath>
 #include <cstdlib>
Index: clang/lib/Headers/cuda_wrappers/__tuple
===================================================================
--- /dev/null
+++ clang/lib/Headers/cuda_wrappers/__tuple
@@ -0,0 +1,24 @@
+/*===---- __tuple - CUDA/HIP wrapper for <__tuple> -------------------------===
+ *
+ * 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___TUPLE
+#define __CLANG_CUDA_WRAPPERS___TUPLE
+
+#if __HIP_USE_LIBCPP
+#pragma push_macro("_LIBCPP_NO_EXCEPTIONS")
+#define _LIBCPP_NO_EXCEPTIONS
+#pragma clang force_cuda_host_device begin
+#include_next <__tuple>
+#pragma clang force_cuda_host_device end
+#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS")
+#else
+#include_next <__tuple>
+#endif // __HIP_USE_LIBCPP
+
+#endif // __CLANG_CUDA_WRAPPERS___TUPLE
Index: clang/lib/Headers/cuda_wrappers/array
===================================================================
--- /dev/null
+++ clang/lib/Headers/cuda_wrappers/array
@@ -0,0 +1,24 @@
+/*===---- array - CUDA/HIP wrapper for <array> -----------------------------===
+ *
+ * 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_ARRAY
+#define __CLANG_CUDA_WRAPPERS_ARRAY
+
+#if __HIP_USE_LIBCPP
+#pragma push_macro("_LIBCPP_NO_EXCEPTIONS")
+#define _LIBCPP_NO_EXCEPTIONS
+#pragma clang force_cuda_host_device begin
+#include_next <array>
+#pragma clang force_cuda_host_device end
+#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS")
+#else
+#include_next <array>
+#endif // __HIP_USE_LIBCPP
+
+#endif // __CLANG_CUDA_WRAPPERS_ARRAY
Index: clang/lib/Headers/cuda_wrappers/functional
===================================================================
--- /dev/null
+++ clang/lib/Headers/cuda_wrappers/functional
@@ -0,0 +1,34 @@
+/*===---- functional - CUDA/HIP wrapper for <functional> -------------------===
+ *
+ * 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_FUNCTIONAL
+#define __CLANG_CUDA_WRAPPERS_FUNCTIONAL
+
+#if __HIP_USE_LIBCPP
+#pragma push_macro("_LIBCPP_NO_EXCEPTIONS")
+#define _LIBCPP_NO_EXCEPTIONS
+
+extern __device__ void abort() __attribute__ ((__noreturn__));
+namespace std {
+namespace __1 {
+  inline __device__ void abort() {
+    return ::abort();
+  }
+}
+}
+
+#pragma clang force_cuda_host_device begin
+#include_next <functional>
+#pragma clang force_cuda_host_device end
+#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS")
+#else
+#include_next <functional>
+#endif // __HIP_USE_LIBCPP
+
+#endif // __CLANG_CUDA_WRAPPERS_FUNCTIONAL
Index: clang/lib/Headers/cuda_wrappers/tuple
===================================================================
--- /dev/null
+++ clang/lib/Headers/cuda_wrappers/tuple
@@ -0,0 +1,24 @@
+/*===---- tuple - CUDA/HIP wrapper for <tuple> -----------------------------===
+ *
+ * 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_TUPLE
+#define __CLANG_CUDA_WRAPPERS_TUPLE
+
+#if __HIP_USE_LIBCPP
+#pragma push_macro("_LIBCPP_NO_EXCEPTIONS")
+#define _LIBCPP_NO_EXCEPTIONS
+#pragma clang force_cuda_host_device begin
+#include_next <tuple>
+#pragma clang force_cuda_host_device end
+#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS")
+#else
+#include_next <tuple>
+#endif // __HIP_USE_LIBCPP
+
+#endif // __CLANG_CUDA_WRAPPERS_TUPLE
Index: clang/lib/Headers/cuda_wrappers/type_traits
===================================================================
--- /dev/null
+++ clang/lib/Headers/cuda_wrappers/type_traits
@@ -0,0 +1,24 @@
+/*===---- type_traits - CUDA/HIP wrapper for <type_traits> -----------------===
+ *
+ * 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_TYPE_TRAITS
+#define __CLANG_CUDA_WRAPPERS_TYPE_TRAITS
+
+#if __HIP_USE_LIBCPP
+#pragma push_macro("_LIBCPP_NO_EXCEPTIONS")
+#define _LIBCPP_NO_EXCEPTIONS
+#pragma clang force_cuda_host_device begin
+#include_next <type_traits>
+#pragma clang force_cuda_host_device end
+#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS")
+#else
+#include_next <type_traits>
+#endif // __HIP_USE_LIBCPP
+
+#endif // __CLANG_CUDA_WRAPPERS_TYPE_TRAITS
Index: clang/lib/Headers/cuda_wrappers/utility
===================================================================
--- /dev/null
+++ clang/lib/Headers/cuda_wrappers/utility
@@ -0,0 +1,24 @@
+/*===---- utility - CUDA/HIP wrapper for <utility> -------------------------===
+ *
+ * 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_UTILITY
+#define __CLANG_CUDA_WRAPPERS_UTILITY
+
+#if __HIP_USE_LIBCPP
+#pragma push_macro("_LIBCPP_NO_EXCEPTIONS")
+#define _LIBCPP_NO_EXCEPTIONS
+#pragma clang force_cuda_host_device begin
+#include_next <utility>
+#pragma clang force_cuda_host_device end
+#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS")
+#else
+#include_next <utility>
+#endif // __HIP_USE_LIBCPP
+
+#endif // __CLANG_CUDA_WRAPPERS_UTILITY