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 <complex>
7
8 // CHECK: define weak {{.*}} @__muldc3
9 // CHECK-DAG: call i32 @__nv_isnand(
10 // CHECK-DAG: call i32 @__nv_isinfd(
11 // CHECK-DAG: call double @__nv_copysign(
12
13 // CHECK: define weak {{.*}} @__mulsc3
14 // CHECK-DAG: call i32 @__nv_isnanf(
15 // CHECK-DAG: call i32 @__nv_isinff(
16 // CHECK-DAG: call float @__nv_copysignf(
17
18 // CHECK: define weak {{.*}} @__divdc3
19 // CHECK-DAG: call i32 @__nv_isnand(
20 // CHECK-DAG: call i32 @__nv_isinfd(
21 // CHECK-DAG: call i32 @__nv_isfinited(
22 // CHECK-DAG: call double @__nv_copysign(
23 // CHECK-DAG: call double @__nv_scalbn(
24 // CHECK-DAG: call double @__nv_fabs(
25 // CHECK-DAG: call double @__nv_logb(
26
27 // CHECK: define weak {{.*}} @__divsc3
28 // CHECK-DAG: call i32 @__nv_isnanf(
29 // CHECK-DAG: call i32 @__nv_isinff(
30 // CHECK-DAG: call i32 @__nv_finitef(
31 // CHECK-DAG: call float @__nv_copysignf(
32 // CHECK-DAG: call float @__nv_scalbnf(
33 // CHECK-DAG: call float @__nv_fabsf(
34 // CHECK-DAG: call float @__nv_logbf(
35
test_scmplx(std::complex<float> a)36 void test_scmplx(std::complex<float> a) {
37 #pragma omp target
38 {
39 (void)(a * (a / a));
40 }
41 }
42
test_dcmplx(std::complex<double> a)43 void test_dcmplx(std::complex<double> a) {
44 #pragma omp target
45 {
46 (void)(a * (a / a));
47 }
48 }
49