[mlir] [libc] [llvm] [mlir][Tensor] Remove unnecessary `#include` from TestTensorTransforms.cpp (PR #80383)
Justin Fargnoli via llvm-commits
llvm-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 llvm-commits
mailing list