1 // REQUIRES: nvptx-registered-target
2 // RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
3 // RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 %s --check-prefix=SLOW
4 // RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast
5 // RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -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 - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=FAST
6 // expected-no-diagnostics
7
8 #include <cmath>
9
math(float f,double d)10 double math(float f, double d) {
11 double r = 0;
12 // SLOW: call float @__nv_sinf(float
13 // FAST: call fast float @__nv_fast_sinf(float
14 r += sin(f);
15 // SLOW: call double @__nv_sin(double
16 // FAST: call fast double @__nv_sin(double
17 r += sin(d);
18 return r;
19 }
20
foo(float f,double d,long double ld)21 long double foo(float f, double d, long double ld) {
22 double r = ld;
23 r += math(f, d);
24 #pragma omp target map(r)
25 { r += math(f, d); }
26 return r;
27 }
28