Skip to content

Commit 5ea0ac7

Browse files
wzssyqawzssyqa
committed
fix some tests
1 parent a251c90 commit 5ea0ac7

File tree

4 files changed

+24
-26
lines changed

4 files changed

+24
-26
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -4418,8 +4418,7 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
44184418
} else {
44194419
FastMathFlags FMF;
44204420
FMF.setNoSignedZeros(true);
4421-
FMFSource FMFSrc(FMF);
4422-
Result = Builder.CreateMaxNum(Op0, Op1, /*FMFSource=*/FMFSrc, "elt.max");
4421+
Result = Builder.CreateMaxNum(Op0, Op1, /*FMFSource=*/FMF, "elt.max");
44234422
}
44244423
return RValue::get(Result);
44254424
}
@@ -4438,8 +4437,7 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
44384437
} else {
44394438
FastMathFlags FMF;
44404439
FMF.setNoSignedZeros(true);
4441-
FMFSource FMFSrc(FMF);
4442-
Result = Builder.CreateMinNum(Op0, Op1, /*FMFSource=*/FMFSrc, "elt.min");
4440+
Result = Builder.CreateMinNum(Op0, Op1, /*FMFSource=*/FMF, "elt.min");
44434441
}
44444442
return RValue::get(Result);
44454443
}

‎clang/test/CodeGen/builtins-elementwise-math.c‎

Lines changed: 12 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -347,21 +347,21 @@ void test_builtin_elementwise_max(float f1, float f2, double d1, double d2,
347347
// CHECK-LABEL: define void @test_builtin_elementwise_max(
348348
// CHECK: [[F1:%.+]] = load float, ptr %f1.addr, align 4
349349
// CHECK-NEXT: [[F2:%.+]] = load float, ptr %f2.addr, align 4
350-
// CHECK-NEXT: call float @llvm.maxnum.f32(float [[F1]], float [[F2]])
350+
// CHECK-NEXT: call nsz float @llvm.maxnum.f32(float [[F1]], float [[F2]])
351351
f1 = __builtin_elementwise_max(f1, f2);
352352

353353
// CHECK: [[D1:%.+]] = load double, ptr %d1.addr, align 8
354354
// CHECK-NEXT: [[D2:%.+]] = load double, ptr %d2.addr, align 8
355-
// CHECK-NEXT: call double @llvm.maxnum.f64(double [[D1]], double [[D2]])
355+
// CHECK-NEXT: call nsz double @llvm.maxnum.f64(double [[D1]], double [[D2]])
356356
d1 = __builtin_elementwise_max(d1, d2);
357357

358358
// CHECK: [[D2:%.+]] = load double, ptr %d2.addr, align 8
359-
// CHECK-NEXT: call double @llvm.maxnum.f64(double 2.000000e+01, double [[D2]])
359+
// CHECK-NEXT: call nsz double @llvm.maxnum.f64(double 2.000000e+01, double [[D2]])
360360
d1 = __builtin_elementwise_max(20.0, d2);
361361

362362
// CHECK: [[VF1:%.+]] = load <4 x float>, ptr %vf1.addr, align 16
363363
// CHECK-NEXT: [[VF2:%.+]] = load <4 x float>, ptr %vf2.addr, align 16
364-
// CHECK-NEXT: call <4 x float> @llvm.maxnum.v4f32(<4 x float> [[VF1]], <4 x float> [[VF2]])
364+
// CHECK-NEXT: call nsz <4 x float> @llvm.maxnum.v4f32(<4 x float> [[VF1]], <4 x float> [[VF2]])
365365
vf1 = __builtin_elementwise_max(vf1, vf2);
366366

367367
// CHECK: [[I1:%.+]] = load i64, ptr %i1.addr, align 8
@@ -404,13 +404,13 @@ void test_builtin_elementwise_max(float f1, float f2, double d1, double d2,
404404

405405
// CHECK: [[CVF1:%.+]] = load <4 x float>, ptr %cvf1, align 16
406406
// CHECK-NEXT: [[VF2:%.+]] = load <4 x float>, ptr %vf2.addr, align 16
407-
// CHECK-NEXT: call <4 x float> @llvm.maxnum.v4f32(<4 x float> [[CVF1]], <4 x float> [[VF2]])
407+
// CHECK-NEXT: call nsz <4 x float> @llvm.maxnum.v4f32(<4 x float> [[CVF1]], <4 x float> [[VF2]])
408408
const float4 cvf1 = vf1;
409409
vf1 = __builtin_elementwise_max(cvf1, vf2);
410410

411411
// CHECK: [[VF2:%.+]] = load <4 x float>, ptr %vf2.addr, align 16
412412
// CHECK-NEXT: [[CVF1:%.+]] = load <4 x float>, ptr %cvf1, align 16
413-
// CHECK-NEXT: call <4 x float> @llvm.maxnum.v4f32(<4 x float> [[VF2]], <4 x float> [[CVF1]])
413+
// CHECK-NEXT: call nsz <4 x float> @llvm.maxnum.v4f32(<4 x float> [[VF2]], <4 x float> [[CVF1]])
414414
vf1 = __builtin_elementwise_max(vf2, cvf1);
415415

416416
// CHECK: [[IAS1:%.+]] = load i32, ptr addrspace(1) @int_as_one, align 4
@@ -431,21 +431,21 @@ void test_builtin_elementwise_min(float f1, float f2, double d1, double d2,
431431
// CHECK-LABEL: define void @test_builtin_elementwise_min(
432432
// CHECK: [[F1:%.+]] = load float, ptr %f1.addr, align 4
433433
// CHECK-NEXT: [[F2:%.+]] = load float, ptr %f2.addr, align 4
434-
// CHECK-NEXT: call float @llvm.minnum.f32(float [[F1]], float [[F2]])
434+
// CHECK-NEXT: call nsz float @llvm.minnum.f32(float [[F1]], float [[F2]])
435435
f1 = __builtin_elementwise_min(f1, f2);
436436

437437
// CHECK: [[D1:%.+]] = load double, ptr %d1.addr, align 8
438438
// CHECK-NEXT: [[D2:%.+]] = load double, ptr %d2.addr, align 8
439-
// CHECK-NEXT: call double @llvm.minnum.f64(double [[D1]], double [[D2]])
439+
// CHECK-NEXT: call nsz double @llvm.minnum.f64(double [[D1]], double [[D2]])
440440
d1 = __builtin_elementwise_min(d1, d2);
441441

442442
// CHECK: [[D1:%.+]] = load double, ptr %d1.addr, align 8
443-
// CHECK-NEXT: call double @llvm.minnum.f64(double [[D1]], double 2.000000e+00)
443+
// CHECK-NEXT: call nsz double @llvm.minnum.f64(double [[D1]], double 2.000000e+00)
444444
d1 = __builtin_elementwise_min(d1, 2.0);
445445

446446
// CHECK: [[VF1:%.+]] = load <4 x float>, ptr %vf1.addr, align 16
447447
// CHECK-NEXT: [[VF2:%.+]] = load <4 x float>, ptr %vf2.addr, align 16
448-
// CHECK-NEXT: call <4 x float> @llvm.minnum.v4f32(<4 x float> [[VF1]], <4 x float> [[VF2]])
448+
// CHECK-NEXT: call nsz <4 x float> @llvm.minnum.v4f32(<4 x float> [[VF1]], <4 x float> [[VF2]])
449449
vf1 = __builtin_elementwise_min(vf1, vf2);
450450

451451
// CHECK: [[I1:%.+]] = load i64, ptr %i1.addr, align 8
@@ -495,13 +495,13 @@ void test_builtin_elementwise_min(float f1, float f2, double d1, double d2,
495495

496496
// CHECK: [[CVF1:%.+]] = load <4 x float>, ptr %cvf1, align 16
497497
// CHECK-NEXT: [[VF2:%.+]] = load <4 x float>, ptr %vf2.addr, align 16
498-
// CHECK-NEXT: call <4 x float> @llvm.minnum.v4f32(<4 x float> [[CVF1]], <4 x float> [[VF2]])
498+
// CHECK-NEXT: call nsz <4 x float> @llvm.minnum.v4f32(<4 x float> [[CVF1]], <4 x float> [[VF2]])
499499
const float4 cvf1 = vf1;
500500
vf1 = __builtin_elementwise_min(cvf1, vf2);
501501

502502
// CHECK: [[VF2:%.+]] = load <4 x float>, ptr %vf2.addr, align 16
503503
// CHECK-NEXT: [[CVF1:%.+]] = load <4 x float>, ptr %cvf1, align 16
504-
// CHECK-NEXT: call <4 x float> @llvm.minnum.v4f32(<4 x float> [[VF2]], <4 x float> [[CVF1]])
504+
// CHECK-NEXT: call nsz <4 x float> @llvm.minnum.v4f32(<4 x float> [[VF2]], <4 x float> [[CVF1]])
505505
vf1 = __builtin_elementwise_min(vf2, cvf1);
506506

507507
// CHECK: [[IAS1:%.+]] = load i32, ptr addrspace(1) @int_as_one, align 4

‎clang/test/CodeGen/strictfp-elementwise-builtins.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ float4 strict_elementwise_abs(float4 a) {
3030
// CHECK-LABEL: define dso_local noundef <4 x float> @_Z22strict_elementwise_maxDv4_fS_
3131
// CHECK-SAME: (<4 x float> noundef [[A:%.*]], <4 x float> noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] {
3232
// CHECK-NEXT: entry:
33-
// CHECK-NEXT: [[ELT_MAX:%.*]] = tail call <4 x float> @llvm.experimental.constrained.maxnum.v4f32(<4 x float> [[A]], <4 x float> [[B]], metadata !"fpexcept.strict") #[[ATTR4]]
33+
// CHECK-NEXT: [[ELT_MAX:%.*]] = tail call nsz <4 x float> @llvm.experimental.constrained.maxnum.v4f32(<4 x float> [[A]], <4 x float> [[B]], metadata !"fpexcept.strict") #[[ATTR4]]
3434
// CHECK-NEXT: ret <4 x float> [[ELT_MAX]]
3535
//
3636
float4 strict_elementwise_max(float4 a, float4 b) {
@@ -40,7 +40,7 @@ float4 strict_elementwise_max(float4 a, float4 b) {
4040
// CHECK-LABEL: define dso_local noundef <4 x float> @_Z22strict_elementwise_minDv4_fS_
4141
// CHECK-SAME: (<4 x float> noundef [[A:%.*]], <4 x float> noundef [[B:%.*]]) local_unnamed_addr #[[ATTR0]] {
4242
// CHECK-NEXT: entry:
43-
// CHECK-NEXT: [[ELT_MIN:%.*]] = tail call <4 x float> @llvm.experimental.constrained.minnum.v4f32(<4 x float> [[A]], <4 x float> [[B]], metadata !"fpexcept.strict") #[[ATTR4]]
43+
// CHECK-NEXT: [[ELT_MIN:%.*]] = tail call nsz <4 x float> @llvm.experimental.constrained.minnum.v4f32(<4 x float> [[A]], <4 x float> [[B]], metadata !"fpexcept.strict") #[[ATTR4]]
4444
// CHECK-NEXT: ret <4 x float> [[ELT_MIN]]
4545
//
4646
float4 strict_elementwise_min(float4 a, float4 b) {

‎clang/test/Headers/__clang_hip_math.hip‎

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1716,7 +1716,7 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
17161716
//
17171717
// AMDGCNSPIRV-LABEL: @test_fmaxf(
17181718
// AMDGCNSPIRV-NEXT: entry:
1719-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
1719+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
17201720
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
17211721
//
17221722
extern "C" __device__ float test_fmaxf(float x, float y) {
@@ -1740,7 +1740,7 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
17401740
//
17411741
// AMDGCNSPIRV-LABEL: @test_fmax(
17421742
// AMDGCNSPIRV-NEXT: entry:
1743-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
1743+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
17441744
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
17451745
//
17461746
extern "C" __device__ double test_fmax(double x, double y) {
@@ -1764,7 +1764,7 @@ extern "C" __device__ double test_fmax(double x, double y) {
17641764
//
17651765
// AMDGCNSPIRV-LABEL: @test_fminf(
17661766
// AMDGCNSPIRV-NEXT: entry:
1767-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
1767+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
17681768
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
17691769
//
17701770
extern "C" __device__ float test_fminf(float x, float y) {
@@ -1788,7 +1788,7 @@ extern "C" __device__ float test_fminf(float x, float y) {
17881788
//
17891789
// AMDGCNSPIRV-LABEL: @test_fmin(
17901790
// AMDGCNSPIRV-NEXT: entry:
1791-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
1791+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
17921792
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
17931793
//
17941794
extern "C" __device__ double test_fmin(double x, double y) {
@@ -6689,7 +6689,7 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
66896689
//
66906690
// AMDGCNSPIRV-LABEL: @test_float_min(
66916691
// AMDGCNSPIRV-NEXT: entry:
6692-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
6692+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
66936693
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
66946694
//
66956695
extern "C" __device__ float test_float_min(float x, float y) {
@@ -6713,7 +6713,7 @@ extern "C" __device__ float test_float_min(float x, float y) {
67136713
//
67146714
// AMDGCNSPIRV-LABEL: @test_float_max(
67156715
// AMDGCNSPIRV-NEXT: entry:
6716-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
6716+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
67176717
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
67186718
//
67196719
extern "C" __device__ float test_float_max(float x, float y) {
@@ -6737,7 +6737,7 @@ extern "C" __device__ float test_float_max(float x, float y) {
67376737
//
67386738
// AMDGCNSPIRV-LABEL: @test_double_min(
67396739
// AMDGCNSPIRV-NEXT: entry:
6740-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
6740+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
67416741
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
67426742
//
67436743
extern "C" __device__ double test_double_min(double x, double y) {
@@ -6761,7 +6761,7 @@ extern "C" __device__ double test_double_min(double x, double y) {
67616761
//
67626762
// AMDGCNSPIRV-LABEL: @test_double_max(
67636763
// AMDGCNSPIRV-NEXT: entry:
6764-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
6764+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
67656765
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
67666766
//
67676767
extern "C" __device__ double test_double_max(double x, double y) {

0 commit comments

Comments
 (0)