Skip to content

Commit 77267f8

Browse files
wzssyqawzssyqa
committed
fix some tests
1 parent ce675ee commit 77267f8

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
@@ -4387,8 +4387,7 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
43874387
} else {
43884388
FastMathFlags FMF;
43894389
FMF.setNoSignedZeros(true);
4390-
FMFSource FMFSrc(FMF);
4391-
Result = Builder.CreateMaxNum(Op0, Op1, /*FMFSource=*/FMFSrc, "elt.max");
4390+
Result = Builder.CreateMaxNum(Op0, Op1, /*FMFSource=*/FMF, "elt.max");
43924391
}
43934392
return RValue::get(Result);
43944393
}
@@ -4407,8 +4406,7 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
44074406
} else {
44084407
FastMathFlags FMF;
44094408
FMF.setNoSignedZeros(true);
4410-
FMFSource FMFSrc(FMF);
4411-
Result = Builder.CreateMinNum(Op0, Op1, /*FMFSource=*/FMFSrc, "elt.min");
4409+
Result = Builder.CreateMinNum(Op0, Op1, /*FMFSource=*/FMF, "elt.min");
44124410
}
44134411
return RValue::get(Result);
44144412
}

‎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-bulitins.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
@@ -1717,7 +1717,7 @@ extern "C" __device__ double test_fma_rn(double x, double y, double z) {
17171717
//
17181718
// AMDGCNSPIRV-LABEL: @test_fmaxf(
17191719
// AMDGCNSPIRV-NEXT: entry:
1720-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
1720+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
17211721
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
17221722
//
17231723
extern "C" __device__ float test_fmaxf(float x, float y) {
@@ -1741,7 +1741,7 @@ extern "C" __device__ float test_fmaxf(float x, float y) {
17411741
//
17421742
// AMDGCNSPIRV-LABEL: @test_fmax(
17431743
// AMDGCNSPIRV-NEXT: entry:
1744-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
1744+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
17451745
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
17461746
//
17471747
extern "C" __device__ double test_fmax(double x, double y) {
@@ -1766,7 +1766,7 @@ extern "C" __device__ double test_fmax(double x, double y) {
17661766
//
17671767
// AMDGCNSPIRV-LABEL: @test_fminf(
17681768
// AMDGCNSPIRV-NEXT: entry:
1769-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
1769+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
17701770
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
17711771
//
17721772
extern "C" __device__ float test_fminf(float x, float y) {
@@ -1790,7 +1790,7 @@ extern "C" __device__ float test_fminf(float x, float y) {
17901790
//
17911791
// AMDGCNSPIRV-LABEL: @test_fmin(
17921792
// AMDGCNSPIRV-NEXT: entry:
1793-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
1793+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
17941794
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
17951795
//
17961796
extern "C" __device__ double test_fmin(double x, double y) {
@@ -6755,7 +6755,7 @@ extern "C" __device__ double test__fma_rn(double x, double y, double z) {
67556755
//
67566756
// AMDGCNSPIRV-LABEL: @test_float_min(
67576757
// AMDGCNSPIRV-NEXT: entry:
6758-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
6758+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.minnum.f32(float [[X:%.*]], float [[Y:%.*]])
67596759
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
67606760
//
67616761
extern "C" __device__ float test_float_min(float x, float y) {
@@ -6779,7 +6779,7 @@ extern "C" __device__ float test_float_min(float x, float y) {
67796779
//
67806780
// AMDGCNSPIRV-LABEL: @test_float_max(
67816781
// AMDGCNSPIRV-NEXT: entry:
6782-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
6782+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) float @llvm.maxnum.f32(float [[X:%.*]], float [[Y:%.*]])
67836783
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
67846784
//
67856785
extern "C" __device__ float test_float_max(float x, float y) {
@@ -6803,7 +6803,7 @@ extern "C" __device__ float test_float_max(float x, float y) {
68036803
//
68046804
// AMDGCNSPIRV-LABEL: @test_double_min(
68056805
// AMDGCNSPIRV-NEXT: entry:
6806-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
6806+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.minnum.f64(double [[X:%.*]], double [[Y:%.*]])
68076807
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
68086808
//
68096809
extern "C" __device__ double test_double_min(double x, double y) {
@@ -6827,7 +6827,7 @@ extern "C" __device__ double test_double_min(double x, double y) {
68276827
//
68286828
// AMDGCNSPIRV-LABEL: @test_double_max(
68296829
// AMDGCNSPIRV-NEXT: entry:
6830-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
6830+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call nsz contract noundef addrspace(4) double @llvm.maxnum.f64(double [[X:%.*]], double [[Y:%.*]])
68316831
// AMDGCNSPIRV-NEXT: ret double [[TMP0]]
68326832
//
68336833
extern "C" __device__ double test_double_max(double x, double y) {

0 commit comments

Comments
 (0)