8 files changed, 108 insertions, 0 deletions
diff --git a/clang/test/Driver/openmp-offload-gpu.c b/clang/test/Driver/openmp-offload-gpu.c
index 7a4dd95e541..3d2ac4525f6 100644
--- a/clang/test/Driver/openmp-offload-gpu.c
+++ b/clang/test/Driver/openmp-offload-gpu.c
@@ -278,3 +278,8 @@
 // RUN:   | FileCheck -check-prefix=CUDA_RED_RECS %s
 // CUDA_RED_RECS: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
 // CUDA_RED_RECS-SAME: "-fopenmp-cuda-teams-reduction-recs-num=2048"
+
+// RUN:   %clang -### -no-canonical-prefixes -fopenmp=libomp -fopenmp-targets=nvptx64-nvidia-cuda %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=OPENMP_NVPTX_WRAPPERS %s
+// OPENMP_NVPTX_WRAPPERS: clang{{.*}}"-cc1"{{.*}}"-triple" "nvptx64-nvidia-cuda"
+// OPENMP_NVPTX_WRAPPERS-SAME: "-internal-isystem" "{{.*}}openmp_wrappers"
diff --git a/clang/test/Headers/Inputs/include/cmath b/clang/test/Headers/Inputs/include/cmath
new file mode 100644
index 00000000000..4ba17951378
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/cmath
@@ -0,0 +1,5 @@
+#pragma once
+
+double sqrt(double);
+double pow(double, double);
+double modf(double, double*);
diff --git a/clang/test/Headers/Inputs/include/limits b/clang/test/Headers/Inputs/include/limits
new file mode 100644
index 00000000000..fbee11ef118
--- /dev/null
+++ b/clang/test/Headers/Inputs/include/limits
@@ -0,0 +1,10 @@
+#pragma once
+
+namespace std
+{
+struct __numeric_limits_base
+  {};
+template<typename _Tp>
+  struct numeric_limits : public __numeric_limits_base
+    {};
+}
diff --git a/clang/test/Headers/Inputs/include/math.h b/clang/test/Headers/Inputs/include/math.h
index 6f70f09beec..4ba17951378 100644
--- a/clang/test/Headers/Inputs/include/math.h
+++ b/clang/test/Headers/Inputs/include/math.h
@@ -1 +1,5 @@
 #pragma once
+
+double sqrt(double);
+double pow(double, double);
+double modf(double, double*);
diff --git a/clang/test/Headers/nvptx_device_cmath_functions.c b/clang/test/Headers/nvptx_device_cmath_functions.c
new file mode 100644
index 00000000000..aa55c1eb652
--- /dev/null
+++ b/clang/test/Headers/nvptx_device_cmath_functions.c
@@ -0,0 +1,21 @@
+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include cmath -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <cmath>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}
diff --git a/clang/test/Headers/nvptx_device_cmath_functions.cpp b/clang/test/Headers/nvptx_device_cmath_functions.cpp
new file mode 100644
index 00000000000..a5b43774133
--- /dev/null
+++ b/clang/test/Headers/nvptx_device_cmath_functions.cpp
@@ -0,0 +1,21 @@
+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include cmath -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include cmath -internal-isystem %S/Inputs/include -include stdlib.h -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <cmath>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}
diff --git a/clang/test/Headers/nvptx_device_math_functions.c b/clang/test/Headers/nvptx_device_math_functions.c
new file mode 100644
index 00000000000..733ad52bd1d
--- /dev/null
+++ b/clang/test/Headers/nvptx_device_math_functions.c
@@ -0,0 +1,21 @@
+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include math.h -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <math.h>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}
diff --git a/clang/test/Headers/nvptx_device_math_functions.cpp b/clang/test/Headers/nvptx_device_math_functions.cpp
new file mode 100644
index 00000000000..97530112436
--- /dev/null
+++ b/clang/test/Headers/nvptx_device_math_functions.cpp
@@ -0,0 +1,21 @@
+// Test calling of device math functions.
+///==========================================================================///
+
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -internal-isystem %S/Inputs/include -include math.h -x c++ -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -internal-isystem %S/../../lib/Headers/openmp_wrappers -include math.h -internal-isystem %S/Inputs/include -include stdlib.h -include limits -x c++ -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck -check-prefix CHECK-YES %s
+
+#include <math.h>
+
+void test_sqrt(double a1) {
+  #pragma omp target
+  {
+    // CHECK-YES: call double @__nv_sqrt(double
+    double l1 = sqrt(a1);
+    // CHECK-YES: call double @__nv_pow(double
+    double l2 = pow(a1, a1);
+    // CHECK-YES: call double @__nv_modf(double
+    double l3 = modf(a1 + 3.5, &a1);
+  }
+}