shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

Clang: Add nsz to llvm.minnum and llvm.maxnum emitted from fmin and fmax (#113133)

Browse Source 
See: https://github.com/llvm/llvm-project/pull/112852

We will define llvm.minnum and llvm.maxnum with +0.0>-0.0, by default,
while libc doesn't require it.
This commit is contained in:
YunQiang Su 2026-02-11 08:33:29 +08:00 committed by GitHub
parent 0f6ee50e4e
commit 7e734da346
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
9 changed files with 98 additions and 84 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
clang/lib/CodeGen/CGBuiltin.cpp
View File

@ -2877,10 +2877,13 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
case Builtin::BI__builtin_fmaxf:
case Builtin::BI__builtin_fmaxf16:
case Builtin::BI__builtin_fmaxl:
case Builtin::BI__builtin_fmaxf128:
return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(*this, E,
Intrinsic::maxnum,
Intrinsic::experimental_constrained_maxnum));
case Builtin::BI__builtin_fmaxf128: {
IRBuilder<>::FastMathFlagGuard FMFGuard(Builder);
Builder.getFastMathFlags().setNoSignedZeros();
return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(
*this, E, Intrinsic::maxnum,
Intrinsic::experimental_constrained_maxnum));
}
case Builtin::BIfmin:
case Builtin::BIfminf:
@ -2889,10 +2892,13 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
case Builtin::BI__builtin_fminf:
case Builtin::BI__builtin_fminf16:
case Builtin::BI__builtin_fminl:
case Builtin::BI__builtin_fminf128:
return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(*this, E,
Intrinsic::minnum,
Intrinsic::experimental_constrained_minnum));
case Builtin::BI__builtin_fminf128: {
IRBuilder<>::FastMathFlagGuard FMFGuard(Builder);
Builder.getFastMathFlags().setNoSignedZeros();
return RValue::get(emitBinaryMaybeConstrainedFPBuiltin(
*this, E, Intrinsic::minnum,
Intrinsic::experimental_constrained_minnum));
}
case Builtin::BIfmaximum_num:
case Builtin::BIfmaximum_numf:

24
clang/test/CodeGen/RISCV/math-builtins.c
View File

@ -134,22 +134,22 @@ long double truncl(long double);
// RV32-NEXT: [[TMP44:%.*]] = call fp128 @llvm.floor.f128(fp128 [[TMP43]])
// RV32-NEXT: [[TMP45:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV32-NEXT: [[TMP46:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV32-NEXT: [[TMP47:%.*]] = call float @llvm.maxnum.f32(float [[TMP45]], float [[TMP46]])
// RV32-NEXT: [[TMP47:%.*]] = call nsz float @llvm.maxnum.f32(float [[TMP45]], float [[TMP46]])
// RV32-NEXT: [[TMP48:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV32-NEXT: [[TMP49:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV32-NEXT: [[TMP50:%.*]] = call double @llvm.maxnum.f64(double [[TMP48]], double [[TMP49]])
// RV32-NEXT: [[TMP50:%.*]] = call nsz double @llvm.maxnum.f64(double [[TMP48]], double [[TMP49]])
// RV32-NEXT: [[TMP51:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV32-NEXT: [[TMP52:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV32-NEXT: [[TMP53:%.*]] = call fp128 @llvm.maxnum.f128(fp128 [[TMP51]], fp128 [[TMP52]])
// RV32-NEXT: [[TMP53:%.*]] = call nsz fp128 @llvm.maxnum.f128(fp128 [[TMP51]], fp128 [[TMP52]])
// RV32-NEXT: [[TMP54:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV32-NEXT: [[TMP55:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV32-NEXT: [[TMP56:%.*]] = call float @llvm.minnum.f32(float [[TMP54]], float [[TMP55]])
// RV32-NEXT: [[TMP56:%.*]] = call nsz float @llvm.minnum.f32(float [[TMP54]], float [[TMP55]])
// RV32-NEXT: [[TMP57:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV32-NEXT: [[TMP58:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV32-NEXT: [[TMP59:%.*]] = call double @llvm.minnum.f64(double [[TMP57]], double [[TMP58]])
// RV32-NEXT: [[TMP59:%.*]] = call nsz double @llvm.minnum.f64(double [[TMP57]], double [[TMP58]])
// RV32-NEXT: [[TMP60:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV32-NEXT: [[TMP61:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV32-NEXT: [[TMP62:%.*]] = call fp128 @llvm.minnum.f128(fp128 [[TMP60]], fp128 [[TMP61]])
// RV32-NEXT: [[TMP62:%.*]] = call nsz fp128 @llvm.minnum.f128(fp128 [[TMP60]], fp128 [[TMP61]])
// RV32-NEXT: [[TMP63:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV32-NEXT: [[TMP64:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV32-NEXT: [[FMOD:%.*]] = frem float [[TMP63]], [[TMP64]]
@ -310,22 +310,22 @@ long double truncl(long double);
// RV64-NEXT: [[TMP44:%.*]] = call fp128 @llvm.floor.f128(fp128 [[TMP43]])
// RV64-NEXT: [[TMP45:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV64-NEXT: [[TMP46:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV64-NEXT: [[TMP47:%.*]] = call float @llvm.maxnum.f32(float [[TMP45]], float [[TMP46]])
// RV64-NEXT: [[TMP47:%.*]] = call nsz float @llvm.maxnum.f32(float [[TMP45]], float [[TMP46]])
// RV64-NEXT: [[TMP48:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV64-NEXT: [[TMP49:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV64-NEXT: [[TMP50:%.*]] = call double @llvm.maxnum.f64(double [[TMP48]], double [[TMP49]])
// RV64-NEXT: [[TMP50:%.*]] = call nsz double @llvm.maxnum.f64(double [[TMP48]], double [[TMP49]])
// RV64-NEXT: [[TMP51:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV64-NEXT: [[TMP52:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV64-NEXT: [[TMP53:%.*]] = call fp128 @llvm.maxnum.f128(fp128 [[TMP51]], fp128 [[TMP52]])
// RV64-NEXT: [[TMP53:%.*]] = call nsz fp128 @llvm.maxnum.f128(fp128 [[TMP51]], fp128 [[TMP52]])
// RV64-NEXT: [[TMP54:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV64-NEXT: [[TMP55:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV64-NEXT: [[TMP56:%.*]] = call float @llvm.minnum.f32(float [[TMP54]], float [[TMP55]])
// RV64-NEXT: [[TMP56:%.*]] = call nsz float @llvm.minnum.f32(float [[TMP54]], float [[TMP55]])
// RV64-NEXT: [[TMP57:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV64-NEXT: [[TMP58:%.*]] = load double, ptr [[DARG_ADDR]], align 8
// RV64-NEXT: [[TMP59:%.*]] = call double @llvm.minnum.f64(double [[TMP57]], double [[TMP58]])
// RV64-NEXT: [[TMP59:%.*]] = call nsz double @llvm.minnum.f64(double [[TMP57]], double [[TMP58]])
// RV64-NEXT: [[TMP60:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV64-NEXT: [[TMP61:%.*]] = load fp128, ptr [[LDARG_ADDR]], align 16
// RV64-NEXT: [[TMP62:%.*]] = call fp128 @llvm.minnum.f128(fp128 [[TMP60]], fp128 [[TMP61]])
// RV64-NEXT: [[TMP62:%.*]] = call nsz fp128 @llvm.minnum.f128(fp128 [[TMP60]], fp128 [[TMP61]])
// RV64-NEXT: [[TMP63:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV64-NEXT: [[TMP64:%.*]] = load float, ptr [[FARG_ADDR]], align 4
// RV64-NEXT: [[FMOD:%.*]] = frem float [[TMP63]], [[TMP64]]

12
clang/test/CodeGen/builtins.c
View File

@ -372,22 +372,22 @@ void test_float_builtin_ops(float F, double D, long double LD, int I) {
// CHECK: call [[LDTYPE]] @llvm.canonicalize.[[LDLLVMTY]]([[LDTYPE]]
resf = __builtin_fminf(F, F);
// CHECK: call float @llvm.minnum.f32
// CHECK: call nsz float @llvm.minnum.f32
resd = __builtin_fmin(D, D);
// CHECK: call double @llvm.minnum.f64
// CHECK: call nsz double @llvm.minnum.f64
resld = __builtin_fminl(LD, LD);
// CHECK: call [[LDTYPE]] @llvm.minnum.[[LDLLVMTY]]
// CHECK: call nsz [[LDTYPE]] @llvm.minnum.[[LDLLVMTY]]
resf = __builtin_fmaxf(F, F);
// CHECK: call float @llvm.maxnum.f32
// CHECK: call nsz float @llvm.maxnum.f32
resd = __builtin_fmax(D, D);
// CHECK: call double @llvm.maxnum.f64
// CHECK: call nsz double @llvm.maxnum.f64
resld = __builtin_fmaxl(LD, LD);
// CHECK: call [[LDTYPE]] @llvm.maxnum.[[LDLLVMTY]]
// CHECK: call nsz [[LDTYPE]] @llvm.maxnum.[[LDLLVMTY]]
resf = __builtin_fminimum_numf(F, F);
// CHECK: call float @llvm.minimumnum.f32

16
clang/test/CodeGen/math-builtins-long.c
View File

@ -148,16 +148,16 @@ void foo(long double f, long double *l, int *i, const char *c) {
// PPCF128: call fp128 @llvm.floor.f128(fp128 %{{.+}})
__builtin_floorl(f);
// F80: call x86_fp80 @llvm.maxnum.f80(x86_fp80 %{{.+}}, x86_fp80 %{{.+}})
// PPC: call ppc_fp128 @llvm.maxnum.ppcf128(ppc_fp128 %{{.+}}, ppc_fp128 %{{.+}})
// X86F128: call fp128 @llvm.maxnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
// PPCF128: call fp128 @llvm.maxnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
// F80: call nsz x86_fp80 @llvm.maxnum.f80(x86_fp80 %{{.+}}, x86_fp80 %{{.+}})
// PPC: call nsz ppc_fp128 @llvm.maxnum.ppcf128(ppc_fp128 %{{.+}}, ppc_fp128 %{{.+}})
// X86F128: call nsz fp128 @llvm.maxnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
// PPCF128: call nsz fp128 @llvm.maxnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
__builtin_fmaxl(f,f);
// F80: call x86_fp80 @llvm.minnum.f80(x86_fp80 %{{.+}}, x86_fp80 %{{.+}})
// PPC: call ppc_fp128 @llvm.minnum.ppcf128(ppc_fp128 %{{.+}}, ppc_fp128 %{{.+}})
// X86F128: call fp128 @llvm.minnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
// PPCF128: call fp128 @llvm.minnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
// F80: call nsz x86_fp80 @llvm.minnum.f80(x86_fp80 %{{.+}}, x86_fp80 %{{.+}})
// PPC: call nsz ppc_fp128 @llvm.minnum.ppcf128(ppc_fp128 %{{.+}}, ppc_fp128 %{{.+}})
// X86F128: call nsz fp128 @llvm.minnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
// PPCF128: call nsz fp128 @llvm.minnum.f128(fp128 %{{.+}}, fp128 %{{.+}})
__builtin_fminl(f,f);
// F80: call x86_fp80 @llvm.nearbyint.f80(x86_fp80 %{{.+}})

4
clang/test/CodeGenOpenCL/builtins-f16.cl
View File

@ -48,10 +48,10 @@ void test_half_builtins(half h0, half h1, half h2, int i0) {
// CHECK: call half @llvm.fma.f16(half %h0, half %h1, half %h2)
res = __builtin_fmaf16(h0, h1 ,h2);
// CHECK: call half @llvm.maxnum.f16(half %h0, half %h1)
// CHECK: call nsz half @llvm.maxnum.f16(half %h0, half %h1)
res = __builtin_fmaxf16(h0, h1);
// CHECK: call half @llvm.minnum.f16(half %h0, half %h1)
// CHECK: call nsz half @llvm.minnum.f16(half %h0, half %h1)
res = __builtin_fminf16(h0, h1);
// CHECK: frem half %h0, %h1

80
clang/test/Headers/__clang_hip_math.hip
View File

@ -2218,31 +2218,31 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: define dso_local noundef float @test_fmaxf(
// DEFAULT-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_fmaxf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_fmaxf(
// APPROX-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: ret float [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef float @test_fmaxf(
// NCRDIV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: ret float [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_fmaxf(
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fmaxf(float x, float y) {
@ -2252,31 +2252,31 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
// DEFAULT-LABEL: define dso_local noundef double @test_fmax(
// DEFAULT-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_fmax(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_fmax(
// APPROX-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: ret double [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef double @test_fmax(
// NCRDIV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: ret double [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef double @test_fmax(
// AMDGCNSPIRV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fmax(double x, double y) {
@ -2286,31 +2286,31 @@ extern "C" __device__ double test_fmax(double x, double y) {
// DEFAULT-LABEL: define dso_local noundef float @test_fminf(
// DEFAULT-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_fminf(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_fminf(
// APPROX-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: ret float [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef float @test_fminf(
// NCRDIV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: ret float [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_fminf(
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.minnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_fminf(float x, float y) {
@ -2320,31 +2320,31 @@ extern "C" __device__ float test_fminf(float x, float y) {
// DEFAULT-LABEL: define dso_local noundef double @test_fmin(
// DEFAULT-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_fmin(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_fmin(
// APPROX-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: ret double [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef double @test_fmin(
// NCRDIV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: ret double [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef double @test_fmin(
// AMDGCNSPIRV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.minnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_fmin(double x, double y) {
@ -9170,31 +9170,31 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
// DEFAULT-LABEL: define dso_local noundef float @test_float_min(
// DEFAULT-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_float_min(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) float @llvm.minnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_float_min(
// APPROX-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: ret float [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef float @test_float_min(
// NCRDIV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.minnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: ret float [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_float_min(
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.minnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_float_min(float x, float y) {
@ -9204,31 +9204,31 @@ extern "C" __device__ float test_float_min(float x, float y) {
// DEFAULT-LABEL: define dso_local noundef float @test_float_max(
// DEFAULT-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// DEFAULT-NEXT: ret float [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_float_max(
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]], float noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) float @llvm.maxnum.f32(float nofpclass(nan inf) [[X]], float nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret float [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef float @test_float_max(
// APPROX-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// APPROX-NEXT: ret float [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef float @test_float_max(
// NCRDIV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef float @llvm.maxnum.f32(float [[X]], float [[Y]])
// NCRDIV-NEXT: ret float [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_float_max(
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]], float noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X]], float [[Y]])
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
//
extern "C" __device__ float test_float_max(float x, float y) {
@ -9238,31 +9238,31 @@ extern "C" __device__ float test_float_max(float x, float y) {
// DEFAULT-LABEL: define dso_local noundef double @test_double_min(
// DEFAULT-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_double_min(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) double @llvm.minnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_double_min(
// APPROX-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: ret double [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef double @test_double_min(
// NCRDIV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.minnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: ret double [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef double @test_double_min(
// AMDGCNSPIRV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.minnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_double_min(double x, double y) {
@ -9272,31 +9272,31 @@ extern "C" __device__ double test_double_min(double x, double y) {
// DEFAULT-LABEL: define dso_local noundef double @test_double_max(
// DEFAULT-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// DEFAULT-NEXT: [[ENTRY:.*:]]
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// DEFAULT-NEXT: ret double [[TMP0]]
//
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) double @test_double_max(
// FINITEONLY-SAME: double noundef nofpclass(nan inf) [[X:%.*]], double noundef nofpclass(nan inf) [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// FINITEONLY-NEXT: [[ENTRY:.*:]]
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef nofpclass(nan inf) double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf nsz contract noundef nofpclass(nan inf) double @llvm.maxnum.f64(double nofpclass(nan inf) [[X]], double nofpclass(nan inf) [[Y]])
// FINITEONLY-NEXT: ret double [[TMP0]]
//
// APPROX-LABEL: define dso_local noundef double @test_double_max(
// APPROX-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// APPROX-NEXT: [[ENTRY:.*:]]
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// APPROX-NEXT: ret double [[TMP0]]
//
// NCRDIV-LABEL: define dso_local noundef double @test_double_max(
// NCRDIV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR3]] {
// NCRDIV-NEXT: [[ENTRY:.*:]]
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef double @llvm.maxnum.f64(double [[X]], double [[Y]])
// NCRDIV-NEXT: ret double [[TMP0]]
//
// AMDGCNSPIRV-LABEL: define spir_func noundef double @test_double_max(
// AMDGCNSPIRV-SAME: double noundef [[X:%.*]], double noundef [[Y:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X]], double [[Y]])
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
//
extern "C" __device__ double test_double_max(double x, double y) {

8
clang/test/Headers/amdgcn_openmp_device_math_constexpr.cpp
View File

@ -182,7 +182,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: [[TMP2:%.*]] = call nsz noundef float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL17constexpr_min_f32 to ptr), align 4
// CHECK-NEXT: ret void
//
@ -198,7 +198,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: [[TMP2:%.*]] = call nsz noundef float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL17constexpr_max_f32 to ptr), align 4
// CHECK-NEXT: ret void
//
@ -230,7 +230,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: [[TMP2:%.*]] = call nsz noundef float @llvm.minnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL19constexpr_fminf_f32 to ptr), align 4
// CHECK-NEXT: ret void
//
@ -246,7 +246,7 @@ const float constexpr_fmaxf_f32 = fmaxf(2.0f, -4.0f);
// CHECK-NEXT: store float -4.000000e+00, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[__X_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[__Y_ADDR_ASCAST_I]], align 4
// CHECK-NEXT: [[TMP2:%.*]] = call noundef float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: [[TMP2:%.*]] = call nsz noundef float @llvm.maxnum.f32(float [[TMP0]], float [[TMP1]])
// CHECK-NEXT: store float [[TMP2]], ptr addrspacecast (ptr addrspace(1) @_ZL19constexpr_fmaxf_f32 to ptr), align 4
// CHECK-NEXT: ret void
//

8
clang/test/Headers/nvptx_device_math_complex.c
View File

@ -17,7 +17,9 @@
// CHECK-DAG: call i1 @llvm.is.fpclass.f32(float %{{.+}}, i32 504)
// CHECK-DAG: call float @llvm.copysign.f32(
// CHECK-DAG: call float @__nv_scalbnf(
// CHECK-DAG: call float @llvm.fabs.f32(
// CHECK-DAG: call nsz float @llvm.fabs.f32(
// CHECK-DAG: call nsz float @llvm.fabs.f32(
// CHECK-DAG: call nsz float @llvm.maxnum.f32(
// CHECK-DAG: call float @__nv_logbf(
// CHECK: define weak {{.*}} @__mulsc3
@ -31,7 +33,9 @@
// CHECK-DAG: call i1 @llvm.is.fpclass.f64(double %{{.+}}, i32 516)
// CHECK-DAG: call double @llvm.copysign.f64(
// CHECK-DAG: call double @__nv_scalbn(
// CHECK-DAG: call double @llvm.fabs.f64(
// CHECK-DAG: call nsz double @llvm.fabs.f64(
// CHECK-DAG: call nsz double @llvm.fabs.f64(
// CHECK-DAG: call nsz double @llvm.maxnum.f64(
// CHECK-DAG: call double @__nv_logb(
// CHECK: define weak {{.*}} @__muldc3

8
clang/test/Headers/nvptx_device_math_complex.cpp
View File

@ -22,7 +22,9 @@
// CHECK-DAG: call i1 @llvm.is.fpclass.f64(double %{{.+}}, i32 504)
// CHECK-DAG: call double @llvm.copysign.f64(
// CHECK-DAG: call double @__nv_scalbn(
// CHECK-DAG: call double @llvm.fabs.f64(
// CHECK-DAG: call nsz double @llvm.fabs.f64(
// CHECK-DAG: call nsz double @llvm.fabs.f64(
// CHECK-DAG: call nsz double @llvm.maxnum.f64(
// CHECK-DAG: call double @__nv_logb(
// CHECK: define weak {{.*}} @__divsc3
@ -31,7 +33,9 @@
// CHECK-DAG: call i1 @llvm.is.fpclass.f32(float %{{.+}}, i32 504)
// CHECK-DAG: call float @llvm.copysign.f32(
// CHECK-DAG: call float @__nv_scalbnf(
// CHECK-DAG: call float @llvm.fabs.f32(
// CHECK-DAG: call nsz float @llvm.fabs.f32(
// CHECK-DAG: call nsz float @llvm.fabs.f32(
// CHECK-DAG: call nsz float @llvm.maxnum.f32(
// CHECK-DAG: call float @__nv_logbf(
// We actually check that there are no declarations of non-OpenMP functions.