diff --git a/clang/lib/Headers/openmp_wrappers/complex b/clang/lib/Headers/openmp_wrappers/complex --- a/clang/lib/Headers/openmp_wrappers/complex +++ b/clang/lib/Headers/openmp_wrappers/complex @@ -45,7 +45,7 @@ #ifndef _LIBCPP_STD_VER #pragma omp begin declare variant match( \ - device = {arch(nvptx, nvptx64)}, \ + device = {arch(amdgcn, nvptx, nvptx64)}, \ implementation = {extension(match_any, allow_templates)}) #include diff --git a/clang/test/Headers/amdgcn-openmp-device-math-complex.cpp b/clang/test/Headers/amdgcn-openmp-device-math-complex.cpp new file mode 100644 --- /dev/null +++ b/clang/test/Headers/amdgcn-openmp-device-math-complex.cpp @@ -0,0 +1,85 @@ +// RUN: %clang_cc1 -verify -internal-isystem %S/Inputs/include -fopenmp -x c++ -triple x86_64-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-x86-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 amdgcn-amd-amdhsa -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -aux-triple x86_64-unknown-unknown -o - | FileCheck %s +// expected-no-diagnostics + +#include +#include + +// CHECK: define weak {{.*}} @__muldc3 +// CHECK-DAG: call i32 @__ocml_isnan_f64( +// CHECK-DAG: call i32 @__ocml_isinf_f64( + +// CHECK: define weak {{.*}} @__mulsc3 +// CHECK-DAG: call i32 @__ocml_isnan_f32( +// CHECK-DAG: call i32 @__ocml_isinf_f32( +// CHECK-DAG: call float @__ocml_copysign_f32( + +// CHECK: define weak {{.*}} @__divdc3 +// CHECK-DAG: call i32 @__ocml_isnan_f64( +// CHECK-DAG: call i32 @__ocml_isinf_f64( +// CHECK-DAG: call i32 @__ocml_isfinite_f64( +// CHECK-DAG: call double @__ocml_copysign_f64( +// CHECK-DAG: call double @__ocml_scalbn_f64( +// CHECK-DAG: call double @__ocml_fabs_f64( +// CHECK-DAG: call double @__ocml_logb_f64( + +// CHECK: define weak {{.*}} @__divsc3 +// CHECK-DAG: call i32 @__ocml_isnan_f32( +// CHECK-DAG: call i32 @__ocml_isinf_f32( +// CHECK-DAG: call i32 @__ocml_isfinite_f32( +// CHECK-DAG: call float @__ocml_copysign_f32( +// CHECK-DAG: call float @__ocml_scalbn_f32( +// CHECK-DAG: call float @__ocml_fabs_f32( +// CHECK-DAG: call float @__ocml_logb_f32( + +// 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 a) { +#pragma omp target + { + (void)(a * (a / a)); + } +} + +void test_dcmplx(std::complex a) { +#pragma omp target + { + (void)(a * (a / a)); + } +} + +template +std::complex test_template_math_calls(std::complex 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 test_scall(std::complex a) { + decltype(a) r; +#pragma omp target + { + r = std::sin(a); + } + return test_template_math_calls(r); +} + +std::complex test_dcall(std::complex a) { + decltype(a) r; +#pragma omp target + { + r = std::exp(a); + } + return test_template_math_calls(r); +}