clang/test/Headers/nvptx_device_math_sin.c

   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-target-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-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fgpu-approx-transcendentals -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=FAST
   6 // expected-no-diagnostics
   7
   8 #include <math.h>
   9
  10 double math(float f, double d) {
  11   double r = 0;
  12 // SLOW:  call float @__nv_sinf(float
  13 // FAST:  call fast nofpclass(nan inf) float @__nv_fast_sinf(float
  14   r += sinf(f);
  15 // SLOW:  call double @__nv_sin(double
  16 // FAST:  call fast nofpclass(nan inf) double @__nv_sin(double
  17   r += sin(d);
  18   return r;
  19 }
  20
  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 }