1d999cbc9SJohannes Doerfert // REQUIRES: nvptx-registered-target 2d999cbc9SJohannes Doerfert // 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*63ca93c7SSergio Afonso // 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-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -aux-triple powerpc64le-unknown-unknown -o - | FileCheck %s 4d999cbc9SJohannes Doerfert // expected-no-diagnostics 5d999cbc9SJohannes Doerfert 656069b5cSJohannes Doerfert #include <cmath> 7d999cbc9SJohannes Doerfert #include <complex> 8d999cbc9SJohannes Doerfert 9b5667d00SJohannes Doerfert // CHECK: define weak {{.*}} @__muldc3 10b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isnand( 11b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isinfd( 12b5667d00SJohannes Doerfert // CHECK-DAG: call double @__nv_copysign( 13d999cbc9SJohannes Doerfert 14b5667d00SJohannes Doerfert // CHECK: define weak {{.*}} @__mulsc3 15b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isnanf( 16b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isinff( 17b5667d00SJohannes Doerfert // CHECK-DAG: call float @__nv_copysignf( 18b5667d00SJohannes Doerfert 19b5667d00SJohannes Doerfert // CHECK: define weak {{.*}} @__divdc3 20b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isnand( 21b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isinfd( 22b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isfinited( 23b5667d00SJohannes Doerfert // CHECK-DAG: call double @__nv_copysign( 24b5667d00SJohannes Doerfert // CHECK-DAG: call double @__nv_scalbn( 25b5667d00SJohannes Doerfert // CHECK-DAG: call double @__nv_fabs( 26b5667d00SJohannes Doerfert // CHECK-DAG: call double @__nv_logb( 27b5667d00SJohannes Doerfert 28b5667d00SJohannes Doerfert // CHECK: define weak {{.*}} @__divsc3 29b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isnanf( 30b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_isinff( 31b5667d00SJohannes Doerfert // CHECK-DAG: call i32 @__nv_finitef( 32b5667d00SJohannes Doerfert // CHECK-DAG: call float @__nv_copysignf( 33d999cbc9SJohannes Doerfert // CHECK-DAG: call float @__nv_scalbnf( 34b5667d00SJohannes Doerfert // CHECK-DAG: call float @__nv_fabsf( 35b5667d00SJohannes Doerfert // CHECK-DAG: call float @__nv_logbf( 36b5667d00SJohannes Doerfert 3756069b5cSJohannes Doerfert // We actually check that there are no declarations of non-OpenMP functions. 3856069b5cSJohannes Doerfert // That is, as long as we don't call an unkown function with a name that 3956069b5cSJohannes Doerfert // doesn't start with '__' we are good :) 4056069b5cSJohannes Doerfert 4156069b5cSJohannes Doerfert // CHECK-NOT: declare.*@[^_] 4256069b5cSJohannes Doerfert test_scmplx(std::complex<float> a)43d999cbc9SJohannes Doerfertvoid test_scmplx(std::complex<float> a) { 44d999cbc9SJohannes Doerfert #pragma omp target 45d999cbc9SJohannes Doerfert { 46d999cbc9SJohannes Doerfert (void)(a * (a / a)); 47d999cbc9SJohannes Doerfert } 48d999cbc9SJohannes Doerfert } 49d999cbc9SJohannes Doerfert test_dcmplx(std::complex<double> a)50d999cbc9SJohannes Doerfertvoid test_dcmplx(std::complex<double> a) { 51d999cbc9SJohannes Doerfert #pragma omp target 52d999cbc9SJohannes Doerfert { 53d999cbc9SJohannes Doerfert (void)(a * (a / a)); 54d999cbc9SJohannes Doerfert } 55d999cbc9SJohannes Doerfert } 5656069b5cSJohannes Doerfert 5756069b5cSJohannes Doerfert template <typename T> test_template_math_calls(std::complex<T> a)5856069b5cSJohannes Doerfertstd::complex<T> test_template_math_calls(std::complex<T> a) { 5956069b5cSJohannes Doerfert decltype(a) r = a; 6056069b5cSJohannes Doerfert #pragma omp target 6156069b5cSJohannes Doerfert { 6256069b5cSJohannes Doerfert r = std::sin(r); 6356069b5cSJohannes Doerfert r = std::cos(r); 6456069b5cSJohannes Doerfert r = std::exp(r); 6556069b5cSJohannes Doerfert r = std::atan(r); 6656069b5cSJohannes Doerfert r = std::acos(r); 6756069b5cSJohannes Doerfert } 6856069b5cSJohannes Doerfert return r; 6956069b5cSJohannes Doerfert } 7056069b5cSJohannes Doerfert test_scall(std::complex<float> a)7156069b5cSJohannes Doerfertstd::complex<float> test_scall(std::complex<float> a) { 7256069b5cSJohannes Doerfert decltype(a) r; 7356069b5cSJohannes Doerfert #pragma omp target 7456069b5cSJohannes Doerfert { 7556069b5cSJohannes Doerfert r = std::sin(a); 7656069b5cSJohannes Doerfert } 7756069b5cSJohannes Doerfert return test_template_math_calls(r); 7856069b5cSJohannes Doerfert } 7956069b5cSJohannes Doerfert test_dcall(std::complex<double> a)8056069b5cSJohannes Doerfertstd::complex<double> test_dcall(std::complex<double> a) { 8156069b5cSJohannes Doerfert decltype(a) r; 8256069b5cSJohannes Doerfert #pragma omp target 8356069b5cSJohannes Doerfert { 8456069b5cSJohannes Doerfert r = std::exp(a); 8556069b5cSJohannes Doerfert } 8656069b5cSJohannes Doerfert return test_template_math_calls(r); 8756069b5cSJohannes Doerfert } 88