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/openmp_device_math_isnan.cpp

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

43 lines
5.6 KiB
C++
Raw Normal View History

[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=BOOL_RETURN
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=AMD_BOOL_RETURN
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=BOOL_RETURN
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=AMD_BOOL_RETURN
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -DUSE_ISNAN_WITH_INT_RETURN
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc -DUSE_ISNAN_WITH_INT_RETURN
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=INT_RETURN
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=AMD_INT_RETURN
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=INT_RETURN
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple amdgcn-amd-amdhsa -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=amdgcn-amd-amdhsa -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast -DUSE_ISNAN_WITH_INT_RETURN | FileCheck %s --check-prefix=AMD_INT_RETURN
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// expected-no-diagnostics
#include <cmath>
double math(float f, double d) {
double r = 0;
// INT_RETURN: call i32 @__nv_isnanf(float
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// AMD_INT_RETURN: call i32 @_{{.*}}isnanf(float
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// BOOL_RETURN: call i32 @__nv_isnanf(float
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnanf(float
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
r += std::isnan(f);
// INT_RETURN: call i32 @__nv_isnand(double
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// AMD_INT_RETURN: call i32 @_{{.*}}isnand(double
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
// BOOL_RETURN: call i32 @__nv_isnand(double
[OpenMP][AMDGCN] Apply fix for isnan, isinf and isfinite for amdgcn. This fixes issues with various return types(bool/int) and was already in place for nvptx headers, adjusted to work for amdgcn. This does not affect hip as the change is guarded with OPENMP_AMDGCN. Similar to D85879. Reviewed By: jdoerfert, JonChesterfield, yaxunl Differential Revision: https://reviews.llvm.org/D104677
2021-06-23 22:25:00 +08:00
// AMD_BOOL_RETURN: call zeroext i1 @_{{.*}}isnand(double
[OpenMP] Overload `std::isnan` and friends multiple times for the GPU `std::isnan` and friends can be found in two variants in the wild, one returns `bool`, as the standard defines it, one returns `int`, as the C macros do. So far we kinda hoped the system versions of these functions will work for people, e.g. they are definitions that can be compiled for the target. We know that is not the case always so we leverage the `disable_implicit_base` OpenMP context extension to specialize both versions of these functions without causing an invalid redeclaration. Reviewed By: JonChesterfield, tra Differential Revision: https://reviews.llvm.org/D85879
2020-08-13 05:45:46 +08:00
r += std::isnan(d);
return r;
}
long double foo(float f, double d, long double ld) {
double r = ld;
r += math(f, d);
#pragma omp target map(r)
{ r += math(f, d); }
return r;
}