diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -1249,7 +1249,7 @@
   // 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.
diff --git a/clang/lib/Headers/__clang_hip_cmath.h b/clang/lib/Headers/__clang_hip_cmath.h
--- a/clang/lib/Headers/__clang_hip_cmath.h
+++ b/clang/lib/Headers/__clang_hip_cmath.h
@@ -10,7 +10,7 @@
 #ifndef __CLANG_HIP_CMATH_H__
 #define __CLANG_HIP_CMATH_H__
 
-#if !defined(__HIP__)
+#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
 #endif
 
@@ -25,7 +25,13 @@
 #endif // !defined(__HIPCC_RTC__)
 
 #pragma push_macro("__DEVICE__")
+#ifdef __OPENMP_AMDGCN__
+#define __DEVICE__ static constexpr __attribute__((always_inline,nothrow))
+#define __constant__ __attribute__((constant))
+#else
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#endif // __OPENMP_AMDGCN__
+
 
 // Start with functions that cannot be defined by DEF macros below.
 #if defined(__cplusplus)
diff --git a/clang/lib/Headers/__clang_hip_math.h b/clang/lib/Headers/__clang_hip_math.h
--- a/clang/lib/Headers/__clang_hip_math.h
+++ b/clang/lib/Headers/__clang_hip_math.h
@@ -9,7 +9,7 @@
 #ifndef __CLANG_HIP_MATH_H__
 #define __CLANG_HIP_MATH_H__
 
-#if !defined(__HIP__)
+#if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
 #endif
 
@@ -19,18 +19,27 @@
 #endif
 #include <limits.h>
 #include <stdint.h>
-#endif // __HIPCC_RTC__
+#endif // !defined(__HIPCC_RTC__)
 
 #pragma push_macro("__DEVICE__")
+
+#ifdef __OPENMP_AMDGCN__
+#define __DEVICE__ static inline __attribute__((always_inline,nothrow))
+#else
 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
+#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 // __OPENMP_AMDGCN__
 
 #if defined (__cplusplus) && __cplusplus < 201103L
 // emulate static_assert on type sizes
@@ -1262,7 +1271,7 @@
 __DEVICE__
 double min(double __x, double __y) { return fmin(__x, __y); }
 
-#if !defined(__HIPCC_RTC__)
+#if !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
 __host__ inline static int min(int __arg1, int __arg2) {
   return std::min(__arg1, __arg2);
 }
@@ -1270,7 +1279,7 @@
 __host__ inline static int max(int __arg1, int __arg2) {
   return std::max(__arg1, __arg2);
 }
-#endif // __HIPCC_RTC__
+#endif // !defined(__HIPCC_RTC__) && !defined(__OPENMP_AMDGCN__)
 #endif
 
 #pragma pop_macro("__DEVICE__")
diff --git a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
--- a/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
+++ b/clang/lib/Headers/openmp_wrappers/__clang_openmp_device_functions.h
@@ -81,4 +81,34 @@
 #pragma pop_macro("OPENMP_NOEXCEPT")
 #endif
 
+#pragma omp begin declare variant match(                                       \
+    device = {arch(amdgcn)}, implementation = {extension(match_any)})
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#define __OPENMP_AMDGCN__
+
+#define __device__ __attribute__((device))
+
+// Import types which will be used by __clang_hip_libdevice_declares.h
+#ifndef __cplusplus
+#include <stdbool.h>
+#include <stdint.h>
+#endif
+
+/// Include declarations for libdevice functions.
+#include <__clang_hip_libdevice_declares.h>
+
+#undef __device__
+#undef __OPENMP_AMDGCN__
+
+#ifdef __cplusplus
+} // extern "C"
+
+#endif
+
+#pragma omp end declare variant
+
 #endif
diff --git a/clang/lib/Headers/openmp_wrappers/cmath b/clang/lib/Headers/openmp_wrappers/cmath
--- a/clang/lib/Headers/openmp_wrappers/cmath
+++ b/clang/lib/Headers/openmp_wrappers/cmath
@@ -75,4 +75,14 @@
 
 #pragma omp end declare variant
 
+#ifdef __AMDGCN__
+#pragma omp begin declare variant match( \
+  device = {arch(amdgcn)}, implementation = {extension(match_any, allow_templates)})
+#define __OPENMP_AMDGCN__
+#include <__clang_hip_cmath.h>
+#undef __OPENMP_AMDGCN__
+
+
+#pragma omp end declare variant
+#endif // __AMDGCN__
 #endif
diff --git a/clang/lib/Headers/openmp_wrappers/math.h b/clang/lib/Headers/openmp_wrappers/math.h
--- a/clang/lib/Headers/openmp_wrappers/math.h
+++ b/clang/lib/Headers/openmp_wrappers/math.h
@@ -48,4 +48,13 @@
 
 #pragma omp end declare variant
 
+#pragma omp begin declare variant match(                                       \
+    device = {arch(amdgcn)}, implementation = {extension(match_any)})
+
+#define __OPENMP_AMDGCN__
+#include <__clang_hip_math.h>
+#undef __OPENMP_AMDGCN__
+
+#pragma omp end declare variant
+
 #endif
diff --git a/clang/test/Headers/Inputs/include/algorithm b/clang/test/Headers/Inputs/include/algorithm
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/algorithm
@@ -0,0 +1,6 @@
+#pragma once
+
+namespace std {
+ template<class T> constexpr const T& min(const T& a, const T& b);
+ template<class T> constexpr const T& max(const T& a, const T& b);
+}
\ No newline at end of file
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
@@ -12,6 +12,7 @@
 extern float fabs (float __x) __attribute__ ((__const__)) ;
 #endif
 
+#ifndef __AMDGCN__
 namespace std
 {
 
@@ -29,3 +30,5 @@
 double abs(double __x) { return fabs(__x); }
 
 }
+
+#endif
diff --git a/clang/test/Headers/Inputs/include/utility b/clang/test/Headers/Inputs/include/utility
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/utility
@@ -0,0 +1,2 @@
+#pragma once
+
diff --git a/clang/test/Headers/amdgcn_openmp_device_math.c b/clang/test/Headers/amdgcn_openmp_device_math.c
new file mode 100644
--- /dev/null
+++ b/clang/test/Headers/amdgcn_openmp_device_math.c
@@ -0,0 +1,51 @@
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK-C,CHECK
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -x c++ -fopenmp -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -o - | FileCheck %s --check-prefixes=CHECK-CPP,CHECK
+
+#ifdef __cplusplus
+#include <cmath>
+#else
+#include <math.h>
+#endif
+
+void test_math_f64(double x) {
+  // CHECK-LABEL: define {{.*}}test_math_f64
+  #pragma omp target
+  {
+    // CHECK: call double @__ocml_sin_f64
+    double l1 = sin(x);
+    // CHECK: call double @__ocml_cos_f64
+    double l2 = cos(x);
+    // CHECK: call double @__ocml_fabs_f64
+    double l3 = fabs(x);
+  }
+}
+
+void test_math_f32(float x) {
+  // CHECK-LABEL: define {{.*}}test_math_f32
+  #pragma omp target
+  {
+    // CHECK-C: call double @__ocml_sin_f64
+    // CHECK-CPP: call float @__ocml_sin_f32
+    float l1 = sin(x);
+    // CHECK-C: call double @__ocml_cos_f64
+    // CHECK-CPP: call float @__ocml_cos_f32
+    float l2 = cos(x);
+    // CHECK-C: call double @__ocml_fabs_f64
+    // CHECK-CPP: call float @__ocml_fabs_f32
+    float l3 = fabs(x);
+  }
+}
+void test_math_f32_suffix(float x) {
+  // CHECK-LABEL: define {{.*}}test_math_f32_suffix
+  #pragma omp target
+  {
+    // CHECK: call float @__ocml_sin_f32
+    float l1 = sinf(x);
+    // CHECK: call float @__ocml_cos_f32
+    float l2 = cosf(x);
+    // CHECK: call float @__ocml_fabs_f32
+    float l3 = fabsf(x);
+  }
+}
diff --git a/clang/test/Headers/openmp_device_math_isnan.cpp b/clang/test/Headers/openmp_device_math_isnan.cpp
--- a/clang/test/Headers/openmp_device_math_isnan.cpp
+++ b/clang/test/Headers/openmp_device_math_isnan.cpp
@@ -21,14 +21,14 @@
 double math(float f, double d) {
   double r = 0;
   // INT_RETURN: call i32 @__nv_isnanf(float
-  // AMD_INT_RETURN: call i32 @_{{.*}}isnanf(float
+  // AMD_INT_RETURN: call i32 @__ocml_isnan_f32(float
   // BOOL_RETURN: call i32 @__nv_isnanf(float
-  // AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnanf(float
+  // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f32(float
   r += std::isnan(f);
   // INT_RETURN: call i32 @__nv_isnand(double
-  // AMD_INT_RETURN: call i32 @_{{.*}}isnand(double
+  // AMD_INT_RETURN: call i32 @__ocml_isnan_f64(double
   // BOOL_RETURN: call i32 @__nv_isnand(double
-  // AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnand(double
+  // AMD_BOOL_RETURN: call i32 @__ocml_isnan_f64(double
   r += std::isnan(d);
   return r;
 }