// Test calling of device math functions. ///==========================================================================/// // REQUIRES: nvptx-registered-target // 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 // RUN: %clang_cc1 -x c -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -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=CHECK-C // 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 // RUN: %clang_cc1 -x c++ -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -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=CHECK-CXX #ifdef __cplusplus #include #include #else #include #include #endif void test_sqrt(double a1) { #pragma omp target { // CHECK-C: call double @__nv_sqrt(double // CHECK-CXX: call noundef double @__nv_sqrt(double double l1 = sqrt(a1); // CHECK-C: call double @__nv_pow(double // CHECK-CXX: call noundef double @__nv_pow(double double l2 = pow(a1, a1); // CHECK-C: call double @__nv_modf(double // CHECK-CXX: call noundef double @__nv_modf(double double l3 = modf(a1 + 3.5, &a1); // CHECK-C: call double @__nv_fabs(double // CHECK-CXX: call noundef double @__nv_fabs(double double l4 = fabs(a1); // CHECK-C: call i32 @__nv_abs(i32 // CHECK-CXX: call noundef i32 @__nv_abs(i32 double l5 = abs((int)a1); } }