diff --git a/clang/test/Headers/Inputs/include/crt/device_double_functions.hpp b/clang/test/Headers/Inputs/include/crt/device_double_functions.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/device_double_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/crt/device_functions.hpp b/clang/test/Headers/Inputs/include/crt/device_functions.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/device_functions.hpp
@@ -0,0 +1,3 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+__device__ void                   __brkpt();
diff --git a/clang/test/Headers/Inputs/include/crt/device_runtime.h b/clang/test/Headers/Inputs/include/crt/device_runtime.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/device_runtime.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/crt/host_runtime.h b/clang/test/Headers/Inputs/include/crt/host_runtime.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/host_runtime.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/crt/math_functions.hpp b/clang/test/Headers/Inputs/include/crt/math_functions.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/math_functions.hpp
@@ -0,0 +1,12 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+__device__ int                    __isinff(float);
+__device__ int                    __isinf(double);
+__device__ int                    __finitef(float);
+__device__ int                    __isfinited(double);
+__device__ int                    __isnanf(float);
+__device__ int                    __isnan(double);
+__device__ int                    __signbitf(float);
+__device__ int                    __signbitd(double);
+__device__ double                 max(double, double);
+__device__ float                  max(float, float);
diff --git a/clang/test/Headers/Inputs/include/crt/sm_70_rt.hpp b/clang/test/Headers/Inputs/include/crt/sm_70_rt.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/crt/sm_70_rt.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/cstdlib b/clang/test/Headers/Inputs/include/cstdlib
--- a/clang/test/Headers/Inputs/include/cstdlib
+++ b/clang/test/Headers/Inputs/include/cstdlib
@@ -5,11 +5,9 @@
 #if __cplusplus >= 201703L
 extern int abs (int __x) throw()  __attribute__ ((__const__)) ;
 extern long int labs (long int __x) throw() __attribute__ ((__const__)) ;
-extern float fabs (float __x) throw() __attribute__ ((__const__)) ;
 #else
 extern int abs (int __x) __attribute__ ((__const__)) ;
 extern long int labs (long int __x) __attribute__ ((__const__)) ;
-extern float fabs (float __x) __attribute__ ((__const__)) ;
 #endif
 
 namespace std
diff --git a/clang/test/Headers/Inputs/include/cuda.h b/clang/test/Headers/Inputs/include/cuda.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/cuda.h
@@ -0,0 +1,127 @@
+/* Minimal declarations for CUDA support.  Testing purposes only. */
+#pragma once
+
+#include <stddef.h>
+
+// Make this file work with nvcc, for testing compatibility.
+
+#ifndef __NVCC__
+#define __constant__ __attribute__((constant))
+#define __device__ __attribute__((device))
+#define __global__ __attribute__((global))
+#define __host__ __attribute__((host))
+#define __shared__ __attribute__((shared))
+#define __managed__ __attribute__((managed))
+#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
+
+struct dim3 {
+  unsigned x, y, z;
+  __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
+};
+
+// Host- and device-side placement new overloads.
+void *operator new(__SIZE_TYPE__, void *p) { return p; }
+void *operator new[](__SIZE_TYPE__, void *p) { return p; }
+__device__ void *operator new(__SIZE_TYPE__, void *p) { return p; }
+__device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; }
+
+#define CUDA_VERSION 10100
+
+struct char2 {
+  char x, y;
+  __host__ __device__ char2(char x = 0, char y = 0) : x(x), y(y) {}
+};
+struct char4 {
+  char x, y, z, w;
+  __host__ __device__ char4(char x = 0, char y = 0, char z = 0, char w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct uchar2 {
+  unsigned char x, y;
+  __host__ __device__ uchar2(unsigned char x = 0, unsigned char y = 0) : x(x), y(y) {}
+};
+struct uchar4 {
+  unsigned char x, y, z, w;
+  __host__ __device__ uchar4(unsigned char x = 0, unsigned char y = 0, unsigned char z = 0, unsigned char w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct short2 {
+  short x, y;
+  __host__ __device__ short2(short x = 0, short y = 0) : x(x), y(y) {}
+};
+struct short4 {
+  short x, y, z, w;
+  __host__ __device__ short4(short x = 0, short y = 0, short z = 0, short w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct ushort2 {
+  unsigned short x, y;
+  __host__ __device__ ushort2(unsigned short x = 0, unsigned short y = 0) : x(x), y(y) {}
+};
+struct ushort4 {
+  unsigned short x, y, z, w;
+  __host__ __device__ ushort4(unsigned short x = 0, unsigned short y = 0, unsigned short z = 0, unsigned short w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct int2 {
+  int x, y;
+  __host__ __device__ int2(int x = 0, int y = 0) : x(x), y(y) {}
+};
+struct int4 {
+  int x, y, z, w;
+  __host__ __device__ int4(int x = 0, int y = 0, int z = 0, int w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct uint2 {
+  unsigned x, y;
+  __host__ __device__ uint2(unsigned x = 0, unsigned y = 0) : x(x), y(y) {}
+};
+struct uint3 {
+  unsigned x, y, z;
+  __host__ __device__ uint3(unsigned x = 0, unsigned y = 0, unsigned z = 0) : x(x), y(y), z(z) {}
+};
+struct uint4 {
+  unsigned x, y, z, w;
+  __host__ __device__ uint4(unsigned x = 0, unsigned y = 0, unsigned z = 0, unsigned w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+
+struct longlong2 {
+  long long x, y;
+  __host__ __device__ longlong2(long long x = 0, long long y = 0) : x(x), y(y) {}
+};
+struct longlong4 {
+  long long x, y, z, w;
+  __host__ __device__ longlong4(long long x = 0, long long y = 0, long long z = 0, long long w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct ulonglong2 {
+  unsigned long long x, y;
+  __host__ __device__ ulonglong2(unsigned long long x = 0, unsigned long long y = 0) : x(x), y(y) {}
+};
+struct ulonglong4 {
+  unsigned long long x, y, z, w;
+  __host__ __device__ ulonglong4(unsigned long long x = 0, unsigned long long y = 0, unsigned long long z = 0, unsigned long long w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+
+struct float2 {
+  float x, y;
+  __host__ __device__ float2(float x = 0, float y = 0) : x(x), y(y) {}
+};
+struct float4 {
+  float x, y, z, w;
+  __host__ __device__ float4(float x = 0, float y = 0, float z = 0, float w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+struct double2 {
+  double x, y;
+  __host__ __device__ double2(double x = 0, double y = 0) : x(x), y(y) {}
+};
+struct double4 {
+  double x, y, z, w;
+  __host__ __device__ double4(double x = 0, double y = 0, double z = 0, double w = 0) : x(x), y(y), z(z), w(w) {}
+};
+
+
+#endif // !__NVCC__
diff --git a/clang/test/Headers/Inputs/include/cuda_runtime.h b/clang/test/Headers/Inputs/include/cuda_runtime.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/cuda_runtime.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/curand_mtgp32_kernel.h b/clang/test/Headers/Inputs/include/curand_mtgp32_kernel.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/curand_mtgp32_kernel.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/device_atomic_functions.h b/clang/test/Headers/Inputs/include/device_atomic_functions.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/device_atomic_functions.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/device_atomic_functions.hpp b/clang/test/Headers/Inputs/include/device_atomic_functions.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/device_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/device_double_functions.h b/clang/test/Headers/Inputs/include/device_double_functions.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/device_double_functions.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/driver_types.h b/clang/test/Headers/Inputs/include/driver_types.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/driver_types.h
@@ -0,0 +1,4 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+
+#include <limits.h>
diff --git a/clang/test/Headers/Inputs/include/host_config.h b/clang/test/Headers/Inputs/include/host_config.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/host_config.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/host_defines.h b/clang/test/Headers/Inputs/include/host_defines.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/host_defines.h
@@ -0,0 +1,3 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+#define __forceinline__ 
diff --git a/clang/test/Headers/Inputs/include/math_functions_dbl_ptx3.hpp b/clang/test/Headers/Inputs/include/math_functions_dbl_ptx3.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/math_functions_dbl_ptx3.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/new b/clang/test/Headers/Inputs/include/new
--- a/clang/test/Headers/Inputs/include/new
+++ b/clang/test/Headers/Inputs/include/new
@@ -1,3 +1,4 @@
+#pragma once
 
 namespace std
 {
diff --git a/clang/test/Headers/Inputs/include/sm_20_atomic_functions.hpp b/clang/test/Headers/Inputs/include/sm_20_atomic_functions.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_20_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_20_intrinsics.hpp b/clang/test/Headers/Inputs/include/sm_20_intrinsics.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_20_intrinsics.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_32_atomic_functions.hpp b/clang/test/Headers/Inputs/include/sm_32_atomic_functions.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_32_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_60_atomic_functions.hpp b/clang/test/Headers/Inputs/include/sm_60_atomic_functions.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_60_atomic_functions.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/sm_61_intrinsics.hpp b/clang/test/Headers/Inputs/include/sm_61_intrinsics.hpp
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/sm_61_intrinsics.hpp
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/Inputs/include/string.h b/clang/test/Headers/Inputs/include/string.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/string.h
@@ -0,0 +1,3 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
+void* memcpy(void* dst, const void* src, size_t num);
diff --git a/clang/test/Headers/Inputs/include/texture_indirect_functions.h b/clang/test/Headers/Inputs/include/texture_indirect_functions.h
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/texture_indirect_functions.h
@@ -0,0 +1,2 @@
+// required for __clang_cuda_runtime_wrapper.h tests
+#pragma once
diff --git a/clang/test/Headers/cuda_with_openmp.cu b/clang/test/Headers/cuda_with_openmp.cu
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/cuda_with_openmp.cu
@@ -0,0 +1,8 @@
+// Test using -x cuda -fopenmp does not clash integrated headers.
+// Reported in https://bugs.llvm.org/show_bug.cgi?id=48014
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang -x cuda -fopenmp -c %s -o - --cuda-path=%S/../Driver/Inputs/CUDA/usr/local/cuda -nocudalib -isystem %S/Inputs/include -isystem %S/../../lib/Headers -fsyntax-only
+