diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -154,6 +154,7 @@ openmp_wrappers/complex.h openmp_wrappers/complex openmp_wrappers/__clang_openmp_device_functions.h + openmp_wrappers/complex_cmath.h openmp_wrappers/new ) 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 @@ -23,3 +23,28 @@ // Grab the host header too. #include_next + + +#ifdef __cplusplus + +// If we are compiling against libc++, the macro _LIBCPP_STD_VER should be set +// after including above. Since the complex header we use is a +// simplified version of the libc++, we don't need it in this case. If we +// compile against libstdc++, or any other standard library, we will overload +// the (hopefully template) functions in the header with the ones we +// got from libc++ which decomposes math functions, like `std::sin`, into +// arithmetic and calls to non-complex functions, all of which we can then +// handle. +#ifndef _LIBCPP_STD_VER + +#pragma omp begin declare variant match( \ + device = {arch(nvptx, nvptx64)}, \ + implementation = {extension(match_any, allow_templates)}) + +#include + +#pragma omp end declare variant + +#endif + +#endif diff --git a/clang/lib/Headers/openmp_wrappers/complex_cmath.h b/clang/lib/Headers/openmp_wrappers/complex_cmath.h new file mode 100644 --- /dev/null +++ b/clang/lib/Headers/openmp_wrappers/complex_cmath.h @@ -0,0 +1,388 @@ +//===------------------------- __complex_cmath.h --------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// std::complex header copied from the libcxx source and simplified for use in +// OpenMP target offload regions. +// +//===----------------------------------------------------------------------===// + +#ifndef _OPENMP +#error "This file is for OpenMP compilation only." +#endif + +#ifndef __cplusplus +#error "This file is for C++ compilation only." +#endif + +#ifndef _LIBCPP_COMPLEX +#define _LIBCPP_COMPLEX + +#include +#include + +#define __DEVICE__ static constexpr __attribute__((nothrow)) + +namespace std { + +// abs + +template __DEVICE__ _Tp abs(const std::complex<_Tp> &__c) { + return hypot(__c.real(), __c.imag()); +} + +// arg + +template __DEVICE__ _Tp arg(const std::complex<_Tp> &__c) { + return atan2(__c.imag(), __c.real()); +} + +template +typename enable_if::value || is_same<_Tp, double>::value, + double>::type +arg(_Tp __re) { + return atan2(0., __re); +} + +template +typename enable_if::value, float>::type arg(_Tp __re) { + return atan2f(0.F, __re); +} + +// norm + +template __DEVICE__ _Tp norm(const std::complex<_Tp> &__c) { + if (std::isinf(__c.real())) + return abs(__c.real()); + if (std::isinf(__c.imag())) + return abs(__c.imag()); + return __c.real() * __c.real() + __c.imag() * __c.imag(); +} + +// conj + +template std::complex<_Tp> conj(const std::complex<_Tp> &__c) { + return std::complex<_Tp>(__c.real(), -__c.imag()); +} + +// proj + +template std::complex<_Tp> proj(const std::complex<_Tp> &__c) { + std::complex<_Tp> __r = __c; + if (std::isinf(__c.real()) || std::isinf(__c.imag())) + __r = std::complex<_Tp>(INFINITY, copysign(_Tp(0), __c.imag())); + return __r; +} + +// polar + +template +complex<_Tp> polar(const _Tp &__rho, const _Tp &__theta = _Tp()) { + if (std::isnan(__rho) || signbit(__rho)) + return std::complex<_Tp>(_Tp(NAN), _Tp(NAN)); + if (std::isnan(__theta)) { + if (std::isinf(__rho)) + return std::complex<_Tp>(__rho, __theta); + return std::complex<_Tp>(__theta, __theta); + } + if (std::isinf(__theta)) { + if (std::isinf(__rho)) + return std::complex<_Tp>(__rho, _Tp(NAN)); + return std::complex<_Tp>(_Tp(NAN), _Tp(NAN)); + } + _Tp __x = __rho * cos(__theta); + if (std::isnan(__x)) + __x = 0; + _Tp __y = __rho * sin(__theta); + if (std::isnan(__y)) + __y = 0; + return std::complex<_Tp>(__x, __y); +} + +// log + +template std::complex<_Tp> log(const std::complex<_Tp> &__x) { + return std::complex<_Tp>(log(abs(__x)), arg(__x)); +} + +// log10 + +template std::complex<_Tp> log10(const std::complex<_Tp> &__x) { + return log(__x) / log(_Tp(10)); +} + +// sqrt + +template +__DEVICE__ std::complex<_Tp> sqrt(const std::complex<_Tp> &__x) { + if (std::isinf(__x.imag())) + return std::complex<_Tp>(_Tp(INFINITY), __x.imag()); + if (std::isinf(__x.real())) { + if (__x.real() > _Tp(0)) + return std::complex<_Tp>(__x.real(), std::isnan(__x.imag()) + ? __x.imag() + : copysign(_Tp(0), __x.imag())); + return std::complex<_Tp>(std::isnan(__x.imag()) ? __x.imag() : _Tp(0), + copysign(__x.real(), __x.imag())); + } + return polar(sqrt(abs(__x)), arg(__x) / _Tp(2)); +} + +// exp + +template +__DEVICE__ std::complex<_Tp> exp(const std::complex<_Tp> &__x) { + _Tp __i = __x.imag(); + if (std::isinf(__x.real())) { + if (__x.real() < _Tp(0)) { + if (!std::isfinite(__i)) + __i = _Tp(1); + } else if (__i == 0 || !std::isfinite(__i)) { + if (std::isinf(__i)) + __i = _Tp(NAN); + return std::complex<_Tp>(__x.real(), __i); + } + } else if (std::isnan(__x.real()) && __x.imag() == 0) + return __x; + _Tp __e = exp(__x.real()); + return std::complex<_Tp>(__e * cos(__i), __e * sin(__i)); +} + +// pow + +template +std::complex<_Tp> pow(const std::complex<_Tp> &__x, + const std::complex<_Tp> &__y) { + return exp(__y * log(__x)); +} + +// __sqr, computes pow(x, 2) + +template std::complex<_Tp> __sqr(const std::complex<_Tp> &__x) { + return std::complex<_Tp>((__x.real() - __x.imag()) * + (__x.real() + __x.imag()), + _Tp(2) * __x.real() * __x.imag()); +} + +// asinh + +template +__DEVICE__ std::complex<_Tp> asinh(const std::complex<_Tp> &__x) { + const _Tp __pi(atan2(+0., -0.)); + if (std::isinf(__x.real())) { + if (std::isnan(__x.imag())) + return __x; + if (std::isinf(__x.imag())) + return std::complex<_Tp>(__x.real(), + copysign(__pi * _Tp(0.25), __x.imag())); + return std::complex<_Tp>(__x.real(), copysign(_Tp(0), __x.imag())); + } + if (std::isnan(__x.real())) { + if (std::isinf(__x.imag())) + return std::complex<_Tp>(__x.imag(), __x.real()); + if (__x.imag() == 0) + return __x; + return std::complex<_Tp>(__x.real(), __x.real()); + } + if (std::isinf(__x.imag())) + return std::complex<_Tp>(copysign(__x.imag(), __x.real()), + copysign(__pi / _Tp(2), __x.imag())); + std::complex<_Tp> __z = log(__x + sqrt(__sqr(__x) + _Tp(1))); + return std::complex<_Tp>(copysign(__z.real(), __x.real()), + copysign(__z.imag(), __x.imag())); +} + +// acosh + +template +__DEVICE__ std::complex<_Tp> acosh(const std::complex<_Tp> &__x) { + const _Tp __pi(atan2(+0., -0.)); + if (std::isinf(__x.real())) { + if (std::isnan(__x.imag())) + return std::complex<_Tp>(abs(__x.real()), __x.imag()); + if (std::isinf(__x.imag())) { + if (__x.real() > 0) + return std::complex<_Tp>(__x.real(), + copysign(__pi * _Tp(0.25), __x.imag())); + else + return std::complex<_Tp>(-__x.real(), + copysign(__pi * _Tp(0.75), __x.imag())); + } + if (__x.real() < 0) + return std::complex<_Tp>(-__x.real(), copysign(__pi, __x.imag())); + return std::complex<_Tp>(__x.real(), copysign(_Tp(0), __x.imag())); + } + if (std::isnan(__x.real())) { + if (std::isinf(__x.imag())) + return std::complex<_Tp>(abs(__x.imag()), __x.real()); + return std::complex<_Tp>(__x.real(), __x.real()); + } + if (std::isinf(__x.imag())) + return std::complex<_Tp>(abs(__x.imag()), + copysign(__pi / _Tp(2), __x.imag())); + std::complex<_Tp> __z = log(__x + sqrt(__sqr(__x) - _Tp(1))); + return std::complex<_Tp>(copysign(__z.real(), _Tp(0)), + copysign(__z.imag(), __x.imag())); +} + +// atanh + +template +__DEVICE__ std::complex<_Tp> atanh(const std::complex<_Tp> &__x) { + const _Tp __pi(atan2(+0., -0.)); + if (std::isinf(__x.imag())) { + return std::complex<_Tp>(copysign(_Tp(0), __x.real()), + copysign(__pi / _Tp(2), __x.imag())); + } + if (std::isnan(__x.imag())) { + if (std::isinf(__x.real()) || __x.real() == 0) + return std::complex<_Tp>(copysign(_Tp(0), __x.real()), __x.imag()); + return std::complex<_Tp>(__x.imag(), __x.imag()); + } + if (std::isnan(__x.real())) { + return std::complex<_Tp>(__x.real(), __x.real()); + } + if (std::isinf(__x.real())) { + return std::complex<_Tp>(copysign(_Tp(0), __x.real()), + copysign(__pi / _Tp(2), __x.imag())); + } + if (abs(__x.real()) == _Tp(1) && __x.imag() == _Tp(0)) { + return std::complex<_Tp>(copysign(_Tp(INFINITY), __x.real()), + copysign(_Tp(0), __x.imag())); + } + std::complex<_Tp> __z = log((_Tp(1) + __x) / (_Tp(1) - __x)) / _Tp(2); + return std::complex<_Tp>(copysign(__z.real(), __x.real()), + copysign(__z.imag(), __x.imag())); +} + +// sinh + +template +__DEVICE__ std::complex<_Tp> sinh(const std::complex<_Tp> &__x) { + if (std::isinf(__x.real()) && !std::isfinite(__x.imag())) + return std::complex<_Tp>(__x.real(), _Tp(NAN)); + if (__x.real() == 0 && !std::isfinite(__x.imag())) + return std::complex<_Tp>(__x.real(), _Tp(NAN)); + if (__x.imag() == 0 && !std::isfinite(__x.real())) + return __x; + return std::complex<_Tp>(sinh(__x.real()) * cos(__x.imag()), + cosh(__x.real()) * sin(__x.imag())); +} + +// cosh + +template +__DEVICE__ std::complex<_Tp> cosh(const std::complex<_Tp> &__x) { + if (std::isinf(__x.real()) && !std::isfinite(__x.imag())) + return std::complex<_Tp>(abs(__x.real()), _Tp(NAN)); + if (__x.real() == 0 && !std::isfinite(__x.imag())) + return std::complex<_Tp>(_Tp(NAN), __x.real()); + if (__x.real() == 0 && __x.imag() == 0) + return std::complex<_Tp>(_Tp(1), __x.imag()); + if (__x.imag() == 0 && !std::isfinite(__x.real())) + return std::complex<_Tp>(abs(__x.real()), __x.imag()); + return std::complex<_Tp>(cosh(__x.real()) * cos(__x.imag()), + sinh(__x.real()) * sin(__x.imag())); +} + +// tanh + +template +__DEVICE__ std::complex<_Tp> tanh(const std::complex<_Tp> &__x) { + if (std::isinf(__x.real())) { + if (!std::isfinite(__x.imag())) + return std::complex<_Tp>(_Tp(1), _Tp(0)); + return std::complex<_Tp>(_Tp(1), + copysign(_Tp(0), sin(_Tp(2) * __x.imag()))); + } + if (std::isnan(__x.real()) && __x.imag() == 0) + return __x; + _Tp __2r(_Tp(2) * __x.real()); + _Tp __2i(_Tp(2) * __x.imag()); + _Tp __d(cosh(__2r) + cos(__2i)); + _Tp __2rsh(sinh(__2r)); + if (std::isinf(__2rsh) && std::isinf(__d)) + return std::complex<_Tp>(__2rsh > _Tp(0) ? _Tp(1) : _Tp(-1), + __2i > _Tp(0) ? _Tp(0) : _Tp(-0.)); + return std::complex<_Tp>(__2rsh / __d, sin(__2i) / __d); +} + +// asin + +template +__DEVICE__ std::complex<_Tp> asin(const std::complex<_Tp> &__x) { + std::complex<_Tp> __z = asinh(complex<_Tp>(-__x.imag(), __x.real())); + return std::complex<_Tp>(__z.imag(), -__z.real()); +} + +// acos + +template +__DEVICE__ std::complex<_Tp> acos(const std::complex<_Tp> &__x) { + const _Tp __pi(atan2(+0., -0.)); + if (std::isinf(__x.real())) { + if (std::isnan(__x.imag())) + return std::complex<_Tp>(__x.imag(), __x.real()); + if (std::isinf(__x.imag())) { + if (__x.real() < _Tp(0)) + return std::complex<_Tp>(_Tp(0.75) * __pi, -__x.imag()); + return std::complex<_Tp>(_Tp(0.25) * __pi, -__x.imag()); + } + if (__x.real() < _Tp(0)) + return std::complex<_Tp>(__pi, + signbit(__x.imag()) ? -__x.real() : __x.real()); + return std::complex<_Tp>(_Tp(0), + signbit(__x.imag()) ? __x.real() : -__x.real()); + } + if (std::isnan(__x.real())) { + if (std::isinf(__x.imag())) + return std::complex<_Tp>(__x.real(), -__x.imag()); + return std::complex<_Tp>(__x.real(), __x.real()); + } + if (std::isinf(__x.imag())) + return std::complex<_Tp>(__pi / _Tp(2), -__x.imag()); + if (__x.real() == 0 && (__x.imag() == 0 || isnan(__x.imag()))) + return std::complex<_Tp>(__pi / _Tp(2), -__x.imag()); + std::complex<_Tp> __z = log(__x + sqrt(__sqr(__x) - _Tp(1))); + if (signbit(__x.imag())) + return std::complex<_Tp>(abs(__z.imag()), abs(__z.real())); + return std::complex<_Tp>(abs(__z.imag()), -abs(__z.real())); +} + +// atan + +template +__DEVICE__ std::complex<_Tp> atan(const std::complex<_Tp> &__x) { + std::complex<_Tp> __z = atanh(complex<_Tp>(-__x.imag(), __x.real())); + return std::complex<_Tp>(__z.imag(), -__z.real()); +} + +// sin + +template +__DEVICE__ std::complex<_Tp> sin(const std::complex<_Tp> &__x) { + std::complex<_Tp> __z = sinh(complex<_Tp>(-__x.imag(), __x.real())); + return std::complex<_Tp>(__z.imag(), -__z.real()); +} + +// cos + +template std::complex<_Tp> cos(const std::complex<_Tp> &__x) { + return cosh(complex<_Tp>(-__x.imag(), __x.real())); +} + +// tan + +template +__DEVICE__ std::complex<_Tp> tan(const std::complex<_Tp> &__x) { + std::complex<_Tp> __z = tanh(complex<_Tp>(-__x.imag(), __x.real())); + return std::complex<_Tp>(__z.imag(), -__z.real()); +} + +} // namespace std + +#endif diff --git a/clang/test/Headers/Inputs/include/complex b/clang/test/Headers/Inputs/include/complex --- a/clang/test/Headers/Inputs/include/complex +++ b/clang/test/Headers/Inputs/include/complex @@ -3,6 +3,7 @@ #include #define INFINITY (__builtin_inff()) +#define NAN (__builtin_nanf ("")) namespace std { @@ -298,4 +299,114 @@ return !(__x == __y); } +template _Tp abs(const std::complex<_Tp> &__c); + +// arg + +template _Tp arg(const std::complex<_Tp> &__c); + +// norm + +template _Tp norm(const std::complex<_Tp> &__c); + +// conj + +template std::complex<_Tp> conj(const std::complex<_Tp> &__c); + +// proj + +template std::complex<_Tp> proj(const std::complex<_Tp> &__c); + +// polar + +template +complex<_Tp> polar(const _Tp &__rho, const _Tp &__theta = _Tp()); + +// log + +template std::complex<_Tp> log(const std::complex<_Tp> &__x); + +// log10 + +template std::complex<_Tp> log10(const std::complex<_Tp> &__x); + +// sqrt + +template +std::complex<_Tp> sqrt(const std::complex<_Tp> &__x); + +// exp + +template +std::complex<_Tp> exp(const std::complex<_Tp> &__x); + +// pow + +template +std::complex<_Tp> pow(const std::complex<_Tp> &__x, + const std::complex<_Tp> &__y); + +// __sqr, computes pow(x, 2) + +template std::complex<_Tp> __sqr(const std::complex<_Tp> &__x); + +// asinh + +template +std::complex<_Tp> asinh(const std::complex<_Tp> &__x); + +// acosh + +template +std::complex<_Tp> acosh(const std::complex<_Tp> &__x); + +// atanh + +template +std::complex<_Tp> atanh(const std::complex<_Tp> &__x); + +// sinh + +template +std::complex<_Tp> sinh(const std::complex<_Tp> &__x); + +// cosh + +template +std::complex<_Tp> cosh(const std::complex<_Tp> &__x); + +// tanh + +template +std::complex<_Tp> tanh(const std::complex<_Tp> &__x); + +// asin + +template +std::complex<_Tp> asin(const std::complex<_Tp> &__x); + +// acos + +template +std::complex<_Tp> acos(const std::complex<_Tp> &__x); + +// atan + +template +std::complex<_Tp> atan(const std::complex<_Tp> &__x); + +// sin + +template +std::complex<_Tp> sin(const std::complex<_Tp> &__x); + +// cos + +template std::complex<_Tp> cos(const std::complex<_Tp> &__x); + +// tan + +template +std::complex<_Tp> tan(const std::complex<_Tp> &__x); + } // namespace std diff --git a/clang/test/Headers/nvptx_device_math_complex.c b/clang/test/Headers/nvptx_device_math_complex.c --- a/clang/test/Headers/nvptx_device_math_complex.c +++ b/clang/test/Headers/nvptx_device_math_complex.c @@ -1,4 +1,3 @@ -// 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 // 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 diff --git a/clang/test/Headers/nvptx_device_math_complex.cpp b/clang/test/Headers/nvptx_device_math_complex.cpp --- a/clang/test/Headers/nvptx_device_math_complex.cpp +++ b/clang/test/Headers/nvptx_device_math_complex.cpp @@ -1,8 +1,8 @@ -// 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 #include // CHECK: define weak {{.*}} @__muldc3 @@ -33,6 +33,38 @@ // CHECK-DAG: call float @__nv_fabsf( // CHECK-DAG: call float @__nv_logbf( +// We might actually want to check that there are no declarations of non-OpenMP functions instead. +// That is, as long as we don't call an unkown function with a name that doesn't start with '__' we are good :) + +// CHECK-DAG: call %"struct.std::complex" @"_ZStL71sin$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL72sinh$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL71exp$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL71sin$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZSt71cos$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL71exp$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL72atan$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL72acos$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL72cosh$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL73atanh$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZSt73__sqr$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZStL72sqrt$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZSt71log$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL71sin$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZSt71cos$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL71exp$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL72atan$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL72acos$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL72sinh$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL72cosh$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL73atanh$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZSt73__sqr$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZStL72sqrt$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZSt71log$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex" @"_ZSt71log$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS2_"(%"struct.std::complex"* +// CHECK-DAG: call %"struct.std::complex" @"_ZSt73polar$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIfESt7complexIT_ERKS1_S4_"(float* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZSt71log$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS2_"(%"struct.std::complex.0"* +// CHECK-DAG: call %"struct.std::complex.0" @"_ZSt73polar$ompvariant$S2$s8$Pnvptx$Pnvptx64$S3$s10$Pmatch_any$Pallow_templatesIdESt7complexIT_ERKS1_S4_"(double* + void test_scmplx(std::complex a) { #pragma omp target { @@ -46,3 +78,35 @@ (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); +}