[llvm] Reland "[NVPTX] Prefer prmt.b32 over bfi.b32" (PR #114326)
Justin Fargnoli via llvm-commits
llvm-commits at lists.llvm.org
Wed Oct 30 16:18:32 PDT 2024
https://github.com/justinfargnoli created https://github.com/llvm/llvm-project/pull/114326
None
>From a3e21a33e28f627cecc76a52b6a37536376c596b Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Tue, 24 Sep 2024 17:08:14 -0700
Subject: [PATCH 1/6] Initial commit
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 31 +-
llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 382 +++++++++----------
2 files changed, 208 insertions(+), 205 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 57bc5fe0ac361c..1677e29aac87f0 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2332,20 +2332,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 96a4359d0ec43e..fdc25cf95d06af 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -101,38 +101,38 @@ 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: 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
@@ -143,29 +143,29 @@ 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: prmt.b32 %r12, %r9, %r11, 16912;
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12;
; CHECK-NEXT: ret;
%r = add <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a
@@ -176,29 +176,29 @@ 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: prmt.b32 %r12, %r9, %r11, 16912;
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12;
; CHECK-NEXT: ret;
%r = add <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4>
@@ -209,38 +209,38 @@ 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: 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
@@ -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,27 +262,27 @@ 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: 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
@@ -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,19 +305,19 @@ 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: 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
@@ -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,27 +340,27 @@ 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: 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
@@ -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,19 +383,19 @@ 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: 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
@@ -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,23 +419,23 @@ 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: 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
@@ -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,23 +459,23 @@ 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: 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
@@ -487,38 +487,38 @@ 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: 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
@@ -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+0], %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
@@ -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+0], %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
@@ -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+0], %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
@@ -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,27 +939,27 @@ 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: 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
@@ -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,23 +1015,23 @@ 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: 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 {
@@ -1044,13 +1044,13 @@ 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: 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>
@@ -1060,19 +1060,19 @@ 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: 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>
@@ -1189,9 +1189,9 @@ 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: 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
@@ -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,18 +1264,18 @@ 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: 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>
@@ -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,18 +1295,18 @@ 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: 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>
@@ -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 940efe91a900bf6d0e85c5fd751976d6e067e307 Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Tue, 24 Sep 2024 17:09:06 -0700
Subject: [PATCH 2/6] Update sext-setcc.ll
---
llvm/test/CodeGen/NVPTX/sext-setcc.ll | 12 ++++++------
1 file changed, 6 insertions(+), 6 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/sext-setcc.ll b/llvm/test/CodeGen/NVPTX/sext-setcc.ll
index f471d47077cf0d..f4622f462dbce3 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,21 +47,21 @@ 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: prmt.b32 %r12, %r10, %r11, 16912;
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r12;
; CHECK-NEXT: ret;
entry:
>From 56e7755afa3bf8a731340cc60b5188eebb86a21b Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Fri, 4 Oct 2024 12:51:04 -0700
Subject: [PATCH 3/6] 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 1677e29aac87f0..7b162f4a7bf3c9 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2338,15 +2338,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 3f5cd57c328020773b4dfd18dbc25b23fa307495 Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Fri, 4 Oct 2024 12:56:38 -0700
Subject: [PATCH 4/6] 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 7b162f4a7bf3c9..d95f8f214be557 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2341,13 +2341,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 e9086459ec2c41a3f65a511d7ad02a1f18e1d0bd Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Fri, 4 Oct 2024 15:09:49 -0700
Subject: [PATCH 5/6] 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 fdc25cf95d06af..84dde539ce4c47 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> <i8 1, i8 2, i8 3, i8 4>, %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, <i8 1, i8 2, i8 3, i8 4>
@@ -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 f4622f462dbce3..8b7e5235443f05 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 9fe8fdb65004e6cd43cb030ab0834bb80b4ea2fb Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <jfargnoli at nvidia.com>
Date: Tue, 29 Oct 2024 12:15:47 -0700
Subject: [PATCH 6/6] 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 d95f8f214be557..47e8c06656a083 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2319,38 +2319,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<ConstantSDNode>(Operand) ||
isa<ConstantFPSDNode>(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 84dde539ce4c47..359eb06fae9f71 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 8b7e5235443f05..7c8462b457296f 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;
More information about the llvm-commits
mailing list