Skip to content

Commit 50ed841

Browse files
Prince781Prince781
committed
[tests] update f32x2-instructions selp.b16 and sink moves
Change selp.u16 -> selp.b16 and
1 parent 955d811 commit 50ed841

File tree

1 file changed

+29
-29
lines changed

1 file changed

+29
-29
lines changed

‎llvm/test/CodeGen/NVPTX/f32x2-instructions.ll‎

Lines changed: 29 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -780,9 +780,9 @@ define <2 x i1> @test_fcmp_une(<2 x float> %a, <2 x float> %b) #0 {
780780
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
781781
; CHECK-NEXT: setp.neu.f32 %p1, %f4, %f2;
782782
; CHECK-NEXT: setp.neu.f32 %p2, %f3, %f1;
783-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
783+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
784784
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
785-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
785+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
786786
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
787787
; CHECK-NEXT: ret;
788788
%r = fcmp une <2 x float> %a, %b
@@ -804,9 +804,9 @@ define <2 x i1> @test_fcmp_ueq(<2 x float> %a, <2 x float> %b) #0 {
804804
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
805805
; CHECK-NEXT: setp.equ.f32 %p1, %f4, %f2;
806806
; CHECK-NEXT: setp.equ.f32 %p2, %f3, %f1;
807-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
807+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
808808
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
809-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
809+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
810810
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
811811
; CHECK-NEXT: ret;
812812
%r = fcmp ueq <2 x float> %a, %b
@@ -828,9 +828,9 @@ define <2 x i1> @test_fcmp_ugt(<2 x float> %a, <2 x float> %b) #0 {
828828
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
829829
; CHECK-NEXT: setp.gtu.f32 %p1, %f4, %f2;
830830
; CHECK-NEXT: setp.gtu.f32 %p2, %f3, %f1;
831-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
831+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
832832
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
833-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
833+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
834834
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
835835
; CHECK-NEXT: ret;
836836
%r = fcmp ugt <2 x float> %a, %b
@@ -852,9 +852,9 @@ define <2 x i1> @test_fcmp_uge(<2 x float> %a, <2 x float> %b) #0 {
852852
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
853853
; CHECK-NEXT: setp.geu.f32 %p1, %f4, %f2;
854854
; CHECK-NEXT: setp.geu.f32 %p2, %f3, %f1;
855-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
855+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
856856
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
857-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
857+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
858858
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
859859
; CHECK-NEXT: ret;
860860
%r = fcmp uge <2 x float> %a, %b
@@ -876,9 +876,9 @@ define <2 x i1> @test_fcmp_ult(<2 x float> %a, <2 x float> %b) #0 {
876876
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
877877
; CHECK-NEXT: setp.ltu.f32 %p1, %f4, %f2;
878878
; CHECK-NEXT: setp.ltu.f32 %p2, %f3, %f1;
879-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
879+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
880880
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
881-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
881+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
882882
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
883883
; CHECK-NEXT: ret;
884884
%r = fcmp ult <2 x float> %a, %b
@@ -900,9 +900,9 @@ define <2 x i1> @test_fcmp_ule(<2 x float> %a, <2 x float> %b) #0 {
900900
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
901901
; CHECK-NEXT: setp.leu.f32 %p1, %f4, %f2;
902902
; CHECK-NEXT: setp.leu.f32 %p2, %f3, %f1;
903-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
903+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
904904
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
905-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
905+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
906906
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
907907
; CHECK-NEXT: ret;
908908
%r = fcmp ule <2 x float> %a, %b
@@ -924,9 +924,9 @@ define <2 x i1> @test_fcmp_uno(<2 x float> %a, <2 x float> %b) #0 {
924924
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
925925
; CHECK-NEXT: setp.nan.f32 %p1, %f4, %f2;
926926
; CHECK-NEXT: setp.nan.f32 %p2, %f3, %f1;
927-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
927+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
928928
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
929-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
929+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
930930
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
931931
; CHECK-NEXT: ret;
932932
%r = fcmp uno <2 x float> %a, %b
@@ -948,9 +948,9 @@ define <2 x i1> @test_fcmp_one(<2 x float> %a, <2 x float> %b) #0 {
948948
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
949949
; CHECK-NEXT: setp.ne.f32 %p1, %f4, %f2;
950950
; CHECK-NEXT: setp.ne.f32 %p2, %f3, %f1;
951-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
951+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
952952
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
953-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
953+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
954954
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
955955
; CHECK-NEXT: ret;
956956
%r = fcmp one <2 x float> %a, %b
@@ -972,9 +972,9 @@ define <2 x i1> @test_fcmp_oeq(<2 x float> %a, <2 x float> %b) #0 {
972972
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
973973
; CHECK-NEXT: setp.eq.f32 %p1, %f4, %f2;
974974
; CHECK-NEXT: setp.eq.f32 %p2, %f3, %f1;
975-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
975+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
976976
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
977-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
977+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
978978
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
979979
; CHECK-NEXT: ret;
980980
%r = fcmp oeq <2 x float> %a, %b
@@ -996,9 +996,9 @@ define <2 x i1> @test_fcmp_ogt(<2 x float> %a, <2 x float> %b) #0 {
996996
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
997997
; CHECK-NEXT: setp.gt.f32 %p1, %f4, %f2;
998998
; CHECK-NEXT: setp.gt.f32 %p2, %f3, %f1;
999-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
999+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
10001000
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
1001-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
1001+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
10021002
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
10031003
; CHECK-NEXT: ret;
10041004
%r = fcmp ogt <2 x float> %a, %b
@@ -1020,9 +1020,9 @@ define <2 x i1> @test_fcmp_oge(<2 x float> %a, <2 x float> %b) #0 {
10201020
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
10211021
; CHECK-NEXT: setp.ge.f32 %p1, %f4, %f2;
10221022
; CHECK-NEXT: setp.ge.f32 %p2, %f3, %f1;
1023-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
1023+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
10241024
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
1025-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
1025+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
10261026
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
10271027
; CHECK-NEXT: ret;
10281028
%r = fcmp oge <2 x float> %a, %b
@@ -1044,9 +1044,9 @@ define <2 x i1> @test_fcmp_olt(<2 x float> %a, <2 x float> %b) #0 {
10441044
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
10451045
; CHECK-NEXT: setp.lt.f32 %p1, %f4, %f2;
10461046
; CHECK-NEXT: setp.lt.f32 %p2, %f3, %f1;
1047-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
1047+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
10481048
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
1049-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
1049+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
10501050
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
10511051
; CHECK-NEXT: ret;
10521052
%r = fcmp olt <2 x float> %a, %b
@@ -1068,9 +1068,9 @@ define <2 x i1> @test_fcmp_ole(<2 x float> %a, <2 x float> %b) #0 {
10681068
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
10691069
; CHECK-NEXT: setp.le.f32 %p1, %f4, %f2;
10701070
; CHECK-NEXT: setp.le.f32 %p2, %f3, %f1;
1071-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
1071+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
10721072
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
1073-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
1073+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
10741074
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
10751075
; CHECK-NEXT: ret;
10761076
%r = fcmp ole <2 x float> %a, %b
@@ -1092,9 +1092,9 @@ define <2 x i1> @test_fcmp_ord(<2 x float> %a, <2 x float> %b) #0 {
10921092
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
10931093
; CHECK-NEXT: setp.num.f32 %p1, %f4, %f2;
10941094
; CHECK-NEXT: setp.num.f32 %p2, %f3, %f1;
1095-
; CHECK-NEXT: selp.u16 %rs1, -1, 0, %p2;
1095+
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
10961096
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
1097-
; CHECK-NEXT: selp.u16 %rs2, -1, 0, %p1;
1097+
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
10981098
; CHECK-NEXT: st.param.b8 [func_retval0+1], %rs2;
10991099
; CHECK-NEXT: ret;
11001100
%r = fcmp ord <2 x float> %a, %b
@@ -1298,7 +1298,7 @@ define <2 x i32> @test_bitcast_2xfloat_to_2xi32(<2 x float> %a) #0 {
12981298
; CHECK-EMPTY:
12991299
; CHECK-NEXT: // %bb.0:
13001300
; CHECK-NEXT: ld.param.u64 %rd2, [test_bitcast_2xfloat_to_2xi32_param_0];
1301-
; CHECK-NEXT: { .reg .b32 tmp; mov.b64 {tmp, %r1}, %rd2; }
1301+
; CHECK-NEXT: mov.b64 {_, %r1}, %rd2;
13021302
; CHECK-NEXT: cvt.u32.u64 %r2, %rd2;
13031303
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};
13041304
; CHECK-NEXT: ret;

0 commit comments

Comments
 (0)