[OpenMP][Clang] Support for target math functions

Summary:
In this patch we propose a temporary solution to resolving math functions for the NVPTX toolchain, temporary until OpenMP variant is supported by Clang.

We intercept the inclusion of math.h and cmath headers and if we are in the OpenMP-NVPTX case, we re-use CUDA's math function resolution mechanism.

Authors:
@gtbercea
@jdoerfert

Reviewers: hfinkel, caomhin, ABataev, tra

Reviewed By: hfinkel, ABataev, tra

Subscribers: mgorny, guansong, cfe-commits, jdoerfert

Tags: #clang

Differential Revision: https://reviews.llvm.org/D61399

llvm-svn: 360063

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);
+  }
+}