diff options
author | Gheorghe-Teodor Bercea <gheorghe-teod.bercea@ibm.com> | 2019-05-06 18:19:15 +0000 |
---|---|---|
committer | Gheorghe-Teodor Bercea <gheorghe-teod.bercea@ibm.com> | 2019-05-06 18:19:15 +0000 |
commit | 1c40858fc28adb9b87e3f8f8bd5b4319defa9624 (patch) | |
tree | 2a06151ae9647e5aef344f6098584a13615818c3 /test | |
parent | 336ca528612aa91867d963f22b4c3c6a6d61874f (diff) |
[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
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@360063 91177308-0d34-0410-b5e6-96231b3b80d8
Diffstat (limited to 'test')
-rw-r--r-- | test/Driver/openmp-offload-gpu.c | 5 | ||||
-rw-r--r-- | test/Headers/Inputs/include/cmath | 5 | ||||
-rw-r--r-- | test/Headers/Inputs/include/limits | 10 | ||||
-rw-r--r-- | test/Headers/Inputs/include/math.h | 4 | ||||
-rw-r--r-- | test/Headers/nvptx_device_cmath_functions.c | 21 | ||||
-rw-r--r-- | test/Headers/nvptx_device_cmath_functions.cpp | 21 | ||||
-rw-r--r-- | test/Headers/nvptx_device_math_functions.c | 21 | ||||
-rw-r--r-- | test/Headers/nvptx_device_math_functions.cpp | 21 |
8 files changed, 108 insertions, 0 deletions
diff --git a/test/Driver/openmp-offload-gpu.c b/test/Driver/openmp-offload-gpu.c index 7a4dd95e54..3d2ac4525f 100644 --- a/test/Driver/openmp-offload-gpu.c +++ b/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/test/Headers/Inputs/include/cmath b/test/Headers/Inputs/include/cmath new file mode 100644 index 0000000000..4ba1795137 --- /dev/null +++ b/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/test/Headers/Inputs/include/limits b/test/Headers/Inputs/include/limits new file mode 100644 index 0000000000..fbee11ef11 --- /dev/null +++ b/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/test/Headers/Inputs/include/math.h b/test/Headers/Inputs/include/math.h index 6f70f09bee..4ba1795137 100644 --- a/test/Headers/Inputs/include/math.h +++ b/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/test/Headers/nvptx_device_cmath_functions.c b/test/Headers/nvptx_device_cmath_functions.c new file mode 100644 index 0000000000..aa55c1eb65 --- /dev/null +++ b/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/test/Headers/nvptx_device_cmath_functions.cpp b/test/Headers/nvptx_device_cmath_functions.cpp new file mode 100644 index 0000000000..a5b4377413 --- /dev/null +++ b/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/test/Headers/nvptx_device_math_functions.c b/test/Headers/nvptx_device_math_functions.c new file mode 100644 index 0000000000..733ad52bd1 --- /dev/null +++ b/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/test/Headers/nvptx_device_math_functions.cpp b/test/Headers/nvptx_device_math_functions.cpp new file mode 100644 index 0000000000..9753011243 --- /dev/null +++ b/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); + } +} |