xref: /llvm-project/clang/test/Headers/nvptx_device_math_complex.cpp (revision 63ca93c7d1d1ee91281ff7ccdbd7014151319324)
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 Doerfert void 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 Doerfert void 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 Doerfert std::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 Doerfert std::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 Doerfert std::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