From 06a4e5caf624c80d46e3f601084ff568dbc75e16 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Tue, 24 Sep 2024 17:08:14 -0700 Subject: [PATCH 1/9] Initial commit --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 31 +- llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 494 +++++++++---------- 2 files changed, 264 insertions(+), 261 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 01abf9591e342..7b4c6a8d34048 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -2328,20 +2328,23 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, // Lower non-const v4i8 vector as byte-wise constructed i32, which allows us // to optimize calculation of constant parts. if (VT == MVT::v4i8) { - SDValue C8 = DAG.getConstant(8, DL, MVT::i32); - SDValue E01 = DAG.getNode( - NVPTXISD::BFI, DL, MVT::i32, - DAG.getAnyExtOrTrunc(Op->getOperand(1), DL, MVT::i32), - DAG.getAnyExtOrTrunc(Op->getOperand(0), DL, MVT::i32), C8, C8); - SDValue E012 = - DAG.getNode(NVPTXISD::BFI, DL, MVT::i32, - DAG.getAnyExtOrTrunc(Op->getOperand(2), DL, MVT::i32), - E01, DAG.getConstant(16, DL, MVT::i32), C8); - SDValue E0123 = - DAG.getNode(NVPTXISD::BFI, DL, MVT::i32, - DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32), - E012, DAG.getConstant(24, DL, MVT::i32), C8); - return DAG.getNode(ISD::BITCAST, DL, VT, E0123); + SDValue PRMT__10 = DAG.getNode( + NVPTXISD::PRMT, DL, MVT::v4i8, + {DAG.getAnyExtOrTrunc(Op->getOperand(0), DL, MVT::i32), + DAG.getAnyExtOrTrunc(Op->getOperand(1), DL, MVT::i32), + DAG.getConstant(0x3340, DL, MVT::i32), + DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); + SDValue PRMT_210 = DAG.getNode( + NVPTXISD::PRMT, DL, MVT::v4i8, + {PRMT__10, DAG.getAnyExtOrTrunc(Op->getOperand(2), DL, MVT::i32), + DAG.getConstant(0x3410, DL, MVT::i32), + DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); + SDValue PRMT3210 = DAG.getNode( + NVPTXISD::PRMT, DL, MVT::v4i8, + {PRMT_210, DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32), + DAG.getConstant(0x4210, DL, MVT::i32), + DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); + return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210); } return Op; } diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll index 5b5662a1eea76..fdc25cf95d06a 100644 --- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll @@ -18,7 +18,7 @@ define <4 x i8> @test_ret_const() #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: mov.b32 %r1, -66911489; -; CHECK-NEXT: st.param.b32 [func_retval0], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; ; CHECK-NEXT: ret; ret <4 x i8> } @@ -31,7 +31,7 @@ define i8 @test_extract_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_0_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 0, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 0 ret i8 %e @@ -45,7 +45,7 @@ define i8 @test_extract_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_1_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 8, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 1 ret i8 %e @@ -59,7 +59,7 @@ define i8 @test_extract_2(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_2_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 16, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 2 ret i8 %e @@ -73,7 +73,7 @@ define i8 @test_extract_3(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_3_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 3 ret i8 %e @@ -91,7 +91,7 @@ define i8 @test_extract_i(<4 x i8> %a, i64 %idx) #0 { ; CHECK-NEXT: cvt.u32.u64 %r2, %rd1; ; CHECK-NEXT: shl.b32 %r3, %r2, 3; ; CHECK-NEXT: bfe.u32 %r4, %r1, %r3, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r4; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i64 %idx ret i8 %e @@ -101,39 +101,39 @@ define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-LABEL: test_add( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<19>; +; CHECK-NEXT: .reg .b32 %r<18>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_add_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_add_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; ; CHECK-NEXT: add.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.u32 %r6, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: add.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.u32 %r10, %r2, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: bfe.u32 %r11, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; ; CHECK-NEXT: add.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; +; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; ; CHECK-NEXT: bfe.u32 %r14, %r2, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; ; CHECK-NEXT: bfe.u32 %r15, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; ; CHECK-NEXT: add.s16 %rs12, %rs11, %rs10; ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; +; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %r = add <4 x i8> %a, %b ret <4 x i8> %r @@ -143,30 +143,30 @@ define <4 x i8> @test_add_imm_0(<4 x i8> %a) #0 { ; CHECK-LABEL: test_add_imm_0( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<9>; -; CHECK-NEXT: .reg .b32 %r<14>; +; CHECK-NEXT: .reg .b32 %r<13>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_add_imm_0_param_0]; -; CHECK-NEXT: bfe.u32 %r2, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r2, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; -; CHECK-NEXT: add.s16 %rs2, %rs1, 1; +; CHECK-NEXT: add.s16 %rs2, %rs1, 2; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs2; -; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; -; CHECK-NEXT: add.s16 %rs4, %rs3, 2; +; CHECK-NEXT: add.s16 %rs4, %rs3, 1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; -; CHECK-NEXT: bfi.b32 %r6, %r5, %r3, 8, 8; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120; ; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: add.s16 %rs6, %rs5, 3; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: bfi.b32 %r9, %r8, %r6, 16, 8; +; CHECK-NEXT: prmt.b32 %r9, %r6, %r8, 13328; ; CHECK-NEXT: bfe.u32 %r10, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: add.s16 %rs8, %rs7, 4; ; CHECK-NEXT: cvt.u32.u16 %r11, %rs8; -; CHECK-NEXT: bfi.b32 %r12, %r11, %r9, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r12; +; CHECK-NEXT: prmt.b32 %r12, %r9, %r11, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; ; CHECK-NEXT: ret; %r = add <4 x i8> , %a ret <4 x i8> %r @@ -176,30 +176,30 @@ define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 { ; CHECK-LABEL: test_add_imm_1( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<9>; -; CHECK-NEXT: .reg .b32 %r<14>; +; CHECK-NEXT: .reg .b32 %r<13>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_add_imm_1_param_0]; -; CHECK-NEXT: bfe.u32 %r2, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r2, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; -; CHECK-NEXT: add.s16 %rs2, %rs1, 1; +; CHECK-NEXT: add.s16 %rs2, %rs1, 2; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs2; -; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; -; CHECK-NEXT: add.s16 %rs4, %rs3, 2; +; CHECK-NEXT: add.s16 %rs4, %rs3, 1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; -; CHECK-NEXT: bfi.b32 %r6, %r5, %r3, 8, 8; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120; ; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: add.s16 %rs6, %rs5, 3; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: bfi.b32 %r9, %r8, %r6, 16, 8; +; CHECK-NEXT: prmt.b32 %r9, %r6, %r8, 13328; ; CHECK-NEXT: bfe.u32 %r10, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: add.s16 %rs8, %rs7, 4; ; CHECK-NEXT: cvt.u32.u16 %r11, %rs8; -; CHECK-NEXT: bfi.b32 %r12, %r11, %r9, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r12; +; CHECK-NEXT: prmt.b32 %r12, %r9, %r11, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; ; CHECK-NEXT: ret; %r = add <4 x i8> %a, ret <4 x i8> %r @@ -209,39 +209,39 @@ define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-LABEL: test_sub( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<19>; +; CHECK-NEXT: .reg .b32 %r<18>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_sub_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_sub_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; ; CHECK-NEXT: sub.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.u32 %r6, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: sub.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.u32 %r10, %r2, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: bfe.u32 %r11, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; ; CHECK-NEXT: sub.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; +; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; ; CHECK-NEXT: bfe.u32 %r14, %r2, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; ; CHECK-NEXT: bfe.u32 %r15, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; ; CHECK-NEXT: sub.s16 %rs12, %rs11, %rs10; ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; +; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %r = sub <4 x i8> %a, %b ret <4 x i8> %r @@ -251,7 +251,7 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-LABEL: test_smax( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<27>; +; CHECK-NEXT: .reg .b32 %r<26>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_smax_param_1]; @@ -262,28 +262,28 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: bfe.s32 %r5, %r2, 16, 8; ; CHECK-NEXT: bfe.s32 %r6, %r1, 16, 8; ; CHECK-NEXT: setp.gt.s32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.s32 %r7, %r2, 8, 8; -; CHECK-NEXT: bfe.s32 %r8, %r1, 8, 8; +; CHECK-NEXT: bfe.s32 %r7, %r2, 0, 8; +; CHECK-NEXT: bfe.s32 %r8, %r1, 0, 8; ; CHECK-NEXT: setp.gt.s32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.s32 %r9, %r2, 0, 8; -; CHECK-NEXT: bfe.s32 %r10, %r1, 0, 8; +; CHECK-NEXT: bfe.s32 %r9, %r2, 8, 8; +; CHECK-NEXT: bfe.s32 %r10, %r1, 8, 8; ; CHECK-NEXT: setp.gt.s32 %p4, %r10, %r9; ; CHECK-NEXT: bfe.u32 %r11, %r1, 24, 8; ; CHECK-NEXT: bfe.u32 %r12, %r1, 16, 8; -; CHECK-NEXT: bfe.u32 %r13, %r1, 8, 8; -; CHECK-NEXT: bfe.u32 %r14, %r1, 0, 8; -; CHECK-NEXT: bfe.u32 %r15, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r13, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r14, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r15, %r2, 8, 8; ; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4; -; CHECK-NEXT: bfe.u32 %r17, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r17, %r2, 0, 8; ; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3; -; CHECK-NEXT: bfi.b32 %r19, %r18, %r16, 8, 8; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120; ; CHECK-NEXT: bfe.u32 %r20, %r2, 16, 8; ; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2; -; CHECK-NEXT: bfi.b32 %r22, %r21, %r19, 16, 8; +; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 13328; ; CHECK-NEXT: bfe.u32 %r23, %r2, 24, 8; ; CHECK-NEXT: selp.b32 %r24, %r11, %r23, %p1; -; CHECK-NEXT: bfi.b32 %r25, %r24, %r22, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r25; +; CHECK-NEXT: prmt.b32 %r25, %r22, %r24, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; ; CHECK-NEXT: ret; %cmp = icmp sgt <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -294,7 +294,7 @@ define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-LABEL: test_umax( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<19>; +; CHECK-NEXT: .reg .b32 %r<18>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_umax_param_1]; @@ -305,20 +305,20 @@ define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: bfe.u32 %r5, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r6, %r1, 16, 8; ; CHECK-NEXT: setp.hi.u32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.u32 %r7, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r8, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r8, %r1, 0, 8; ; CHECK-NEXT: setp.hi.u32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r10, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r9, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r10, %r1, 8, 8; ; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; ; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; ; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; +; CHECK-NEXT: prmt.b32 %r15, %r13, %r14, 13328; ; CHECK-NEXT: selp.b32 %r16, %r4, %r3, %p1; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; +; CHECK-NEXT: prmt.b32 %r17, %r15, %r16, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %cmp = icmp ugt <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -329,7 +329,7 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-LABEL: test_smin( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<27>; +; CHECK-NEXT: .reg .b32 %r<26>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_smin_param_1]; @@ -340,28 +340,28 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: bfe.s32 %r5, %r2, 16, 8; ; CHECK-NEXT: bfe.s32 %r6, %r1, 16, 8; ; CHECK-NEXT: setp.le.s32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.s32 %r7, %r2, 8, 8; -; CHECK-NEXT: bfe.s32 %r8, %r1, 8, 8; +; CHECK-NEXT: bfe.s32 %r7, %r2, 0, 8; +; CHECK-NEXT: bfe.s32 %r8, %r1, 0, 8; ; CHECK-NEXT: setp.le.s32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.s32 %r9, %r2, 0, 8; -; CHECK-NEXT: bfe.s32 %r10, %r1, 0, 8; +; CHECK-NEXT: bfe.s32 %r9, %r2, 8, 8; +; CHECK-NEXT: bfe.s32 %r10, %r1, 8, 8; ; CHECK-NEXT: setp.le.s32 %p4, %r10, %r9; ; CHECK-NEXT: bfe.u32 %r11, %r1, 24, 8; ; CHECK-NEXT: bfe.u32 %r12, %r1, 16, 8; -; CHECK-NEXT: bfe.u32 %r13, %r1, 8, 8; -; CHECK-NEXT: bfe.u32 %r14, %r1, 0, 8; -; CHECK-NEXT: bfe.u32 %r15, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r13, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r14, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r15, %r2, 8, 8; ; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4; -; CHECK-NEXT: bfe.u32 %r17, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r17, %r2, 0, 8; ; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3; -; CHECK-NEXT: bfi.b32 %r19, %r18, %r16, 8, 8; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120; ; CHECK-NEXT: bfe.u32 %r20, %r2, 16, 8; ; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2; -; CHECK-NEXT: bfi.b32 %r22, %r21, %r19, 16, 8; +; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 13328; ; CHECK-NEXT: bfe.u32 %r23, %r2, 24, 8; ; CHECK-NEXT: selp.b32 %r24, %r11, %r23, %p1; -; CHECK-NEXT: bfi.b32 %r25, %r24, %r22, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r25; +; CHECK-NEXT: prmt.b32 %r25, %r22, %r24, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; ; CHECK-NEXT: ret; %cmp = icmp sle <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -372,7 +372,7 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-LABEL: test_umin( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<19>; +; CHECK-NEXT: .reg .b32 %r<18>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_umin_param_1]; @@ -383,20 +383,20 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: bfe.u32 %r5, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r6, %r1, 16, 8; ; CHECK-NEXT: setp.ls.u32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.u32 %r7, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r8, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r8, %r1, 0, 8; ; CHECK-NEXT: setp.ls.u32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r10, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r9, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r10, %r1, 8, 8; ; CHECK-NEXT: setp.ls.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; ; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; ; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; +; CHECK-NEXT: prmt.b32 %r15, %r13, %r14, 13328; ; CHECK-NEXT: selp.b32 %r16, %r4, %r3, %p1; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; +; CHECK-NEXT: prmt.b32 %r17, %r15, %r16, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %cmp = icmp ule <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -407,7 +407,7 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-LABEL: test_eq( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<24>; +; CHECK-NEXT: .reg .b32 %r<23>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r3, [test_eq_param_2]; @@ -419,24 +419,24 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: bfe.u32 %r6, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; ; CHECK-NEXT: setp.eq.u32 %p2, %r7, %r6; -; CHECK-NEXT: bfe.u32 %r8, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r9, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r8, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r9, %r1, 0, 8; ; CHECK-NEXT: setp.eq.u32 %p3, %r9, %r8; -; CHECK-NEXT: bfe.u32 %r10, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r11, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; ; CHECK-NEXT: setp.eq.u32 %p4, %r11, %r10; -; CHECK-NEXT: bfe.u32 %r12, %r3, 0, 8; +; CHECK-NEXT: bfe.u32 %r12, %r3, 8, 8; ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; -; CHECK-NEXT: bfe.u32 %r14, %r3, 8, 8; +; CHECK-NEXT: bfe.u32 %r14, %r3, 0, 8; ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: bfi.b32 %r16, %r15, %r13, 8, 8; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120; ; CHECK-NEXT: bfe.u32 %r17, %r3, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; -; CHECK-NEXT: bfi.b32 %r19, %r18, %r16, 16, 8; +; CHECK-NEXT: prmt.b32 %r19, %r16, %r18, 13328; ; CHECK-NEXT: bfe.u32 %r20, %r3, 24, 8; ; CHECK-NEXT: selp.b32 %r21, %r5, %r20, %p1; -; CHECK-NEXT: bfi.b32 %r22, %r21, %r19, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r22; +; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; ; CHECK-NEXT: ret; %cmp = icmp eq <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c @@ -447,7 +447,7 @@ define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-LABEL: test_ne( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<24>; +; CHECK-NEXT: .reg .b32 %r<23>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r3, [test_ne_param_2]; @@ -459,24 +459,24 @@ define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: bfe.u32 %r6, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; ; CHECK-NEXT: setp.ne.u32 %p2, %r7, %r6; -; CHECK-NEXT: bfe.u32 %r8, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r9, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r8, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r9, %r1, 0, 8; ; CHECK-NEXT: setp.ne.u32 %p3, %r9, %r8; -; CHECK-NEXT: bfe.u32 %r10, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r11, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; ; CHECK-NEXT: setp.ne.u32 %p4, %r11, %r10; -; CHECK-NEXT: bfe.u32 %r12, %r3, 0, 8; +; CHECK-NEXT: bfe.u32 %r12, %r3, 8, 8; ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; -; CHECK-NEXT: bfe.u32 %r14, %r3, 8, 8; +; CHECK-NEXT: bfe.u32 %r14, %r3, 0, 8; ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: bfi.b32 %r16, %r15, %r13, 8, 8; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120; ; CHECK-NEXT: bfe.u32 %r17, %r3, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; -; CHECK-NEXT: bfi.b32 %r19, %r18, %r16, 16, 8; +; CHECK-NEXT: prmt.b32 %r19, %r16, %r18, 13328; ; CHECK-NEXT: bfe.u32 %r20, %r3, 24, 8; ; CHECK-NEXT: selp.b32 %r21, %r5, %r20, %p1; -; CHECK-NEXT: bfi.b32 %r22, %r21, %r19, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r22; +; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; ; CHECK-NEXT: ret; %cmp = icmp ne <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c @@ -487,39 +487,39 @@ define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-LABEL: test_mul( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<19>; +; CHECK-NEXT: .reg .b32 %r<18>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_mul_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_mul_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; ; CHECK-NEXT: mul.lo.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.u32 %r6, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: mul.lo.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.u32 %r10, %r2, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: bfe.u32 %r11, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; ; CHECK-NEXT: mul.lo.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; +; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; ; CHECK-NEXT: bfe.u32 %r14, %r2, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; ; CHECK-NEXT: bfe.u32 %r15, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; ; CHECK-NEXT: mul.lo.s16 %rs12, %rs11, %rs10; ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r17; +; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %r = mul <4 x i8> %a, %b ret <4 x i8> %r @@ -534,7 +534,7 @@ define <4 x i8> @test_or(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_or_param_1]; ; CHECK-NEXT: ld.param.u32 %r4, [test_or_param_0]; ; CHECK-NEXT: or.b32 %r5, %r4, %r3; -; CHECK-NEXT: st.param.b32 [func_retval0], %r5; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; ; CHECK-NEXT: ret; %r = or <4 x i8> %a, %b ret <4 x i8> %r @@ -544,17 +544,17 @@ define <4 x i8> @test_or_computed(i8 %a) { ; CHECK-LABEL: test_or_computed( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<9>; +; CHECK-NEXT: .reg .b32 %r<8>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_or_computed_param_0]; ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: bfi.b32 %r2, 0, %r1, 8, 8; -; CHECK-NEXT: bfi.b32 %r3, 0, %r2, 16, 8; -; CHECK-NEXT: bfi.b32 %r4, 0, %r3, 24, 8; -; CHECK-NEXT: bfi.b32 %r6, 5, %r4, 8, 8; -; CHECK-NEXT: or.b32 %r8, %r6, %r4; -; CHECK-NEXT: st.param.b32 [func_retval0], %r8; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; +; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 13328; +; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 16912; +; CHECK-NEXT: bfi.b32 %r5, 5, %r4, 8, 8; +; CHECK-NEXT: or.b32 %r7, %r5, %r4; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -570,7 +570,7 @@ define <4 x i8> @test_or_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_or_imm_0_param_0]; ; CHECK-NEXT: or.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %r = or <4 x i8> , %a ret <4 x i8> %r @@ -584,7 +584,7 @@ define <4 x i8> @test_or_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_or_imm_1_param_0]; ; CHECK-NEXT: or.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %r = or <4 x i8> %a, ret <4 x i8> %r @@ -599,7 +599,7 @@ define <4 x i8> @test_xor(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_xor_param_1]; ; CHECK-NEXT: ld.param.u32 %r4, [test_xor_param_0]; ; CHECK-NEXT: xor.b32 %r5, %r4, %r3; -; CHECK-NEXT: st.param.b32 [func_retval0], %r5; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; ; CHECK-NEXT: ret; %r = xor <4 x i8> %a, %b ret <4 x i8> %r @@ -609,17 +609,17 @@ define <4 x i8> @test_xor_computed(i8 %a) { ; CHECK-LABEL: test_xor_computed( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<9>; +; CHECK-NEXT: .reg .b32 %r<8>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_xor_computed_param_0]; ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: bfi.b32 %r2, 0, %r1, 8, 8; -; CHECK-NEXT: bfi.b32 %r3, 0, %r2, 16, 8; -; CHECK-NEXT: bfi.b32 %r4, 0, %r3, 24, 8; -; CHECK-NEXT: bfi.b32 %r6, 5, %r4, 8, 8; -; CHECK-NEXT: xor.b32 %r8, %r6, %r4; -; CHECK-NEXT: st.param.b32 [func_retval0], %r8; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; +; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 13328; +; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 16912; +; CHECK-NEXT: bfi.b32 %r5, 5, %r4, 8, 8; +; CHECK-NEXT: xor.b32 %r7, %r5, %r4; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -635,7 +635,7 @@ define <4 x i8> @test_xor_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_xor_imm_0_param_0]; ; CHECK-NEXT: xor.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %r = xor <4 x i8> , %a ret <4 x i8> %r @@ -649,7 +649,7 @@ define <4 x i8> @test_xor_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_xor_imm_1_param_0]; ; CHECK-NEXT: xor.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %r = xor <4 x i8> %a, ret <4 x i8> %r @@ -664,7 +664,7 @@ define <4 x i8> @test_and(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_and_param_1]; ; CHECK-NEXT: ld.param.u32 %r4, [test_and_param_0]; ; CHECK-NEXT: and.b32 %r5, %r4, %r3; -; CHECK-NEXT: st.param.b32 [func_retval0], %r5; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; ; CHECK-NEXT: ret; %r = and <4 x i8> %a, %b ret <4 x i8> %r @@ -674,17 +674,17 @@ define <4 x i8> @test_and_computed(i8 %a) { ; CHECK-LABEL: test_and_computed( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<9>; +; CHECK-NEXT: .reg .b32 %r<8>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_and_computed_param_0]; ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: bfi.b32 %r2, 0, %r1, 8, 8; -; CHECK-NEXT: bfi.b32 %r3, 0, %r2, 16, 8; -; CHECK-NEXT: bfi.b32 %r4, 0, %r3, 24, 8; -; CHECK-NEXT: bfi.b32 %r6, 5, %r4, 8, 8; -; CHECK-NEXT: and.b32 %r8, %r6, %r4; -; CHECK-NEXT: st.param.b32 [func_retval0], %r8; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; +; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 13328; +; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 16912; +; CHECK-NEXT: bfi.b32 %r5, 5, %r4, 8, 8; +; CHECK-NEXT: and.b32 %r7, %r5, %r4; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -700,7 +700,7 @@ define <4 x i8> @test_and_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_and_imm_0_param_0]; ; CHECK-NEXT: and.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %r = and <4 x i8> , %a ret <4 x i8> %r @@ -714,7 +714,7 @@ define <4 x i8> @test_and_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_and_imm_1_param_0]; ; CHECK-NEXT: and.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %r = and <4 x i8> %a, ret <4 x i8> %r @@ -828,9 +828,9 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_call_param_0]; ; CHECK-NEXT: { // callseq 0, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0], %r1; +; CHECK-NEXT: st.param.b32 [param0+0], %r1; ; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1], %r2; +; CHECK-NEXT: st.param.b32 [param1+0], %r2; ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; ; CHECK-NEXT: call.uni (retval0), ; CHECK-NEXT: test_callee, @@ -838,9 +838,9 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: param0, ; CHECK-NEXT: param1 ; CHECK-NEXT: ); -; CHECK-NEXT: ld.param.b32 %r3, [retval0]; +; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; ; CHECK-NEXT: } // callseq 0 -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; ; CHECK-NEXT: ret; %r = call <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) ret <4 x i8> %r @@ -856,9 +856,9 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_call_flipped_param_0]; ; CHECK-NEXT: { // callseq 1, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0], %r2; +; CHECK-NEXT: st.param.b32 [param0+0], %r2; ; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1], %r1; +; CHECK-NEXT: st.param.b32 [param1+0], %r1; ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; ; CHECK-NEXT: call.uni (retval0), ; CHECK-NEXT: test_callee, @@ -866,9 +866,9 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: param0, ; CHECK-NEXT: param1 ; CHECK-NEXT: ); -; CHECK-NEXT: ld.param.b32 %r3, [retval0]; +; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; ; CHECK-NEXT: } // callseq 1 -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; ; CHECK-NEXT: ret; %r = call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) ret <4 x i8> %r @@ -884,9 +884,9 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_tailcall_flipped_param_0]; ; CHECK-NEXT: { // callseq 2, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0], %r2; +; CHECK-NEXT: st.param.b32 [param0+0], %r2; ; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1], %r1; +; CHECK-NEXT: st.param.b32 [param1+0], %r1; ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; ; CHECK-NEXT: call.uni (retval0), ; CHECK-NEXT: test_callee, @@ -894,9 +894,9 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: param0, ; CHECK-NEXT: param1 ; CHECK-NEXT: ); -; CHECK-NEXT: ld.param.b32 %r3, [retval0]; +; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; ; CHECK-NEXT: } // callseq 2 -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; ; CHECK-NEXT: ret; %r = tail call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) ret <4 x i8> %r @@ -916,7 +916,7 @@ define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 { ; CHECK-NEXT: ld.param.u32 %r2, [test_select_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_select_param_0]; ; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; ; CHECK-NEXT: ret; %r = select i1 %c, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r @@ -926,7 +926,7 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> ; CHECK-LABEL: test_select_cc( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<29>; +; CHECK-NEXT: .reg .b32 %r<28>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r4, [test_select_cc_param_3]; @@ -939,28 +939,28 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> ; CHECK-NEXT: bfe.u32 %r7, %r4, 16, 8; ; CHECK-NEXT: bfe.u32 %r8, %r3, 16, 8; ; CHECK-NEXT: setp.ne.u32 %p2, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r4, 8, 8; -; CHECK-NEXT: bfe.u32 %r10, %r3, 8, 8; +; CHECK-NEXT: bfe.u32 %r9, %r4, 0, 8; +; CHECK-NEXT: bfe.u32 %r10, %r3, 0, 8; ; CHECK-NEXT: setp.ne.u32 %p3, %r10, %r9; -; CHECK-NEXT: bfe.u32 %r11, %r4, 0, 8; -; CHECK-NEXT: bfe.u32 %r12, %r3, 0, 8; +; CHECK-NEXT: bfe.u32 %r11, %r4, 8, 8; +; CHECK-NEXT: bfe.u32 %r12, %r3, 8, 8; ; CHECK-NEXT: setp.ne.u32 %p4, %r12, %r11; -; CHECK-NEXT: bfe.u32 %r13, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r14, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r13, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r14, %r1, 8, 8; ; CHECK-NEXT: selp.b32 %r15, %r14, %r13, %p4; -; CHECK-NEXT: bfe.u32 %r16, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r17, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r16, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r17, %r1, 0, 8; ; CHECK-NEXT: selp.b32 %r18, %r17, %r16, %p3; -; CHECK-NEXT: bfi.b32 %r19, %r18, %r15, 8, 8; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 13120; ; CHECK-NEXT: bfe.u32 %r20, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r21, %r1, 16, 8; ; CHECK-NEXT: selp.b32 %r22, %r21, %r20, %p2; -; CHECK-NEXT: bfi.b32 %r23, %r22, %r19, 16, 8; +; CHECK-NEXT: prmt.b32 %r23, %r19, %r22, 13328; ; CHECK-NEXT: bfe.u32 %r24, %r2, 24, 8; ; CHECK-NEXT: bfe.u32 %r25, %r1, 24, 8; ; CHECK-NEXT: selp.b32 %r26, %r25, %r24, %p1; -; CHECK-NEXT: bfi.b32 %r27, %r26, %r23, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r27; +; CHECK-NEXT: prmt.b32 %r27, %r23, %r26, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r27; ; CHECK-NEXT: ret; %cc = icmp ne <4 x i8> %c, %d %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b @@ -994,7 +994,7 @@ define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b, ; CHECK-NEXT: selp.b32 %r20, %r3, %r7, %p3; ; CHECK-NEXT: selp.b32 %r21, %r2, %r6, %p2; ; CHECK-NEXT: selp.b32 %r22, %r1, %r5, %p1; -; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19}; +; CHECK-NEXT: st.param.v4.b32 [func_retval0+0], {%r22, %r21, %r20, %r19}; ; CHECK-NEXT: ret; <4 x i8> %c, <4 x i8> %d) #0 { %cc = icmp ne <4 x i8> %c, %d @@ -1006,7 +1006,7 @@ define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, ; CHECK-LABEL: test_select_cc_i8_i32( ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; -; CHECK-NEXT: .reg .b32 %r<27>; +; CHECK-NEXT: .reg .b32 %r<26>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v4.u32 {%r7, %r8, %r9, %r10}, [test_select_cc_i8_i32_param_3]; @@ -1015,24 +1015,24 @@ define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, ; CHECK-NEXT: ld.param.u32 %r1, [test_select_cc_i8_i32_param_0]; ; CHECK-NEXT: setp.ne.s32 %p1, %r6, %r10; ; CHECK-NEXT: setp.ne.s32 %p2, %r5, %r9; -; CHECK-NEXT: setp.ne.s32 %p3, %r4, %r8; -; CHECK-NEXT: setp.ne.s32 %p4, %r3, %r7; -; CHECK-NEXT: bfe.u32 %r11, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r12, %r1, 0, 8; +; CHECK-NEXT: setp.ne.s32 %p3, %r3, %r7; +; CHECK-NEXT: setp.ne.s32 %p4, %r4, %r8; +; CHECK-NEXT: bfe.u32 %r11, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r12, %r1, 8, 8; ; CHECK-NEXT: selp.b32 %r13, %r12, %r11, %p4; -; CHECK-NEXT: bfe.u32 %r14, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r15, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r14, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r15, %r1, 0, 8; ; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p3; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 8, 8; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 13120; ; CHECK-NEXT: bfe.u32 %r18, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r19, %r1, 16, 8; ; CHECK-NEXT: selp.b32 %r20, %r19, %r18, %p2; -; CHECK-NEXT: bfi.b32 %r21, %r20, %r17, 16, 8; +; CHECK-NEXT: prmt.b32 %r21, %r17, %r20, 13328; ; CHECK-NEXT: bfe.u32 %r22, %r2, 24, 8; ; CHECK-NEXT: bfe.u32 %r23, %r1, 24, 8; ; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p1; -; CHECK-NEXT: bfi.b32 %r25, %r24, %r21, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r25; +; CHECK-NEXT: prmt.b32 %r25, %r21, %r24, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; ; CHECK-NEXT: ret; <4 x i32> %c, <4 x i32> %d) #0 { %cc = icmp ne <4 x i32> %c, %d @@ -1044,14 +1044,14 @@ define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 { ; CHECK-LABEL: test_trunc_2xi32( ; CHECK: { -; CHECK-NEXT: .reg .b32 %r<9>; +; CHECK-NEXT: .reg .b32 %r<8>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_trunc_2xi32_param_0]; -; CHECK-NEXT: bfi.b32 %r5, %r2, %r1, 8, 8; -; CHECK-NEXT: bfi.b32 %r6, %r3, %r5, 16, 8; -; CHECK-NEXT: bfi.b32 %r7, %r4, %r6, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r7; +; CHECK-NEXT: prmt.b32 %r5, %r1, %r2, 13120; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13328; +; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; ; CHECK-NEXT: ret; %r = trunc <4 x i32> %a to <4 x i8> ret <4 x i8> %r @@ -1060,20 +1060,20 @@ define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 { define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 { ; CHECK-LABEL: test_trunc_2xi64( ; CHECK: { -; CHECK-NEXT: .reg .b32 %r<9>; +; CHECK-NEXT: .reg .b32 %r<8>; ; CHECK-NEXT: .reg .b64 %rd<5>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v2.u64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16]; ; CHECK-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [test_trunc_2xi64_param_0]; -; CHECK-NEXT: cvt.u32.u64 %r1, %rd1; -; CHECK-NEXT: cvt.u32.u64 %r2, %rd2; -; CHECK-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; +; CHECK-NEXT: cvt.u32.u64 %r1, %rd2; +; CHECK-NEXT: cvt.u32.u64 %r2, %rd1; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 13120; ; CHECK-NEXT: cvt.u32.u64 %r4, %rd3; -; CHECK-NEXT: bfi.b32 %r5, %r4, %r3, 16, 8; +; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 13328; ; CHECK-NEXT: cvt.u32.u64 %r6, %rd4; -; CHECK-NEXT: bfi.b32 %r7, %r6, %r5, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r7; +; CHECK-NEXT: prmt.b32 %r7, %r5, %r6, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; ; CHECK-NEXT: ret; %r = trunc <4 x i64> %a to <4 x i8> ret <4 x i8> %r @@ -1090,7 +1090,7 @@ define <4 x i32> @test_zext_2xi32(<4 x i8> %a) #0 { ; CHECK-NEXT: bfe.u32 %r3, %r1, 16, 8; ; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; ; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; -; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r5, %r4, %r3, %r2}; +; CHECK-NEXT: st.param.v4.b32 [func_retval0+0], {%r5, %r4, %r3, %r2}; ; CHECK-NEXT: ret; %r = zext <4 x i8> %a to <4 x i32> ret <4 x i32> %r @@ -1116,7 +1116,7 @@ define <4 x i64> @test_zext_2xi64(<4 x i8> %a) #0 { ; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; ; CHECK-NEXT: cvt.u64.u32 %rd7, %r5; ; CHECK-NEXT: and.b64 %rd8, %rd7, 255; -; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd8, %rd6}; +; CHECK-NEXT: st.param.v2.b64 [func_retval0+0], {%rd8, %rd6}; ; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd4, %rd2}; ; CHECK-NEXT: ret; %r = zext <4 x i8> %a to <4 x i64> @@ -1130,7 +1130,7 @@ define <4 x i8> @test_bitcast_i32_to_4xi8(i32 %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_bitcast_i32_to_4xi8_param_0]; -; CHECK-NEXT: st.param.b32 [func_retval0], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; ; CHECK-NEXT: ret; %r = bitcast i32 %a to <4 x i8> ret <4 x i8> %r @@ -1145,7 +1145,7 @@ define <4 x i8> @test_bitcast_float_to_4xi8(float %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.f32 %f1, [test_bitcast_float_to_4xi8_param_0]; ; CHECK-NEXT: mov.b32 %r1, %f1; -; CHECK-NEXT: st.param.b32 [func_retval0], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; ; CHECK-NEXT: ret; %r = bitcast float %a to <4 x i8> ret <4 x i8> %r @@ -1158,7 +1158,7 @@ define i32 @test_bitcast_4xi8_to_i32(<4 x i8> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_bitcast_4xi8_to_i32_param_0]; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %r = bitcast <4 x i8> %a to i32 ret i32 %r @@ -1173,7 +1173,7 @@ define float @test_bitcast_4xi8_to_float(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_bitcast_4xi8_to_float_param_0]; ; CHECK-NEXT: mov.b32 %f1, %r2; -; CHECK-NEXT: st.param.f32 [func_retval0], %f1; +; CHECK-NEXT: st.param.f32 [func_retval0+0], %f1; ; CHECK-NEXT: ret; %r = bitcast <4 x i8> %a to float ret float %r @@ -1189,10 +1189,10 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0]; ; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: bfi.b32 %r2, 5, %r1, 8, 8; -; CHECK-NEXT: bfi.b32 %r3, 6, %r2, 16, 8; -; CHECK-NEXT: bfi.b32 %r4, 7, %r3, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r4; +; CHECK-NEXT: prmt.b32 %r2, %r1, 5, 13120; +; CHECK-NEXT: prmt.b32 %r3, %r2, 6, 13328; +; CHECK-NEXT: prmt.b32 %r4, %r3, 7, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> undef, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -1212,7 +1212,7 @@ define <4 x i8> @test_shufflevector(<4 x i8> %a) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_shufflevector_param_0]; ; CHECK-NEXT: // implicit-def: %r3 ; CHECK-NEXT: prmt.b32 %r2, %r1, %r3, 291; -; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; ; CHECK-NEXT: ret; %s = shufflevector <4 x i8> %a, <4 x i8> undef, <4 x i32> ret <4 x i8> %s @@ -1227,7 +1227,7 @@ define <4 x i8> @test_shufflevector_2(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r2, [test_shufflevector_2_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_shufflevector_2_param_0]; ; CHECK-NEXT: prmt.b32 %r3, %r1, %r2, 9527; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; ; CHECK-NEXT: ret; %s = shufflevector <4 x i8> %a, <4 x i8> %b, <4 x i32> ret <4 x i8> %s @@ -1245,7 +1245,7 @@ define <4 x i8> @test_insertelement(<4 x i8> %a, i8 %x) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_insertelement_param_0]; ; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; ; CHECK-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; ; CHECK-NEXT: ret; %i = insertelement <4 x i8> %a, i8 %x, i64 1 ret <4 x i8> %i @@ -1255,7 +1255,7 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-LABEL: test_fptosi_4xhalf_to_4xi8( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<15>; +; CHECK-NEXT: .reg .b32 %r<14>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_fptosi_4xhalf_to_4xi8_param_0]; @@ -1264,19 +1264,19 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-NEXT: cvt.rzi.s16.f16 %rs4, %rs1; ; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3}; ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; -; CHECK-NEXT: cvt.u32.u16 %r6, %rs5; -; CHECK-NEXT: cvt.u32.u16 %r7, %rs6; -; CHECK-NEXT: bfi.b32 %r8, %r7, %r6, 8, 8; +; CHECK-NEXT: cvt.u32.u16 %r6, %rs6; +; CHECK-NEXT: cvt.u32.u16 %r7, %rs5; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; ; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r4; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs9, %rs8; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs10, %rs7; ; CHECK-NEXT: mov.b32 %r9, {%rs10, %rs9}; ; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r9; ; CHECK-NEXT: cvt.u32.u16 %r10, %rs11; -; CHECK-NEXT: bfi.b32 %r11, %r10, %r8, 16, 8; +; CHECK-NEXT: prmt.b32 %r11, %r8, %r10, 13328; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs12; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r13; +; CHECK-NEXT: prmt.b32 %r13, %r11, %r12, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; ; CHECK-NEXT: ret; %r = fptosi <4 x half> %a to <4 x i8> ret <4 x i8> %r @@ -1286,7 +1286,7 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-LABEL: test_fptoui_4xhalf_to_4xi8( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<13>; -; CHECK-NEXT: .reg .b32 %r<15>; +; CHECK-NEXT: .reg .b32 %r<14>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_fptoui_4xhalf_to_4xi8_param_0]; @@ -1295,19 +1295,19 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-NEXT: cvt.rzi.u16.f16 %rs4, %rs1; ; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3}; ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; -; CHECK-NEXT: cvt.u32.u16 %r6, %rs5; -; CHECK-NEXT: cvt.u32.u16 %r7, %rs6; -; CHECK-NEXT: bfi.b32 %r8, %r7, %r6, 8, 8; +; CHECK-NEXT: cvt.u32.u16 %r6, %rs6; +; CHECK-NEXT: cvt.u32.u16 %r7, %rs5; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; ; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r4; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs9, %rs8; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs10, %rs7; ; CHECK-NEXT: mov.b32 %r9, {%rs10, %rs9}; ; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r9; ; CHECK-NEXT: cvt.u32.u16 %r10, %rs11; -; CHECK-NEXT: bfi.b32 %r11, %r10, %r8, 16, 8; +; CHECK-NEXT: prmt.b32 %r11, %r8, %r10, 13328; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs12; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r13; +; CHECK-NEXT: prmt.b32 %r13, %r11, %r12, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; ; CHECK-NEXT: ret; %r = fptoui <4 x half> %a to <4 x i8> ret <4 x i8> %r @@ -1326,33 +1326,33 @@ define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: ld.param.u64 %rd1, [test_srem_v4i8_param_0]; ; CHECK-NEXT: ld.u32 %r1, [%rd1]; ; CHECK-NEXT: ld.u32 %r2, [%rd2]; -; CHECK-NEXT: bfe.s32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.s32 %r3, %r2, 8, 8; ; CHECK-NEXT: cvt.s8.s32 %rs1, %r3; -; CHECK-NEXT: bfe.s32 %r4, %r1, 0, 8; +; CHECK-NEXT: bfe.s32 %r4, %r1, 8, 8; ; CHECK-NEXT: cvt.s8.s32 %rs2, %r4; ; CHECK-NEXT: rem.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.s32 %r6, %r2, 8, 8; +; CHECK-NEXT: bfe.s32 %r6, %r2, 0, 8; ; CHECK-NEXT: cvt.s8.s32 %rs4, %r6; -; CHECK-NEXT: bfe.s32 %r7, %r1, 8, 8; +; CHECK-NEXT: bfe.s32 %r7, %r1, 0, 8; ; CHECK-NEXT: cvt.s8.s32 %rs5, %r7; ; CHECK-NEXT: rem.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: bfi.b32 %r9, %r8, %r5, 8, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.s32 %r10, %r2, 16, 8; ; CHECK-NEXT: cvt.s8.s32 %rs7, %r10; ; CHECK-NEXT: bfe.s32 %r11, %r1, 16, 8; ; CHECK-NEXT: cvt.s8.s32 %rs8, %r11; ; CHECK-NEXT: rem.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r9, 16, 8; +; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; ; CHECK-NEXT: bfe.s32 %r14, %r2, 24, 8; ; CHECK-NEXT: cvt.s8.s32 %rs10, %r14; ; CHECK-NEXT: bfe.s32 %r15, %r1, 24, 8; ; CHECK-NEXT: cvt.s8.s32 %rs11, %r15; ; CHECK-NEXT: rem.s16 %rs12, %rs11, %rs10; ; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r13, 24, 8; +; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; ; CHECK-NEXT: st.u32 [%rd3], %r17; ; CHECK-NEXT: ret; entry: @@ -1373,7 +1373,7 @@ define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) { ; CHECK-LABEL: test_srem_v3i8( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<20>; -; CHECK-NEXT: .reg .b32 %r<16>; +; CHECK-NEXT: .reg .b32 %r<17>; ; CHECK-NEXT: .reg .b64 %rd<4>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: // %entry @@ -1392,25 +1392,25 @@ define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: or.b16 %rs9, %rs8, %rs6; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs9; ; CHECK-NEXT: ld.s8 %rs10, [%rd2+2]; -; CHECK-NEXT: bfe.s32 %r5, %r3, 0, 8; +; CHECK-NEXT: bfe.s32 %r5, %r3, 8, 8; ; CHECK-NEXT: cvt.s8.s32 %rs11, %r5; -; CHECK-NEXT: bfe.s32 %r6, %r1, 0, 8; +; CHECK-NEXT: bfe.s32 %r6, %r1, 8, 8; ; CHECK-NEXT: cvt.s8.s32 %rs12, %r6; ; CHECK-NEXT: rem.s16 %rs13, %rs12, %rs11; ; CHECK-NEXT: cvt.u32.u16 %r7, %rs13; -; CHECK-NEXT: bfe.s32 %r8, %r3, 8, 8; +; CHECK-NEXT: bfe.s32 %r8, %r3, 0, 8; ; CHECK-NEXT: cvt.s8.s32 %rs14, %r8; -; CHECK-NEXT: bfe.s32 %r9, %r1, 8, 8; +; CHECK-NEXT: bfe.s32 %r9, %r1, 0, 8; ; CHECK-NEXT: cvt.s8.s32 %rs15, %r9; ; CHECK-NEXT: rem.s16 %rs16, %rs15, %rs14; ; CHECK-NEXT: cvt.u32.u16 %r10, %rs16; -; CHECK-NEXT: bfi.b32 %r11, %r10, %r7, 8, 8; +; CHECK-NEXT: prmt.b32 %r11, %r10, %r7, 13120; ; CHECK-NEXT: // implicit-def: %r13 -; CHECK-NEXT: bfi.b32 %r12, %r13, %r11, 16, 8; +; CHECK-NEXT: prmt.b32 %r12, %r11, %r13, 13328; ; CHECK-NEXT: // implicit-def: %r15 -; CHECK-NEXT: bfi.b32 %r14, %r15, %r12, 24, 8; +; CHECK-NEXT: prmt.b32 %r14, %r12, %r15, 16912; ; CHECK-NEXT: rem.s16 %rs17, %rs5, %rs10; -; CHECK-NEXT: cvt.u16.u32 %rs18, %r14; +; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs18, tmp}, %r14; } ; CHECK-NEXT: st.u8 [%rd3], %rs18; ; CHECK-NEXT: shr.u16 %rs19, %rs18, 8; ; CHECK-NEXT: st.u8 [%rd3+1], %rs19; @@ -1443,19 +1443,19 @@ define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: bfe.u32 %r5, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r6, %r1, 16, 8; ; CHECK-NEXT: setp.hi.u32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.u32 %r7, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r8, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r8, %r1, 0, 8; ; CHECK-NEXT: setp.hi.u32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r10, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r9, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r10, %r1, 8, 8; ; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.s32 %r11, -1, 0, %p4; ; CHECK-NEXT: selp.s32 %r12, -1, 0, %p3; -; CHECK-NEXT: bfi.b32 %r13, %r12, %r11, 8, 8; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; ; CHECK-NEXT: selp.s32 %r14, -1, 0, %p2; -; CHECK-NEXT: bfi.b32 %r15, %r14, %r13, 16, 8; +; CHECK-NEXT: prmt.b32 %r15, %r13, %r14, 13328; ; CHECK-NEXT: selp.s32 %r16, -1, 0, %p1; -; CHECK-NEXT: bfi.b32 %r17, %r16, %r15, 24, 8; +; CHECK-NEXT: prmt.b32 %r17, %r15, %r16, 16912; ; CHECK-NEXT: st.u32 [%rd3], %r17; ; CHECK-NEXT: ret; entry: From 0d6b6812ecfdefa42bdb798e772165ccaabc1c44 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Tue, 24 Sep 2024 17:09:06 -0700 Subject: [PATCH 2/9] Update sext-setcc.ll --- llvm/test/CodeGen/NVPTX/sext-setcc.ll | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/llvm/test/CodeGen/NVPTX/sext-setcc.ll b/llvm/test/CodeGen/NVPTX/sext-setcc.ll index 0cb0c1ba8c6bd..7c22b0020596a 100644 --- a/llvm/test/CodeGen/NVPTX/sext-setcc.ll +++ b/llvm/test/CodeGen/NVPTX/sext-setcc.ll @@ -33,7 +33,7 @@ define <4 x i8> @sext_setcc_v4i1_to_v4i8(ptr %p) { ; CHECK: { ; CHECK-NEXT: .reg .pred %p<5>; ; CHECK-NEXT: .reg .b16 %rs<9>; -; CHECK-NEXT: .reg .b32 %r<14>; +; CHECK-NEXT: .reg .b32 %r<13>; ; CHECK-NEXT: .reg .b64 %rd<2>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: // %entry @@ -47,22 +47,22 @@ define <4 x i8> @sext_setcc_v4i1_to_v4i8(ptr %p) { ; CHECK-NEXT: cvt.u16.u32 %rs3, %r3; ; CHECK-NEXT: and.b16 %rs4, %rs3, 255; ; CHECK-NEXT: setp.eq.s16 %p2, %rs4, 0; -; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r4; ; CHECK-NEXT: and.b16 %rs6, %rs5, 255; ; CHECK-NEXT: setp.eq.s16 %p3, %rs6, 0; -; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r5, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r5; ; CHECK-NEXT: and.b16 %rs8, %rs7, 255; ; CHECK-NEXT: setp.eq.s16 %p4, %rs8, 0; ; CHECK-NEXT: selp.s32 %r6, -1, 0, %p4; ; CHECK-NEXT: selp.s32 %r7, -1, 0, %p3; -; CHECK-NEXT: bfi.b32 %r8, %r7, %r6, 8, 8; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; ; CHECK-NEXT: selp.s32 %r9, -1, 0, %p2; -; CHECK-NEXT: bfi.b32 %r10, %r9, %r8, 16, 8; +; CHECK-NEXT: prmt.b32 %r10, %r8, %r9, 13328; ; CHECK-NEXT: selp.s32 %r11, -1, 0, %p1; -; CHECK-NEXT: bfi.b32 %r12, %r11, %r10, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0], %r12; +; CHECK-NEXT: prmt.b32 %r12, %r10, %r11, 16912; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; ; CHECK-NEXT: ret; entry: %v = load <4 x i8>, ptr %p, align 4 From 42cc30ca5ca060d1f85395723fc51a5c463fc991 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Fri, 4 Oct 2024 12:51:04 -0700 Subject: [PATCH 3/9] prmt(d, prmt(c, prmt(a,b))) --> prmt(prmt(c,d), prmt(a,b)) --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 7b4c6a8d34048..6638392254f21 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -2334,15 +2334,16 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, DAG.getAnyExtOrTrunc(Op->getOperand(1), DL, MVT::i32), DAG.getConstant(0x3340, DL, MVT::i32), DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); - SDValue PRMT_210 = DAG.getNode( + SDValue PRMT32__ = DAG.getNode( NVPTXISD::PRMT, DL, MVT::v4i8, - {PRMT__10, DAG.getAnyExtOrTrunc(Op->getOperand(2), DL, MVT::i32), - DAG.getConstant(0x3410, DL, MVT::i32), + {DAG.getAnyExtOrTrunc(Op->getOperand(2), DL, MVT::i32), + DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32), + DAG.getConstant(0x4033, DL, MVT::i32), DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); SDValue PRMT3210 = DAG.getNode( NVPTXISD::PRMT, DL, MVT::v4i8, - {PRMT_210, DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32), - DAG.getConstant(0x4210, DL, MVT::i32), + {PRMT__10, PRMT32__, + DAG.getConstant(0x5410, DL, MVT::i32), DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210); } From 506256945f1730500dbba07c0728a3a3d060b134 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Fri, 4 Oct 2024 12:56:38 -0700 Subject: [PATCH 4/9] clang-format --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 6638392254f21..18ec9704a5b02 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -2337,13 +2337,12 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, SDValue PRMT32__ = DAG.getNode( NVPTXISD::PRMT, DL, MVT::v4i8, {DAG.getAnyExtOrTrunc(Op->getOperand(2), DL, MVT::i32), - DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32), + DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32), DAG.getConstant(0x4033, DL, MVT::i32), DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); SDValue PRMT3210 = DAG.getNode( NVPTXISD::PRMT, DL, MVT::v4i8, - {PRMT__10, PRMT32__, - DAG.getConstant(0x5410, DL, MVT::i32), + {PRMT__10, PRMT32__, DAG.getConstant(0x5410, DL, MVT::i32), DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210); } From 6563cbef363384ad9160ed1b2931318843621612 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Fri, 4 Oct 2024 15:09:49 -0700 Subject: [PATCH 5/9] Update tests --- llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 572 ++++++++++--------- llvm/test/CodeGen/NVPTX/sext-setcc.ll | 16 +- 2 files changed, 296 insertions(+), 292 deletions(-) diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll index fdc25cf95d06a..84dde539ce4c4 100644 --- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll @@ -106,33 +106,33 @@ define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_add_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_add_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; ; CHECK-NEXT: add.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.u32 %r6, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: bfe.u32 %r7, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: add.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; -; CHECK-NEXT: bfe.u32 %r10, %r2, 16, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: bfe.u32 %r11, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; ; CHECK-NEXT: add.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; -; CHECK-NEXT: bfe.u32 %r14, %r2, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; -; CHECK-NEXT: bfe.u32 %r15, %r1, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; +; CHECK-NEXT: bfe.u32 %r13, %r2, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs10, %r13; +; CHECK-NEXT: bfe.u32 %r14, %r1, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs11, %r14; ; CHECK-NEXT: add.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; +; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 13120; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %r = add <4 x i8> %a, %b @@ -147,25 +147,25 @@ define <4 x i8> @test_add_imm_0(<4 x i8> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_add_imm_0_param_0]; -; CHECK-NEXT: bfe.u32 %r2, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r2, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; -; CHECK-NEXT: add.s16 %rs2, %rs1, 2; +; CHECK-NEXT: add.s16 %rs2, %rs1, 4; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs2; -; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; -; CHECK-NEXT: add.s16 %rs4, %rs3, 1; +; CHECK-NEXT: add.s16 %rs4, %rs3, 3; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120; -; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 16435; +; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; -; CHECK-NEXT: add.s16 %rs6, %rs5, 3; +; CHECK-NEXT: add.s16 %rs6, %rs5, 2; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r6, %r8, 13328; -; CHECK-NEXT: bfe.u32 %r10, %r1, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: add.s16 %rs8, %rs7, 4; -; CHECK-NEXT: cvt.u32.u16 %r11, %rs8; -; CHECK-NEXT: prmt.b32 %r12, %r9, %r11, 16912; +; CHECK-NEXT: bfe.u32 %r9, %r1, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs7, %r9; +; CHECK-NEXT: add.s16 %rs8, %rs7, 1; +; CHECK-NEXT: cvt.u32.u16 %r10, %rs8; +; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 13120; +; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; ; CHECK-NEXT: ret; %r = add <4 x i8> , %a @@ -180,25 +180,25 @@ define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_add_imm_1_param_0]; -; CHECK-NEXT: bfe.u32 %r2, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r2, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; -; CHECK-NEXT: add.s16 %rs2, %rs1, 2; +; CHECK-NEXT: add.s16 %rs2, %rs1, 4; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs2; -; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; -; CHECK-NEXT: add.s16 %rs4, %rs3, 1; +; CHECK-NEXT: add.s16 %rs4, %rs3, 3; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120; -; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 16435; +; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; -; CHECK-NEXT: add.s16 %rs6, %rs5, 3; +; CHECK-NEXT: add.s16 %rs6, %rs5, 2; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r6, %r8, 13328; -; CHECK-NEXT: bfe.u32 %r10, %r1, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: add.s16 %rs8, %rs7, 4; -; CHECK-NEXT: cvt.u32.u16 %r11, %rs8; -; CHECK-NEXT: prmt.b32 %r12, %r9, %r11, 16912; +; CHECK-NEXT: bfe.u32 %r9, %r1, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs7, %r9; +; CHECK-NEXT: add.s16 %rs8, %rs7, 1; +; CHECK-NEXT: cvt.u32.u16 %r10, %rs8; +; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 13120; +; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; ; CHECK-NEXT: ret; %r = add <4 x i8> %a, @@ -214,33 +214,33 @@ define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_sub_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_sub_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; ; CHECK-NEXT: sub.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.u32 %r6, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: bfe.u32 %r7, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: sub.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; -; CHECK-NEXT: bfe.u32 %r10, %r2, 16, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: bfe.u32 %r11, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; ; CHECK-NEXT: sub.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; -; CHECK-NEXT: bfe.u32 %r14, %r2, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; -; CHECK-NEXT: bfe.u32 %r15, %r1, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; +; CHECK-NEXT: bfe.u32 %r13, %r2, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs10, %r13; +; CHECK-NEXT: bfe.u32 %r14, %r1, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs11, %r14; ; CHECK-NEXT: sub.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; +; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 13120; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %r = sub <4 x i8> %a, %b @@ -256,33 +256,33 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_smax_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_smax_param_0]; -; CHECK-NEXT: bfe.s32 %r3, %r2, 24, 8; -; CHECK-NEXT: bfe.s32 %r4, %r1, 24, 8; +; CHECK-NEXT: bfe.s32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.s32 %r4, %r1, 0, 8; ; CHECK-NEXT: setp.gt.s32 %p1, %r4, %r3; -; CHECK-NEXT: bfe.s32 %r5, %r2, 16, 8; -; CHECK-NEXT: bfe.s32 %r6, %r1, 16, 8; +; CHECK-NEXT: bfe.s32 %r5, %r2, 8, 8; +; CHECK-NEXT: bfe.s32 %r6, %r1, 8, 8; ; CHECK-NEXT: setp.gt.s32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.s32 %r7, %r2, 0, 8; -; CHECK-NEXT: bfe.s32 %r8, %r1, 0, 8; +; CHECK-NEXT: bfe.s32 %r7, %r2, 16, 8; +; CHECK-NEXT: bfe.s32 %r8, %r1, 16, 8; ; CHECK-NEXT: setp.gt.s32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.s32 %r9, %r2, 8, 8; -; CHECK-NEXT: bfe.s32 %r10, %r1, 8, 8; +; CHECK-NEXT: bfe.s32 %r9, %r2, 24, 8; +; CHECK-NEXT: bfe.s32 %r10, %r1, 24, 8; ; CHECK-NEXT: setp.gt.s32 %p4, %r10, %r9; -; CHECK-NEXT: bfe.u32 %r11, %r1, 24, 8; -; CHECK-NEXT: bfe.u32 %r12, %r1, 16, 8; -; CHECK-NEXT: bfe.u32 %r13, %r1, 0, 8; -; CHECK-NEXT: bfe.u32 %r14, %r1, 8, 8; -; CHECK-NEXT: bfe.u32 %r15, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r12, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r13, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r14, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r15, %r2, 24, 8; ; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4; -; CHECK-NEXT: bfe.u32 %r17, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r17, %r2, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3; -; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120; -; CHECK-NEXT: bfe.u32 %r20, %r2, 16, 8; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 16435; +; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8; ; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2; -; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 13328; -; CHECK-NEXT: bfe.u32 %r23, %r2, 24, 8; -; CHECK-NEXT: selp.b32 %r24, %r11, %r23, %p1; -; CHECK-NEXT: prmt.b32 %r25, %r22, %r24, 16912; +; CHECK-NEXT: bfe.u32 %r22, %r2, 0, 8; +; CHECK-NEXT: selp.b32 %r23, %r11, %r22, %p1; +; CHECK-NEXT: prmt.b32 %r24, %r23, %r21, 13120; +; CHECK-NEXT: prmt.b32 %r25, %r24, %r19, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; ; CHECK-NEXT: ret; %cmp = icmp sgt <4 x i8> %a, %b @@ -299,25 +299,25 @@ define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_umax_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_umax_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 24, 8; -; CHECK-NEXT: bfe.u32 %r4, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; ; CHECK-NEXT: setp.hi.u32 %p1, %r4, %r3; -; CHECK-NEXT: bfe.u32 %r5, %r2, 16, 8; -; CHECK-NEXT: bfe.u32 %r6, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r5, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r6, %r1, 8, 8; ; CHECK-NEXT: setp.hi.u32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.u32 %r7, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r8, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r7, %r2, 16, 8; +; CHECK-NEXT: bfe.u32 %r8, %r1, 16, 8; ; CHECK-NEXT: setp.hi.u32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r10, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r9, %r2, 24, 8; +; CHECK-NEXT: bfe.u32 %r10, %r1, 24, 8; ; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; ; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435; ; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: prmt.b32 %r15, %r13, %r14, 13328; -; CHECK-NEXT: selp.b32 %r16, %r4, %r3, %p1; -; CHECK-NEXT: prmt.b32 %r17, %r15, %r16, 16912; +; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %cmp = icmp ugt <4 x i8> %a, %b @@ -334,33 +334,33 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_smin_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_smin_param_0]; -; CHECK-NEXT: bfe.s32 %r3, %r2, 24, 8; -; CHECK-NEXT: bfe.s32 %r4, %r1, 24, 8; +; CHECK-NEXT: bfe.s32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.s32 %r4, %r1, 0, 8; ; CHECK-NEXT: setp.le.s32 %p1, %r4, %r3; -; CHECK-NEXT: bfe.s32 %r5, %r2, 16, 8; -; CHECK-NEXT: bfe.s32 %r6, %r1, 16, 8; +; CHECK-NEXT: bfe.s32 %r5, %r2, 8, 8; +; CHECK-NEXT: bfe.s32 %r6, %r1, 8, 8; ; CHECK-NEXT: setp.le.s32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.s32 %r7, %r2, 0, 8; -; CHECK-NEXT: bfe.s32 %r8, %r1, 0, 8; +; CHECK-NEXT: bfe.s32 %r7, %r2, 16, 8; +; CHECK-NEXT: bfe.s32 %r8, %r1, 16, 8; ; CHECK-NEXT: setp.le.s32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.s32 %r9, %r2, 8, 8; -; CHECK-NEXT: bfe.s32 %r10, %r1, 8, 8; +; CHECK-NEXT: bfe.s32 %r9, %r2, 24, 8; +; CHECK-NEXT: bfe.s32 %r10, %r1, 24, 8; ; CHECK-NEXT: setp.le.s32 %p4, %r10, %r9; -; CHECK-NEXT: bfe.u32 %r11, %r1, 24, 8; -; CHECK-NEXT: bfe.u32 %r12, %r1, 16, 8; -; CHECK-NEXT: bfe.u32 %r13, %r1, 0, 8; -; CHECK-NEXT: bfe.u32 %r14, %r1, 8, 8; -; CHECK-NEXT: bfe.u32 %r15, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r12, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r13, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r14, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r15, %r2, 24, 8; ; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4; -; CHECK-NEXT: bfe.u32 %r17, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r17, %r2, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3; -; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120; -; CHECK-NEXT: bfe.u32 %r20, %r2, 16, 8; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 16435; +; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8; ; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2; -; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 13328; -; CHECK-NEXT: bfe.u32 %r23, %r2, 24, 8; -; CHECK-NEXT: selp.b32 %r24, %r11, %r23, %p1; -; CHECK-NEXT: prmt.b32 %r25, %r22, %r24, 16912; +; CHECK-NEXT: bfe.u32 %r22, %r2, 0, 8; +; CHECK-NEXT: selp.b32 %r23, %r11, %r22, %p1; +; CHECK-NEXT: prmt.b32 %r24, %r23, %r21, 13120; +; CHECK-NEXT: prmt.b32 %r25, %r24, %r19, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; ; CHECK-NEXT: ret; %cmp = icmp sle <4 x i8> %a, %b @@ -377,25 +377,25 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_umin_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_umin_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 24, 8; -; CHECK-NEXT: bfe.u32 %r4, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; ; CHECK-NEXT: setp.ls.u32 %p1, %r4, %r3; -; CHECK-NEXT: bfe.u32 %r5, %r2, 16, 8; -; CHECK-NEXT: bfe.u32 %r6, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r5, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r6, %r1, 8, 8; ; CHECK-NEXT: setp.ls.u32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.u32 %r7, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r8, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r7, %r2, 16, 8; +; CHECK-NEXT: bfe.u32 %r8, %r1, 16, 8; ; CHECK-NEXT: setp.ls.u32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r10, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r9, %r2, 24, 8; +; CHECK-NEXT: bfe.u32 %r10, %r1, 24, 8; ; CHECK-NEXT: setp.ls.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; ; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435; ; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; -; CHECK-NEXT: prmt.b32 %r15, %r13, %r14, 13328; -; CHECK-NEXT: selp.b32 %r16, %r4, %r3, %p1; -; CHECK-NEXT: prmt.b32 %r17, %r15, %r16, 16912; +; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %cmp = icmp ule <4 x i8> %a, %b @@ -413,29 +413,29 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_eq_param_2]; ; CHECK-NEXT: ld.param.u32 %r2, [test_eq_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_eq_param_0]; -; CHECK-NEXT: bfe.u32 %r4, %r2, 24, 8; -; CHECK-NEXT: bfe.u32 %r5, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r4, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; ; CHECK-NEXT: setp.eq.u32 %p1, %r5, %r4; -; CHECK-NEXT: bfe.u32 %r6, %r2, 16, 8; -; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; ; CHECK-NEXT: setp.eq.u32 %p2, %r7, %r6; -; CHECK-NEXT: bfe.u32 %r8, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r9, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r8, %r2, 16, 8; +; CHECK-NEXT: bfe.u32 %r9, %r1, 16, 8; ; CHECK-NEXT: setp.eq.u32 %p3, %r9, %r8; -; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r10, %r2, 24, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 24, 8; ; CHECK-NEXT: setp.eq.u32 %p4, %r11, %r10; -; CHECK-NEXT: bfe.u32 %r12, %r3, 8, 8; +; CHECK-NEXT: bfe.u32 %r12, %r3, 24, 8; ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; -; CHECK-NEXT: bfe.u32 %r14, %r3, 0, 8; +; CHECK-NEXT: bfe.u32 %r14, %r3, 16, 8; ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120; -; CHECK-NEXT: bfe.u32 %r17, %r3, 16, 8; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 16435; +; CHECK-NEXT: bfe.u32 %r17, %r3, 8, 8; ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; -; CHECK-NEXT: prmt.b32 %r19, %r16, %r18, 13328; -; CHECK-NEXT: bfe.u32 %r20, %r3, 24, 8; -; CHECK-NEXT: selp.b32 %r21, %r5, %r20, %p1; -; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 16912; +; CHECK-NEXT: bfe.u32 %r19, %r3, 0, 8; +; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1; +; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 13120; +; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; ; CHECK-NEXT: ret; %cmp = icmp eq <4 x i8> %a, %b @@ -453,29 +453,29 @@ define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_ne_param_2]; ; CHECK-NEXT: ld.param.u32 %r2, [test_ne_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_ne_param_0]; -; CHECK-NEXT: bfe.u32 %r4, %r2, 24, 8; -; CHECK-NEXT: bfe.u32 %r5, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r4, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; ; CHECK-NEXT: setp.ne.u32 %p1, %r5, %r4; -; CHECK-NEXT: bfe.u32 %r6, %r2, 16, 8; -; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; ; CHECK-NEXT: setp.ne.u32 %p2, %r7, %r6; -; CHECK-NEXT: bfe.u32 %r8, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r9, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r8, %r2, 16, 8; +; CHECK-NEXT: bfe.u32 %r9, %r1, 16, 8; ; CHECK-NEXT: setp.ne.u32 %p3, %r9, %r8; -; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r10, %r2, 24, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 24, 8; ; CHECK-NEXT: setp.ne.u32 %p4, %r11, %r10; -; CHECK-NEXT: bfe.u32 %r12, %r3, 8, 8; +; CHECK-NEXT: bfe.u32 %r12, %r3, 24, 8; ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; -; CHECK-NEXT: bfe.u32 %r14, %r3, 0, 8; +; CHECK-NEXT: bfe.u32 %r14, %r3, 16, 8; ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120; -; CHECK-NEXT: bfe.u32 %r17, %r3, 16, 8; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 16435; +; CHECK-NEXT: bfe.u32 %r17, %r3, 8, 8; ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; -; CHECK-NEXT: prmt.b32 %r19, %r16, %r18, 13328; -; CHECK-NEXT: bfe.u32 %r20, %r3, 24, 8; -; CHECK-NEXT: selp.b32 %r21, %r5, %r20, %p1; -; CHECK-NEXT: prmt.b32 %r22, %r19, %r21, 16912; +; CHECK-NEXT: bfe.u32 %r19, %r3, 0, 8; +; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1; +; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 13120; +; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; ; CHECK-NEXT: ret; %cmp = icmp ne <4 x i8> %a, %b @@ -492,33 +492,33 @@ define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_mul_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_mul_param_0]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r3; -; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs2, %r4; ; CHECK-NEXT: mul.lo.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.u32 %r6, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r6, %r2, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs4, %r6; -; CHECK-NEXT: bfe.u32 %r7, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r7, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: mul.lo.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; -; CHECK-NEXT: bfe.u32 %r10, %r2, 16, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; -; CHECK-NEXT: bfe.u32 %r11, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs8, %r11; ; CHECK-NEXT: mul.lo.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; -; CHECK-NEXT: bfe.u32 %r14, %r2, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs10, %r14; -; CHECK-NEXT: bfe.u32 %r15, %r1, 24, 8; -; CHECK-NEXT: cvt.u16.u32 %rs11, %r15; +; CHECK-NEXT: bfe.u32 %r13, %r2, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs10, %r13; +; CHECK-NEXT: bfe.u32 %r14, %r1, 0, 8; +; CHECK-NEXT: cvt.u16.u32 %rs11, %r14; ; CHECK-NEXT: mul.lo.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; +; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 13120; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; ; CHECK-NEXT: ret; %r = mul <4 x i8> %a, %b @@ -544,17 +544,18 @@ define <4 x i8> @test_or_computed(i8 %a) { ; CHECK-LABEL: test_or_computed( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<8>; +; CHECK-NEXT: .reg .b32 %r<9>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_or_computed_param_0]; -; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 13328; -; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 16912; -; CHECK-NEXT: bfi.b32 %r5, 5, %r4, 8, 8; -; CHECK-NEXT: or.b32 %r7, %r5, %r4; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; +; CHECK-NEXT: mov.b32 %r1, 0; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; +; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; +; CHECK-NEXT: or.b32 %r8, %r6, %r5; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -609,17 +610,18 @@ define <4 x i8> @test_xor_computed(i8 %a) { ; CHECK-LABEL: test_xor_computed( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<8>; +; CHECK-NEXT: .reg .b32 %r<9>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_xor_computed_param_0]; -; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 13328; -; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 16912; -; CHECK-NEXT: bfi.b32 %r5, 5, %r4, 8, 8; -; CHECK-NEXT: xor.b32 %r7, %r5, %r4; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; +; CHECK-NEXT: mov.b32 %r1, 0; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; +; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; +; CHECK-NEXT: xor.b32 %r8, %r6, %r5; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -674,17 +676,18 @@ define <4 x i8> @test_and_computed(i8 %a) { ; CHECK-LABEL: test_and_computed( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<8>; +; CHECK-NEXT: .reg .b32 %r<9>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_and_computed_param_0]; -; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; -; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 13328; -; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 16912; -; CHECK-NEXT: bfi.b32 %r5, 5, %r4, 8, 8; -; CHECK-NEXT: and.b32 %r7, %r5, %r4; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; +; CHECK-NEXT: mov.b32 %r1, 0; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; +; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; +; CHECK-NEXT: and.b32 %r8, %r6, %r5; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -933,33 +936,33 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> ; CHECK-NEXT: ld.param.u32 %r3, [test_select_cc_param_2]; ; CHECK-NEXT: ld.param.u32 %r2, [test_select_cc_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_select_cc_param_0]; -; CHECK-NEXT: bfe.u32 %r5, %r4, 24, 8; -; CHECK-NEXT: bfe.u32 %r6, %r3, 24, 8; +; CHECK-NEXT: bfe.u32 %r5, %r4, 0, 8; +; CHECK-NEXT: bfe.u32 %r6, %r3, 0, 8; ; CHECK-NEXT: setp.ne.u32 %p1, %r6, %r5; -; CHECK-NEXT: bfe.u32 %r7, %r4, 16, 8; -; CHECK-NEXT: bfe.u32 %r8, %r3, 16, 8; +; CHECK-NEXT: bfe.u32 %r7, %r4, 8, 8; +; CHECK-NEXT: bfe.u32 %r8, %r3, 8, 8; ; CHECK-NEXT: setp.ne.u32 %p2, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r4, 0, 8; -; CHECK-NEXT: bfe.u32 %r10, %r3, 0, 8; +; CHECK-NEXT: bfe.u32 %r9, %r4, 16, 8; +; CHECK-NEXT: bfe.u32 %r10, %r3, 16, 8; ; CHECK-NEXT: setp.ne.u32 %p3, %r10, %r9; -; CHECK-NEXT: bfe.u32 %r11, %r4, 8, 8; -; CHECK-NEXT: bfe.u32 %r12, %r3, 8, 8; +; CHECK-NEXT: bfe.u32 %r11, %r4, 24, 8; +; CHECK-NEXT: bfe.u32 %r12, %r3, 24, 8; ; CHECK-NEXT: setp.ne.u32 %p4, %r12, %r11; -; CHECK-NEXT: bfe.u32 %r13, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r14, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r13, %r2, 24, 8; +; CHECK-NEXT: bfe.u32 %r14, %r1, 24, 8; ; CHECK-NEXT: selp.b32 %r15, %r14, %r13, %p4; -; CHECK-NEXT: bfe.u32 %r16, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r17, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r16, %r2, 16, 8; +; CHECK-NEXT: bfe.u32 %r17, %r1, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r17, %r16, %p3; -; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 13120; -; CHECK-NEXT: bfe.u32 %r20, %r2, 16, 8; -; CHECK-NEXT: bfe.u32 %r21, %r1, 16, 8; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 16435; +; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r21, %r1, 8, 8; ; CHECK-NEXT: selp.b32 %r22, %r21, %r20, %p2; -; CHECK-NEXT: prmt.b32 %r23, %r19, %r22, 13328; -; CHECK-NEXT: bfe.u32 %r24, %r2, 24, 8; -; CHECK-NEXT: bfe.u32 %r25, %r1, 24, 8; -; CHECK-NEXT: selp.b32 %r26, %r25, %r24, %p1; -; CHECK-NEXT: prmt.b32 %r27, %r23, %r26, 16912; +; CHECK-NEXT: bfe.u32 %r23, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r24, %r1, 0, 8; +; CHECK-NEXT: selp.b32 %r25, %r24, %r23, %p1; +; CHECK-NEXT: prmt.b32 %r26, %r25, %r22, 13120; +; CHECK-NEXT: prmt.b32 %r27, %r26, %r19, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r27; ; CHECK-NEXT: ret; %cc = icmp ne <4 x i8> %c, %d @@ -1013,25 +1016,25 @@ define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, ; CHECK-NEXT: ld.param.v4.u32 {%r3, %r4, %r5, %r6}, [test_select_cc_i8_i32_param_2]; ; CHECK-NEXT: ld.param.u32 %r2, [test_select_cc_i8_i32_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_select_cc_i8_i32_param_0]; -; CHECK-NEXT: setp.ne.s32 %p1, %r6, %r10; -; CHECK-NEXT: setp.ne.s32 %p2, %r5, %r9; -; CHECK-NEXT: setp.ne.s32 %p3, %r3, %r7; -; CHECK-NEXT: setp.ne.s32 %p4, %r4, %r8; -; CHECK-NEXT: bfe.u32 %r11, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r12, %r1, 8, 8; +; CHECK-NEXT: setp.ne.s32 %p1, %r3, %r7; +; CHECK-NEXT: setp.ne.s32 %p2, %r4, %r8; +; CHECK-NEXT: setp.ne.s32 %p3, %r5, %r9; +; CHECK-NEXT: setp.ne.s32 %p4, %r6, %r10; +; CHECK-NEXT: bfe.u32 %r11, %r2, 24, 8; +; CHECK-NEXT: bfe.u32 %r12, %r1, 24, 8; ; CHECK-NEXT: selp.b32 %r13, %r12, %r11, %p4; -; CHECK-NEXT: bfe.u32 %r14, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r15, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r14, %r2, 16, 8; +; CHECK-NEXT: bfe.u32 %r15, %r1, 16, 8; ; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 13120; -; CHECK-NEXT: bfe.u32 %r18, %r2, 16, 8; -; CHECK-NEXT: bfe.u32 %r19, %r1, 16, 8; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 16435; +; CHECK-NEXT: bfe.u32 %r18, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r19, %r1, 8, 8; ; CHECK-NEXT: selp.b32 %r20, %r19, %r18, %p2; -; CHECK-NEXT: prmt.b32 %r21, %r17, %r20, 13328; -; CHECK-NEXT: bfe.u32 %r22, %r2, 24, 8; -; CHECK-NEXT: bfe.u32 %r23, %r1, 24, 8; -; CHECK-NEXT: selp.b32 %r24, %r23, %r22, %p1; -; CHECK-NEXT: prmt.b32 %r25, %r21, %r24, 16912; +; CHECK-NEXT: bfe.u32 %r21, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r22, %r1, 0, 8; +; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p1; +; CHECK-NEXT: prmt.b32 %r24, %r23, %r20, 13120; +; CHECK-NEXT: prmt.b32 %r25, %r24, %r17, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; ; CHECK-NEXT: ret; <4 x i32> %c, <4 x i32> %d) #0 { @@ -1048,9 +1051,9 @@ define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_trunc_2xi32_param_0]; -; CHECK-NEXT: prmt.b32 %r5, %r1, %r2, 13120; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13328; -; CHECK-NEXT: prmt.b32 %r7, %r6, %r4, 16912; +; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 16435; +; CHECK-NEXT: prmt.b32 %r6, %r1, %r2, 13120; +; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; ; CHECK-NEXT: ret; %r = trunc <4 x i32> %a to <4 x i8> @@ -1066,13 +1069,13 @@ define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v2.u64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16]; ; CHECK-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [test_trunc_2xi64_param_0]; -; CHECK-NEXT: cvt.u32.u64 %r1, %rd2; -; CHECK-NEXT: cvt.u32.u64 %r2, %rd1; -; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 13120; -; CHECK-NEXT: cvt.u32.u64 %r4, %rd3; -; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 13328; -; CHECK-NEXT: cvt.u32.u64 %r6, %rd4; -; CHECK-NEXT: prmt.b32 %r7, %r5, %r6, 16912; +; CHECK-NEXT: cvt.u32.u64 %r1, %rd4; +; CHECK-NEXT: cvt.u32.u64 %r2, %rd3; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 16435; +; CHECK-NEXT: cvt.u32.u64 %r4, %rd2; +; CHECK-NEXT: cvt.u32.u64 %r5, %rd1; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 13120; +; CHECK-NEXT: prmt.b32 %r7, %r6, %r3, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; ; CHECK-NEXT: ret; %r = trunc <4 x i64> %a to <4 x i8> @@ -1184,15 +1187,16 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 { ; CHECK-LABEL: test_bitcast_4xi8_to_2xhalf( ; CHECK: { ; CHECK-NEXT: .reg .b16 %rs<2>; -; CHECK-NEXT: .reg .b32 %r<6>; +; CHECK-NEXT: .reg .b32 %r<7>; ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0]; -; CHECK-NEXT: cvt.u32.u16 %r1, %rs1; -; CHECK-NEXT: prmt.b32 %r2, %r1, 5, 13120; -; CHECK-NEXT: prmt.b32 %r3, %r2, 6, 13328; -; CHECK-NEXT: prmt.b32 %r4, %r3, 7, 16912; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; +; CHECK-NEXT: mov.b32 %r1, 6; +; CHECK-NEXT: prmt.b32 %r2, %r1, 7, 16435; +; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; +; CHECK-NEXT: prmt.b32 %r4, %r3, 5, 13120; +; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; +; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> undef, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -1259,23 +1263,23 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_fptosi_4xhalf_to_4xi8_param_0]; -; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r3; +; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs3, %rs2; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs4, %rs1; ; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3}; ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; ; CHECK-NEXT: cvt.u32.u16 %r6, %rs6; ; CHECK-NEXT: cvt.u32.u16 %r7, %rs5; -; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; -; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r4; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435; +; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs9, %rs8; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs10, %rs7; ; CHECK-NEXT: mov.b32 %r9, {%rs10, %rs9}; ; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r9; -; CHECK-NEXT: cvt.u32.u16 %r10, %rs11; -; CHECK-NEXT: prmt.b32 %r11, %r8, %r10, 13328; -; CHECK-NEXT: cvt.u32.u16 %r12, %rs12; -; CHECK-NEXT: prmt.b32 %r13, %r11, %r12, 16912; +; CHECK-NEXT: cvt.u32.u16 %r10, %rs12; +; CHECK-NEXT: cvt.u32.u16 %r11, %rs11; +; CHECK-NEXT: prmt.b32 %r12, %r11, %r10, 13120; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r8, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; ; CHECK-NEXT: ret; %r = fptosi <4 x half> %a to <4 x i8> @@ -1290,23 +1294,23 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v2.u32 {%r3, %r4}, [test_fptoui_4xhalf_to_4xi8_param_0]; -; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r3; +; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r4; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs3, %rs2; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs4, %rs1; ; CHECK-NEXT: mov.b32 %r5, {%rs4, %rs3}; ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; ; CHECK-NEXT: cvt.u32.u16 %r6, %rs6; ; CHECK-NEXT: cvt.u32.u16 %r7, %rs5; -; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; -; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r4; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435; +; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs9, %rs8; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs10, %rs7; ; CHECK-NEXT: mov.b32 %r9, {%rs10, %rs9}; ; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r9; -; CHECK-NEXT: cvt.u32.u16 %r10, %rs11; -; CHECK-NEXT: prmt.b32 %r11, %r8, %r10, 13328; -; CHECK-NEXT: cvt.u32.u16 %r12, %rs12; -; CHECK-NEXT: prmt.b32 %r13, %r11, %r12, 16912; +; CHECK-NEXT: cvt.u32.u16 %r10, %rs12; +; CHECK-NEXT: cvt.u32.u16 %r11, %rs11; +; CHECK-NEXT: prmt.b32 %r12, %r11, %r10, 13120; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r8, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; ; CHECK-NEXT: ret; %r = fptoui <4 x half> %a to <4 x i8> @@ -1326,33 +1330,33 @@ define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: ld.param.u64 %rd1, [test_srem_v4i8_param_0]; ; CHECK-NEXT: ld.u32 %r1, [%rd1]; ; CHECK-NEXT: ld.u32 %r2, [%rd2]; -; CHECK-NEXT: bfe.s32 %r3, %r2, 8, 8; +; CHECK-NEXT: bfe.s32 %r3, %r2, 24, 8; ; CHECK-NEXT: cvt.s8.s32 %rs1, %r3; -; CHECK-NEXT: bfe.s32 %r4, %r1, 8, 8; +; CHECK-NEXT: bfe.s32 %r4, %r1, 24, 8; ; CHECK-NEXT: cvt.s8.s32 %rs2, %r4; ; CHECK-NEXT: rem.s16 %rs3, %rs2, %rs1; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs3; -; CHECK-NEXT: bfe.s32 %r6, %r2, 0, 8; +; CHECK-NEXT: bfe.s32 %r6, %r2, 16, 8; ; CHECK-NEXT: cvt.s8.s32 %rs4, %r6; -; CHECK-NEXT: bfe.s32 %r7, %r1, 0, 8; +; CHECK-NEXT: bfe.s32 %r7, %r1, 16, 8; ; CHECK-NEXT: cvt.s8.s32 %rs5, %r7; ; CHECK-NEXT: rem.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; -; CHECK-NEXT: bfe.s32 %r10, %r2, 16, 8; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: bfe.s32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.s8.s32 %rs7, %r10; -; CHECK-NEXT: bfe.s32 %r11, %r1, 16, 8; +; CHECK-NEXT: bfe.s32 %r11, %r1, 8, 8; ; CHECK-NEXT: cvt.s8.s32 %rs8, %r11; ; CHECK-NEXT: rem.s16 %rs9, %rs8, %rs7; ; CHECK-NEXT: cvt.u32.u16 %r12, %rs9; -; CHECK-NEXT: prmt.b32 %r13, %r9, %r12, 13328; -; CHECK-NEXT: bfe.s32 %r14, %r2, 24, 8; -; CHECK-NEXT: cvt.s8.s32 %rs10, %r14; -; CHECK-NEXT: bfe.s32 %r15, %r1, 24, 8; -; CHECK-NEXT: cvt.s8.s32 %rs11, %r15; +; CHECK-NEXT: bfe.s32 %r13, %r2, 0, 8; +; CHECK-NEXT: cvt.s8.s32 %rs10, %r13; +; CHECK-NEXT: bfe.s32 %r14, %r1, 0, 8; +; CHECK-NEXT: cvt.s8.s32 %rs11, %r14; ; CHECK-NEXT: rem.s16 %rs12, %rs11, %rs10; -; CHECK-NEXT: cvt.u32.u16 %r16, %rs12; -; CHECK-NEXT: prmt.b32 %r17, %r13, %r16, 16912; +; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 13120; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 21520; ; CHECK-NEXT: st.u32 [%rd3], %r17; ; CHECK-NEXT: ret; entry: @@ -1406,11 +1410,11 @@ define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: cvt.u32.u16 %r10, %rs16; ; CHECK-NEXT: prmt.b32 %r11, %r10, %r7, 13120; ; CHECK-NEXT: // implicit-def: %r13 -; CHECK-NEXT: prmt.b32 %r12, %r11, %r13, 13328; -; CHECK-NEXT: // implicit-def: %r15 -; CHECK-NEXT: prmt.b32 %r14, %r12, %r15, 16912; +; CHECK-NEXT: // implicit-def: %r14 +; CHECK-NEXT: prmt.b32 %r12, %r13, %r14, 16435; +; CHECK-NEXT: prmt.b32 %r15, %r11, %r12, 21520; ; CHECK-NEXT: rem.s16 %rs17, %rs5, %rs10; -; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs18, tmp}, %r14; } +; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs18, tmp}, %r15; } ; CHECK-NEXT: st.u8 [%rd3], %rs18; ; CHECK-NEXT: shr.u16 %rs19, %rs18, 8; ; CHECK-NEXT: st.u8 [%rd3+1], %rs19; @@ -1437,25 +1441,25 @@ define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: ld.param.u64 %rd1, [test_sext_v4i1_to_v4i8_param_0]; ; CHECK-NEXT: ld.u32 %r1, [%rd1]; ; CHECK-NEXT: ld.u32 %r2, [%rd2]; -; CHECK-NEXT: bfe.u32 %r3, %r2, 24, 8; -; CHECK-NEXT: bfe.u32 %r4, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r3, %r2, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; ; CHECK-NEXT: setp.hi.u32 %p1, %r4, %r3; -; CHECK-NEXT: bfe.u32 %r5, %r2, 16, 8; -; CHECK-NEXT: bfe.u32 %r6, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r5, %r2, 8, 8; +; CHECK-NEXT: bfe.u32 %r6, %r1, 8, 8; ; CHECK-NEXT: setp.hi.u32 %p2, %r6, %r5; -; CHECK-NEXT: bfe.u32 %r7, %r2, 0, 8; -; CHECK-NEXT: bfe.u32 %r8, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r7, %r2, 16, 8; +; CHECK-NEXT: bfe.u32 %r8, %r1, 16, 8; ; CHECK-NEXT: setp.hi.u32 %p3, %r8, %r7; -; CHECK-NEXT: bfe.u32 %r9, %r2, 8, 8; -; CHECK-NEXT: bfe.u32 %r10, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r9, %r2, 24, 8; +; CHECK-NEXT: bfe.u32 %r10, %r1, 24, 8; ; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.s32 %r11, -1, 0, %p4; ; CHECK-NEXT: selp.s32 %r12, -1, 0, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435; ; CHECK-NEXT: selp.s32 %r14, -1, 0, %p2; -; CHECK-NEXT: prmt.b32 %r15, %r13, %r14, 13328; -; CHECK-NEXT: selp.s32 %r16, -1, 0, %p1; -; CHECK-NEXT: prmt.b32 %r17, %r15, %r16, 16912; +; CHECK-NEXT: selp.s32 %r15, -1, 0, %p1; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 21520; ; CHECK-NEXT: st.u32 [%rd3], %r17; ; CHECK-NEXT: ret; entry: diff --git a/llvm/test/CodeGen/NVPTX/sext-setcc.ll b/llvm/test/CodeGen/NVPTX/sext-setcc.ll index 7c22b0020596a..7ace2c08c7193 100644 --- a/llvm/test/CodeGen/NVPTX/sext-setcc.ll +++ b/llvm/test/CodeGen/NVPTX/sext-setcc.ll @@ -39,29 +39,29 @@ define <4 x i8> @sext_setcc_v4i1_to_v4i8(ptr %p) { ; CHECK-NEXT: // %bb.0: // %entry ; CHECK-NEXT: ld.param.u64 %rd1, [sext_setcc_v4i1_to_v4i8_param_0]; ; CHECK-NEXT: ld.u32 %r1, [%rd1]; -; CHECK-NEXT: bfe.u32 %r2, %r1, 24, 8; +; CHECK-NEXT: bfe.u32 %r2, %r1, 0, 8; ; CHECK-NEXT: cvt.u16.u32 %rs1, %r2; ; CHECK-NEXT: and.b16 %rs2, %rs1, 255; ; CHECK-NEXT: setp.eq.s16 %p1, %rs2, 0; -; CHECK-NEXT: bfe.u32 %r3, %r1, 16, 8; +; CHECK-NEXT: bfe.u32 %r3, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs3, %r3; ; CHECK-NEXT: and.b16 %rs4, %rs3, 255; ; CHECK-NEXT: setp.eq.s16 %p2, %rs4, 0; -; CHECK-NEXT: bfe.u32 %r4, %r1, 0, 8; +; CHECK-NEXT: bfe.u32 %r4, %r1, 16, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r4; ; CHECK-NEXT: and.b16 %rs6, %rs5, 255; ; CHECK-NEXT: setp.eq.s16 %p3, %rs6, 0; -; CHECK-NEXT: bfe.u32 %r5, %r1, 8, 8; +; CHECK-NEXT: bfe.u32 %r5, %r1, 24, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r5; ; CHECK-NEXT: and.b16 %rs8, %rs7, 255; ; CHECK-NEXT: setp.eq.s16 %p4, %rs8, 0; ; CHECK-NEXT: selp.s32 %r6, -1, 0, %p4; ; CHECK-NEXT: selp.s32 %r7, -1, 0, %p3; -; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435; ; CHECK-NEXT: selp.s32 %r9, -1, 0, %p2; -; CHECK-NEXT: prmt.b32 %r10, %r8, %r9, 13328; -; CHECK-NEXT: selp.s32 %r11, -1, 0, %p1; -; CHECK-NEXT: prmt.b32 %r12, %r10, %r11, 16912; +; CHECK-NEXT: selp.s32 %r10, -1, 0, %p1; +; CHECK-NEXT: prmt.b32 %r11, %r10, %r9, 13120; +; CHECK-NEXT: prmt.b32 %r12, %r11, %r8, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; ; CHECK-NEXT: ret; entry: From b31dfc19e9701688f0c617b28e888af7db565832 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Tue, 29 Oct 2024 12:15:47 -0700 Subject: [PATCH 6/9] Fix selection value for PRMT__32 --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 40 ++++++++-------- llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 48 ++++++++++---------- llvm/test/CodeGen/NVPTX/sext-setcc.ll | 2 +- 3 files changed, 44 insertions(+), 46 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 18ec9704a5b02..049c424d49018 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -2315,38 +2315,36 @@ NVPTXTargetLowering::LowerCONCAT_VECTORS(SDValue Op, SelectionDAG &DAG) const { // mov.b32 %r2, 0x40003C00 SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, SelectionDAG &DAG) const { + SDLoc DL(Op); EVT VT = Op->getValueType(0); if (!(Isv2x16VT(VT) || VT == MVT::v4i8)) return Op; - SDLoc DL(Op); - if (!llvm::all_of(Op->ops(), [](SDValue Operand) { return Operand->isUndef() || isa(Operand) || isa(Operand); })) { + if (VT != MVT::v4i8) + return Op; // Lower non-const v4i8 vector as byte-wise constructed i32, which allows us // to optimize calculation of constant parts. - if (VT == MVT::v4i8) { - SDValue PRMT__10 = DAG.getNode( - NVPTXISD::PRMT, DL, MVT::v4i8, - {DAG.getAnyExtOrTrunc(Op->getOperand(0), DL, MVT::i32), - DAG.getAnyExtOrTrunc(Op->getOperand(1), DL, MVT::i32), - DAG.getConstant(0x3340, DL, MVT::i32), - DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); - SDValue PRMT32__ = DAG.getNode( - NVPTXISD::PRMT, DL, MVT::v4i8, - {DAG.getAnyExtOrTrunc(Op->getOperand(2), DL, MVT::i32), - DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32), - DAG.getConstant(0x4033, DL, MVT::i32), - DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); - SDValue PRMT3210 = DAG.getNode( - NVPTXISD::PRMT, DL, MVT::v4i8, - {PRMT__10, PRMT32__, DAG.getConstant(0x5410, DL, MVT::i32), + auto GetPRMT = [&](const SDValue Left, const SDValue Right, bool Cast, + uint64_t SelectionValue) -> SDValue { + SDValue L = Left; + SDValue R = Right; + if (Cast) { + L = DAG.getAnyExtOrTrunc(L, DL, MVT::i32); + R = DAG.getAnyExtOrTrunc(R, DL, MVT::i32); + } + return DAG.getNode( + NVPTXISD::PRMT, DL, MVT::v4i8, {L, R, + DAG.getConstant(SelectionValue, DL, MVT::i32), DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); - return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210); - } - return Op; + }; + auto PRMT__10 = GetPRMT(Op->getOperand(0), Op->getOperand(1), true, 0x3340); + auto PRMT__32 = GetPRMT(Op->getOperand(2), Op->getOperand(3), true, 0x3340); + auto PRMT3210 = GetPRMT(PRMT__10, PRMT__32, false, 0x5410); + return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210); } // Get value or the Nth operand as an APInt(32). Undef values treated as 0. diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll index 84dde539ce4c4..359eb06fae9f7 100644 --- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll @@ -118,7 +118,7 @@ define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: add.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; @@ -155,7 +155,7 @@ define <4 x i8> @test_add_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; ; CHECK-NEXT: add.s16 %rs4, %rs3, 3; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 16435; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120; ; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: add.s16 %rs6, %rs5, 2; @@ -188,7 +188,7 @@ define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: cvt.u16.u32 %rs3, %r4; ; CHECK-NEXT: add.s16 %rs4, %rs3, 3; ; CHECK-NEXT: cvt.u32.u16 %r5, %rs4; -; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 16435; +; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120; ; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: add.s16 %rs6, %rs5, 2; @@ -226,7 +226,7 @@ define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: sub.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; @@ -276,7 +276,7 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4; ; CHECK-NEXT: bfe.u32 %r17, %r2, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3; -; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 16435; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120; ; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8; ; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2; ; CHECK-NEXT: bfe.u32 %r22, %r2, 0, 8; @@ -313,7 +313,7 @@ define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; ; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; ; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; ; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; @@ -354,7 +354,7 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4; ; CHECK-NEXT: bfe.u32 %r17, %r2, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3; -; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 16435; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120; ; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8; ; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2; ; CHECK-NEXT: bfe.u32 %r22, %r2, 0, 8; @@ -391,7 +391,7 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: setp.ls.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4; ; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; ; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2; ; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; @@ -429,7 +429,7 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; ; CHECK-NEXT: bfe.u32 %r14, %r3, 16, 8; ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 16435; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120; ; CHECK-NEXT: bfe.u32 %r17, %r3, 8, 8; ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; ; CHECK-NEXT: bfe.u32 %r19, %r3, 0, 8; @@ -469,7 +469,7 @@ define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4; ; CHECK-NEXT: bfe.u32 %r14, %r3, 16, 8; ; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 16435; +; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120; ; CHECK-NEXT: bfe.u32 %r17, %r3, 8, 8; ; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2; ; CHECK-NEXT: bfe.u32 %r19, %r3, 0, 8; @@ -504,7 +504,7 @@ define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: cvt.u16.u32 %rs5, %r7; ; CHECK-NEXT: mul.lo.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.u16.u32 %rs7, %r10; ; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8; @@ -549,7 +549,7 @@ define <4 x i8> @test_or_computed(i8 %a) { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_or_computed_param_0]; ; CHECK-NEXT: mov.b32 %r1, 0; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; ; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120; ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; @@ -615,7 +615,7 @@ define <4 x i8> @test_xor_computed(i8 %a) { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_xor_computed_param_0]; ; CHECK-NEXT: mov.b32 %r1, 0; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; ; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120; ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; @@ -681,7 +681,7 @@ define <4 x i8> @test_and_computed(i8 %a) { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_and_computed_param_0]; ; CHECK-NEXT: mov.b32 %r1, 0; -; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435; +; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; ; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120; ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; @@ -954,7 +954,7 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> ; CHECK-NEXT: bfe.u32 %r16, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r17, %r1, 16, 8; ; CHECK-NEXT: selp.b32 %r18, %r17, %r16, %p3; -; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 16435; +; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 13120; ; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8; ; CHECK-NEXT: bfe.u32 %r21, %r1, 8, 8; ; CHECK-NEXT: selp.b32 %r22, %r21, %r20, %p2; @@ -1026,7 +1026,7 @@ define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, ; CHECK-NEXT: bfe.u32 %r14, %r2, 16, 8; ; CHECK-NEXT: bfe.u32 %r15, %r1, 16, 8; ; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p3; -; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 16435; +; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 13120; ; CHECK-NEXT: bfe.u32 %r18, %r2, 8, 8; ; CHECK-NEXT: bfe.u32 %r19, %r1, 8, 8; ; CHECK-NEXT: selp.b32 %r20, %r19, %r18, %p2; @@ -1051,7 +1051,7 @@ define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_trunc_2xi32_param_0]; -; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 16435; +; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 13120; ; CHECK-NEXT: prmt.b32 %r6, %r1, %r2, 13120; ; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 21520; ; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; @@ -1071,7 +1071,7 @@ define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 { ; CHECK-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [test_trunc_2xi64_param_0]; ; CHECK-NEXT: cvt.u32.u64 %r1, %rd4; ; CHECK-NEXT: cvt.u32.u64 %r2, %rd3; -; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 16435; +; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 13120; ; CHECK-NEXT: cvt.u32.u64 %r4, %rd2; ; CHECK-NEXT: cvt.u32.u64 %r5, %rd1; ; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 13120; @@ -1192,7 +1192,7 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0]; ; CHECK-NEXT: mov.b32 %r1, 6; -; CHECK-NEXT: prmt.b32 %r2, %r1, 7, 16435; +; CHECK-NEXT: prmt.b32 %r2, %r1, 7, 13120; ; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; ; CHECK-NEXT: prmt.b32 %r4, %r3, 5, 13120; ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; @@ -1270,7 +1270,7 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; ; CHECK-NEXT: cvt.u32.u16 %r6, %rs6; ; CHECK-NEXT: cvt.u32.u16 %r7, %rs5; -; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; ; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs9, %rs8; ; CHECK-NEXT: cvt.rzi.s16.f16 %rs10, %rs7; @@ -1301,7 +1301,7 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5; ; CHECK-NEXT: cvt.u32.u16 %r6, %rs6; ; CHECK-NEXT: cvt.u32.u16 %r7, %rs5; -; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; ; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs9, %rs8; ; CHECK-NEXT: cvt.rzi.u16.f16 %rs10, %rs7; @@ -1342,7 +1342,7 @@ define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: cvt.s8.s32 %rs5, %r7; ; CHECK-NEXT: rem.s16 %rs6, %rs5, %rs4; ; CHECK-NEXT: cvt.u32.u16 %r8, %rs6; -; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435; +; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120; ; CHECK-NEXT: bfe.s32 %r10, %r2, 8, 8; ; CHECK-NEXT: cvt.s8.s32 %rs7, %r10; ; CHECK-NEXT: bfe.s32 %r11, %r1, 8, 8; @@ -1411,7 +1411,7 @@ define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: prmt.b32 %r11, %r10, %r7, 13120; ; CHECK-NEXT: // implicit-def: %r13 ; CHECK-NEXT: // implicit-def: %r14 -; CHECK-NEXT: prmt.b32 %r12, %r13, %r14, 16435; +; CHECK-NEXT: prmt.b32 %r12, %r13, %r14, 13120; ; CHECK-NEXT: prmt.b32 %r15, %r11, %r12, 21520; ; CHECK-NEXT: rem.s16 %rs17, %rs5, %rs10; ; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs18, tmp}, %r15; } @@ -1455,7 +1455,7 @@ define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) { ; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9; ; CHECK-NEXT: selp.s32 %r11, -1, 0, %p4; ; CHECK-NEXT: selp.s32 %r12, -1, 0, %p3; -; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435; +; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120; ; CHECK-NEXT: selp.s32 %r14, -1, 0, %p2; ; CHECK-NEXT: selp.s32 %r15, -1, 0, %p1; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; diff --git a/llvm/test/CodeGen/NVPTX/sext-setcc.ll b/llvm/test/CodeGen/NVPTX/sext-setcc.ll index 7ace2c08c7193..63fb759f31886 100644 --- a/llvm/test/CodeGen/NVPTX/sext-setcc.ll +++ b/llvm/test/CodeGen/NVPTX/sext-setcc.ll @@ -57,7 +57,7 @@ define <4 x i8> @sext_setcc_v4i1_to_v4i8(ptr %p) { ; CHECK-NEXT: setp.eq.s16 %p4, %rs8, 0; ; CHECK-NEXT: selp.s32 %r6, -1, 0, %p4; ; CHECK-NEXT: selp.s32 %r7, -1, 0, %p3; -; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435; +; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120; ; CHECK-NEXT: selp.s32 %r9, -1, 0, %p2; ; CHECK-NEXT: selp.s32 %r10, -1, 0, %p1; ; CHECK-NEXT: prmt.b32 %r11, %r10, %r9, 13120; From 03f7ba8e185a257c6f77a4e81d8c566f27c4e6e9 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Wed, 30 Oct 2024 16:23:17 -0700 Subject: [PATCH 7/9] Cleanup use of DL --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 049c424d49018..4b4ad7c8e2c4c 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -2315,10 +2315,10 @@ NVPTXTargetLowering::LowerCONCAT_VECTORS(SDValue Op, SelectionDAG &DAG) const { // mov.b32 %r2, 0x40003C00 SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, SelectionDAG &DAG) const { - SDLoc DL(Op); EVT VT = Op->getValueType(0); if (!(Isv2x16VT(VT) || VT == MVT::v4i8)) return Op; + SDLoc DL(Op); if (!llvm::all_of(Op->ops(), [](SDValue Operand) { return Operand->isUndef() || isa(Operand) || @@ -2375,8 +2375,8 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, } else { llvm_unreachable("Unsupported type"); } - SDValue Const = DAG.getConstant(Value, SDLoc(Op), MVT::i32); - return DAG.getNode(ISD::BITCAST, SDLoc(Op), Op->getValueType(0), Const); + SDValue Const = DAG.getConstant(Value, DL, MVT::i32); + return DAG.getNode(ISD::BITCAST, DL, Op->getValueType(0), Const); } SDValue NVPTXTargetLowering::LowerEXTRACT_VECTOR_ELT(SDValue Op, From c7134115251790f412f1fc3bb796d9afb69f696e Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Wed, 30 Oct 2024 16:51:21 -0700 Subject: [PATCH 8/9] Eliminate +0 in tests --- llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 118 +++++++++---------- llvm/test/CodeGen/NVPTX/sext-setcc.ll | 2 +- 2 files changed, 60 insertions(+), 60 deletions(-) diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll index 359eb06fae9f7..a16a5b435962d 100644 --- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll +++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll @@ -18,7 +18,7 @@ define <4 x i8> @test_ret_const() #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: mov.b32 %r1, -66911489; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r1; ; CHECK-NEXT: ret; ret <4 x i8> } @@ -31,7 +31,7 @@ define i8 @test_extract_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_0_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 0, 8; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 0 ret i8 %e @@ -45,7 +45,7 @@ define i8 @test_extract_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_1_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 8, 8; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 1 ret i8 %e @@ -59,7 +59,7 @@ define i8 @test_extract_2(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_2_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 16, 8; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 2 ret i8 %e @@ -73,7 +73,7 @@ define i8 @test_extract_3(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_extract_3_param_0]; ; CHECK-NEXT: bfe.u32 %r2, %r1, 24, 8; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i32 3 ret i8 %e @@ -91,7 +91,7 @@ define i8 @test_extract_i(<4 x i8> %a, i64 %idx) #0 { ; CHECK-NEXT: cvt.u32.u64 %r2, %rd1; ; CHECK-NEXT: shl.b32 %r3, %r2, 3; ; CHECK-NEXT: bfe.u32 %r4, %r1, %r3, 8; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r4; +; CHECK-NEXT: st.param.b32 [func_retval0], %r4; ; CHECK-NEXT: ret; %e = extractelement <4 x i8> %a, i64 %idx ret i8 %e @@ -133,7 +133,7 @@ define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 13120; ; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; +; CHECK-NEXT: st.param.b32 [func_retval0], %r17; ; CHECK-NEXT: ret; %r = add <4 x i8> %a, %b ret <4 x i8> %r @@ -166,7 +166,7 @@ define <4 x i8> @test_add_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: cvt.u32.u16 %r10, %rs8; ; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 13120; ; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; +; CHECK-NEXT: st.param.b32 [func_retval0], %r12; ; CHECK-NEXT: ret; %r = add <4 x i8> , %a ret <4 x i8> %r @@ -199,7 +199,7 @@ define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: cvt.u32.u16 %r10, %rs8; ; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 13120; ; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; +; CHECK-NEXT: st.param.b32 [func_retval0], %r12; ; CHECK-NEXT: ret; %r = add <4 x i8> %a, ret <4 x i8> %r @@ -241,7 +241,7 @@ define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 13120; ; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; +; CHECK-NEXT: st.param.b32 [func_retval0], %r17; ; CHECK-NEXT: ret; %r = sub <4 x i8> %a, %b ret <4 x i8> %r @@ -283,7 +283,7 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: selp.b32 %r23, %r11, %r22, %p1; ; CHECK-NEXT: prmt.b32 %r24, %r23, %r21, 13120; ; CHECK-NEXT: prmt.b32 %r25, %r24, %r19, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; +; CHECK-NEXT: st.param.b32 [func_retval0], %r25; ; CHECK-NEXT: ret; %cmp = icmp sgt <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -318,7 +318,7 @@ define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; ; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; +; CHECK-NEXT: st.param.b32 [func_retval0], %r17; ; CHECK-NEXT: ret; %cmp = icmp ugt <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -361,7 +361,7 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: selp.b32 %r23, %r11, %r22, %p1; ; CHECK-NEXT: prmt.b32 %r24, %r23, %r21, 13120; ; CHECK-NEXT: prmt.b32 %r25, %r24, %r19, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; +; CHECK-NEXT: st.param.b32 [func_retval0], %r25; ; CHECK-NEXT: ret; %cmp = icmp sle <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -396,7 +396,7 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120; ; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; +; CHECK-NEXT: st.param.b32 [func_retval0], %r17; ; CHECK-NEXT: ret; %cmp = icmp ule <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b @@ -436,7 +436,7 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1; ; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 13120; ; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; +; CHECK-NEXT: st.param.b32 [func_retval0], %r22; ; CHECK-NEXT: ret; %cmp = icmp eq <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c @@ -476,7 +476,7 @@ define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 { ; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1; ; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 13120; ; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r22; +; CHECK-NEXT: st.param.b32 [func_retval0], %r22; ; CHECK-NEXT: ret; %cmp = icmp ne <4 x i8> %a, %b %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c @@ -519,7 +519,7 @@ define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: cvt.u32.u16 %r15, %rs12; ; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 13120; ; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r17; +; CHECK-NEXT: st.param.b32 [func_retval0], %r17; ; CHECK-NEXT: ret; %r = mul <4 x i8> %a, %b ret <4 x i8> %r @@ -534,7 +534,7 @@ define <4 x i8> @test_or(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_or_param_1]; ; CHECK-NEXT: ld.param.u32 %r4, [test_or_param_0]; ; CHECK-NEXT: or.b32 %r5, %r4, %r3; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %r = or <4 x i8> %a, %b ret <4 x i8> %r @@ -555,7 +555,7 @@ define <4 x i8> @test_or_computed(i8 %a) { ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; ; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; ; CHECK-NEXT: or.b32 %r8, %r6, %r5; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; +; CHECK-NEXT: st.param.b32 [func_retval0], %r8; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -571,7 +571,7 @@ define <4 x i8> @test_or_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_or_imm_0_param_0]; ; CHECK-NEXT: or.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %r = or <4 x i8> , %a ret <4 x i8> %r @@ -585,7 +585,7 @@ define <4 x i8> @test_or_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_or_imm_1_param_0]; ; CHECK-NEXT: or.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %r = or <4 x i8> %a, ret <4 x i8> %r @@ -600,7 +600,7 @@ define <4 x i8> @test_xor(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_xor_param_1]; ; CHECK-NEXT: ld.param.u32 %r4, [test_xor_param_0]; ; CHECK-NEXT: xor.b32 %r5, %r4, %r3; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %r = xor <4 x i8> %a, %b ret <4 x i8> %r @@ -621,7 +621,7 @@ define <4 x i8> @test_xor_computed(i8 %a) { ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; ; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; ; CHECK-NEXT: xor.b32 %r8, %r6, %r5; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; +; CHECK-NEXT: st.param.b32 [func_retval0], %r8; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -637,7 +637,7 @@ define <4 x i8> @test_xor_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_xor_imm_0_param_0]; ; CHECK-NEXT: xor.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %r = xor <4 x i8> , %a ret <4 x i8> %r @@ -651,7 +651,7 @@ define <4 x i8> @test_xor_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_xor_imm_1_param_0]; ; CHECK-NEXT: xor.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %r = xor <4 x i8> %a, ret <4 x i8> %r @@ -666,7 +666,7 @@ define <4 x i8> @test_and(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r3, [test_and_param_1]; ; CHECK-NEXT: ld.param.u32 %r4, [test_and_param_0]; ; CHECK-NEXT: and.b32 %r5, %r4, %r3; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %r = and <4 x i8> %a, %b ret <4 x i8> %r @@ -687,7 +687,7 @@ define <4 x i8> @test_and_computed(i8 %a) { ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; ; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8; ; CHECK-NEXT: and.b32 %r8, %r6, %r5; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r8; +; CHECK-NEXT: st.param.b32 [func_retval0], %r8; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -703,7 +703,7 @@ define <4 x i8> @test_and_imm_0(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_and_imm_0_param_0]; ; CHECK-NEXT: and.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %r = and <4 x i8> , %a ret <4 x i8> %r @@ -717,7 +717,7 @@ define <4 x i8> @test_and_imm_1(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_and_imm_1_param_0]; ; CHECK-NEXT: and.b32 %r2, %r1, 67305985; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %r = and <4 x i8> %a, ret <4 x i8> %r @@ -831,9 +831,9 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_call_param_0]; ; CHECK-NEXT: { // callseq 0, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0+0], %r1; +; CHECK-NEXT: st.param.b32 [param0], %r1; ; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1+0], %r2; +; CHECK-NEXT: st.param.b32 [param1], %r2; ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; ; CHECK-NEXT: call.uni (retval0), ; CHECK-NEXT: test_callee, @@ -841,9 +841,9 @@ define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: param0, ; CHECK-NEXT: param1 ; CHECK-NEXT: ); -; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; +; CHECK-NEXT: ld.param.b32 %r3, [retval0]; ; CHECK-NEXT: } // callseq 0 -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %r = call <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) ret <4 x i8> %r @@ -859,9 +859,9 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_call_flipped_param_0]; ; CHECK-NEXT: { // callseq 1, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0+0], %r2; +; CHECK-NEXT: st.param.b32 [param0], %r2; ; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1+0], %r1; +; CHECK-NEXT: st.param.b32 [param1], %r1; ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; ; CHECK-NEXT: call.uni (retval0), ; CHECK-NEXT: test_callee, @@ -869,9 +869,9 @@ define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: param0, ; CHECK-NEXT: param1 ; CHECK-NEXT: ); -; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; +; CHECK-NEXT: ld.param.b32 %r3, [retval0]; ; CHECK-NEXT: } // callseq 1 -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %r = call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) ret <4 x i8> %r @@ -887,9 +887,9 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_tailcall_flipped_param_0]; ; CHECK-NEXT: { // callseq 2, 0 ; CHECK-NEXT: .param .align 4 .b8 param0[4]; -; CHECK-NEXT: st.param.b32 [param0+0], %r2; +; CHECK-NEXT: st.param.b32 [param0], %r2; ; CHECK-NEXT: .param .align 4 .b8 param1[4]; -; CHECK-NEXT: st.param.b32 [param1+0], %r1; +; CHECK-NEXT: st.param.b32 [param1], %r1; ; CHECK-NEXT: .param .align 4 .b8 retval0[4]; ; CHECK-NEXT: call.uni (retval0), ; CHECK-NEXT: test_callee, @@ -897,9 +897,9 @@ define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: param0, ; CHECK-NEXT: param1 ; CHECK-NEXT: ); -; CHECK-NEXT: ld.param.b32 %r3, [retval0+0]; +; CHECK-NEXT: ld.param.b32 %r3, [retval0]; ; CHECK-NEXT: } // callseq 2 -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %r = tail call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a) ret <4 x i8> %r @@ -919,7 +919,7 @@ define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 { ; CHECK-NEXT: ld.param.u32 %r2, [test_select_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_select_param_0]; ; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %r = select i1 %c, <4 x i8> %a, <4 x i8> %b ret <4 x i8> %r @@ -963,7 +963,7 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> ; CHECK-NEXT: selp.b32 %r25, %r24, %r23, %p1; ; CHECK-NEXT: prmt.b32 %r26, %r25, %r22, 13120; ; CHECK-NEXT: prmt.b32 %r27, %r26, %r19, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r27; +; CHECK-NEXT: st.param.b32 [func_retval0], %r27; ; CHECK-NEXT: ret; %cc = icmp ne <4 x i8> %c, %d %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b @@ -997,7 +997,7 @@ define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b, ; CHECK-NEXT: selp.b32 %r20, %r3, %r7, %p3; ; CHECK-NEXT: selp.b32 %r21, %r2, %r6, %p2; ; CHECK-NEXT: selp.b32 %r22, %r1, %r5, %p1; -; CHECK-NEXT: st.param.v4.b32 [func_retval0+0], {%r22, %r21, %r20, %r19}; +; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19}; ; CHECK-NEXT: ret; <4 x i8> %c, <4 x i8> %d) #0 { %cc = icmp ne <4 x i8> %c, %d @@ -1035,7 +1035,7 @@ define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b, ; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p1; ; CHECK-NEXT: prmt.b32 %r24, %r23, %r20, 13120; ; CHECK-NEXT: prmt.b32 %r25, %r24, %r17, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r25; +; CHECK-NEXT: st.param.b32 [func_retval0], %r25; ; CHECK-NEXT: ret; <4 x i32> %c, <4 x i32> %d) #0 { %cc = icmp ne <4 x i32> %c, %d @@ -1054,7 +1054,7 @@ define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 { ; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 13120; ; CHECK-NEXT: prmt.b32 %r6, %r1, %r2, 13120; ; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; +; CHECK-NEXT: st.param.b32 [func_retval0], %r7; ; CHECK-NEXT: ret; %r = trunc <4 x i32> %a to <4 x i8> ret <4 x i8> %r @@ -1076,7 +1076,7 @@ define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 { ; CHECK-NEXT: cvt.u32.u64 %r5, %rd1; ; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 13120; ; CHECK-NEXT: prmt.b32 %r7, %r6, %r3, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7; +; CHECK-NEXT: st.param.b32 [func_retval0], %r7; ; CHECK-NEXT: ret; %r = trunc <4 x i64> %a to <4 x i8> ret <4 x i8> %r @@ -1093,7 +1093,7 @@ define <4 x i32> @test_zext_2xi32(<4 x i8> %a) #0 { ; CHECK-NEXT: bfe.u32 %r3, %r1, 16, 8; ; CHECK-NEXT: bfe.u32 %r4, %r1, 8, 8; ; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; -; CHECK-NEXT: st.param.v4.b32 [func_retval0+0], {%r5, %r4, %r3, %r2}; +; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r5, %r4, %r3, %r2}; ; CHECK-NEXT: ret; %r = zext <4 x i8> %a to <4 x i32> ret <4 x i32> %r @@ -1119,7 +1119,7 @@ define <4 x i64> @test_zext_2xi64(<4 x i8> %a) #0 { ; CHECK-NEXT: bfe.u32 %r5, %r1, 0, 8; ; CHECK-NEXT: cvt.u64.u32 %rd7, %r5; ; CHECK-NEXT: and.b64 %rd8, %rd7, 255; -; CHECK-NEXT: st.param.v2.b64 [func_retval0+0], {%rd8, %rd6}; +; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd8, %rd6}; ; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd4, %rd2}; ; CHECK-NEXT: ret; %r = zext <4 x i8> %a to <4 x i64> @@ -1133,7 +1133,7 @@ define <4 x i8> @test_bitcast_i32_to_4xi8(i32 %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r1, [test_bitcast_i32_to_4xi8_param_0]; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r1; ; CHECK-NEXT: ret; %r = bitcast i32 %a to <4 x i8> ret <4 x i8> %r @@ -1148,7 +1148,7 @@ define <4 x i8> @test_bitcast_float_to_4xi8(float %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.f32 %f1, [test_bitcast_float_to_4xi8_param_0]; ; CHECK-NEXT: mov.b32 %r1, %f1; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r1; ; CHECK-NEXT: ret; %r = bitcast float %a to <4 x i8> ret <4 x i8> %r @@ -1161,7 +1161,7 @@ define i32 @test_bitcast_4xi8_to_i32(<4 x i8> %a) #0 { ; CHECK-EMPTY: ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_bitcast_4xi8_to_i32_param_0]; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %r = bitcast <4 x i8> %a to i32 ret i32 %r @@ -1176,7 +1176,7 @@ define float @test_bitcast_4xi8_to_float(<4 x i8> %a) #0 { ; CHECK-NEXT: // %bb.0: ; CHECK-NEXT: ld.param.u32 %r2, [test_bitcast_4xi8_to_float_param_0]; ; CHECK-NEXT: mov.b32 %f1, %r2; -; CHECK-NEXT: st.param.f32 [func_retval0+0], %f1; +; CHECK-NEXT: st.param.f32 [func_retval0], %f1; ; CHECK-NEXT: ret; %r = bitcast <4 x i8> %a to float ret float %r @@ -1196,7 +1196,7 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 { ; CHECK-NEXT: cvt.u32.u16 %r3, %rs1; ; CHECK-NEXT: prmt.b32 %r4, %r3, 5, 13120; ; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r5; +; CHECK-NEXT: st.param.b32 [func_retval0], %r5; ; CHECK-NEXT: ret; %ins.0 = insertelement <4 x i8> undef, i8 %a, i32 0 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1 @@ -1216,7 +1216,7 @@ define <4 x i8> @test_shufflevector(<4 x i8> %a) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_shufflevector_param_0]; ; CHECK-NEXT: // implicit-def: %r3 ; CHECK-NEXT: prmt.b32 %r2, %r1, %r3, 291; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r2; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; ; CHECK-NEXT: ret; %s = shufflevector <4 x i8> %a, <4 x i8> undef, <4 x i32> ret <4 x i8> %s @@ -1231,7 +1231,7 @@ define <4 x i8> @test_shufflevector_2(<4 x i8> %a, <4 x i8> %b) #0 { ; CHECK-NEXT: ld.param.u32 %r2, [test_shufflevector_2_param_1]; ; CHECK-NEXT: ld.param.u32 %r1, [test_shufflevector_2_param_0]; ; CHECK-NEXT: prmt.b32 %r3, %r1, %r2, 9527; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %s = shufflevector <4 x i8> %a, <4 x i8> %b, <4 x i32> ret <4 x i8> %s @@ -1249,7 +1249,7 @@ define <4 x i8> @test_insertelement(<4 x i8> %a, i8 %x) #0 { ; CHECK-NEXT: ld.param.u32 %r1, [test_insertelement_param_0]; ; CHECK-NEXT: cvt.u32.u16 %r2, %rs1; ; CHECK-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r3; +; CHECK-NEXT: st.param.b32 [func_retval0], %r3; ; CHECK-NEXT: ret; %i = insertelement <4 x i8> %a, i8 %x, i64 1 ret <4 x i8> %i @@ -1280,7 +1280,7 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-NEXT: cvt.u32.u16 %r11, %rs11; ; CHECK-NEXT: prmt.b32 %r12, %r11, %r10, 13120; ; CHECK-NEXT: prmt.b32 %r13, %r12, %r8, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; +; CHECK-NEXT: st.param.b32 [func_retval0], %r13; ; CHECK-NEXT: ret; %r = fptosi <4 x half> %a to <4 x i8> ret <4 x i8> %r @@ -1311,7 +1311,7 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 { ; CHECK-NEXT: cvt.u32.u16 %r11, %rs11; ; CHECK-NEXT: prmt.b32 %r12, %r11, %r10, 13120; ; CHECK-NEXT: prmt.b32 %r13, %r12, %r8, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r13; +; CHECK-NEXT: st.param.b32 [func_retval0], %r13; ; CHECK-NEXT: ret; %r = fptoui <4 x half> %a to <4 x i8> ret <4 x i8> %r diff --git a/llvm/test/CodeGen/NVPTX/sext-setcc.ll b/llvm/test/CodeGen/NVPTX/sext-setcc.ll index 63fb759f31886..0a9fc724945c5 100644 --- a/llvm/test/CodeGen/NVPTX/sext-setcc.ll +++ b/llvm/test/CodeGen/NVPTX/sext-setcc.ll @@ -62,7 +62,7 @@ define <4 x i8> @sext_setcc_v4i1_to_v4i8(ptr %p) { ; CHECK-NEXT: selp.s32 %r10, -1, 0, %p1; ; CHECK-NEXT: prmt.b32 %r11, %r10, %r9, 13120; ; CHECK-NEXT: prmt.b32 %r12, %r11, %r8, 21520; -; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12; +; CHECK-NEXT: st.param.b32 [func_retval0], %r12; ; CHECK-NEXT: ret; entry: %v = load <4 x i8>, ptr %p, align 4 From a175d762c73d546d315a100096262a386e67d599 Mon Sep 17 00:00:00 2001 From: Justin Fargnoli Date: Wed, 30 Oct 2024 17:26:14 -0700 Subject: [PATCH 9/9] Format --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 4b4ad7c8e2c4c..d32be9cd9bf7c 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -2337,8 +2337,8 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op, R = DAG.getAnyExtOrTrunc(R, DL, MVT::i32); } return DAG.getNode( - NVPTXISD::PRMT, DL, MVT::v4i8, {L, R, - DAG.getConstant(SelectionValue, DL, MVT::i32), + NVPTXISD::PRMT, DL, MVT::v4i8, + {L, R, DAG.getConstant(SelectionValue, DL, MVT::i32), DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)}); }; auto PRMT__10 = GetPRMT(Op->getOperand(0), Op->getOperand(1), true, 0x3340);