[AMDGPU][OPENMP] OpenMP AMDGCN Header Support

[ashi1]

Summary

Both nvptx and amdgcn will use the autoheader
openmp_wrappers/__clang_openmp_device_functions.h.
Use the clang headers for math functions. Still need a
libm bitcode file to build the FORTRAN device. The code
in __clang_hip_math.h can beu sed to build a libm
device library for FORTRAN. Add log2 and pow operation
with int arg.

[JonChesterfield]

clang/lib/Headers/__clang_cuda_math.h

29 ↗

(On Diff #298715)

I recognise this from the fortran libs build, but I'm not sure it's applicable to trunk as we don't have any amdgpu/openmp fortran support in trunk.

clang/lib/Headers/__clang_hip_cmath.h

98

What are these _i suffix functions for? They aren't in libm. Writing log2(42) usually works without extra overloads

clang/lib/Headers/__clang_hip_libdevice_declares.h

60

I don't think the return type of this library function can depend on whether this header is compiled as c++/openmp or not. ocml is very macro themed, but it looks to me like the declaration there is an instance of extern __attribute__((const)) int OCML_MANGLE_F32(N)(float); which suggests 'int' was right.

clang/lib/Headers/__clang_hip_math.h

273–277

This is a bit weird. __tmp is a stack variable, which I thought would be in address_space(5) already. Declaration is extern float OCML_MANGLE_F32(frexp)(float, __private int *);.

So I guess I can see why either of these branches would work, but don't understand why either of them don't.

442

I don't recognise the symbol powii (or powi) from libm

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

70

Maybe use this define instead of attribute((address_space(5))) throughout the above?

clang/lib/Headers/openmp_wrappers/math.h

30

Why is <new> needed here?

[ashi1]

Updated based on review comments.

clang/lib/Headers/__clang_hip_cmath.h

98

okay, I will remove these new functions, and a future patch can update these log2 and pow functions.

clang/lib/Headers/__clang_hip_math.h

273–277

I will update this to use the static version for HIP too. That path works correctly, and the HIP path was dropping the addrspace qualifier.

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

70

I will update the HIP and OpenMP to share the same usage of this address space (rather than the 2 paths above). Since HIP doesn't define __private, I may skip this change.

[ashi1]

clang/lib/Headers/__clang_hip_math.h

273–277

Using static seems to cause a HIP error that is not observed on OpenMP.
[2020-10-20T16:18:06.080Z] /opt/rocm-3.9.0-3703/llvm/lib/clang/12.0.0/include/__clang_hip_math.h:280:48: error: within a device function, only shared variables or const variables without device memory qualifier may be marked 'static'
[2020-10-20T16:18:06.080Z] static attribute((address_space(5))) int __tmp;
[2020-10-20T16:18:06.080Z] ^

Should we change this patch to using #ifdef, or we can wait until this patch lands (it will fix this):
https://reviews.llvm.org/D88345

[JonChesterfield]

Patch looks broadly reasonable to me. @jdoerfert does this align with what you want from these headers?

clang/lib/Headers/__clang_hip_math.h

273–277

I guess we ifdef until D88345 lands. Have you opened a task for HIP dropping the addrspace qualifier?

[ashi1]

Added back the #ifdef for few functions that need static with private for OpenMP.

[jdoerfert]

First off, this is the right direction, thanks for working on this!

Then, we need to split this. There are various things happening here and we can easily do them one by one.
Please reach out if you don't know where to split but basically try to make the patches "as small as possible" such that you write a test against them.
So for example, math and complex stuff are separate. New functions are separate, changing the attributes is separate, ...

I added a bunch of comments and tried to be exhaustive, though once we split it in the separate components I will do another review.

Finally, where are the tests? You only added the test/Inputs stuff.

clang/lib/Headers/CMakeLists.txt

163

I'm not sold on the hip/hip... idea here but ok.
I guess I would have put the stuff in __clang_openmp_device_functions instead.

Now that I write this, where is hip_runtime included?

clang/lib/Headers/__clang_cuda_complex_builtins.h

103

What is the difference to the block above? We should also revisit the need for this in the first place. We have template overload support through begin/end declare variant in clang now. Best case, none of this is needed and we simply call the C++ or C standard library functions here. That would be great because they have to resolve to the right nv and ocml functions anyway so why not call them ;)

298

The max stuff you can commit ahead of time. LGTM on that NFC part.

clang/lib/Headers/__clang_hip_cmath.h

31

Nit: Hm, the comment talks about static but we use that either way.

clang/lib/Headers/__clang_hip_libdevice_declares.h

308

This I don't get, why is the signature different between C++ and C? That seems wrong.

clang/lib/Headers/__clang_hip_math.h

39

What is the reason you need constexpr in the first place?

51

So I fought this problem long and hard and in the end realized that this way cannot work.
If you fix the type you will always break some systems. The OpenMP declare variant provides therefore an extension that allows to "overload" on the return type. Take a look at
clang/lib/Headers/__clang_cuda_cmath.h:72 which was introduced D85879.

I don't require you to change this (for this commit or in general) but I would urge you to consider it or we will most certainly get bug reports for HIP and OpenMP offload for AMD at some point.

65

This should probably go to the declaration of the macros not their first use.

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

48

Neither the ifdef here nor the one above are needed.

51

Please elaborate on this and clean it up if we *really* need it. Why is "optimized macro definitions" bad?

70

I think this is problematic but let's talk about this once the stuff is split off.

77

The extern C looks dubious to me. What in here needs it and why isn't it in the included file then?

clang/lib/Headers/openmp_wrappers/cmath

22

If hip cmath needs new, include it there. cmath can arguably rely on math.h but new seems to be hip specific.

79

We don't need this ifdef. Please check also other ifdefs surrounding begin/end declare variant.

90

Same as the ifdef above, not needed. match_any not needed here.

clang/lib/Headers/openmp_wrappers/complex

22

we forgot to undef CUDA here...

25

This is the wrong spot to include libdevice_declares.

48

This I do not understand.

clang/lib/Headers/openmp_wrappers/hip/hip_runtime.h

24

These pragmas are useless and, FWIW, you don't need the extension if you have a single arch here.

undefine the OPENMP_AMDGCN or make sure it is always defined, see also below.

clang/lib/Headers/openmp_wrappers/math.h

63

The extension is not necessary here.

It is dubious that we have to define things here without undefining them.
Please check if these things are set before or not. If not, define and undefine them as needed, I don't want to clutter the user code differently when math is included vs. when it's not.

[ashi1]

In D89584#2345823, @jdoerfert wrote:

First off, this is the right direction, thanks for working on this!

Then, we need to split this. There are various things happening here and we can easily do them one by one.
Please reach out if you don't know where to split but basically try to make the patches "as small as possible" such that you write a test against them.
So for example, math and complex stuff are separate. New functions are separate, changing the attributes is separate, ...

I added a bunch of comments and tried to be exhaustive, though once we split it in the separate components I will do another review.

Finally, where are the tests? You only added the test/Inputs stuff.

Hi @jdoerfert , thank you for the review. I will discuss with the team, and start opening smaller patches one at a time.

[ashi1]

Dropping this patch in favour of smaller patches.

clang/lib/Driver/ToolChains/Clang.cpp

void Clang::AddPreprocessingOptions(Compilation &C, const JobAction &JA,
   if (JA.isOffloading(Action::OFK_Cuda))
     getToolChain().AddCudaIncludeArgs(Args, CmdArgs);
   if (JA.isOffloading(Action::OFK_HIP))
     getToolChain().AddHIPIncludeArgs(Args, CmdArgs);
 
   // If we are offloading to a target via OpenMP we need to include the
   // openmp_wrappers folder which contains alternative system headers.
   if (JA.isDeviceOffloading(Action::OFK_OpenMP) &&
-      getToolChain().getTriple().isNVPTX()){
+      (getToolChain().getTriple().isNVPTX() ||
+       getToolChain().getTriple().isAMDGCN())) {
     if (!Args.hasArg(options::OPT_nobuiltininc)) {
       // Add openmp_wrappers/* to our system include path.  This lets us wrap
       // standard library headers.
       SmallString<128> P(D.ResourceDir);
       llvm::sys::path::append(P, "include");
       llvm::sys::path::append(P, "openmp_wrappers");
       CmdArgs.push_back("-internal-isystem");
       CmdArgs.push_back(Args.MakeArgString(P));

clang/lib/Headers/CMakeLists.txt

 set(openmp_wrapper_files
   openmp_wrappers/math.h
   openmp_wrappers/cmath
   openmp_wrappers/complex.h
   openmp_wrappers/complex
   openmp_wrappers/__clang_openmp_device_functions.h
   openmp_wrappers/complex_cmath.h
   openmp_wrappers/new
+  openmp_wrappers/hip/hip_runtime.h
 )
 
 set(output_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include)
 set(out_files)
 set(generated_files)
 
 function(copy_header_to_output_dir src_dir file)
   set(src ${src_dir}/${file})

clang/lib/Headers/__clang_cuda_complex_builtins.h

 #else
 #define __DEVICE__ __device__ inline
 #endif
 
 // To make the algorithms available for C and C++ in CUDA and OpenMP we select
 // different but equivalent function versions. TODO: For OpenMP we currently
 // select the native builtins as the overload support for templates is lacking.
 #if !defined(_OPENMP)
+#ifdef __NVPTX__
 #define _ISNANd std::isnan
 #define _ISNANf std::isnan
 #define _ISINFd std::isinf
 #define _ISINFf std::isinf
 #define _ISFINITEd std::isfinite
 #define _ISFINITEf std::isfinite
 #define _COPYSIGNd std::copysign
 #define _COPYSIGNf std::copysign
 #define _SCALBNd std::scalbn
 #define _SCALBNf std::scalbn
 #define _ABSd std::abs
 #define _ABSf std::abs
 #define _LOGBd std::logb
 #define _LOGBf std::logb
+#define _fmaxd max
+#define _fmaxf max
 #else
+#include <__clang_hip_libdevice_declares.h>
+#define _ISNANd __ocml_isnan_f64
+#define _ISNANf __ocml_isnan_f32
+#define _ISINFd __ocml_isinf_f64
+#define _ISINFf __ocml_isinf_f32
+#define _ISFINITEd __ocml_isfinite_f64
+#define _ISFINITEf __ocml_isfinite_f32
+#define _COPYSIGNd __ocml_copysign_f64
+#define _COPYSIGNf __ocml_copysign_f32
+#define _SCALBNd __ocml_scalbn_f64
+#define _SCALBNf __ocml_scalbn_f32
+#define _ABSd __ocml_fabs_f64
+#define _ABSf __ocml_fabs_f32
+#define _LOGBd __ocml_logb_f64
+#define _LOGBf __ocml_logb_f32
+#define _fmaxd __ocml_fmax_f64
+#define _fmaxf __ocml_fmax_f32
+#endif
+#else  // OPENMP
+#ifdef __NVPTX__
 #define _ISNANd __nv_isnand
 #define _ISNANf __nv_isnanf
 #define _ISINFd __nv_isinfd
 #define _ISINFf __nv_isinff
 #define _ISFINITEd __nv_isfinited
 #define _ISFINITEf __nv_finitef
 #define _COPYSIGNd __nv_copysign
 #define _COPYSIGNf __nv_copysignf
 #define _SCALBNd __nv_scalbn
 #define _SCALBNf __nv_scalbnf
 #define _ABSd __nv_fabs
 #define _ABSf __nv_fabsf
 #define _LOGBd __nv_logb
 #define _LOGBf __nv_logbf
+#define _fmaxd __nv_fmax
+#define _fmaxf __nv_fmaxf
+#else
+// OPENMP and __AMDGCN__
+#include <__clang_hip_libdevice_declares.h>
+#define _ISNANd __ocml_isnan_f64
+#define _ISNANf __ocml_isnan_f32
+#define _ISINFd __ocml_isinf_f64
+#define _ISINFf __ocml_isinf_f32
+#define _ISFINITEd __ocml_isfinite_f64
+#define _ISFINITEf __ocml_isfinite_f32
+#define _COPYSIGNd __ocml_copysign_f64
+#define _COPYSIGNf __ocml_copysign_f32
+#define _SCALBNd __ocml_scalbn_f64
+#define _SCALBNf __ocml_scalbn_f32
+#define _ABSd __ocml_fabs_f64
+#define _ABSf __ocml_fabs_f32
+#define _LOGBd __ocml_logb_f64
+#define _LOGBf __ocml_logb_f32
+#define _fmaxd __ocml_fmax_f64
+#define _fmaxf __ocml_fmax_f32
+#endif
 #endif
 
 #if defined(__cplusplus)
 extern "C" {
 #endif
 
 __DEVICE__ double _Complex __muldc3(double __a, double __b, double __c,
                                     double __d) {

 }
 
 __DEVICE__ double _Complex __divdc3(double __a, double __b, double __c,
                                     double __d) {
   int __ilogbw = 0;
   // Can't use std::max, because that's defined in <algorithm>, and we don't
   // want to pull that in for every compile.  The CUDA headers define
   // ::max(float, float) and ::max(double, double), which is sufficient for us.
-  double __logbw = _LOGBd(max(_ABSd(__c), _ABSd(__d)));
+  double __logbw = _LOGBd(_fmaxd(_ABSd(__c), _ABSd(__d)));
   if (_ISFINITEd(__logbw)) {
     __ilogbw = (int)__logbw;
     __c = _SCALBNd(__c, -__ilogbw);
     __d = _SCALBNd(__d, -__ilogbw);
   }
   double __denom = __c * __c + __d * __d;
   double _Complex z;
   __real__(z) = _SCALBNd((__a * __c + __b * __d) / __denom, -__ilogbw);
if ((__denom == 0.0) && (!_ISNANd(__a) || !_ISNANd(__b))) {
       __imag__(z) = 0.0 * (__b * __c - __a * __d);
     }
   }
   return z;
 }
 
 __DEVICE__ float _Complex __divsc3(float __a, float __b, float __c, float __d) {
   int __ilogbw = 0;
-  float __logbw = _LOGBf(max(_ABSf(__c), _ABSf(__d)));
+  float __logbw = _LOGBf(_fmaxf(_ABSf(__c), _ABSf(__d)));
   if (_ISFINITEf(__logbw)) {
     __ilogbw = (int)__logbw;
     __c = _SCALBNf(__c, -__ilogbw);
     __d = _SCALBNf(__d, -__ilogbw);
   }
   float __denom = __c * __c + __d * __d;
   float _Complex z;
   __real__(z) = _SCALBNf((__a * __c + __b * __d) / __denom, -__ilogbw);

 #undef _ISFINITEd
 #undef _ISFINITEf
 #undef _SCALBNd
 #undef _SCALBNf
 #undef _ABSd
 #undef _ABSf
 #undef _LOGBd
 #undef _LOGBf
+#undef _fmaxd
+#undef _fmaxf
 
 #ifdef _OPENMP
 #pragma omp end declare target
 #endif
 
 #pragma pop_macro("__DEVICE__")
 
 #endif // __CLANG_CUDA_COMPLEX_BUILTINS

clang/lib/Headers/__clang_hip_cmath.h

 #endif
 
 #if defined(__cplusplus)
 #include <limits>
 #endif
 #include <limits.h>
 #include <stdint.h>
 
+// __DEVICE__ is a helper macro with common set of attributes for the wrappers
+// we implement in this file. We need static in order to avoid emitting unused
+// functions.
 #pragma push_macro("__DEVICE__")
+#ifdef __OPENMP_AMDGCN__
+#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
+#else
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#endif
 
 // Start with functions that cannot be defined by DEF macros below.
 #if defined(__cplusplus)
 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
 __DEVICE__ long abs(long __n) { return ::labs(__n); }
 __DEVICE__ float fma(float __x, float __y, float __z) {

 __DEVICE__ bool isnormal(double __x) { return __builtin_isnormal(__x); }
 __DEVICE__ bool isunordered(float __x, float __y) {
   return __builtin_isunordered(__x, __y);
 }
 __DEVICE__ bool isunordered(double __x, double __y) {
   return __builtin_isunordered(__x, __y);
 }
 __DEVICE__ float modf(float __x, float *__iptr) { return ::modff(__x, __iptr); }
 __DEVICE__ float pow(float __base, int __iexp) {
   return ::powif(__base, __iexp);
 }
 __DEVICE__ double pow(double __base, int __iexp) {
   return ::powi(__base, __iexp);
 }
 __DEVICE__ float remquo(float __x, float __y, int *__quo) {
   return ::remquof(__x, __y, __quo);
 }

clang/lib/Headers/__clang_hip_libdevice_declares.h

 __device__ __attribute__((const)) __device__ float __ocml_fmod_f32(float,
                                                                    float);
 __device__ float __ocml_frexp_f32(float,
                                   __attribute__((address_space(5))) int *);
 __device__ __attribute__((const)) float __ocml_hypot_f32(float, float);
 __device__ __attribute__((const)) int __ocml_ilogb_f32(float);
 __device__ __attribute__((const)) int __ocml_isfinite_f32(float);
 __device__ __attribute__((const)) int __ocml_isinf_f32(float);
 __device__ __attribute__((const)) int __ocml_isnan_f32(float);
 __device__ float __ocml_j0_f32(float);
 __device__ float __ocml_j1_f32(float);
 __device__ __attribute__((const)) float __ocml_ldexp_f32(float, int);
 __device__ float __ocml_lgamma_f32(float);
 __device__ __attribute__((pure)) float __ocml_log10_f32(float);
 __device__ __attribute__((pure)) float __ocml_native_log10_f32(float);
 __device__ __attribute__((pure)) float __ocml_log1p_f32(float);
 __device__ __attribute__((pure)) float __ocml_log2_f32(float);

 __device__ _Float16 __ocml_sin_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
 __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
 __device__ __attribute__((pure)) _Float16 __ocml_pown_f16(_Float16, int);
 
 typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
 typedef short __2i16 __attribute__((ext_vector_type(2)));
 
+#if defined(__cplusplus)
 __device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b,
                                                      float c, bool s);
+#else
+__device__ __attribute__((const)) float __ockl_fdot2(__2f16 a, __2f16 b,
+                                                     float c, unsigned int s);
+#endif
 __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_fabs_2f16(__2f16);
 __device__ __2f16 __ocml_cos_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
 __device__ __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
 __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
 __device__ __attribute__((const))

clang/lib/Headers/__clang_hip_math.h

 #endif
 
 #if defined(__cplusplus)
 #include <algorithm>
 #endif
 #include <limits.h>
 #include <stdint.h>
 
+// __DEVICE__ is a helper macro with common set of attributes for the wrappers
+// we implement in this file. We need static in order to avoid emitting unused
+// functions and __forceinline__ helps inlining these wrappers at -O1.
 #pragma push_macro("__DEVICE__")
+#pragma push_macro("__DEVICE_NOCE__")
+
+#ifdef __OPENMP_AMDGCN__
+#if defined(__cplusplus)
+#define __DEVICE__ static constexpr __attribute__((always_inline, nothrow))
+#define __DEVICE_NOCE__ static __attribute__((always_inline, nothrow))
+#else // !defined(__cplusplus), c openmp compilation
+#define __DEVICE__ static __attribute__((always_inline, nothrow))
+#define __DEVICE_NOCE__ __DEVICE__
+#endif
+#else // !__OPENMP_AMDGCN__, so this is for HIP-Clang which is always C++.
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#define __DEVICE_NOCE__ __DEVICE__
+#endif
 
 // A few functions return bool type starting only in C++11.
 #pragma push_macro("__RETURN_TYPE")
+#ifdef __OPENMP_AMDGCN__
+#define __RETURN_TYPE int
+#else
 #if defined(__cplusplus)
 #define __RETURN_TYPE bool
 #else
 #define __RETURN_TYPE int
 #endif
+#endif
 
 #if defined (__cplusplus) && __cplusplus < 201103L
 // emulate static_assert on type sizes
 template<bool>
 struct __compare_result{};
 template<>
 struct __compare_result<true> {
   static const bool valid;
 };
 
-__DEVICE__
+// All following c-capable function defs have one of two macro modifiers:
+// __DEVICE__
+// __DEVICE_NOCE__ same as __DEVICE__ but no constexpr for those functions
+//                 that cannot return constexpr in c++.
+
+__DEVICE_NOCE__
 void __suppress_unused_warning(bool b){};
 template <unsigned int S, unsigned int T>
 __DEVICE__ void __static_assert_equal_size() {
   __suppress_unused_warning(__compare_result<S == T>::valid);
 }
 
 #define __static_assert_type_size_equal(A, B) \
   __static_assert_equal_size<A,B>()
 
 #else
 #define __static_assert_type_size_equal(A,B) \
   static_assert((A) == (B), "")
 
 #endif
 
-__DEVICE__
+__DEVICE_NOCE__
 uint64_t __make_mantissa_base8(const char *__tagp) {
   uint64_t __r = 0;
   while (__tagp) {
     char __tmp = *__tagp;
 
     if (__tmp >= '0' && __tmp <= '7')
       __r = (__r * 8u) + __tmp - '0';
     else
       return 0;
 
     ++__tagp;
   }
 
   return __r;
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 uint64_t __make_mantissa_base10(const char *__tagp) {
   uint64_t __r = 0;
   while (__tagp) {
     char __tmp = *__tagp;
 
     if (__tmp >= '0' && __tmp <= '9')
       __r = (__r * 10u) + __tmp - '0';
     else
       return 0;
 
     ++__tagp;
   }
 
   return __r;
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 uint64_t __make_mantissa_base16(const char *__tagp) {
   uint64_t __r = 0;
   while (__tagp) {
     char __tmp = *__tagp;
 
     if (__tmp >= '0' && __tmp <= '9')
       __r = (__r * 16u) + __tmp - '0';
     else if (__tmp >= 'a' && __tmp <= 'f')

 float fmaxf(float __x, float __y) { return __ocml_fmax_f32(__x, __y); }
 
 __DEVICE__
 float fminf(float __x, float __y) { return __ocml_fmin_f32(__x, __y); }
 
 __DEVICE__
 float fmodf(float __x, float __y) { return __ocml_fmod_f32(__x, __y); }
 
-__DEVICE__
+__DEVICE_NOCE__
 float frexpf(float __x, int *__nptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) int __tmp;
+  float __r = __ocml_frexp_f32(__x, &__tmp);
+#else
   int __tmp;
   float __r =
       __ocml_frexp_f32(__x, (__attribute__((address_space(5))) int *)&__tmp);
+#endif
   *__nptr = __tmp;
-
   return __r;
 }
 
 __DEVICE__
 float hypotf(float __x, float __y) { return __ocml_hypot_f32(__x, __y); }
 
 __DEVICE__
 int ilogbf(float __x) { return __ocml_ilogb_f32(__x); }
 
 __DEVICE__
 __RETURN_TYPE __finitef(float __x) { return __ocml_isfinite_f32(__x); }
 
 __DEVICE__
 __RETURN_TYPE __isinff(float __x) { return __ocml_isinf_f32(__x); }
 
 __DEVICE__
 __RETURN_TYPE __isnanf(float __x) { return __ocml_isnan_f32(__x); }
 
 __DEVICE__
 float j0f(float __x) { return __ocml_j0_f32(__x); }
 
 __DEVICE__
 float j1f(float __x) { return __ocml_j1_f32(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 float jnf(int __n, float __x) { // TODO: we could use Ahmes multiplication
                                 // and the Miller & Brown algorithm
   //       for linear recurrences to get O(log n) steps, but it's unclear if
   //       it'd be beneficial in this case.
   if (__n == 0)
     return j0f(__x);
   if (__n == 1)
     return j1f(__x);

 float logf(float __x) { return __ocml_log_f32(__x); }
 
 __DEVICE__
 long int lrintf(float __x) { return __ocml_rint_f32(__x); }
 
 __DEVICE__
 long int lroundf(float __x) { return __ocml_round_f32(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 float modff(float __x, float *__iptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) float __tmp;
+  float __r = __ocml_modf_f32(__x, &__tmp);
+#else
   float __tmp;
   float __r =
       __ocml_modf_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
+#endif
   *__iptr = __tmp;
   return __r;
 }
 
+// FIXME need a c version of nanf
+#if defined(__cplusplus)
 __DEVICE__
 float nanf(const char *__tagp) {
   union {
     float val;
     struct ieee_float {
       unsigned int mantissa : 22;
       unsigned int quiet : 1;
       unsigned int exponent : 8;
       unsigned int sign : 1;
     } bits;
   } __tmp;
   __static_assert_type_size_equal(sizeof(__tmp.val), sizeof(__tmp.bits));
 
   __tmp.bits.sign = 0u;
   __tmp.bits.exponent = ~0u;
   __tmp.bits.quiet = 1u;
   __tmp.bits.mantissa = __make_mantissa(__tagp);
 
   return __tmp.val;
 }
+#endif
 
 __DEVICE__
 float nearbyintf(float __x) { return __ocml_nearbyint_f32(__x); }
 
 __DEVICE__
 float nextafterf(float __x, float __y) {
   return __ocml_nextafter_f32(__x, __y);
 }

 }
 
 __DEVICE__
 float normcdff(float __x) { return __ocml_ncdf_f32(__x); }
 
 __DEVICE__
 float normcdfinvf(float __x) { return __ocml_ncdfinv_f32(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 float normf(int __dim,
             const float *__a) { // TODO: placeholder until OCML adds support.
   float __r = 0;
   while (__dim--) {
     __r += __a[0] * __a[0];
     ++__a;
   }
 
   return __ocml_sqrt_f32(__r);
 }
 
 __DEVICE__
 float powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
 
 __DEVICE__
 float powif(float __x, int __y) { return __ocml_pown_f32(__x, __y); }
 
 __DEVICE__
 float rcbrtf(float __x) { return __ocml_rcbrt_f32(__x); }
 
 __DEVICE__
 float remainderf(float __x, float __y) {
   return __ocml_remainder_f32(__x, __y);
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 float remquof(float __x, float __y, int *__quo) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) int __tmp;
+  float __r = __ocml_remquo_f32( __x, __y, &__tmp);
+#else
   int __tmp;
   float __r = __ocml_remquo_f32(
       __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
+#endif
   *__quo = __tmp;
-
   return __r;
 }
 
 __DEVICE__
 float rhypotf(float __x, float __y) { return __ocml_rhypot_f32(__x, __y); }
 
 __DEVICE__
 float rintf(float __x) { return __ocml_rint_f32(__x); }
 
 __DEVICE__
 float rnorm3df(float __x, float __y, float __z) {
   return __ocml_rlen3_f32(__x, __y, __z);
 }
 
 __DEVICE__
 float rnorm4df(float __x, float __y, float __z, float __w) {
   return __ocml_rlen4_f32(__x, __y, __z, __w);
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 float rnormf(int __dim,
              const float *__a) { // TODO: placeholder until OCML adds support.
   float __r = 0;
   while (__dim--) {
     __r += __a[0] * __a[0];
     ++__a;
   }
 

 }
 
 __DEVICE__
 float scalbnf(float __x, int __n) { return __ocml_scalbn_f32(__x, __n); }
 
 __DEVICE__
 __RETURN_TYPE __signbitf(float __x) { return __ocml_signbit_f32(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 void sincosf(float __x, float *__sinptr, float *__cosptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) float __tmp;
+  *__sinptr = __ocml_sincos_f32(__x, &__tmp);
+#else
   float __tmp;
   *__sinptr =
       __ocml_sincos_f32(__x, (__attribute__((address_space(5))) float *)&__tmp);
+#endif
   *__cosptr = __tmp;
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 void sincospif(float __x, float *__sinptr, float *__cosptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) float __tmp;
+  *__sinptr = __ocml_sincospi_f32(__x, &__tmp);
+#else
   float __tmp;
   *__sinptr = __ocml_sincospi_f32(
       __x, (__attribute__((address_space(5))) float *)&__tmp);
+#endif
   *__cosptr = __tmp;
 }
 
 __DEVICE__
 float sinf(float __x) { return __ocml_sin_f32(__x); }
 
 __DEVICE__
 float sinhf(float __x) { return __ocml_sinh_f32(__x); }

 float truncf(float __x) { return __ocml_trunc_f32(__x); }
 
 __DEVICE__
 float y0f(float __x) { return __ocml_y0_f32(__x); }
 
 __DEVICE__
 float y1f(float __x) { return __ocml_y1_f32(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 float ynf(int __n, float __x) { // TODO: we could use Ahmes multiplication
                                 // and the Miller & Brown algorithm
   //       for linear recurrences to get O(log n) steps, but it's unclear if
   //       it'd be beneficial in this case. Placeholder until OCML adds
   //       support.
   if (__n == 0)
     return y0f(__x);
   if (__n == 1)

 float __logf(float __x) { return __ocml_native_log_f32(__x); }
 
 __DEVICE__
 float __powf(float __x, float __y) { return __ocml_pow_f32(__x, __y); }
 
 __DEVICE__
 float __saturatef(float __x) { return (__x < 0) ? 0 : ((__x > 1) ? 1 : __x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 void __sincosf(float __x, float *__sinptr, float *__cosptr) {
   *__sinptr = __ocml_native_sin_f32(__x);
   *__cosptr = __ocml_native_cos_f32(__x);
 }
 
 __DEVICE__
 float __sinf(float __x) { return __ocml_native_sin_f32(__x); }
 

 double fmax(double __x, double __y) { return __ocml_fmax_f64(__x, __y); }
 
 __DEVICE__
 double fmin(double __x, double __y) { return __ocml_fmin_f64(__x, __y); }
 
 __DEVICE__
 double fmod(double __x, double __y) { return __ocml_fmod_f64(__x, __y); }
 
-__DEVICE__
+__DEVICE_NOCE__
 double frexp(double __x, int *__nptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) int __tmp;
+  double __r = __ocml_frexp_f64(__x, &__tmp);
+#else
   int __tmp;
   double __r =
       __ocml_frexp_f64(__x, (__attribute__((address_space(5))) int *)&__tmp);
+#endif
   *__nptr = __tmp;
   return __r;
 }
 
 __DEVICE__
 double hypot(double __x, double __y) { return __ocml_hypot_f64(__x, __y); }
 
 __DEVICE__

 __RETURN_TYPE __isnan(double __x) { return __ocml_isnan_f64(__x); }
 
 __DEVICE__
 double j0(double __x) { return __ocml_j0_f64(__x); }
 
 __DEVICE__
 double j1(double __x) { return __ocml_j1_f64(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 double jn(int __n, double __x) { // TODO: we could use Ahmes multiplication
                                  // and the Miller & Brown algorithm
   //       for linear recurrences to get O(log n) steps, but it's unclear if
   //       it'd be beneficial in this case. Placeholder until OCML adds
   //       support.
   if (__n == 0)
     return j0(__x);
   if (__n == 1)

 double logb(double __x) { return __ocml_logb_f64(__x); }
 
 __DEVICE__
 long int lrint(double __x) { return __ocml_rint_f64(__x); }
 
 __DEVICE__
 long int lround(double __x) { return __ocml_round_f64(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 double modf(double __x, double *__iptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) double __tmp;
+  double __r = __ocml_modf_f64(__x, &__tmp);
+#else
   double __tmp;
   double __r =
       __ocml_modf_f64(__x, (__attribute__((address_space(5))) double *)&__tmp);
+#endif
   *__iptr = __tmp;
-
   return __r;
 }
 
+// FIXME need a c version of nan
+#if defined(__cplusplus)
 __DEVICE__
 double nan(const char *__tagp) {
 #if !_WIN32
   union {
     double val;
     struct ieee_double {
       uint64_t mantissa : 51;
       uint32_t quiet : 1;
#if !_WIN32
   return __tmp.val;
 #else
   __static_assert_type_size_equal(sizeof(uint64_t), sizeof(double));
   uint64_t __val = __make_mantissa(__tagp);
   __val |= 0xFFF << 51;
   return *reinterpret_cast<double *>(&__val);
 #endif
 }
+#endif
 
 __DEVICE__
 double nearbyint(double __x) { return __ocml_nearbyint_f64(__x); }
 
 __DEVICE__
 double nextafter(double __x, double __y) {
   return __ocml_nextafter_f64(__x, __y);
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 double norm(int __dim,
             const double *__a) { // TODO: placeholder until OCML adds support.
   double __r = 0;
   while (__dim--) {
     __r += __a[0] * __a[0];
     ++__a;
   }
 

 __DEVICE__
 double rcbrt(double __x) { return __ocml_rcbrt_f64(__x); }
 
 __DEVICE__
 double remainder(double __x, double __y) {
   return __ocml_remainder_f64(__x, __y);
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 double remquo(double __x, double __y, int *__quo) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) int __tmp;
+  double __r = __ocml_remquo_f64(__x, __y, &__tmp);
+#else
   int __tmp;
   double __r = __ocml_remquo_f64(
       __x, __y, (__attribute__((address_space(5))) int *)&__tmp);
+#endif
   *__quo = __tmp;
-
   return __r;
 }
 
 __DEVICE__
 double rhypot(double __x, double __y) { return __ocml_rhypot_f64(__x, __y); }
 
 __DEVICE__
 double rint(double __x) { return __ocml_rint_f64(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 double rnorm(int __dim,
              const double *__a) { // TODO: placeholder until OCML adds support.
   double __r = 0;
   while (__dim--) {
     __r += __a[0] * __a[0];
     ++__a;
   }
 

 double scalbn(double __x, int __n) { return __ocml_scalbn_f64(__x, __n); }
 
 __DEVICE__
 __RETURN_TYPE __signbit(double __x) { return __ocml_signbit_f64(__x); }
 
 __DEVICE__
 double sin(double __x) { return __ocml_sin_f64(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 void sincos(double __x, double *__sinptr, double *__cosptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) double __tmp;
+  *__sinptr = __ocml_sincos_f64(__x, &__tmp);
+#else
   double __tmp;
   *__sinptr = __ocml_sincos_f64(
       __x, (__attribute__((address_space(5))) double *)&__tmp);
+#endif
   *__cosptr = __tmp;
 }
 
-__DEVICE__
+__DEVICE_NOCE__
 void sincospi(double __x, double *__sinptr, double *__cosptr) {
+#ifdef __OPENMP_AMDGCN__
+  static __attribute__((address_space(5))) double __tmp;
+  *__sinptr = __ocml_sincospi_f64(__x, &__tmp);
+#else
   double __tmp;
   *__sinptr = __ocml_sincospi_f64(
       __x, (__attribute__((address_space(5))) double *)&__tmp);
+#endif
   *__cosptr = __tmp;
 }
 
 __DEVICE__
 double sinh(double __x) { return __ocml_sinh_f64(__x); }
 
 __DEVICE__
 double sinpi(double __x) { return __ocml_sinpi_f64(__x); }

 double trunc(double __x) { return __ocml_trunc_f64(__x); }
 
 __DEVICE__
 double y0(double __x) { return __ocml_y0_f64(__x); }
 
 __DEVICE__
 double y1(double __x) { return __ocml_y1_f64(__x); }
 
-__DEVICE__
+__DEVICE_NOCE__
 double yn(int __n, double __x) { // TODO: we could use Ahmes multiplication
                                  // and the Miller & Brown algorithm
   //       for linear recurrences to get O(log n) steps, but it's unclear if
   //       it'd be beneficial in this case. Placeholder until OCML adds
   //       support.
   if (__n == 0)
     return y0(__x);
   if (__n == 1)

 double max(double __x, double __y) { return fmax(__x, __y); }
 
 __DEVICE__
 float min(float __x, float __y) { return fminf(__x, __y); }
 
 __DEVICE__
 double min(double __x, double __y) { return fmin(__x, __y); }
 
+#ifndef __OPENMP_AMDGCN__
 __host__ inline static int min(int __arg1, int __arg2) {
   return std::min(__arg1, __arg2);
 }
 
 __host__ inline static int max(int __arg1, int __arg2) {
   return std::max(__arg1, __arg2);
 }
 #endif
+#endif // defined(__cplusplus)
 
+#pragma pop_macro("__DEVICE_NOCE__")
 #pragma pop_macro("__DEVICE__")
 #pragma pop_macro("__RETURN_TYPE")
 
 #endif // __CLANG_HIP_MATH_H__

clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h

 /*===- __clang_openmp_device_functions.h - OpenMP device function declares -===
  *
  * 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_OPENMP_DEVICE_FUNCTIONS_H__
 #define __CLANG_OPENMP_DEVICE_FUNCTIONS_H__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
+#ifdef __NVPTX__
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
 #ifdef __cplusplus
+#include <cstdint>
 extern "C" {
+#else
+#include <stdint.h>
 #endif
 
 #define __CUDA__
 #define __OPENMP_NVPTX__
 
 /// Include declarations for libdevice functions.
 #include <__clang_cuda_libdevice_declares.h>
 
 /// Provide definitions for these functions.
 #include <__clang_cuda_device_functions.h>
 
 #undef __OPENMP_NVPTX__
 #undef __CUDA__
 
 #ifdef __cplusplus
 } // extern "C"
 #endif
 
 #pragma omp end declare variant
 
+#endif // __NVPTX__
+
+#ifdef __AMDGCN__
+
+// __NO_INLINE__ prevents some x86 optimized macro definitions in system headers
+#define __NO_INLINE__ 1
+#pragma omp begin declare variant match(                                       \
+    device = {arch(amdgcn)}, implementation = {extension(match_any)})
+
+#ifdef __cplusplus
+#include <cstdint>
+extern "C" {
+#else
+#include <stdint.h>
+#endif
+
+#define __HIP__
+#define __OPENMP_AMDGCN__
+
+#define __host__ __attribute__((host))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __shared__ __attribute__((shared))
+#define __constant__ __attribute__((constant))
+#define __private __attribute__((address_space(5)))
+
+/// Include declarations for libdevice functions.
+#include <__clang_hip_libdevice_declares.h>
+
+#ifdef __cplusplus
+} // extern "C"
 #endif
+
+#pragma omp end declare variant
+
+#endif // __AMDGCN__
+
+#endif // __CLANG_OPENMP_DEVICE_FUNCTIONS_H__

clang/lib/Headers/openmp_wrappers/cmath

 #define __CLANG_OPENMP_CMATH_H__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 #include_next <cmath>
 
-// Make sure we include our math.h overlay, it probably happend already but we
-// need to be sure.
+// Make sure we include our new and math.h overlays, it probably happened already
+// but we need to be sure.
+#include <new>
 #include <math.h>
 
 // We (might) need cstdlib because __clang_cuda_cmath.h below declares `abs`
 // which might live in cstdlib.
 #include <cstdlib>
 
+#ifdef __NVPTX__
 // We need limits because __clang_cuda_cmath.h below uses `std::numeric_limit`.
 #include <limits>
 
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any, allow_templates)})
 
 #define __CUDA__
 #define __OPENMP_NVPTX__
__DEVICE__ float scalbln(float __x, long int __y) {
   return ::scalblnf(__x, __y);
 }
 __DEVICE__ float scalbn(float __x, int __y) { return ::scalbnf(__x, __y); }
 __DEVICE__ float tgamma(float __x) { return ::tgammaf(__x); }
 
 #undef __DEVICE__
 
 #pragma omp end declare variant
+#endif // __NVPTX__
 
-#endif
+#ifdef __AMDGCN__
+#pragma omp begin declare variant match(                                       \
+    device = {arch(amdgcn)}, implementation = {extension(match_any, allow_templates)})
+#define __HIP__
+#define __OPENMP_AMDGCN__
+#include <__clang_hip_cmath.h>
+#undef __HIP__
+
+#pragma omp end declare variant
+#endif // __AMDGCN__
+
+#endif // __CLANG_OPENMP_CMATH_H__

clang/lib/Headers/openmp_wrappers/complex

 #ifndef __CLANG_OPENMP_COMPLEX__
 #define __CLANG_OPENMP_COMPLEX__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 // We require std::math functions in the complex builtins below.
+#ifdef __NVPTX__
 #include <cmath>
-
 #define __CUDA__
+#define __ARCHTYPES__ nvptx,nvptx64
+#endif // __NVPTX__
+
+#ifdef __AMDGCN__
+#include <__clang_hip_libdevice_declares.h>
+#define __HIP__
+#define __ARCHTYPES__ amdgcn
+#endif // __AMDGCN__
+
 #include <__clang_cuda_complex_builtins.h>
-#endif
+#endif //__CLANG_OPENMP_COMPLEX__
 
 // Grab the host header too.
 #include_next <complex>
 
 
 #ifdef __cplusplus
 
 // If we are compiling against libc++, the macro _LIBCPP_STD_VER should be set
 // after including <cmath> above. Since the complex header we use is a
 // simplified version of the libc++, we don't need it in this case. If we
 // compile against libstdc++, or any other standard library, we will overload
 // the (hopefully template) functions in the <complex> header with the ones we
 // got from libc++ which decomposes math functions, like `std::sin`, into
 // arithmetic and calls to non-complex functions, all of which we can then
 // handle.
 #ifndef _LIBCPP_STD_VER
+#ifndef _GLIBCXX_COMPLEX
 
 #pragma omp begin declare variant match(                                       \
-    device = {arch(nvptx, nvptx64)},                                           \
+    device = {arch(__ARCHTYPES__)},                                            \
     implementation = {extension(match_any, allow_templates)})
 
 #include <complex_cmath.h>
 
 #pragma omp end declare variant
 
-#endif
+#endif // _GLIBCXX_COMPLEX
+#endif // _LIBCPP_STD_VER
 
-#endif
+#endif // __cplusplus

clang/lib/Headers/openmp_wrappers/hip/hip_runtime.h

+/*===-- hip_runtime - OpenMP hip_runtime.h wrapper for target regions ------===
+ *
+ * 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_OPENMP_HIP_RUNTIME_H__
+#define __CLANG_OPENMP_HIP_RUNTIME_H__
+
+#ifndef _OPENMP
+#error "This file is for OpenMP compilation only."
+#endif
+
+#include <time.h>
+#define __OPENMP_AMDGCN__
+#include_next <hip/hip_runtime.h>
+
+#pragma omp begin declare variant match(                                       \
+    device = {arch(amdgcn)}, implementation = {extension(match_any)})
+
+#pragma omp end declare variant
+
+// Now get the actual hip headers
+
+#endif // __CLANG_OPENMP_HIP_RUNTIME_H__

clang/lib/Headers/openmp_wrappers/math.h

 
 #ifndef __CLANG_OPENMP_MATH_H__
 #define __CLANG_OPENMP_MATH_H__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 #include_next <math.h>
 
 // We need limits.h for __clang_cuda_math.h below and because it should not hurt
 // we include it eagerly here.
 #include <limits.h>
 
 // We need stdlib.h because (for now) __clang_cuda_math.h below declares `abs`
 // which should live in stdlib.h.
 #include <stdlib.h>
 
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
 #define __CUDA__
 #define __OPENMP_NVPTX__
 #include <__clang_cuda_math.h>
 #undef __OPENMP_NVPTX__
 #undef __CUDA__
 
 #pragma omp end declare variant
 
+#pragma omp begin declare variant match(                                       \
+    device = {arch(amdgcn)}, implementation = {extension(match_any)})
+
+#ifndef __OPENMP_AMDGCN__
+#define __OPENMP_AMDGCN__
+#endif
+
+#ifndef __HIP__
+#define __HIP__
 #endif
+
+#include <__clang_hip_math.h>
+#pragma omp end declare variant
+
+#endif // __CLANG_OPENMP_MATH_H__

clang/lib/Headers/openmp_wrappers/new

 //===--------- new - OPENMP wrapper for <new> ------------------------------===
 //
 // 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_OPENMP_WRAPPERS_NEW
 #define __CLANG_OPENMP_WRAPPERS_NEW
 
 #include_next <new>
 
-#if defined(__NVPTX__) && defined(_OPENMP)
+#if (defined(__NVPTX__) || defined(__AMDGCN__))  && defined(_OPENMP)
 
 #include <cstdlib>
 
 #pragma push_macro("OPENMP_NOEXCEPT")
 #if __cplusplus >= 201103L
 #define OPENMP_NOEXCEPT noexcept
 #else
 #define OPENMP_NOEXCEPT

clang/test/Headers/Inputs/include/algorithm

+#pragma once
+
+// Copied from libcxx
+
+namespace std {
+
+template <class T>
+    const T&
+    max(const T& a, const T& b); // constexpr in C++14
+template <class T, class Compare>
+    const T&
+    max(const T& a, const T& b, Compare comp);  // constexpr in C++14
+
+template <class T>
+    const T&
+    min(const T& a, const T& b);  // constexpr in C++14
+template <class T, class Compare>
+    const T&
+    min(const T& a, const T& b, Compare comp);  // constexpr in C++14
+
+}

clang/test/Headers/Inputs/include/cstdint

+#pragma once
+
+#include <stdint.h>
+
+namespace std {
+#ifdef __INT32_TYPE__
+using ::uint32_t;
+#endif
+
+#ifdef __INT64_TYPE__
+using ::uint64_t;
+#endif
+
+#ifdef __INTPTR_TYPE__
+using ::intptr_t;
+using ::uintptr_t;
+#else
+#error Every target should have __INTPTR_TYPE__
+#endif
+
+} // namespace std

clang/test/Headers/Inputs/include/cstdlib

 extern long int labs (long int __x) __attribute__ ((__const__)) ;
 extern float fabs (float __x) __attribute__ ((__const__)) ;
 #endif
 
 namespace std
 {
 
 using ::abs;
+using ::size_t;
 
 inline long
 abs(long __i) { return __builtin_labs(__i); }
 
 inline long long
 abs(long long __x) { return __builtin_llabs (__x); }
 
 float fabs(float __x) { return __builtin_fabs(__x); }
 
 float abs(float __x) { return fabs(__x); }
 double abs(double __x) { return fabs(__x); }
 
 }

clang/test/Headers/Inputs/include/exception

+#pragma once
+
+// Copied from libcxx
+
+namespace std {
+
+class exception
+{
+public:
+    exception() noexcept;
+    exception(const exception&) noexcept;
+    exception& operator=(const exception&) noexcept;
+    virtual ~exception() noexcept;
+    virtual const char* what() const noexcept;
+};
+
+}

clang/test/Headers/Inputs/include/new

+#pragma once
+
+#include <cstdlib>
+#include <exception>
+
+// Copied from libcxx
+
+namespace std {
+
+struct nothrow_t { explicit nothrow_t() = default; };
+extern const nothrow_t nothrow;
+
+class bad_alloc
+    : public exception
+{
+public:
+    bad_alloc();
+    virtual ~bad_alloc();
+    virtual const char* what() const;
+};
+
+class bad_array_new_length
+    : public bad_alloc
+{
+public:
+    bad_array_new_length();
+    virtual ~bad_array_new_length();
+    virtual const char* what() const;
+};
+
+typedef void (*new_handler)();
+new_handler set_new_handler(new_handler);
+new_handler get_new_handler();
+
+}
+
+void* operator new(std::size_t __sz);
+void* operator new(std::size_t __sz, const std::nothrow_t&);
+void  operator delete(void* __p);
+void  operator delete(void* __p, const std::nothrow_t&);
+
+void* operator new[](std::size_t __sz);
+void* operator new[](std::size_t __sz, const std::nothrow_t&);
+void  operator delete[](void* __p);
+void  operator delete[](void* __p, const std::nothrow_t&);
+
+inline void* operator new  (std::size_t, void* __p) {return __p;}
+inline void* operator new[](std::size_t, void* __p) {return __p;}
+inline void  operator delete  (void*, void*) {}
+inline void  operator delete[](void*, void*) {}

clang/test/Headers/Inputs/include/stdlib.h

 #pragma once
 typedef __SIZE_TYPE__ size_t;
 
 #ifndef __cplusplus
 extern int abs(int __x) __attribute__((__const__));
 #endif
+
+void free(void* ptr);
+void* malloc(size_t size);