1 // REQUIRES: nvptx-registered-target 2 // RUN: %clang_cc1 -verify -internal-isystem %S/Inputs/include -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc 3 // RUN: %clang_cc1 -verify -internal-isystem %S/../../lib/Headers/openmp_wrappers -include __clang_openmp_device_functions.h -internal-isystem %S/Inputs/include -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -aux-triple powerpc64le-unknown-unknown -o - | FileCheck %s 4 // expected-no-diagnostics 5 6 #include <cmath> 7 #include <complex> 8 9 // CHECK: define weak {{.*}} @__muldc3 10 // CHECK-DAG: call i32 @__nv_isnand( 11 // CHECK-DAG: call i32 @__nv_isinfd( 12 // CHECK-DAG: call double @__nv_copysign( 13 14 // CHECK: define weak {{.*}} @__mulsc3 15 // CHECK-DAG: call i32 @__nv_isnanf( 16 // CHECK-DAG: call i32 @__nv_isinff( 17 // CHECK-DAG: call float @__nv_copysignf( 18 19 // CHECK: define weak {{.*}} @__divdc3 20 // CHECK-DAG: call i32 @__nv_isnand( 21 // CHECK-DAG: call i32 @__nv_isinfd( 22 // CHECK-DAG: call i32 @__nv_isfinited( 23 // CHECK-DAG: call double @__nv_copysign( 24 // CHECK-DAG: call double @__nv_scalbn( 25 // CHECK-DAG: call double @__nv_fabs( 26 // CHECK-DAG: call double @__nv_logb( 27 28 // CHECK: define weak {{.*}} @__divsc3 29 // CHECK-DAG: call i32 @__nv_isnanf( 30 // CHECK-DAG: call i32 @__nv_isinff( 31 // CHECK-DAG: call i32 @__nv_finitef( 32 // CHECK-DAG: call float @__nv_copysignf( 33 // CHECK-DAG: call float @__nv_scalbnf( 34 // CHECK-DAG: call float @__nv_fabsf( 35 // CHECK-DAG: call float @__nv_logbf( 36 37 // We actually check that there are no declarations of non-OpenMP functions. 38 // That is, as long as we don't call an unkown function with a name that 39 // doesn't start with '__' we are good :) 40 41 // CHECK-NOT: declare.*@[^_] 42 test_scmplx(std::complex<float> a)43void test_scmplx(std::complex<float> a) { 44 #pragma omp target 45 { 46 (void)(a * (a / a)); 47 } 48 } 49 test_dcmplx(std::complex<double> a)50void test_dcmplx(std::complex<double> a) { 51 #pragma omp target 52 { 53 (void)(a * (a / a)); 54 } 55 } 56 57 template <typename T> test_template_math_calls(std::complex<T> a)58std::complex<T> test_template_math_calls(std::complex<T> a) { 59 decltype(a) r = a; 60 #pragma omp target 61 { 62 r = std::sin(r); 63 r = std::cos(r); 64 r = std::exp(r); 65 r = std::atan(r); 66 r = std::acos(r); 67 } 68 return r; 69 } 70 test_scall(std::complex<float> a)71std::complex<float> test_scall(std::complex<float> a) { 72 decltype(a) r; 73 #pragma omp target 74 { 75 r = std::sin(a); 76 } 77 return test_template_math_calls(r); 78 } 79 test_dcall(std::complex<double> a)80std::complex<double> test_dcall(std::complex<double> a) { 81 decltype(a) r; 82 #pragma omp target 83 { 84 r = std::exp(a); 85 } 86 return test_template_math_calls(r); 87 } 88