maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/Headers/nvptx_device_math_complex.cpp

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

49 lines
1.7 KiB
C++
Raw Normal View History

[OpenMP] Initial support for std::complex in target regions This simply follows the scheme we have for other wrappers. It resolves the current link problem, e.g., `__muldc3 not found`, when std::complex operations are used on a device. This will not allow complex make math function calls to work properly, e.g., sin, but that is more complex (pan intended) anyway. Reviewed By: tra, JonChesterfield Differential Revision: https://reviews.llvm.org/D80897
2020-04-12 15:10:28 +08:00
// 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 <complex>
[OpenMP][CUDA] Fix std::complex in GPU regions The old way worked to some degree for C++-mode but in C mode we actually tried to introduce variants of macros (e.g., isinf). To make both modes work reliably we get rid of those extra variants and directly use NVIDIA intrinsics in the complex implementation. While this has to be revisited as we add other GPU targets which want to reuse the code, it should be fine for now. Reviewed By: tra, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D83591
2020-07-11 05:45:02 +08:00
// CHECK: define weak {{.*}} @__muldc3
// CHECK-DAG: call i32 @__nv_isnand(
// CHECK-DAG: call i32 @__nv_isinfd(
// CHECK-DAG: call double @__nv_copysign(
[OpenMP] Initial support for std::complex in target regions This simply follows the scheme we have for other wrappers. It resolves the current link problem, e.g., `__muldc3 not found`, when std::complex operations are used on a device. This will not allow complex make math function calls to work properly, e.g., sin, but that is more complex (pan intended) anyway. Reviewed By: tra, JonChesterfield Differential Revision: https://reviews.llvm.org/D80897
2020-04-12 15:10:28 +08:00
[OpenMP][CUDA] Fix std::complex in GPU regions The old way worked to some degree for C++-mode but in C mode we actually tried to introduce variants of macros (e.g., isinf). To make both modes work reliably we get rid of those extra variants and directly use NVIDIA intrinsics in the complex implementation. While this has to be revisited as we add other GPU targets which want to reuse the code, it should be fine for now. Reviewed By: tra, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D83591
2020-07-11 05:45:02 +08:00
// 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(
[OpenMP] Initial support for std::complex in target regions This simply follows the scheme we have for other wrappers. It resolves the current link problem, e.g., `__muldc3 not found`, when std::complex operations are used on a device. This will not allow complex make math function calls to work properly, e.g., sin, but that is more complex (pan intended) anyway. Reviewed By: tra, JonChesterfield Differential Revision: https://reviews.llvm.org/D80897
2020-04-12 15:10:28 +08:00
// CHECK-DAG: call float @__nv_scalbnf(
[OpenMP][CUDA] Fix std::complex in GPU regions The old way worked to some degree for C++-mode but in C mode we actually tried to introduce variants of macros (e.g., isinf). To make both modes work reliably we get rid of those extra variants and directly use NVIDIA intrinsics in the complex implementation. While this has to be revisited as we add other GPU targets which want to reuse the code, it should be fine for now. Reviewed By: tra, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D83591
2020-07-11 05:45:02 +08:00
// CHECK-DAG: call float @__nv_fabsf(
// CHECK-DAG: call float @__nv_logbf(
[OpenMP] Initial support for std::complex in target regions This simply follows the scheme we have for other wrappers. It resolves the current link problem, e.g., `__muldc3 not found`, when std::complex operations are used on a device. This will not allow complex make math function calls to work properly, e.g., sin, but that is more complex (pan intended) anyway. Reviewed By: tra, JonChesterfield Differential Revision: https://reviews.llvm.org/D80897
2020-04-12 15:10:28 +08:00
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));
}
}