88 lines · cpp
1// REQUIRES: nvptx-registered-target2// 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.bc3// 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 %s4// expected-no-diagnostics5 6#include <cmath>7#include <complex>8 9// CHECK: define weak {{.*}} @__muldc310// 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 {{.*}} @__mulsc315// 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 {{.*}} @__divdc320// 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 {{.*}} @__divsc329// 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 that39// doesn't start with '__' we are good :)40 41// CHECK-NOT: declare.*@[^_]42 43void test_scmplx(std::complex<float> a) {44#pragma omp target45 {46 (void)(a * (a / a));47 }48}49 50void test_dcmplx(std::complex<double> a) {51#pragma omp target52 {53 (void)(a * (a / a));54 }55}56 57template <typename T>58std::complex<T> test_template_math_calls(std::complex<T> a) {59 decltype(a) r = a;60#pragma omp target61 {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 71std::complex<float> test_scall(std::complex<float> a) {72 decltype(a) r;73#pragma omp target74 {75 r = std::sin(a);76 }77 return test_template_math_calls(r);78}79 80std::complex<double> test_dcall(std::complex<double> a) {81 decltype(a) r;82#pragma omp target83 {84 r = std::exp(a);85 }86 return test_template_math_calls(r);87}88