[libc-commits] [mlir] [libc] [llvm] [mlir][Tensor] Remove unnecessary `#include` from TestTensorTransforms.cpp (PR #80383)

Justin Fargnoli via libc-commits libc-commits at lists.llvm.org
Thu Feb 1 20:07:34 PST 2024


https://github.com/justinfargnoli created https://github.com/llvm/llvm-project/pull/80383

None

>From 6c86d89e81f2be7b00cb0f71ebe8015c40dcd0fe Mon Sep 17 00:00:00 2001
From: justinfargnoli <justinfargnoli at gmail.com>
Date: Mon, 29 Jan 2024 16:15:58 -0800
Subject: [PATCH 1/4] Revert "Disable incorrect peephole optimizations"

This reverts commit ff77058141e8026357ca514ad0d45c6c50921290.
---
 llvm/lib/Target/NVPTX/NVPTXInstrInfo.td      |  13 +
 llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 462 +++++++------------
 2 files changed, 184 insertions(+), 291 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 351d8d5371756..e1cced3275449 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2258,6 +2258,19 @@ def: Pat<(setlt (sext_inreg (trunc Int32Regs:$a), i8), (sext_inreg (trunc Int32R
 def: Pat<(setle (sext_inreg (trunc Int32Regs:$a), i8), (sext_inreg (trunc Int32Regs:$b), i8)),
          (SETP_s32rr Int32Regs:$a, Int32Regs:$b, CmpLE)>;
 
+def: Pat<(setugt (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
+         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpHI)>;
+def: Pat<(setuge (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
+         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpHS)>;
+def: Pat<(setult (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
+         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpLO)>;
+def: Pat<(setule (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
+         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpLS)>;
+def: Pat<(seteq (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
+         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpEQ)>;
+def: Pat<(setne (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
+         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpNE)>;
+
 // i1 compare -> i32
 def : Pat<(i32 (setne Int1Regs:$a, Int1Regs:$b)),
           (SELP_u32ii -1, 0, (XORb1rr Int1Regs:$a, Int1Regs:$b))>;
diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index c03081ce68362..1ec68b4a271ba 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -285,51 +285,33 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 {
 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 .b16       %rs<17>;
-; CHECK-NEXT: .reg .b32       %r<19>;
+; CHECK-NEXT:    .reg .pred %p<5>;
+; CHECK-NEXT:    .reg .b32 %r<19>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT: ld.param.u32    %r1, [test_umax_param_0];
-; CHECK-NEXT: ld.param.u32    %r2, [test_umax_param_1];
-; CHECK-NEXT: bfe.s32         %r3, %r2, 24, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs1, %r3;
-; CHECK-NEXT: and.b16         %rs2, %rs1, 255;
-; CHECK-NEXT: bfe.s32         %r4, %r1, 24, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs3, %r4;
-; CHECK-NEXT: and.b16         %rs4, %rs3, 255;
-; CHECK-NEXT: setp.gt.u16     %p1, %rs4, %rs2;
-; CHECK-NEXT: bfe.s32         %r5, %r2, 16, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs5, %r5;
-; CHECK-NEXT: and.b16         %rs6, %rs5, 255;
-; CHECK-NEXT: bfe.s32         %r6, %r1, 16, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs7, %r6;
-; CHECK-NEXT: and.b16         %rs8, %rs7, 255;
-; CHECK-NEXT: setp.gt.u16     %p2, %rs8, %rs6;
-; CHECK-NEXT: bfe.s32         %r7, %r2, 8, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs9, %r7;
-; CHECK-NEXT: and.b16         %rs10, %rs9, 255;
-; CHECK-NEXT: bfe.s32         %r8, %r1, 8, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs11, %r8;
-; CHECK-NEXT: and.b16         %rs12, %rs11, 255;
-; CHECK-NEXT: setp.gt.u16     %p3, %rs12, %rs10;
-; CHECK-NEXT: bfe.s32         %r9, %r2, 0, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs13, %r9;
-; CHECK-NEXT: and.b16         %rs14, %rs13, 255;
-; CHECK-NEXT: bfe.s32         %r10, %r1, 0, 8;
-; CHECK-NEXT: cvt.u16.u32     %rs15, %r10;
-; CHECK-NEXT: and.b16         %rs16, %rs15, 255;
-; CHECK-NEXT: setp.gt.u16     %p4, %rs16, %rs14;
-; 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: selp.b32        %r14, %r6, %r5, %p2;
-; CHECK-NEXT: bfi.b32         %r15, %r14, %r13, 16, 8;
-; CHECK-NEXT: selp.b32        %r16, %r4, %r3, %p1;
-; CHECK-NEXT: bfi.b32         %r17, %r16, %r15, 24, 8;
-; CHECK-NEXT: st.param.b32    [func_retval0+0], %r17;
-; CHECK-NEXT: ret;
-
+; CHECK-NEXT:    ld.param.u32 %r2, [test_umax_param_1];
+; CHECK-NEXT:    ld.param.u32 %r1, [test_umax_param_0];
+; CHECK-NEXT:    bfe.s32 %r3, %r1, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r4, %r2, 24, 8;
+; CHECK-NEXT:    setp.hi.u32 %p1, %r3, %r4;
+; CHECK-NEXT:    bfe.s32 %r5, %r1, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
+; CHECK-NEXT:    setp.hi.u32 %p2, %r5, %r6;
+; CHECK-NEXT:    bfe.s32 %r7, %r1, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
+; CHECK-NEXT:    setp.hi.u32 %p3, %r7, %r8;
+; CHECK-NEXT:    bfe.s32 %r9, %r1, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
+; CHECK-NEXT:    setp.hi.u32 %p4, %r9, %r10;
+; CHECK-NEXT:    selp.b32 %r11, %r9, %r10, %p4;
+; CHECK-NEXT:    selp.b32 %r12, %r7, %r8, %p3;
+; CHECK-NEXT:    bfi.b32 %r13, %r12, %r11, 8, 8;
+; CHECK-NEXT:    selp.b32 %r14, %r5, %r6, %p2;
+; CHECK-NEXT:    bfi.b32 %r15, %r14, %r13, 16, 8;
+; CHECK-NEXT:    selp.b32 %r16, %r3, %r4, %p1;
+; CHECK-NEXT:    bfi.b32 %r17, %r16, %r15, 24, 8;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r17;
+; CHECK-NEXT:    ret;
   %cmp = icmp ugt <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b
   ret <4 x i8> %r
@@ -373,49 +355,32 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 {
 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 .b16       %rs<17>;
-; CHECK-NEXT: .reg .b32       %r<19>;
+; CHECK-NEXT:    .reg .pred %p<5>;
+; CHECK-NEXT:    .reg .b32 %r<19>;
 ; CHECK-EMPTY:
 ; 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.s32         %r3, %r2, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs1, %r3;
-; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
-; CHECK-NEXT:    bfe.s32         %r4, %r1, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs3, %r4;
-; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
-; CHECK-NEXT:    setp.le.u16     %p1, %rs4, %rs2;
-; CHECK-NEXT:    bfe.s32         %r5, %r2, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs5, %r5;
-; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
-; CHECK-NEXT:    bfe.s32         %r6, %r1, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs7, %r6;
-; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
-; CHECK-NEXT:    setp.le.u16     %p2, %rs8, %rs6;
-; CHECK-NEXT:    bfe.s32         %r7, %r2, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs9, %r7;
-; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
-; CHECK-NEXT:    bfe.s32         %r8, %r1, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs11, %r8;
-; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
-; CHECK-NEXT:    setp.le.u16     %p3, %rs12, %rs10;
-; CHECK-NEXT:    bfe.s32         %r9, %r2, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs13, %r9;
-; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
-; CHECK-NEXT:    bfe.s32         %r10, %r1, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs15, %r10;
-; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
-; CHECK-NEXT:    setp.le.u16     %p4, %rs16, %rs14;
-; 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:    selp.b32        %r14, %r6, %r5, %p2;
-; CHECK-NEXT:    bfi.b32         %r15, %r14, %r13, 16, 8;
-; CHECK-NEXT:    selp.b32        %r16, %r4, %r3, %p1;
-; CHECK-NEXT:    bfi.b32         %r17, %r16, %r15, 24, 8;
-; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r17;
+; CHECK-NEXT:    bfe.s32 %r3, %r1, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r4, %r2, 24, 8;
+; CHECK-NEXT:    setp.ls.u32 %p1, %r3, %r4;
+; CHECK-NEXT:    bfe.s32 %r5, %r1, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
+; CHECK-NEXT:    setp.ls.u32 %p2, %r5, %r6;
+; CHECK-NEXT:    bfe.s32 %r7, %r1, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
+; CHECK-NEXT:    setp.ls.u32 %p3, %r7, %r8;
+; CHECK-NEXT:    bfe.s32 %r9, %r1, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
+; CHECK-NEXT:    setp.ls.u32 %p4, %r9, %r10;
+; CHECK-NEXT:    selp.b32 %r11, %r9, %r10, %p4;
+; CHECK-NEXT:    selp.b32 %r12, %r7, %r8, %p3;
+; CHECK-NEXT:    bfi.b32 %r13, %r12, %r11, 8, 8;
+; CHECK-NEXT:    selp.b32 %r14, %r5, %r6, %p2;
+; CHECK-NEXT:    bfi.b32 %r15, %r14, %r13, 16, 8;
+; CHECK-NEXT:    selp.b32 %r16, %r3, %r4, %p1;
+; CHECK-NEXT:    bfi.b32 %r17, %r16, %r15, 24, 8;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r17;
 ; CHECK-NEXT:    ret;
   %cmp = icmp ule <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b
@@ -425,54 +390,37 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 {
 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 .b16       %rs<17>;
-; CHECK-NEXT: .reg .b32       %r<24>;
+; CHECK-NEXT:    .reg .pred %p<5>;
+; CHECK-NEXT:    .reg .b32 %r<24>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.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.s32         %r4, %r2, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs1, %r4;
-; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
-; CHECK-NEXT:    bfe.s32         %r5, %r1, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs3, %r5;
-; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
-; CHECK-NEXT:    setp.eq.s16     %p1, %rs4, %rs2;
-; CHECK-NEXT:    bfe.s32         %r6, %r2, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs5, %r6;
-; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
-; CHECK-NEXT:    bfe.s32         %r7, %r1, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs7, %r7;
-; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
-; CHECK-NEXT:    setp.eq.s16     %p2, %rs8, %rs6;
-; CHECK-NEXT:    bfe.s32         %r8, %r2, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs9, %r8;
-; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
-; CHECK-NEXT:    bfe.s32         %r9, %r1, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs11, %r9;
-; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
-; CHECK-NEXT:    setp.eq.s16     %p3, %rs12, %rs10;
-; CHECK-NEXT:    bfe.s32         %r10, %r2, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs13, %r10;
-; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
-; CHECK-NEXT:    bfe.s32         %r11, %r1, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs15, %r11;
-; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
-; CHECK-NEXT:    setp.eq.s16     %p4, %rs16, %rs14;
-; CHECK-NEXT:    bfe.s32         %r12, %r3, 0, 8;
-; CHECK-NEXT:    selp.b32        %r13, %r11, %r12, %p4;
-; CHECK-NEXT:    bfe.s32         %r14, %r3, 8, 8;
-; CHECK-NEXT:    selp.b32        %r15, %r9, %r14, %p3;
-; CHECK-NEXT:    bfi.b32         %r16, %r15, %r13, 8, 8;
-; CHECK-NEXT:    bfe.s32         %r17, %r3, 16, 8;
-; CHECK-NEXT:    selp.b32        %r18, %r7, %r17, %p2;
-; CHECK-NEXT:    bfi.b32         %r19, %r18, %r16, 16, 8;
-; CHECK-NEXT:    bfe.s32         %r20, %r3, 24, 8;
-; CHECK-NEXT:    selp.b32        %r21, %r5, %r20, %p1;
-; CHECK-NEXT:    bfi.b32         %r22, %r21, %r19, 24, 8;
-; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r22;
+; CHECK-NEXT:    bfe.s32 %r4, %r2, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r5, %r1, 24, 8;
+; CHECK-NEXT:    setp.eq.u32 %p1, %r5, %r4;
+; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r7, %r1, 16, 8;
+; CHECK-NEXT:    setp.eq.u32 %p2, %r7, %r6;
+; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r9, %r1, 8, 8;
+; CHECK-NEXT:    setp.eq.u32 %p3, %r9, %r8;
+; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r11, %r1, 0, 8;
+; CHECK-NEXT:    setp.eq.u32 %p4, %r11, %r10;
+; CHECK-NEXT:    bfe.s32 %r12, %r3, 0, 8;
+; CHECK-NEXT:    selp.b32 %r13, %r11, %r12, %p4;
+; CHECK-NEXT:    bfe.s32 %r14, %r3, 8, 8;
+; CHECK-NEXT:    selp.b32 %r15, %r9, %r14, %p3;
+; CHECK-NEXT:    bfi.b32 %r16, %r15, %r13, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r17, %r3, 16, 8;
+; CHECK-NEXT:    selp.b32 %r18, %r7, %r17, %p2;
+; CHECK-NEXT:    bfi.b32 %r19, %r18, %r16, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r20, %r3, 24, 8;
+; CHECK-NEXT:    selp.b32 %r21, %r5, %r20, %p1;
+; CHECK-NEXT:    bfi.b32 %r22, %r21, %r19, 24, 8;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r22;
 ; CHECK-NEXT:    ret;
   %cmp = icmp eq <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c
@@ -482,54 +430,37 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 {
 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 .b16       %rs<17>;
-; CHECK-NEXT: .reg .b32       %r<24>;
+; CHECK-NEXT:    .reg .pred %p<5>;
+; CHECK-NEXT:    .reg .b32 %r<24>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.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.s32         %r4, %r2, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs1, %r4;
-; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
-; CHECK-NEXT:    bfe.s32         %r5, %r1, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs3, %r5;
-; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
-; CHECK-NEXT:    setp.ne.s16     %p1, %rs4, %rs2;
-; CHECK-NEXT:    bfe.s32         %r6, %r2, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs5, %r6;
-; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
-; CHECK-NEXT:    bfe.s32         %r7, %r1, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs7, %r7;
-; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
-; CHECK-NEXT:    setp.ne.s16     %p2, %rs8, %rs6;
-; CHECK-NEXT:    bfe.s32         %r8, %r2, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs9, %r8;
-; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
-; CHECK-NEXT:    bfe.s32         %r9, %r1, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs11, %r9;
-; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
-; CHECK-NEXT:    setp.ne.s16     %p3, %rs12, %rs10;
-; CHECK-NEXT:    bfe.s32         %r10, %r2, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs13, %r10;
-; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
-; CHECK-NEXT:    bfe.s32         %r11, %r1, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs15, %r11;
-; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
-; CHECK-NEXT:    setp.ne.s16     %p4, %rs16, %rs14;
-; CHECK-NEXT:    bfe.s32         %r12, %r3, 0, 8;
-; CHECK-NEXT:    selp.b32        %r13, %r11, %r12, %p4;
-; CHECK-NEXT:    bfe.s32         %r14, %r3, 8, 8;
-; CHECK-NEXT:    selp.b32        %r15, %r9, %r14, %p3;
-; CHECK-NEXT:    bfi.b32         %r16, %r15, %r13, 8, 8;
-; CHECK-NEXT:    bfe.s32         %r17, %r3, 16, 8;
-; CHECK-NEXT:    selp.b32        %r18, %r7, %r17, %p2;
-; CHECK-NEXT:    bfi.b32         %r19, %r18, %r16, 16, 8;
-; CHECK-NEXT:    bfe.s32         %r20, %r3, 24, 8;
-; CHECK-NEXT:    selp.b32        %r21, %r5, %r20, %p1;
-; CHECK-NEXT:    bfi.b32         %r22, %r21, %r19, 24, 8;
-; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r22;
+; CHECK-NEXT:    bfe.s32 %r4, %r2, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r5, %r1, 24, 8;
+; CHECK-NEXT:    setp.ne.u32 %p1, %r5, %r4;
+; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r7, %r1, 16, 8;
+; CHECK-NEXT:    setp.ne.u32 %p2, %r7, %r6;
+; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r9, %r1, 8, 8;
+; CHECK-NEXT:    setp.ne.u32 %p3, %r9, %r8;
+; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r11, %r1, 0, 8;
+; CHECK-NEXT:    setp.ne.u32 %p4, %r11, %r10;
+; CHECK-NEXT:    bfe.s32 %r12, %r3, 0, 8;
+; CHECK-NEXT:    selp.b32 %r13, %r11, %r12, %p4;
+; CHECK-NEXT:    bfe.s32 %r14, %r3, 8, 8;
+; CHECK-NEXT:    selp.b32 %r15, %r9, %r14, %p3;
+; CHECK-NEXT:    bfi.b32 %r16, %r15, %r13, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r17, %r3, 16, 8;
+; CHECK-NEXT:    selp.b32 %r18, %r7, %r17, %p2;
+; CHECK-NEXT:    bfi.b32 %r19, %r18, %r16, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r20, %r3, 24, 8;
+; CHECK-NEXT:    selp.b32 %r21, %r5, %r20, %p1;
+; CHECK-NEXT:    bfi.b32 %r22, %r21, %r19, 24, 8;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r22;
 ; CHECK-NEXT:    ret;
   %cmp = icmp ne <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c
@@ -981,59 +912,42 @@ define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 {
 define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> %d) #0 {
 ; CHECK-LABEL: test_select_cc(
 ; CHECK:       {
-; CHECK-NEXT: .reg .pred      %p<5>;
-; CHECK-NEXT: .reg .b16       %rs<17>;
-; CHECK-NEXT: .reg .b32       %r<29>;
+; CHECK-NEXT:    .reg .pred %p<5>;
+; CHECK-NEXT:    .reg .b32 %r<29>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.u32 %r4, [test_select_cc_param_3];
 ; 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.s32         %r5, %r4, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs1, %r5;
-; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
-; CHECK-NEXT:    bfe.s32         %r6, %r3, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs3, %r6;
-; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
-; CHECK-NEXT:    setp.ne.s16     %p1, %rs4, %rs2;
-; CHECK-NEXT:    bfe.s32         %r7, %r4, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs5, %r7;
-; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
-; CHECK-NEXT:    bfe.s32         %r8, %r3, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs7, %r8;
-; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
-; CHECK-NEXT:    setp.ne.s16     %p2, %rs8, %rs6;
-; CHECK-NEXT:    bfe.s32         %r9, %r4, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs9, %r9;
-; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
-; CHECK-NEXT:    bfe.s32         %r10, %r3, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs11, %r10;
-; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
-; CHECK-NEXT:    setp.ne.s16     %p3, %rs12, %rs10;
-; CHECK-NEXT:    bfe.s32         %r11, %r4, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs13, %r11;
-; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
-; CHECK-NEXT:    bfe.s32         %r12, %r3, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs15, %r12;
-; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
-; CHECK-NEXT:    setp.ne.s16     %p4, %rs16, %rs14;
-; CHECK-NEXT:    bfe.s32         %r13, %r2, 0, 8;
-; CHECK-NEXT:    bfe.s32         %r14, %r1, 0, 8;
-; CHECK-NEXT:    selp.b32        %r15, %r14, %r13, %p4;
-; CHECK-NEXT:    bfe.s32         %r16, %r2, 8, 8;
-; CHECK-NEXT:    bfe.s32         %r17, %r1, 8, 8;
-; CHECK-NEXT:    selp.b32        %r18, %r17, %r16, %p3;
-; CHECK-NEXT:    bfi.b32         %r19, %r18, %r15, 8, 8;
-; CHECK-NEXT:    bfe.s32         %r20, %r2, 16, 8;
-; CHECK-NEXT:    bfe.s32         %r21, %r1, 16, 8;
-; CHECK-NEXT:    selp.b32        %r22, %r21, %r20, %p2;
-; CHECK-NEXT:    bfi.b32         %r23, %r22, %r19, 16, 8;
-; CHECK-NEXT:    bfe.s32         %r24, %r2, 24, 8;
-; CHECK-NEXT:    bfe.s32         %r25, %r1, 24, 8;
-; CHECK-NEXT:    selp.b32        %r26, %r25, %r24, %p1;
-; CHECK-NEXT:    bfi.b32         %r27, %r26, %r23, 24, 8;
-; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r27;
+; CHECK-NEXT:    bfe.s32 %r5, %r4, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r6, %r3, 24, 8;
+; CHECK-NEXT:    setp.ne.u32 %p1, %r6, %r5;
+; CHECK-NEXT:    bfe.s32 %r7, %r4, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r8, %r3, 16, 8;
+; CHECK-NEXT:    setp.ne.u32 %p2, %r8, %r7;
+; CHECK-NEXT:    bfe.s32 %r9, %r4, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r10, %r3, 8, 8;
+; CHECK-NEXT:    setp.ne.u32 %p3, %r10, %r9;
+; CHECK-NEXT:    bfe.s32 %r11, %r4, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r12, %r3, 0, 8;
+; CHECK-NEXT:    setp.ne.u32 %p4, %r12, %r11;
+; CHECK-NEXT:    bfe.s32 %r13, %r2, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r14, %r1, 0, 8;
+; CHECK-NEXT:    selp.b32 %r15, %r14, %r13, %p4;
+; CHECK-NEXT:    bfe.s32 %r16, %r2, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r17, %r1, 8, 8;
+; CHECK-NEXT:    selp.b32 %r18, %r17, %r16, %p3;
+; CHECK-NEXT:    bfi.b32 %r19, %r18, %r15, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r20, %r2, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r21, %r1, 16, 8;
+; CHECK-NEXT:    selp.b32 %r22, %r21, %r20, %p2;
+; CHECK-NEXT:    bfi.b32 %r23, %r22, %r19, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r24, %r2, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r25, %r1, 24, 8;
+; CHECK-NEXT:    selp.b32 %r26, %r25, %r24, %p1;
+; CHECK-NEXT:    bfi.b32 %r27, %r26, %r23, 24, 8;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r27;
 ; CHECK-NEXT:    ret;
   %cc = icmp ne <4 x i8> %c, %d
   %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b
@@ -1043,48 +957,31 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8>
 define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b,
 ; CHECK-LABEL: test_select_cc_i32_i8(
 ; CHECK:       {
-; CHECK-NEXT: .reg .pred      %p<5>;
-; CHECK-NEXT: .reg .b16       %rs<17>;
-; CHECK-NEXT: .reg .b32       %r<23>;
+; CHECK-NEXT:    .reg .pred %p<5>;
+; CHECK-NEXT:    .reg .b32 %r<23>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.v4.u32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1];
 ; CHECK-NEXT:    ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0];
 ; CHECK-NEXT:    ld.param.u32 %r10, [test_select_cc_i32_i8_param_3];
 ; CHECK-NEXT:    ld.param.u32 %r9, [test_select_cc_i32_i8_param_2];
-; CHECK-NEXT:    bfe.s32         %r11, %r10, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs1, %r11;
-; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
-; CHECK-NEXT:    bfe.s32         %r12, %r9, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs3, %r12;
-; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
-; CHECK-NEXT:    setp.ne.s16     %p1, %rs4, %rs2;
-; CHECK-NEXT:    bfe.s32         %r13, %r10, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs5, %r13;
-; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
-; CHECK-NEXT:    bfe.s32         %r14, %r9, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs7, %r14;
-; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
-; CHECK-NEXT:    setp.ne.s16     %p2, %rs8, %rs6;
-; CHECK-NEXT:    bfe.s32         %r15, %r10, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs9, %r15;
-; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
-; CHECK-NEXT:    bfe.s32         %r16, %r9, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs11, %r16;
-; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
-; CHECK-NEXT:    setp.ne.s16     %p3, %rs12, %rs10;
-; CHECK-NEXT:    bfe.s32         %r17, %r10, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs13, %r17;
-; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
-; CHECK-NEXT:    bfe.s32         %r18, %r9, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs15, %r18;
-; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
-; CHECK-NEXT:    setp.ne.s16     %p4, %rs16, %rs14;
-; CHECK-NEXT:    selp.b32        %r19, %r4, %r8, %p4;
-; CHECK-NEXT:    selp.b32        %r20, %r3, %r7, %p3;
-; CHECK-NEXT:    selp.b32        %r21, %r2, %r6, %p2;
-; CHECK-NEXT:    selp.b32        %r22, %r1, %r5, %p1;
-; CHECK-NEXT:    st.param.v4.b32         [func_retval0+0], {%r22, %r21, %r20, %r19};
+; CHECK-NEXT:    bfe.s32 %r11, %r10, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r12, %r9, 0, 8;
+; CHECK-NEXT:    setp.ne.u32 %p1, %r12, %r11;
+; CHECK-NEXT:    bfe.s32 %r13, %r10, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r14, %r9, 8, 8;
+; CHECK-NEXT:    setp.ne.u32 %p2, %r14, %r13;
+; CHECK-NEXT:    bfe.s32 %r15, %r10, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r16, %r9, 16, 8;
+; CHECK-NEXT:    setp.ne.u32 %p3, %r16, %r15;
+; CHECK-NEXT:    bfe.s32 %r17, %r10, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r18, %r9, 24, 8;
+; CHECK-NEXT:    setp.ne.u32 %p4, %r18, %r17;
+; CHECK-NEXT:    selp.b32 %r19, %r4, %r8, %p4;
+; CHECK-NEXT:    selp.b32 %r20, %r3, %r7, %p3;
+; CHECK-NEXT:    selp.b32 %r21, %r2, %r6, %p2;
+; CHECK-NEXT:    selp.b32 %r22, %r1, %r5, %p1;
+; CHECK-NEXT:    st.param.v4.b32 [func_retval0+0], {%r22, %r21, %r20, %r19};
 ; CHECK-NEXT:    ret;
                                            <4 x i8> %c, <4 x i8> %d) #0 {
   %cc = icmp ne <4 x i8> %c, %d
@@ -1517,10 +1414,9 @@ entry:
 define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) {
 ; CHECK-LABEL: test_sext_v4i1_to_v4i8(
 ; CHECK:       {
-; CHECK-NEXT: .reg .pred      %p<5>;
-; CHECK-NEXT: .reg .b16       %rs<17>;
-; CHECK-NEXT: .reg .b32       %r<18>;
-; CHECK-NEXT: .reg .b64       %rd<4>;
+; CHECK-NEXT:    .reg .pred %p<5>;
+; CHECK-NEXT:    .reg .b32 %r<18>;
+; CHECK-NEXT:    .reg .b64 %rd<4>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0: // %entry
 ; CHECK-NEXT:    ld.param.u64 %rd3, [test_sext_v4i1_to_v4i8_param_2];
@@ -1528,42 +1424,26 @@ 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.s32         %r3, %r2, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs1, %r3;
-; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
-; CHECK-NEXT:    bfe.s32         %r4, %r1, 24, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs3, %r4;
-; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
-; CHECK-NEXT:    setp.gt.u16     %p1, %rs4, %rs2;
-; CHECK-NEXT:    bfe.s32         %r5, %r2, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs5, %r5;
-; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
-; CHECK-NEXT:    bfe.s32         %r6, %r1, 16, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs7, %r6;
-; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
-; CHECK-NEXT:    setp.gt.u16     %p2, %rs8, %rs6;
-; CHECK-NEXT:    bfe.s32         %r7, %r2, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs9, %r7;
-; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
-; CHECK-NEXT:    bfe.s32         %r8, %r1, 8, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs11, %r8;
-; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
-; CHECK-NEXT:    setp.gt.u16     %p3, %rs12, %rs10;
-; CHECK-NEXT:    bfe.s32         %r9, %r2, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs13, %r9;
-; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
-; CHECK-NEXT:    bfe.s32         %r10, %r1, 0, 8;
-; CHECK-NEXT:    cvt.u16.u32     %rs15, %r10;
-; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
-; CHECK-NEXT:    setp.gt.u16     %p4, %rs16, %rs14;
-; 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:    selp.s32        %r14, -1, 0, %p2;
-; CHECK-NEXT:    bfi.b32         %r15, %r14, %r13, 16, 8;
-; CHECK-NEXT:    selp.s32        %r16, -1, 0, %p1;
-; CHECK-NEXT:    bfi.b32         %r17, %r16, %r15, 24, 8;
-; CHECK-NEXT:    st.u32  [%rd3], %r17;
+; CHECK-NEXT:    bfe.s32 %r3, %r2, 24, 8;
+; CHECK-NEXT:    bfe.s32 %r4, %r1, 24, 8;
+; CHECK-NEXT:    setp.hi.u32 %p1, %r4, %r3;
+; CHECK-NEXT:    bfe.s32 %r5, %r2, 16, 8;
+; CHECK-NEXT:    bfe.s32 %r6, %r1, 16, 8;
+; CHECK-NEXT:    setp.hi.u32 %p2, %r6, %r5;
+; CHECK-NEXT:    bfe.s32 %r7, %r2, 8, 8;
+; CHECK-NEXT:    bfe.s32 %r8, %r1, 8, 8;
+; CHECK-NEXT:    setp.hi.u32 %p3, %r8, %r7;
+; CHECK-NEXT:    bfe.s32 %r9, %r2, 0, 8;
+; CHECK-NEXT:    bfe.s32 %r10, %r1, 0, 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:    selp.s32 %r14, -1, 0, %p2;
+; CHECK-NEXT:    bfi.b32 %r15, %r14, %r13, 16, 8;
+; CHECK-NEXT:    selp.s32 %r16, -1, 0, %p1;
+; CHECK-NEXT:    bfi.b32 %r17, %r16, %r15, 24, 8;
+; CHECK-NEXT:    st.u32 [%rd3], %r17;
 ; CHECK-NEXT:    ret;
 entry:
   %t1 = load <4 x i8>, ptr %a, align 4

>From 4487d30fb1009fbaa7c4be049133aec2f304caea Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <34139864+justinfargnoli at users.noreply.github.com>
Date: Mon, 29 Jan 2024 16:34:39 -0800
Subject: [PATCH 2/4] [NVPTX] Disable incorrect peephole optimizations

---
 llvm/lib/Target/NVPTX/NVPTXInstrInfo.td      |  13 -
 llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 464 ++++++++++++-------
 2 files changed, 292 insertions(+), 185 deletions(-)

diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index e1cced3275449..351d8d5371756 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2258,19 +2258,6 @@ def: Pat<(setlt (sext_inreg (trunc Int32Regs:$a), i8), (sext_inreg (trunc Int32R
 def: Pat<(setle (sext_inreg (trunc Int32Regs:$a), i8), (sext_inreg (trunc Int32Regs:$b), i8)),
          (SETP_s32rr Int32Regs:$a, Int32Regs:$b, CmpLE)>;
 
-def: Pat<(setugt (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
-         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpHI)>;
-def: Pat<(setuge (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
-         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpHS)>;
-def: Pat<(setult (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
-         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpLO)>;
-def: Pat<(setule (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
-         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpLS)>;
-def: Pat<(seteq (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
-         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpEQ)>;
-def: Pat<(setne (i16 (and (trunc Int32Regs:$a), 255)), (i16 (and (trunc Int32Regs:$b), 255))),
-         (SETP_u32rr Int32Regs:$a, Int32Regs:$b, CmpNE)>;
-
 // i1 compare -> i32
 def : Pat<(i32 (setne Int1Regs:$a, Int1Regs:$b)),
           (SELP_u32ii -1, 0, (XORb1rr Int1Regs:$a, Int1Regs:$b))>;
diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index 1ec68b4a271ba..ba247c523963b 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -285,33 +285,51 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 {
 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 .pred      %p<5>;
+; CHECK-NEXT: .reg .b16       %rs<17>;
+; CHECK-NEXT: .reg .b32       %r<19>;
 ; CHECK-EMPTY:
 ; 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.s32 %r3, %r1, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r4, %r2, 24, 8;
-; CHECK-NEXT:    setp.hi.u32 %p1, %r3, %r4;
-; CHECK-NEXT:    bfe.s32 %r5, %r1, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
-; CHECK-NEXT:    setp.hi.u32 %p2, %r5, %r6;
-; CHECK-NEXT:    bfe.s32 %r7, %r1, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
-; CHECK-NEXT:    setp.hi.u32 %p3, %r7, %r8;
-; CHECK-NEXT:    bfe.s32 %r9, %r1, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
-; CHECK-NEXT:    setp.hi.u32 %p4, %r9, %r10;
-; CHECK-NEXT:    selp.b32 %r11, %r9, %r10, %p4;
-; CHECK-NEXT:    selp.b32 %r12, %r7, %r8, %p3;
-; CHECK-NEXT:    bfi.b32 %r13, %r12, %r11, 8, 8;
-; CHECK-NEXT:    selp.b32 %r14, %r5, %r6, %p2;
-; CHECK-NEXT:    bfi.b32 %r15, %r14, %r13, 16, 8;
-; CHECK-NEXT:    selp.b32 %r16, %r3, %r4, %p1;
-; CHECK-NEXT:    bfi.b32 %r17, %r16, %r15, 24, 8;
-; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r17;
-; CHECK-NEXT:    ret;
+; CHECK-NEXT: ld.param.u32    %r1, [test_umax_param_0];
+; CHECK-NEXT: ld.param.u32    %r2, [test_umax_param_1];
+; CHECK-NEXT: bfe.s32         %r3, %r2, 24, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs1, %r3;
+; CHECK-NEXT: and.b16         %rs2, %rs1, 255;
+; CHECK-NEXT: bfe.s32         %r4, %r1, 24, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs3, %r4;
+; CHECK-NEXT: and.b16         %rs4, %rs3, 255;
+; CHECK-NEXT: setp.gt.u16     %p1, %rs4, %rs2;
+; CHECK-NEXT: bfe.s32         %r5, %r2, 16, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs5, %r5;
+; CHECK-NEXT: and.b16         %rs6, %rs5, 255;
+; CHECK-NEXT: bfe.s32         %r6, %r1, 16, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs7, %r6;
+; CHECK-NEXT: and.b16         %rs8, %rs7, 255;
+; CHECK-NEXT: setp.gt.u16     %p2, %rs8, %rs6;
+; CHECK-NEXT: bfe.s32         %r7, %r2, 8, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs9, %r7;
+; CHECK-NEXT: and.b16         %rs10, %rs9, 255;
+; CHECK-NEXT: bfe.s32         %r8, %r1, 8, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs11, %r8;
+; CHECK-NEXT: and.b16         %rs12, %rs11, 255;
+; CHECK-NEXT: setp.gt.u16     %p3, %rs12, %rs10;
+; CHECK-NEXT: bfe.s32         %r9, %r2, 0, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs13, %r9;
+; CHECK-NEXT: and.b16         %rs14, %rs13, 255;
+; CHECK-NEXT: bfe.s32         %r10, %r1, 0, 8;
+; CHECK-NEXT: cvt.u16.u32     %rs15, %r10;
+; CHECK-NEXT: and.b16         %rs16, %rs15, 255;
+; CHECK-NEXT: setp.gt.u16     %p4, %rs16, %rs14;
+; 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: selp.b32        %r14, %r6, %r5, %p2;
+; CHECK-NEXT: bfi.b32         %r15, %r14, %r13, 16, 8;
+; CHECK-NEXT: selp.b32        %r16, %r4, %r3, %p1;
+; CHECK-NEXT: bfi.b32         %r17, %r16, %r15, 24, 8;
+; CHECK-NEXT: st.param.b32    [func_retval0+0], %r17;
+; CHECK-NEXT: ret;
+
   %cmp = icmp ugt <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b
   ret <4 x i8> %r
@@ -355,32 +373,49 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 {
 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 .pred      %p<5>;
+; CHECK-NEXT: .reg .b16       %rs<17>;
+; CHECK-NEXT: .reg .b32       %r<19>;
 ; CHECK-EMPTY:
 ; 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.s32 %r3, %r1, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r4, %r2, 24, 8;
-; CHECK-NEXT:    setp.ls.u32 %p1, %r3, %r4;
-; CHECK-NEXT:    bfe.s32 %r5, %r1, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
-; CHECK-NEXT:    setp.ls.u32 %p2, %r5, %r6;
-; CHECK-NEXT:    bfe.s32 %r7, %r1, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
-; CHECK-NEXT:    setp.ls.u32 %p3, %r7, %r8;
-; CHECK-NEXT:    bfe.s32 %r9, %r1, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
-; CHECK-NEXT:    setp.ls.u32 %p4, %r9, %r10;
-; CHECK-NEXT:    selp.b32 %r11, %r9, %r10, %p4;
-; CHECK-NEXT:    selp.b32 %r12, %r7, %r8, %p3;
-; CHECK-NEXT:    bfi.b32 %r13, %r12, %r11, 8, 8;
-; CHECK-NEXT:    selp.b32 %r14, %r5, %r6, %p2;
-; CHECK-NEXT:    bfi.b32 %r15, %r14, %r13, 16, 8;
-; CHECK-NEXT:    selp.b32 %r16, %r3, %r4, %p1;
-; CHECK-NEXT:    bfi.b32 %r17, %r16, %r15, 24, 8;
-; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r17;
+; CHECK-NEXT:    bfe.s32         %r3, %r2, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs1, %r3;
+; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
+; CHECK-NEXT:    bfe.s32         %r4, %r1, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs3, %r4;
+; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
+; CHECK-NEXT:    setp.le.u16     %p1, %rs4, %rs2;
+; CHECK-NEXT:    bfe.s32         %r5, %r2, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs5, %r5;
+; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
+; CHECK-NEXT:    bfe.s32         %r6, %r1, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs7, %r6;
+; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
+; CHECK-NEXT:    setp.le.u16     %p2, %rs8, %rs6;
+; CHECK-NEXT:    bfe.s32         %r7, %r2, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs9, %r7;
+; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
+; CHECK-NEXT:    bfe.s32         %r8, %r1, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs11, %r8;
+; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
+; CHECK-NEXT:    setp.le.u16     %p3, %rs12, %rs10;
+; CHECK-NEXT:    bfe.s32         %r9, %r2, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs13, %r9;
+; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
+; CHECK-NEXT:    bfe.s32         %r10, %r1, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs15, %r10;
+; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
+; CHECK-NEXT:    setp.le.u16     %p4, %rs16, %rs14;
+; 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:    selp.b32        %r14, %r6, %r5, %p2;
+; CHECK-NEXT:    bfi.b32         %r15, %r14, %r13, 16, 8;
+; CHECK-NEXT:    selp.b32        %r16, %r4, %r3, %p1;
+; CHECK-NEXT:    bfi.b32         %r17, %r16, %r15, 24, 8;
+; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r17;
 ; CHECK-NEXT:    ret;
   %cmp = icmp ule <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b
@@ -390,37 +425,54 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 {
 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 .pred      %p<5>;
+; CHECK-NEXT: .reg .b16       %rs<17>;
+; CHECK-NEXT: .reg .b32       %r<24>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.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.s32 %r4, %r2, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r5, %r1, 24, 8;
-; CHECK-NEXT:    setp.eq.u32 %p1, %r5, %r4;
-; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r7, %r1, 16, 8;
-; CHECK-NEXT:    setp.eq.u32 %p2, %r7, %r6;
-; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r9, %r1, 8, 8;
-; CHECK-NEXT:    setp.eq.u32 %p3, %r9, %r8;
-; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r11, %r1, 0, 8;
-; CHECK-NEXT:    setp.eq.u32 %p4, %r11, %r10;
-; CHECK-NEXT:    bfe.s32 %r12, %r3, 0, 8;
-; CHECK-NEXT:    selp.b32 %r13, %r11, %r12, %p4;
-; CHECK-NEXT:    bfe.s32 %r14, %r3, 8, 8;
-; CHECK-NEXT:    selp.b32 %r15, %r9, %r14, %p3;
-; CHECK-NEXT:    bfi.b32 %r16, %r15, %r13, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r17, %r3, 16, 8;
-; CHECK-NEXT:    selp.b32 %r18, %r7, %r17, %p2;
-; CHECK-NEXT:    bfi.b32 %r19, %r18, %r16, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r20, %r3, 24, 8;
-; CHECK-NEXT:    selp.b32 %r21, %r5, %r20, %p1;
-; CHECK-NEXT:    bfi.b32 %r22, %r21, %r19, 24, 8;
-; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r22;
+; CHECK-NEXT:    bfe.s32         %r4, %r2, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs1, %r4;
+; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
+; CHECK-NEXT:    bfe.s32         %r5, %r1, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs3, %r5;
+; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
+; CHECK-NEXT:    setp.eq.s16     %p1, %rs4, %rs2;
+; CHECK-NEXT:    bfe.s32         %r6, %r2, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs5, %r6;
+; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
+; CHECK-NEXT:    bfe.s32         %r7, %r1, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs7, %r7;
+; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
+; CHECK-NEXT:    setp.eq.s16     %p2, %rs8, %rs6;
+; CHECK-NEXT:    bfe.s32         %r8, %r2, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs9, %r8;
+; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
+; CHECK-NEXT:    bfe.s32         %r9, %r1, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs11, %r9;
+; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
+; CHECK-NEXT:    setp.eq.s16     %p3, %rs12, %rs10;
+; CHECK-NEXT:    bfe.s32         %r10, %r2, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs13, %r10;
+; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
+; CHECK-NEXT:    bfe.s32         %r11, %r1, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs15, %r11;
+; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
+; CHECK-NEXT:    setp.eq.s16     %p4, %rs16, %rs14;
+; CHECK-NEXT:    bfe.s32         %r12, %r3, 0, 8;
+; CHECK-NEXT:    selp.b32        %r13, %r11, %r12, %p4;
+; CHECK-NEXT:    bfe.s32         %r14, %r3, 8, 8;
+; CHECK-NEXT:    selp.b32        %r15, %r9, %r14, %p3;
+; CHECK-NEXT:    bfi.b32         %r16, %r15, %r13, 8, 8;
+; CHECK-NEXT:    bfe.s32         %r17, %r3, 16, 8;
+; CHECK-NEXT:    selp.b32        %r18, %r7, %r17, %p2;
+; CHECK-NEXT:    bfi.b32         %r19, %r18, %r16, 16, 8;
+; CHECK-NEXT:    bfe.s32         %r20, %r3, 24, 8;
+; CHECK-NEXT:    selp.b32        %r21, %r5, %r20, %p1;
+; CHECK-NEXT:    bfi.b32         %r22, %r21, %r19, 24, 8;
+; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r22;
 ; CHECK-NEXT:    ret;
   %cmp = icmp eq <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c
@@ -430,37 +482,54 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 {
 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 .pred      %p<5>;
+; CHECK-NEXT: .reg .b16       %rs<17>;
+; CHECK-NEXT: .reg .b32       %r<24>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.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.s32 %r4, %r2, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r5, %r1, 24, 8;
-; CHECK-NEXT:    setp.ne.u32 %p1, %r5, %r4;
-; CHECK-NEXT:    bfe.s32 %r6, %r2, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r7, %r1, 16, 8;
-; CHECK-NEXT:    setp.ne.u32 %p2, %r7, %r6;
-; CHECK-NEXT:    bfe.s32 %r8, %r2, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r9, %r1, 8, 8;
-; CHECK-NEXT:    setp.ne.u32 %p3, %r9, %r8;
-; CHECK-NEXT:    bfe.s32 %r10, %r2, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r11, %r1, 0, 8;
-; CHECK-NEXT:    setp.ne.u32 %p4, %r11, %r10;
-; CHECK-NEXT:    bfe.s32 %r12, %r3, 0, 8;
-; CHECK-NEXT:    selp.b32 %r13, %r11, %r12, %p4;
-; CHECK-NEXT:    bfe.s32 %r14, %r3, 8, 8;
-; CHECK-NEXT:    selp.b32 %r15, %r9, %r14, %p3;
-; CHECK-NEXT:    bfi.b32 %r16, %r15, %r13, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r17, %r3, 16, 8;
-; CHECK-NEXT:    selp.b32 %r18, %r7, %r17, %p2;
-; CHECK-NEXT:    bfi.b32 %r19, %r18, %r16, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r20, %r3, 24, 8;
-; CHECK-NEXT:    selp.b32 %r21, %r5, %r20, %p1;
-; CHECK-NEXT:    bfi.b32 %r22, %r21, %r19, 24, 8;
-; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r22;
+; CHECK-NEXT:    bfe.s32         %r4, %r2, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs1, %r4;
+; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
+; CHECK-NEXT:    bfe.s32         %r5, %r1, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs3, %r5;
+; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
+; CHECK-NEXT:    setp.ne.s16     %p1, %rs4, %rs2;
+; CHECK-NEXT:    bfe.s32         %r6, %r2, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs5, %r6;
+; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
+; CHECK-NEXT:    bfe.s32         %r7, %r1, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs7, %r7;
+; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
+; CHECK-NEXT:    setp.ne.s16     %p2, %rs8, %rs6;
+; CHECK-NEXT:    bfe.s32         %r8, %r2, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs9, %r8;
+; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
+; CHECK-NEXT:    bfe.s32         %r9, %r1, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs11, %r9;
+; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
+; CHECK-NEXT:    setp.ne.s16     %p3, %rs12, %rs10;
+; CHECK-NEXT:    bfe.s32         %r10, %r2, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs13, %r10;
+; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
+; CHECK-NEXT:    bfe.s32         %r11, %r1, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs15, %r11;
+; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
+; CHECK-NEXT:    setp.ne.s16     %p4, %rs16, %rs14;
+; CHECK-NEXT:    bfe.s32         %r12, %r3, 0, 8;
+; CHECK-NEXT:    selp.b32        %r13, %r11, %r12, %p4;
+; CHECK-NEXT:    bfe.s32         %r14, %r3, 8, 8;
+; CHECK-NEXT:    selp.b32        %r15, %r9, %r14, %p3;
+; CHECK-NEXT:    bfi.b32         %r16, %r15, %r13, 8, 8;
+; CHECK-NEXT:    bfe.s32         %r17, %r3, 16, 8;
+; CHECK-NEXT:    selp.b32        %r18, %r7, %r17, %p2;
+; CHECK-NEXT:    bfi.b32         %r19, %r18, %r16, 16, 8;
+; CHECK-NEXT:    bfe.s32         %r20, %r3, 24, 8;
+; CHECK-NEXT:    selp.b32        %r21, %r5, %r20, %p1;
+; CHECK-NEXT:    bfi.b32         %r22, %r21, %r19, 24, 8;
+; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r22;
 ; CHECK-NEXT:    ret;
   %cmp = icmp ne <4 x i8> %a, %b
   %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c
@@ -912,42 +981,59 @@ define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 {
 define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> %d) #0 {
 ; CHECK-LABEL: test_select_cc(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .pred %p<5>;
-; CHECK-NEXT:    .reg .b32 %r<29>;
+; CHECK-NEXT: .reg .pred      %p<5>;
+; CHECK-NEXT: .reg .b16       %rs<17>;
+; CHECK-NEXT: .reg .b32       %r<29>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.u32 %r4, [test_select_cc_param_3];
 ; 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.s32 %r5, %r4, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r6, %r3, 24, 8;
-; CHECK-NEXT:    setp.ne.u32 %p1, %r6, %r5;
-; CHECK-NEXT:    bfe.s32 %r7, %r4, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r8, %r3, 16, 8;
-; CHECK-NEXT:    setp.ne.u32 %p2, %r8, %r7;
-; CHECK-NEXT:    bfe.s32 %r9, %r4, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r10, %r3, 8, 8;
-; CHECK-NEXT:    setp.ne.u32 %p3, %r10, %r9;
-; CHECK-NEXT:    bfe.s32 %r11, %r4, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r12, %r3, 0, 8;
-; CHECK-NEXT:    setp.ne.u32 %p4, %r12, %r11;
-; CHECK-NEXT:    bfe.s32 %r13, %r2, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r14, %r1, 0, 8;
-; CHECK-NEXT:    selp.b32 %r15, %r14, %r13, %p4;
-; CHECK-NEXT:    bfe.s32 %r16, %r2, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r17, %r1, 8, 8;
-; CHECK-NEXT:    selp.b32 %r18, %r17, %r16, %p3;
-; CHECK-NEXT:    bfi.b32 %r19, %r18, %r15, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r20, %r2, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r21, %r1, 16, 8;
-; CHECK-NEXT:    selp.b32 %r22, %r21, %r20, %p2;
-; CHECK-NEXT:    bfi.b32 %r23, %r22, %r19, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r24, %r2, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r25, %r1, 24, 8;
-; CHECK-NEXT:    selp.b32 %r26, %r25, %r24, %p1;
-; CHECK-NEXT:    bfi.b32 %r27, %r26, %r23, 24, 8;
-; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r27;
+; CHECK-NEXT:    bfe.s32         %r5, %r4, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs1, %r5;
+; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
+; CHECK-NEXT:    bfe.s32         %r6, %r3, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs3, %r6;
+; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
+; CHECK-NEXT:    setp.ne.s16     %p1, %rs4, %rs2;
+; CHECK-NEXT:    bfe.s32         %r7, %r4, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs5, %r7;
+; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
+; CHECK-NEXT:    bfe.s32         %r8, %r3, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs7, %r8;
+; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
+; CHECK-NEXT:    setp.ne.s16     %p2, %rs8, %rs6;
+; CHECK-NEXT:    bfe.s32         %r9, %r4, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs9, %r9;
+; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
+; CHECK-NEXT:    bfe.s32         %r10, %r3, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs11, %r10;
+; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
+; CHECK-NEXT:    setp.ne.s16     %p3, %rs12, %rs10;
+; CHECK-NEXT:    bfe.s32         %r11, %r4, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs13, %r11;
+; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
+; CHECK-NEXT:    bfe.s32         %r12, %r3, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs15, %r12;
+; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
+; CHECK-NEXT:    setp.ne.s16     %p4, %rs16, %rs14;
+; CHECK-NEXT:    bfe.s32         %r13, %r2, 0, 8;
+; CHECK-NEXT:    bfe.s32         %r14, %r1, 0, 8;
+; CHECK-NEXT:    selp.b32        %r15, %r14, %r13, %p4;
+; CHECK-NEXT:    bfe.s32         %r16, %r2, 8, 8;
+; CHECK-NEXT:    bfe.s32         %r17, %r1, 8, 8;
+; CHECK-NEXT:    selp.b32        %r18, %r17, %r16, %p3;
+; CHECK-NEXT:    bfi.b32         %r19, %r18, %r15, 8, 8;
+; CHECK-NEXT:    bfe.s32         %r20, %r2, 16, 8;
+; CHECK-NEXT:    bfe.s32         %r21, %r1, 16, 8;
+; CHECK-NEXT:    selp.b32        %r22, %r21, %r20, %p2;
+; CHECK-NEXT:    bfi.b32         %r23, %r22, %r19, 16, 8;
+; CHECK-NEXT:    bfe.s32         %r24, %r2, 24, 8;
+; CHECK-NEXT:    bfe.s32         %r25, %r1, 24, 8;
+; CHECK-NEXT:    selp.b32        %r26, %r25, %r24, %p1;
+; CHECK-NEXT:    bfi.b32         %r27, %r26, %r23, 24, 8;
+; CHECK-NEXT:    st.param.b32    [func_retval0+0], %r27;
 ; CHECK-NEXT:    ret;
   %cc = icmp ne <4 x i8> %c, %d
   %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b
@@ -957,31 +1043,48 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8>
 define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b,
 ; CHECK-LABEL: test_select_cc_i32_i8(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .pred %p<5>;
-; CHECK-NEXT:    .reg .b32 %r<23>;
+; CHECK-NEXT: .reg .pred      %p<5>;
+; CHECK-NEXT: .reg .b16       %rs<17>;
+; CHECK-NEXT: .reg .b32       %r<23>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.v4.u32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1];
 ; CHECK-NEXT:    ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0];
 ; CHECK-NEXT:    ld.param.u32 %r10, [test_select_cc_i32_i8_param_3];
 ; CHECK-NEXT:    ld.param.u32 %r9, [test_select_cc_i32_i8_param_2];
-; CHECK-NEXT:    bfe.s32 %r11, %r10, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r12, %r9, 0, 8;
-; CHECK-NEXT:    setp.ne.u32 %p1, %r12, %r11;
-; CHECK-NEXT:    bfe.s32 %r13, %r10, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r14, %r9, 8, 8;
-; CHECK-NEXT:    setp.ne.u32 %p2, %r14, %r13;
-; CHECK-NEXT:    bfe.s32 %r15, %r10, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r16, %r9, 16, 8;
-; CHECK-NEXT:    setp.ne.u32 %p3, %r16, %r15;
-; CHECK-NEXT:    bfe.s32 %r17, %r10, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r18, %r9, 24, 8;
-; CHECK-NEXT:    setp.ne.u32 %p4, %r18, %r17;
-; CHECK-NEXT:    selp.b32 %r19, %r4, %r8, %p4;
-; CHECK-NEXT:    selp.b32 %r20, %r3, %r7, %p3;
-; CHECK-NEXT:    selp.b32 %r21, %r2, %r6, %p2;
-; CHECK-NEXT:    selp.b32 %r22, %r1, %r5, %p1;
-; CHECK-NEXT:    st.param.v4.b32 [func_retval0+0], {%r22, %r21, %r20, %r19};
+; CHECK-NEXT:    bfe.s32         %r11, %r10, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs1, %r11;
+; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
+; CHECK-NEXT:    bfe.s32         %r12, %r9, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs3, %r12;
+; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
+; CHECK-NEXT:    setp.ne.s16     %p1, %rs4, %rs2;
+; CHECK-NEXT:    bfe.s32         %r13, %r10, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs5, %r13;
+; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
+; CHECK-NEXT:    bfe.s32         %r14, %r9, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs7, %r14;
+; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
+; CHECK-NEXT:    setp.ne.s16     %p2, %rs8, %rs6;
+; CHECK-NEXT:    bfe.s32         %r15, %r10, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs9, %r15;
+; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
+; CHECK-NEXT:    bfe.s32         %r16, %r9, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs11, %r16;
+; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
+; CHECK-NEXT:    setp.ne.s16     %p3, %rs12, %rs10;
+; CHECK-NEXT:    bfe.s32         %r17, %r10, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs13, %r17;
+; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
+; CHECK-NEXT:    bfe.s32         %r18, %r9, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs15, %r18;
+; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
+; CHECK-NEXT:    setp.ne.s16     %p4, %rs16, %rs14;
+; CHECK-NEXT:    selp.b32        %r19, %r4, %r8, %p4;
+; CHECK-NEXT:    selp.b32        %r20, %r3, %r7, %p3;
+; CHECK-NEXT:    selp.b32        %r21, %r2, %r6, %p2;
+; CHECK-NEXT:    selp.b32        %r22, %r1, %r5, %p1;
+; CHECK-NEXT:    st.param.v4.b32         [func_retval0+0], {%r22, %r21, %r20, %r19};
 ; CHECK-NEXT:    ret;
                                            <4 x i8> %c, <4 x i8> %d) #0 {
   %cc = icmp ne <4 x i8> %c, %d
@@ -1414,9 +1517,10 @@ entry:
 define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) {
 ; CHECK-LABEL: test_sext_v4i1_to_v4i8(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .pred %p<5>;
-; CHECK-NEXT:    .reg .b32 %r<18>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .pred      %p<5>;
+; CHECK-NEXT: .reg .b16       %rs<17>;
+; CHECK-NEXT: .reg .b32       %r<18>;
+; CHECK-NEXT: .reg .b64       %rd<4>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0: // %entry
 ; CHECK-NEXT:    ld.param.u64 %rd3, [test_sext_v4i1_to_v4i8_param_2];
@@ -1424,26 +1528,42 @@ 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.s32 %r3, %r2, 24, 8;
-; CHECK-NEXT:    bfe.s32 %r4, %r1, 24, 8;
-; CHECK-NEXT:    setp.hi.u32 %p1, %r4, %r3;
-; CHECK-NEXT:    bfe.s32 %r5, %r2, 16, 8;
-; CHECK-NEXT:    bfe.s32 %r6, %r1, 16, 8;
-; CHECK-NEXT:    setp.hi.u32 %p2, %r6, %r5;
-; CHECK-NEXT:    bfe.s32 %r7, %r2, 8, 8;
-; CHECK-NEXT:    bfe.s32 %r8, %r1, 8, 8;
-; CHECK-NEXT:    setp.hi.u32 %p3, %r8, %r7;
-; CHECK-NEXT:    bfe.s32 %r9, %r2, 0, 8;
-; CHECK-NEXT:    bfe.s32 %r10, %r1, 0, 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:    selp.s32 %r14, -1, 0, %p2;
-; CHECK-NEXT:    bfi.b32 %r15, %r14, %r13, 16, 8;
-; CHECK-NEXT:    selp.s32 %r16, -1, 0, %p1;
-; CHECK-NEXT:    bfi.b32 %r17, %r16, %r15, 24, 8;
-; CHECK-NEXT:    st.u32 [%rd3], %r17;
+; CHECK-NEXT:    bfe.s32         %r3, %r2, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs1, %r3;
+; CHECK-NEXT:    and.b16         %rs2, %rs1, 255;
+; CHECK-NEXT:    bfe.s32         %r4, %r1, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs3, %r4;
+; CHECK-NEXT:    and.b16         %rs4, %rs3, 255;
+; CHECK-NEXT:    setp.gt.u16     %p1, %rs4, %rs2;
+; CHECK-NEXT:    bfe.s32         %r5, %r2, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs5, %r5;
+; CHECK-NEXT:    and.b16         %rs6, %rs5, 255;
+; CHECK-NEXT:    bfe.s32         %r6, %r1, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs7, %r6;
+; CHECK-NEXT:    and.b16         %rs8, %rs7, 255;
+; CHECK-NEXT:    setp.gt.u16     %p2, %rs8, %rs6;
+; CHECK-NEXT:    bfe.s32         %r7, %r2, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs9, %r7;
+; CHECK-NEXT:    and.b16         %rs10, %rs9, 255;
+; CHECK-NEXT:    bfe.s32         %r8, %r1, 8, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs11, %r8;
+; CHECK-NEXT:    and.b16         %rs12, %rs11, 255;
+; CHECK-NEXT:    setp.gt.u16     %p3, %rs12, %rs10;
+; CHECK-NEXT:    bfe.s32         %r9, %r2, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs13, %r9;
+; CHECK-NEXT:    and.b16         %rs14, %rs13, 255;
+; CHECK-NEXT:    bfe.s32         %r10, %r1, 0, 8;
+; CHECK-NEXT:    cvt.u16.u32     %rs15, %r10;
+; CHECK-NEXT:    and.b16         %rs16, %rs15, 255;
+; CHECK-NEXT:    setp.gt.u16     %p4, %rs16, %rs14;
+; 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:    selp.s32        %r14, -1, 0, %p2;
+; CHECK-NEXT:    bfi.b32         %r15, %r14, %r13, 16, 8;
+; CHECK-NEXT:    selp.s32        %r16, -1, 0, %p1;
+; CHECK-NEXT:    bfi.b32         %r17, %r16, %r15, 24, 8;
+; CHECK-NEXT:    st.u32  [%rd3], %r17;
 ; CHECK-NEXT:    ret;
 entry:
   %t1 = load <4 x i8>, ptr %a, align 4
@@ -1454,4 +1574,4 @@ entry:
   ret void
 }
 
-attributes #0 = { nounwind }
+attributes #0 = { nounwind }
\ No newline at end of file

>From 0170bcdf9331ccdb945cb11a5b682e999c2bea30 Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <34139864+justinfargnoli at users.noreply.github.com>
Date: Mon, 29 Jan 2024 16:36:54 -0800
Subject: [PATCH 3/4] Add newline to end of test

---
 llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index ba247c523963b..c03081ce68362 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -1574,4 +1574,4 @@ entry:
   ret void
 }
 
-attributes #0 = { nounwind }
\ No newline at end of file
+attributes #0 = { nounwind }

>From 4b6f7d57942b56525dc40f9e3ed9cd512c8bc89c Mon Sep 17 00:00:00 2001
From: Justin Fargnoli <34139864+justinfargnoli at users.noreply.github.com>
Date: Thu, 1 Feb 2024 20:06:28 -0800
Subject: [PATCH 4/4] Remove unnecessary `#include` from
 TestTensorTransforms.cpp

---
 mlir/test/lib/Dialect/Tensor/TestTensorTransforms.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/mlir/test/lib/Dialect/Tensor/TestTensorTransforms.cpp b/mlir/test/lib/Dialect/Tensor/TestTensorTransforms.cpp
index b907f77e91082..3fd83303891b0 100644
--- a/mlir/test/lib/Dialect/Tensor/TestTensorTransforms.cpp
+++ b/mlir/test/lib/Dialect/Tensor/TestTensorTransforms.cpp
@@ -16,7 +16,6 @@
 #include "mlir/Dialect/Tensor/IR/Tensor.h"
 #include "mlir/Dialect/Tensor/Transforms/TransformUtils.h"
 #include "mlir/Dialect/Tensor/Transforms/Transforms.h"
-#include "mlir/Dialect/Transform/IR/TransformInterfaces.h"
 #include "mlir/Dialect/Transform/IR/TransformOps.h"
 #include "mlir/Pass/Pass.h"
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"



More information about the libc-commits mailing list