Skip to content

Commit c153bb8

Browse files
wzssyqawzssyqa
committed
Fix testcase
1 parent 842f8b6 commit c153bb8

File tree

2 files changed

+64
-199
lines changed

2 files changed

+64
-199
lines changed

‎llvm/test/CodeGen/NVPTX/math-intrins.ll

Lines changed: 30 additions & 100 deletions
Original file line numberDiff line numberDiff line change
@@ -685,25 +685,16 @@ define half @minimum_half(half %a, half %b) {
685685
define float @minimum_float(float %a, float %b) {
686686
; CHECK-NOF16-LABEL: minimum_float(
687687
; CHECK-NOF16: {
688-
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
689-
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
690-
; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
688+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
689+
; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
691690
; CHECK-NOF16-EMPTY:
692691
; CHECK-NOF16-NEXT: // %bb.0:
693692
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
694-
; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
695693
; CHECK-NOF16-NEXT: ld.param.f32 %f2, [minimum_float_param_1];
696694
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f2;
697695
; CHECK-NOF16-NEXT: min.f32 %f3, %f1, %f2;
698696
; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
699-
; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
700-
; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
701-
; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
702-
; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, -2147483648;
703-
; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
704-
; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %f4, 0f00000000;
705-
; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
706-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f7;
697+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f4;
707698
; CHECK-NOF16-NEXT: ret;
708699
;
709700
; CHECK-F16-LABEL: minimum_float(
@@ -734,21 +725,15 @@ define float @minimum_float(float %a, float %b) {
734725
define float @minimum_imm1(float %a) {
735726
; CHECK-NOF16-LABEL: minimum_imm1(
736727
; CHECK-NOF16: {
737-
; CHECK-NOF16-NEXT: .reg .pred %p<4>;
738-
; CHECK-NOF16-NEXT: .reg .b32 %r<2>;
739-
; CHECK-NOF16-NEXT: .reg .b32 %f<6>;
728+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
729+
; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
740730
; CHECK-NOF16-EMPTY:
741731
; CHECK-NOF16-NEXT: // %bb.0:
742732
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
743-
; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
744733
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
745734
; CHECK-NOF16-NEXT: min.f32 %f2, %f1, 0f00000000;
746735
; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
747-
; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
748-
; CHECK-NOF16-NEXT: selp.f32 %f4, %f1, %f3, %p2;
749-
; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %f3, 0f00000000;
750-
; CHECK-NOF16-NEXT: selp.f32 %f5, %f4, %f3, %p3;
751-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f5;
736+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f3;
752737
; CHECK-NOF16-NEXT: ret;
753738
;
754739
; CHECK-F16-LABEL: minimum_imm1(
@@ -777,21 +762,15 @@ define float @minimum_imm1(float %a) {
777762
define float @minimum_imm2(float %a) {
778763
; CHECK-NOF16-LABEL: minimum_imm2(
779764
; CHECK-NOF16: {
780-
; CHECK-NOF16-NEXT: .reg .pred %p<4>;
781-
; CHECK-NOF16-NEXT: .reg .b32 %r<2>;
782-
; CHECK-NOF16-NEXT: .reg .b32 %f<6>;
765+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
766+
; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
783767
; CHECK-NOF16-EMPTY:
784768
; CHECK-NOF16-NEXT: // %bb.0:
785769
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
786-
; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
787770
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
788771
; CHECK-NOF16-NEXT: min.f32 %f2, %f1, 0f00000000;
789772
; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
790-
; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
791-
; CHECK-NOF16-NEXT: selp.f32 %f4, %f1, %f3, %p2;
792-
; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %f3, 0f00000000;
793-
; CHECK-NOF16-NEXT: selp.f32 %f5, %f4, %f3, %p3;
794-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f5;
773+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f3;
795774
; CHECK-NOF16-NEXT: ret;
796775
;
797776
; CHECK-F16-LABEL: minimum_imm2(
@@ -820,25 +799,16 @@ define float @minimum_imm2(float %a) {
820799
define float @minimum_float_ftz(float %a, float %b) #1 {
821800
; CHECK-NOF16-LABEL: minimum_float_ftz(
822801
; CHECK-NOF16: {
823-
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
824-
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
825-
; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
802+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
803+
; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
826804
; CHECK-NOF16-EMPTY:
827805
; CHECK-NOF16-NEXT: // %bb.0:
828806
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
829-
; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
830807
; CHECK-NOF16-NEXT: ld.param.f32 %f2, [minimum_float_ftz_param_1];
831808
; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %f1, %f2;
832809
; CHECK-NOF16-NEXT: min.ftz.f32 %f3, %f1, %f2;
833810
; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
834-
; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
835-
; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
836-
; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
837-
; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, -2147483648;
838-
; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
839-
; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %f4, 0f00000000;
840-
; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
841-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f7;
811+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f4;
842812
; CHECK-NOF16-NEXT: ret;
843813
;
844814
; CHECK-F16-LABEL: minimum_float_ftz(
@@ -869,25 +839,16 @@ define float @minimum_float_ftz(float %a, float %b) #1 {
869839
define double @minimum_double(double %a, double %b) {
870840
; CHECK-LABEL: minimum_double(
871841
; CHECK: {
872-
; CHECK-NEXT: .reg .pred %p<5>;
873-
; CHECK-NEXT: .reg .b64 %rd<3>;
874-
; CHECK-NEXT: .reg .b64 %fd<8>;
842+
; CHECK-NEXT: .reg .pred %p<2>;
843+
; CHECK-NEXT: .reg .b64 %fd<5>;
875844
; CHECK-EMPTY:
876845
; CHECK-NEXT: // %bb.0:
877846
; CHECK-NEXT: ld.param.f64 %fd1, [minimum_double_param_0];
878-
; CHECK-NEXT: mov.b64 %rd1, %fd1;
879847
; CHECK-NEXT: ld.param.f64 %fd2, [minimum_double_param_1];
880848
; CHECK-NEXT: setp.nan.f64 %p1, %fd1, %fd2;
881849
; CHECK-NEXT: min.f64 %fd3, %fd1, %fd2;
882850
; CHECK-NEXT: selp.f64 %fd4, 0d7FF8000000000000, %fd3, %p1;
883-
; CHECK-NEXT: setp.eq.s64 %p2, %rd1, -9223372036854775808;
884-
; CHECK-NEXT: selp.f64 %fd5, %fd1, %fd4, %p2;
885-
; CHECK-NEXT: mov.b64 %rd2, %fd2;
886-
; CHECK-NEXT: setp.eq.s64 %p3, %rd2, -9223372036854775808;
887-
; CHECK-NEXT: selp.f64 %fd6, %fd2, %fd5, %p3;
888-
; CHECK-NEXT: setp.eq.f64 %p4, %fd4, 0d0000000000000000;
889-
; CHECK-NEXT: selp.f64 %fd7, %fd6, %fd4, %p4;
890-
; CHECK-NEXT: st.param.f64 [func_retval0], %fd7;
851+
; CHECK-NEXT: st.param.f64 [func_retval0], %fd4;
891852
; CHECK-NEXT: ret;
892853
%x = call double @llvm.minimum.f64(double %a, double %b)
893854
ret double %x
@@ -1243,17 +1204,15 @@ define half @maximum_half(half %a, half %b) {
12431204
define float @maximum_imm1(float %a) {
12441205
; CHECK-NOF16-LABEL: maximum_imm1(
12451206
; CHECK-NOF16: {
1246-
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
1247-
; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
1207+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
1208+
; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
12481209
; CHECK-NOF16-EMPTY:
12491210
; CHECK-NOF16-NEXT: // %bb.0:
12501211
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm1_param_0];
12511212
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
12521213
; CHECK-NOF16-NEXT: max.f32 %f2, %f1, 0f00000000;
12531214
; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
1254-
; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %f3, 0f00000000;
1255-
; CHECK-NOF16-NEXT: selp.f32 %f4, 0f00000000, %f3, %p2;
1256-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f4;
1215+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f3;
12571216
; CHECK-NOF16-NEXT: ret;
12581217
;
12591218
; CHECK-F16-LABEL: maximum_imm1(
@@ -1282,17 +1241,15 @@ define float @maximum_imm1(float %a) {
12821241
define float @maximum_imm2(float %a) {
12831242
; CHECK-NOF16-LABEL: maximum_imm2(
12841243
; CHECK-NOF16: {
1285-
; CHECK-NOF16-NEXT: .reg .pred %p<3>;
1286-
; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
1244+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
1245+
; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
12871246
; CHECK-NOF16-EMPTY:
12881247
; CHECK-NOF16-NEXT: // %bb.0:
12891248
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm2_param_0];
12901249
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
12911250
; CHECK-NOF16-NEXT: max.f32 %f2, %f1, 0f00000000;
12921251
; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
1293-
; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %f3, 0f00000000;
1294-
; CHECK-NOF16-NEXT: selp.f32 %f4, 0f00000000, %f3, %p2;
1295-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f4;
1252+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f3;
12961253
; CHECK-NOF16-NEXT: ret;
12971254
;
12981255
; CHECK-F16-LABEL: maximum_imm2(
@@ -1321,25 +1278,16 @@ define float @maximum_imm2(float %a) {
13211278
define float @maximum_float(float %a, float %b) {
13221279
; CHECK-NOF16-LABEL: maximum_float(
13231280
; CHECK-NOF16: {
1324-
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
1325-
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
1326-
; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
1281+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
1282+
; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
13271283
; CHECK-NOF16-EMPTY:
13281284
; CHECK-NOF16-NEXT: // %bb.0:
13291285
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
1330-
; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
13311286
; CHECK-NOF16-NEXT: ld.param.f32 %f2, [maximum_float_param_1];
13321287
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f2;
13331288
; CHECK-NOF16-NEXT: max.f32 %f3, %f1, %f2;
13341289
; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
1335-
; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, 0;
1336-
; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
1337-
; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
1338-
; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, 0;
1339-
; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
1340-
; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %f4, 0f00000000;
1341-
; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
1342-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f7;
1290+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f4;
13431291
; CHECK-NOF16-NEXT: ret;
13441292
;
13451293
; CHECK-F16-LABEL: maximum_float(
@@ -1370,25 +1318,16 @@ define float @maximum_float(float %a, float %b) {
13701318
define float @maximum_float_ftz(float %a, float %b) #1 {
13711319
; CHECK-NOF16-LABEL: maximum_float_ftz(
13721320
; CHECK-NOF16: {
1373-
; CHECK-NOF16-NEXT: .reg .pred %p<5>;
1374-
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
1375-
; CHECK-NOF16-NEXT: .reg .b32 %f<8>;
1321+
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
1322+
; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
13761323
; CHECK-NOF16-EMPTY:
13771324
; CHECK-NOF16-NEXT: // %bb.0:
13781325
; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
1379-
; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
13801326
; CHECK-NOF16-NEXT: ld.param.f32 %f2, [maximum_float_ftz_param_1];
13811327
; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %f1, %f2;
13821328
; CHECK-NOF16-NEXT: max.ftz.f32 %f3, %f1, %f2;
13831329
; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
1384-
; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, 0;
1385-
; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
1386-
; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
1387-
; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, 0;
1388-
; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
1389-
; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %f4, 0f00000000;
1390-
; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
1391-
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f7;
1330+
; CHECK-NOF16-NEXT: st.param.f32 [func_retval0], %f4;
13921331
; CHECK-NOF16-NEXT: ret;
13931332
;
13941333
; CHECK-F16-LABEL: maximum_float_ftz(
@@ -1419,25 +1358,16 @@ define float @maximum_float_ftz(float %a, float %b) #1 {
14191358
define double @maximum_double(double %a, double %b) {
14201359
; CHECK-LABEL: maximum_double(
14211360
; CHECK: {
1422-
; CHECK-NEXT: .reg .pred %p<5>;
1423-
; CHECK-NEXT: .reg .b64 %rd<3>;
1424-
; CHECK-NEXT: .reg .b64 %fd<8>;
1361+
; CHECK-NEXT: .reg .pred %p<2>;
1362+
; CHECK-NEXT: .reg .b64 %fd<5>;
14251363
; CHECK-EMPTY:
14261364
; CHECK-NEXT: // %bb.0:
14271365
; CHECK-NEXT: ld.param.f64 %fd1, [maximum_double_param_0];
1428-
; CHECK-NEXT: mov.b64 %rd1, %fd1;
14291366
; CHECK-NEXT: ld.param.f64 %fd2, [maximum_double_param_1];
14301367
; CHECK-NEXT: setp.nan.f64 %p1, %fd1, %fd2;
14311368
; CHECK-NEXT: max.f64 %fd3, %fd1, %fd2;
14321369
; CHECK-NEXT: selp.f64 %fd4, 0d7FF8000000000000, %fd3, %p1;
1433-
; CHECK-NEXT: setp.eq.s64 %p2, %rd1, 0;
1434-
; CHECK-NEXT: selp.f64 %fd5, %fd1, %fd4, %p2;
1435-
; CHECK-NEXT: mov.b64 %rd2, %fd2;
1436-
; CHECK-NEXT: setp.eq.s64 %p3, %rd2, 0;
1437-
; CHECK-NEXT: selp.f64 %fd6, %fd2, %fd5, %p3;
1438-
; CHECK-NEXT: setp.eq.f64 %p4, %fd4, 0d0000000000000000;
1439-
; CHECK-NEXT: selp.f64 %fd7, %fd6, %fd4, %p4;
1440-
; CHECK-NEXT: st.param.f64 [func_retval0], %fd7;
1370+
; CHECK-NEXT: st.param.f64 [func_retval0], %fd4;
14411371
; CHECK-NEXT: ret;
14421372
%x = call double @llvm.maximum.f64(double %a, double %b)
14431373
ret double %x

0 commit comments

Comments
 (0)