nvptx_device_math_complex.cpp 2.54 KB
// REQUIRES: nvptx-registered-target
// 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
// 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
// expected-no-diagnostics

#include <cmath>
#include <complex>

// CHECK: define weak {{.*}} @__muldc3
// CHECK-DAG: call i32 @__nv_isnand(
// CHECK-DAG: call i32 @__nv_isinfd(
// CHECK-DAG: call double @__nv_copysign(

// CHECK: define weak {{.*}} @__mulsc3
// CHECK-DAG: call i32 @__nv_isnanf(
// CHECK-DAG: call i32 @__nv_isinff(
// CHECK-DAG: call float @__nv_copysignf(

// CHECK: define weak {{.*}} @__divdc3
// CHECK-DAG: call i32 @__nv_isnand(
// CHECK-DAG: call i32 @__nv_isinfd(
// CHECK-DAG: call i32 @__nv_isfinited(
// CHECK-DAG: call double @__nv_copysign(
// CHECK-DAG: call double @__nv_scalbn(
// CHECK-DAG: call double @__nv_fabs(
// CHECK-DAG: call double @__nv_logb(

// CHECK: define weak {{.*}} @__divsc3
// CHECK-DAG: call i32 @__nv_isnanf(
// CHECK-DAG: call i32 @__nv_isinff(
// CHECK-DAG: call i32 @__nv_finitef(
// CHECK-DAG: call float @__nv_copysignf(
// CHECK-DAG: call float @__nv_scalbnf(
// CHECK-DAG: call float @__nv_fabsf(
// CHECK-DAG: call float @__nv_logbf(

// We actually check that there are no declarations of non-OpenMP functions.
// That is, as long as we don't call an unkown function with a name that
// doesn't start with '__' we are good :)

// CHECK-NOT: declare.*@[^_]

void test_scmplx(std::complex<float> a) {
#pragma omp target
  {
    (void)(a * (a / a));
  }
}

void test_dcmplx(std::complex<double> a) {
#pragma omp target
  {
    (void)(a * (a / a));
  }
}

template <typename T>
std::complex<T> test_template_math_calls(std::complex<T> a) {
  decltype(a) r = a;
#pragma omp target
  {
    r = std::sin(r);
    r = std::cos(r);
    r = std::exp(r);
    r = std::atan(r);
    r = std::acos(r);
  }
  return r;
}

std::complex<float> test_scall(std::complex<float> a) {
  decltype(a) r;
#pragma omp target
  {
    r = std::sin(a);
  }
  return test_template_math_calls(r);
}

std::complex<double> test_dcall(std::complex<double> a) {
  decltype(a) r;
#pragma omp target
  {
    r = std::exp(a);
  }
  return test_template_math_calls(r);
}