[llvm] [NVPTX] Add PRMT constant folding and cleanup usage of PRMT node (PR #148906)
Alex MacLean via llvm-commits
llvm-commits at lists.llvm.org
Thu Jul 17 09:58:45 PDT 2025
https://github.com/AlexMaclean updated https://github.com/llvm/llvm-project/pull/148906
>From 2cba2415226ee9fd785811567b3027b0ed46cb1e Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Tue, 15 Jul 2025 16:19:10 +0000
Subject: [PATCH 1/3] pre-commit tests
---
llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 2638 ++++++++++++------
1 file changed, 1766 insertions(+), 872 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index 410c0019c7222..728636573a42d 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -1,14 +1,19 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 3
; ## Support i16x2 instructions
-; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx80 \
-; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
-; RUN: | FileCheck -allow-deprecated-dag-overlap %s
-; RUN: %if ptxas %{ \
-; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 \
-; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
-; RUN: | %ptxas-verify -arch=sm_90 \
+; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
+; RUN: | FileCheck %s --check-prefixes=CHECK,O0
+; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -verify-machineinstrs \
+; RUN: | FileCheck %s --check-prefixes=CHECK,O3
+; RUN: %if ptxas %{ \
+; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \
+; RUN: | %ptxas-verify -arch=sm_90 \
+; RUN: %}
+; RUN: %if ptxas %{ \
+; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -verify-machineinstrs \
+; RUN: | %ptxas-verify -arch=sm_90 \
; RUN: %}
+target triple = "nvptx64-nvidia-cuda"
target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
define <4 x i8> @test_ret_const() #0 {
@@ -79,61 +84,111 @@ define i8 @test_extract_3(<4 x i8> %a) #0 {
}
define i8 @test_extract_i(<4 x i8> %a, i64 %idx) #0 {
-; CHECK-LABEL: test_extract_i(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];
-; CHECK-NEXT: cvt.u32.u64 %r2, %rd1;
-; CHECK-NEXT: or.b32 %r3, %r2, 30576;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, %r3;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_extract_i(
+; O0: {
+; O0-NEXT: .reg .b32 %r<5>;
+; O0-NEXT: .reg .b64 %rd<2>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];
+; O0-NEXT: cvt.u32.u64 %r2, %rd1;
+; O0-NEXT: or.b32 %r3, %r2, 30576;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, %r3;
+; O0-NEXT: st.param.b32 [func_retval0], %r4;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_extract_i(
+; O3: {
+; O3-NEXT: .reg .b32 %r<5>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_extract_i_param_1];
+; O3-NEXT: or.b32 %r3, %r2, 30576;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, %r3;
+; O3-NEXT: st.param.b32 [func_retval0], %r4;
+; O3-NEXT: ret;
%e = extractelement <4 x i8> %a, i64 %idx
ret i8 %e
}
define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_add(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<13>;
-; CHECK-NEXT: .reg .b32 %r<18>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_add_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_add_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
-; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
-; CHECK-NEXT: cvt.u16.u32 %rs2, %r4;
-; CHECK-NEXT: add.s16 %rs3, %rs2, %rs1;
-; CHECK-NEXT: cvt.u32.u16 %r5, %rs3;
-; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
-; CHECK-NEXT: cvt.u16.u32 %rs4, %r6;
-; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
-; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
-; CHECK-NEXT: add.s16 %rs6, %rs5, %rs4;
-; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
-; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
-; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;
-; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
-; CHECK-NEXT: cvt.u16.u32 %rs8, %r11;
-; CHECK-NEXT: add.s16 %rs9, %rs8, %rs7;
-; CHECK-NEXT: cvt.u32.u16 %r12, %rs9;
-; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
-; CHECK-NEXT: cvt.u16.u32 %rs10, %r13;
-; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
-; CHECK-NEXT: cvt.u16.u32 %rs11, %r14;
-; CHECK-NEXT: add.s16 %rs12, %rs11, %rs10;
-; CHECK-NEXT: cvt.u32.u16 %r15, %rs12;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_add(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<13>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_add_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_add_param_0];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
+; O0-NEXT: cvt.u16.u32 %rs1, %r3;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
+; O0-NEXT: cvt.u16.u32 %rs2, %r4;
+; O0-NEXT: add.s16 %rs3, %rs2, %rs1;
+; O0-NEXT: cvt.u32.u16 %r5, %rs3;
+; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
+; O0-NEXT: cvt.u16.u32 %rs4, %r6;
+; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
+; O0-NEXT: cvt.u16.u32 %rs5, %r7;
+; O0-NEXT: add.s16 %rs6, %rs5, %rs4;
+; O0-NEXT: cvt.u32.u16 %r8, %rs6;
+; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
+; O0-NEXT: cvt.u16.u32 %rs7, %r10;
+; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
+; O0-NEXT: cvt.u16.u32 %rs8, %r11;
+; O0-NEXT: add.s16 %rs9, %rs8, %rs7;
+; O0-NEXT: cvt.u32.u16 %r12, %rs9;
+; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
+; O0-NEXT: cvt.u16.u32 %rs10, %r13;
+; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
+; O0-NEXT: cvt.u16.u32 %rs11, %r14;
+; O0-NEXT: add.s16 %rs12, %rs11, %rs10;
+; O0-NEXT: cvt.u32.u16 %r15, %rs12;
+; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_add(
+; O3: {
+; O3-NEXT: .reg .b16 %rs<13>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_add_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_add_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
+; O3-NEXT: cvt.u16.u32 %rs1, %r3;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
+; O3-NEXT: cvt.u16.u32 %rs2, %r4;
+; O3-NEXT: add.s16 %rs3, %rs2, %rs1;
+; O3-NEXT: cvt.u32.u16 %r5, %rs3;
+; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
+; O3-NEXT: cvt.u16.u32 %rs4, %r6;
+; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
+; O3-NEXT: cvt.u16.u32 %rs5, %r7;
+; O3-NEXT: add.s16 %rs6, %rs5, %rs4;
+; O3-NEXT: cvt.u32.u16 %r8, %rs6;
+; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
+; O3-NEXT: cvt.u16.u32 %rs7, %r10;
+; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
+; O3-NEXT: cvt.u16.u32 %rs8, %r11;
+; O3-NEXT: add.s16 %rs9, %rs8, %rs7;
+; O3-NEXT: cvt.u32.u16 %r12, %rs9;
+; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
+; O3-NEXT: cvt.u16.u32 %rs10, %r13;
+; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
+; O3-NEXT: cvt.u16.u32 %rs11, %r14;
+; O3-NEXT: add.s16 %rs12, %rs11, %rs10;
+; O3-NEXT: cvt.u32.u16 %r15, %rs12;
+; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r17;
+; O3-NEXT: ret;
%r = add <4 x i8> %a, %b
ret <4 x i8> %r
}
@@ -205,341 +260,633 @@ define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 {
}
define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_sub(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<13>;
-; CHECK-NEXT: .reg .b32 %r<18>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_sub_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_sub_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
-; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
-; CHECK-NEXT: cvt.u16.u32 %rs2, %r4;
-; CHECK-NEXT: sub.s16 %rs3, %rs2, %rs1;
-; CHECK-NEXT: cvt.u32.u16 %r5, %rs3;
-; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
-; CHECK-NEXT: cvt.u16.u32 %rs4, %r6;
-; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
-; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
-; CHECK-NEXT: sub.s16 %rs6, %rs5, %rs4;
-; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
-; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
-; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;
-; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
-; CHECK-NEXT: cvt.u16.u32 %rs8, %r11;
-; CHECK-NEXT: sub.s16 %rs9, %rs8, %rs7;
-; CHECK-NEXT: cvt.u32.u16 %r12, %rs9;
-; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
-; CHECK-NEXT: cvt.u16.u32 %rs10, %r13;
-; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
-; CHECK-NEXT: cvt.u16.u32 %rs11, %r14;
-; CHECK-NEXT: sub.s16 %rs12, %rs11, %rs10;
-; CHECK-NEXT: cvt.u32.u16 %r15, %rs12;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_sub(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<13>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_sub_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_sub_param_0];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
+; O0-NEXT: cvt.u16.u32 %rs1, %r3;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
+; O0-NEXT: cvt.u16.u32 %rs2, %r4;
+; O0-NEXT: sub.s16 %rs3, %rs2, %rs1;
+; O0-NEXT: cvt.u32.u16 %r5, %rs3;
+; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
+; O0-NEXT: cvt.u16.u32 %rs4, %r6;
+; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
+; O0-NEXT: cvt.u16.u32 %rs5, %r7;
+; O0-NEXT: sub.s16 %rs6, %rs5, %rs4;
+; O0-NEXT: cvt.u32.u16 %r8, %rs6;
+; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
+; O0-NEXT: cvt.u16.u32 %rs7, %r10;
+; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
+; O0-NEXT: cvt.u16.u32 %rs8, %r11;
+; O0-NEXT: sub.s16 %rs9, %rs8, %rs7;
+; O0-NEXT: cvt.u32.u16 %r12, %rs9;
+; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
+; O0-NEXT: cvt.u16.u32 %rs10, %r13;
+; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
+; O0-NEXT: cvt.u16.u32 %rs11, %r14;
+; O0-NEXT: sub.s16 %rs12, %rs11, %rs10;
+; O0-NEXT: cvt.u32.u16 %r15, %rs12;
+; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_sub(
+; O3: {
+; O3-NEXT: .reg .b16 %rs<13>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_sub_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_sub_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
+; O3-NEXT: cvt.u16.u32 %rs1, %r3;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
+; O3-NEXT: cvt.u16.u32 %rs2, %r4;
+; O3-NEXT: sub.s16 %rs3, %rs2, %rs1;
+; O3-NEXT: cvt.u32.u16 %r5, %rs3;
+; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
+; O3-NEXT: cvt.u16.u32 %rs4, %r6;
+; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
+; O3-NEXT: cvt.u16.u32 %rs5, %r7;
+; O3-NEXT: sub.s16 %rs6, %rs5, %rs4;
+; O3-NEXT: cvt.u32.u16 %r8, %rs6;
+; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
+; O3-NEXT: cvt.u16.u32 %rs7, %r10;
+; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
+; O3-NEXT: cvt.u16.u32 %rs8, %r11;
+; O3-NEXT: sub.s16 %rs9, %rs8, %rs7;
+; O3-NEXT: cvt.u32.u16 %r12, %rs9;
+; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
+; O3-NEXT: cvt.u16.u32 %rs10, %r13;
+; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
+; O3-NEXT: cvt.u16.u32 %rs11, %r14;
+; O3-NEXT: sub.s16 %rs12, %rs11, %rs10;
+; O3-NEXT: cvt.u32.u16 %r15, %rs12;
+; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r17;
+; O3-NEXT: ret;
%r = sub <4 x i8> %a, %b
ret <4 x i8> %r
}
define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_smax(
-; CHECK: {
-; CHECK-NEXT: .reg .pred %p<5>;
-; CHECK-NEXT: .reg .b32 %r<18>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_smax_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_smax_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
-; CHECK-NEXT: setp.gt.s32 %p1, %r4, %r3;
-; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
-; CHECK-NEXT: setp.gt.s32 %p2, %r6, %r5;
-; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
-; CHECK-NEXT: setp.gt.s32 %p3, %r8, %r7;
-; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
-; CHECK-NEXT: setp.gt.s32 %p4, %r10, %r9;
-; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4;
-; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3;
-; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
-; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2;
-; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_smax(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_smax_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_smax_param_0];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O0-NEXT: setp.gt.s32 %p1, %r4, %r3;
+; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O0-NEXT: setp.gt.s32 %p2, %r6, %r5;
+; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O0-NEXT: setp.gt.s32 %p3, %r8, %r7;
+; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O0-NEXT: setp.gt.s32 %p4, %r10, %r9;
+; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_smax(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_smax_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_smax_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O3-NEXT: setp.gt.s32 %p1, %r4, %r3;
+; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O3-NEXT: setp.gt.s32 %p2, %r6, %r5;
+; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O3-NEXT: setp.gt.s32 %p3, %r8, %r7;
+; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O3-NEXT: setp.gt.s32 %p4, %r10, %r9;
+; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r17;
+; O3-NEXT: ret;
%cmp = icmp sgt <4 x i8> %a, %b
%r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b
ret <4 x i8> %r
}
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<18>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_umax_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_umax_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
-; CHECK-NEXT: setp.gt.u32 %p1, %r4, %r3;
-; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
-; CHECK-NEXT: setp.gt.u32 %p2, %r6, %r5;
-; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
-; CHECK-NEXT: setp.gt.u32 %p3, %r8, %r7;
-; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
-; CHECK-NEXT: setp.gt.u32 %p4, %r10, %r9;
-; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4;
-; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3;
-; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
-; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2;
-; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_umax(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_umax_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_umax_param_0];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O0-NEXT: setp.gt.u32 %p1, %r4, %r3;
+; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O0-NEXT: setp.gt.u32 %p2, %r6, %r5;
+; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O0-NEXT: setp.gt.u32 %p3, %r8, %r7;
+; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O0-NEXT: setp.gt.u32 %p4, %r10, %r9;
+; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_umax(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_umax_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_umax_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O3-NEXT: setp.gt.u32 %p1, %r4, %r3;
+; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O3-NEXT: setp.gt.u32 %p2, %r6, %r5;
+; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O3-NEXT: setp.gt.u32 %p3, %r8, %r7;
+; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O3-NEXT: setp.gt.u32 %p4, %r10, %r9;
+; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r17;
+; O3-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
}
define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_smin(
-; CHECK: {
-; CHECK-NEXT: .reg .pred %p<5>;
-; CHECK-NEXT: .reg .b32 %r<18>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_smin_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_smin_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
-; CHECK-NEXT: setp.le.s32 %p1, %r4, %r3;
-; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
-; CHECK-NEXT: setp.le.s32 %p2, %r6, %r5;
-; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
-; CHECK-NEXT: setp.le.s32 %p3, %r8, %r7;
-; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
-; CHECK-NEXT: setp.le.s32 %p4, %r10, %r9;
-; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4;
-; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3;
-; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
-; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2;
-; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_smin(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_smin_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_smin_param_0];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O0-NEXT: setp.le.s32 %p1, %r4, %r3;
+; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O0-NEXT: setp.le.s32 %p2, %r6, %r5;
+; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O0-NEXT: setp.le.s32 %p3, %r8, %r7;
+; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O0-NEXT: setp.le.s32 %p4, %r10, %r9;
+; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_smin(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_smin_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_smin_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O3-NEXT: setp.le.s32 %p1, %r4, %r3;
+; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O3-NEXT: setp.le.s32 %p2, %r6, %r5;
+; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O3-NEXT: setp.le.s32 %p3, %r8, %r7;
+; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O3-NEXT: setp.le.s32 %p4, %r10, %r9;
+; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r17;
+; O3-NEXT: ret;
%cmp = icmp sle <4 x i8> %a, %b
%r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b
ret <4 x i8> %r
}
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<18>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_umin_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_umin_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
-; CHECK-NEXT: setp.le.u32 %p1, %r4, %r3;
-; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
-; CHECK-NEXT: setp.le.u32 %p2, %r6, %r5;
-; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
-; CHECK-NEXT: setp.le.u32 %p3, %r8, %r7;
-; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
-; CHECK-NEXT: setp.le.u32 %p4, %r10, %r9;
-; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4;
-; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3;
-; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
-; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2;
-; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_umin(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_umin_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_umin_param_0];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O0-NEXT: setp.le.u32 %p1, %r4, %r3;
+; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O0-NEXT: setp.le.u32 %p2, %r6, %r5;
+; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O0-NEXT: setp.le.u32 %p3, %r8, %r7;
+; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O0-NEXT: setp.le.u32 %p4, %r10, %r9;
+; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_umin(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_umin_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_umin_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O3-NEXT: setp.le.u32 %p1, %r4, %r3;
+; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O3-NEXT: setp.le.u32 %p2, %r6, %r5;
+; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O3-NEXT: setp.le.u32 %p3, %r8, %r7;
+; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O3-NEXT: setp.le.u32 %p4, %r10, %r9;
+; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4;
+; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3;
+; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2;
+; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1;
+; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r17;
+; O3-NEXT: ret;
%cmp = icmp ule <4 x i8> %a, %b
%r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b
ret <4 x i8> %r
}
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<23>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r3, [test_eq_param_2];
-; CHECK-NEXT: ld.param.b32 %r2, [test_eq_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_eq_param_0];
-; CHECK-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;
-; CHECK-NEXT: setp.eq.b32 %p1, %r5, %r4;
-; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;
-; CHECK-NEXT: setp.eq.b32 %p2, %r7, %r6;
-; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U;
-; CHECK-NEXT: setp.eq.b32 %p3, %r9, %r8;
-; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U;
-; CHECK-NEXT: setp.eq.b32 %p4, %r11, %r10;
-; CHECK-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;
-; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4;
-; CHECK-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U;
-; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U;
-; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2;
-; CHECK-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U;
-; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1;
-; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r22;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_eq(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<23>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r3, [test_eq_param_2];
+; O0-NEXT: ld.param.b32 %r2, [test_eq_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_eq_param_0];
+; O0-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;
+; O0-NEXT: setp.eq.b32 %p1, %r5, %r4;
+; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;
+; O0-NEXT: setp.eq.b32 %p2, %r7, %r6;
+; O0-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U;
+; O0-NEXT: setp.eq.b32 %p3, %r9, %r8;
+; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U;
+; O0-NEXT: setp.eq.b32 %p4, %r11, %r10;
+; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;
+; O0-NEXT: selp.b32 %r13, %r11, %r12, %p4;
+; O0-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U;
+; O0-NEXT: selp.b32 %r15, %r9, %r14, %p3;
+; O0-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U;
+; O0-NEXT: selp.b32 %r18, %r7, %r17, %p2;
+; O0-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U;
+; O0-NEXT: selp.b32 %r20, %r5, %r19, %p1;
+; O0-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;
+; O0-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r22;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_eq(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<23>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_eq_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_eq_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O3-NEXT: setp.eq.b32 %p1, %r4, %r3;
+; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O3-NEXT: setp.eq.b32 %p2, %r6, %r5;
+; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O3-NEXT: setp.eq.b32 %p3, %r8, %r7;
+; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O3-NEXT: setp.eq.b32 %p4, %r10, %r9;
+; O3-NEXT: ld.param.b32 %r11, [test_eq_param_2];
+; O3-NEXT: prmt.b32 %r12, %r11, 0, 0x7773U;
+; O3-NEXT: selp.b32 %r13, %r10, %r12, %p4;
+; O3-NEXT: prmt.b32 %r14, %r11, 0, 0x7772U;
+; O3-NEXT: selp.b32 %r15, %r8, %r14, %p3;
+; O3-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r11, 0, 0x7771U;
+; O3-NEXT: selp.b32 %r18, %r6, %r17, %p2;
+; O3-NEXT: prmt.b32 %r19, %r11, 0, 0x7770U;
+; O3-NEXT: selp.b32 %r20, %r4, %r19, %p1;
+; O3-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;
+; O3-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r22;
+; O3-NEXT: ret;
%cmp = icmp eq <4 x i8> %a, %b
%r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c
ret <4 x i8> %r
}
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<23>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r3, [test_ne_param_2];
-; CHECK-NEXT: ld.param.b32 %r2, [test_ne_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_ne_param_0];
-; CHECK-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;
-; CHECK-NEXT: setp.ne.b32 %p1, %r5, %r4;
-; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;
-; CHECK-NEXT: setp.ne.b32 %p2, %r7, %r6;
-; CHECK-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U;
-; CHECK-NEXT: setp.ne.b32 %p3, %r9, %r8;
-; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U;
-; CHECK-NEXT: setp.ne.b32 %p4, %r11, %r10;
-; CHECK-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;
-; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4;
-; CHECK-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U;
-; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U;
-; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2;
-; CHECK-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U;
-; CHECK-NEXT: selp.b32 %r20, %r5, %r19, %p1;
-; CHECK-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r22;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_ne(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<23>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r3, [test_ne_param_2];
+; O0-NEXT: ld.param.b32 %r2, [test_ne_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_ne_param_0];
+; O0-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;
+; O0-NEXT: setp.ne.b32 %p1, %r5, %r4;
+; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;
+; O0-NEXT: setp.ne.b32 %p2, %r7, %r6;
+; O0-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U;
+; O0-NEXT: setp.ne.b32 %p3, %r9, %r8;
+; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U;
+; O0-NEXT: setp.ne.b32 %p4, %r11, %r10;
+; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;
+; O0-NEXT: selp.b32 %r13, %r11, %r12, %p4;
+; O0-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U;
+; O0-NEXT: selp.b32 %r15, %r9, %r14, %p3;
+; O0-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U;
+; O0-NEXT: selp.b32 %r18, %r7, %r17, %p2;
+; O0-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U;
+; O0-NEXT: selp.b32 %r20, %r5, %r19, %p1;
+; O0-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;
+; O0-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r22;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_ne(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<23>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_ne_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_ne_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O3-NEXT: setp.ne.b32 %p1, %r4, %r3;
+; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O3-NEXT: setp.ne.b32 %p2, %r6, %r5;
+; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O3-NEXT: setp.ne.b32 %p3, %r8, %r7;
+; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O3-NEXT: setp.ne.b32 %p4, %r10, %r9;
+; O3-NEXT: ld.param.b32 %r11, [test_ne_param_2];
+; O3-NEXT: prmt.b32 %r12, %r11, 0, 0x7773U;
+; O3-NEXT: selp.b32 %r13, %r10, %r12, %p4;
+; O3-NEXT: prmt.b32 %r14, %r11, 0, 0x7772U;
+; O3-NEXT: selp.b32 %r15, %r8, %r14, %p3;
+; O3-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r11, 0, 0x7771U;
+; O3-NEXT: selp.b32 %r18, %r6, %r17, %p2;
+; O3-NEXT: prmt.b32 %r19, %r11, 0, 0x7770U;
+; O3-NEXT: selp.b32 %r20, %r4, %r19, %p1;
+; O3-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;
+; O3-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r22;
+; O3-NEXT: ret;
%cmp = icmp ne <4 x i8> %a, %b
%r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c
ret <4 x i8> %r
}
define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_mul(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<13>;
-; CHECK-NEXT: .reg .b32 %r<18>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_mul_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_mul_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
-; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
-; CHECK-NEXT: cvt.u16.u32 %rs2, %r4;
-; CHECK-NEXT: mul.lo.s16 %rs3, %rs2, %rs1;
-; CHECK-NEXT: cvt.u32.u16 %r5, %rs3;
-; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
-; CHECK-NEXT: cvt.u16.u32 %rs4, %r6;
-; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
-; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
-; CHECK-NEXT: mul.lo.s16 %rs6, %rs5, %rs4;
-; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
-; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
-; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;
-; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
-; CHECK-NEXT: cvt.u16.u32 %rs8, %r11;
-; CHECK-NEXT: mul.lo.s16 %rs9, %rs8, %rs7;
-; CHECK-NEXT: cvt.u32.u16 %r12, %rs9;
-; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
-; CHECK-NEXT: cvt.u16.u32 %rs10, %r13;
-; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
-; CHECK-NEXT: cvt.u16.u32 %rs11, %r14;
-; CHECK-NEXT: mul.lo.s16 %rs12, %rs11, %rs10;
-; CHECK-NEXT: cvt.u32.u16 %r15, %rs12;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_mul(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<13>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_mul_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_mul_param_0];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
+; O0-NEXT: cvt.u16.u32 %rs1, %r3;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
+; O0-NEXT: cvt.u16.u32 %rs2, %r4;
+; O0-NEXT: mul.lo.s16 %rs3, %rs2, %rs1;
+; O0-NEXT: cvt.u32.u16 %r5, %rs3;
+; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
+; O0-NEXT: cvt.u16.u32 %rs4, %r6;
+; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
+; O0-NEXT: cvt.u16.u32 %rs5, %r7;
+; O0-NEXT: mul.lo.s16 %rs6, %rs5, %rs4;
+; O0-NEXT: cvt.u32.u16 %r8, %rs6;
+; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
+; O0-NEXT: cvt.u16.u32 %rs7, %r10;
+; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
+; O0-NEXT: cvt.u16.u32 %rs8, %r11;
+; O0-NEXT: mul.lo.s16 %rs9, %rs8, %rs7;
+; O0-NEXT: cvt.u32.u16 %r12, %rs9;
+; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
+; O0-NEXT: cvt.u16.u32 %rs10, %r13;
+; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
+; O0-NEXT: cvt.u16.u32 %rs11, %r14;
+; O0-NEXT: mul.lo.s16 %rs12, %rs11, %rs10;
+; O0-NEXT: cvt.u32.u16 %r15, %rs12;
+; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_mul(
+; O3: {
+; O3-NEXT: .reg .b16 %rs<13>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_mul_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_mul_param_1];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;
+; O3-NEXT: cvt.u16.u32 %rs1, %r3;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
+; O3-NEXT: cvt.u16.u32 %rs2, %r4;
+; O3-NEXT: mul.lo.s16 %rs3, %rs2, %rs1;
+; O3-NEXT: cvt.u32.u16 %r5, %rs3;
+; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;
+; O3-NEXT: cvt.u16.u32 %rs4, %r6;
+; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;
+; O3-NEXT: cvt.u16.u32 %rs5, %r7;
+; O3-NEXT: mul.lo.s16 %rs6, %rs5, %rs4;
+; O3-NEXT: cvt.u32.u16 %r8, %rs6;
+; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;
+; O3-NEXT: cvt.u16.u32 %rs7, %r10;
+; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;
+; O3-NEXT: cvt.u16.u32 %rs8, %r11;
+; O3-NEXT: mul.lo.s16 %rs9, %rs8, %rs7;
+; O3-NEXT: cvt.u32.u16 %r12, %rs9;
+; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;
+; O3-NEXT: cvt.u16.u32 %rs10, %r13;
+; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;
+; O3-NEXT: cvt.u16.u32 %rs11, %r14;
+; O3-NEXT: mul.lo.s16 %rs12, %rs11, %rs10;
+; O3-NEXT: cvt.u32.u16 %r15, %rs12;
+; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r17;
+; O3-NEXT: ret;
%r = mul <4 x i8> %a, %b
ret <4 x i8> %r
}
define <4 x i8> @test_or(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_or(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_or_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_or_param_0];
-; CHECK-NEXT: or.b32 %r3, %r1, %r2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_or(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_or_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_or_param_0];
+; O0-NEXT: or.b32 %r3, %r1, %r2;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_or(
+; O3: {
+; O3-NEXT: .reg .b32 %r<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_or_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_or_param_1];
+; O3-NEXT: or.b32 %r3, %r1, %r2;
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%r = or <4 x i8> %a, %b
ret <4 x i8> %r
}
define <4 x i8> @test_or_computed(i8 %a) {
-; CHECK-LABEL: test_or_computed(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<8>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b8 %rs1, [test_or_computed_param_0];
-; CHECK-NEXT: mov.b32 %r1, 0;
-; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
-; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
-; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
-; CHECK-NEXT: or.b32 %r7, %r6, %r5;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_or_computed(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<2>;
+; O0-NEXT: .reg .b32 %r<8>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b8 %rs1, [test_or_computed_param_0];
+; O0-NEXT: mov.b32 %r1, 0;
+; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
+; O0-NEXT: cvt.u32.u16 %r3, %rs1;
+; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
+; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
+; O0-NEXT: or.b32 %r7, %r6, %r5;
+; O0-NEXT: st.param.b32 [func_retval0], %r7;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_or_computed(
+; O3: {
+; O3-NEXT: .reg .b32 %r<8>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: mov.b32 %r1, 0;
+; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
+; O3-NEXT: ld.param.b8 %r3, [test_or_computed_param_0];
+; O3-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
+; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O3-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
+; O3-NEXT: or.b32 %r7, %r6, %r5;
+; O3-NEXT: st.param.b32 [func_retval0], %r7;
+; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
%r = or <4 x i8> %ins.1, %ins.0
@@ -575,37 +922,63 @@ define <4 x i8> @test_or_imm_1(<4 x i8> %a) #0 {
}
define <4 x i8> @test_xor(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_xor(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_xor_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_xor_param_0];
-; CHECK-NEXT: xor.b32 %r3, %r1, %r2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_xor(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_xor_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_xor_param_0];
+; O0-NEXT: xor.b32 %r3, %r1, %r2;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_xor(
+; O3: {
+; O3-NEXT: .reg .b32 %r<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_xor_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_xor_param_1];
+; O3-NEXT: xor.b32 %r3, %r1, %r2;
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%r = xor <4 x i8> %a, %b
ret <4 x i8> %r
}
define <4 x i8> @test_xor_computed(i8 %a) {
-; CHECK-LABEL: test_xor_computed(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<8>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b8 %rs1, [test_xor_computed_param_0];
-; CHECK-NEXT: mov.b32 %r1, 0;
-; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
-; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
-; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
-; CHECK-NEXT: xor.b32 %r7, %r6, %r5;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_xor_computed(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<2>;
+; O0-NEXT: .reg .b32 %r<8>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b8 %rs1, [test_xor_computed_param_0];
+; O0-NEXT: mov.b32 %r1, 0;
+; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
+; O0-NEXT: cvt.u32.u16 %r3, %rs1;
+; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
+; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
+; O0-NEXT: xor.b32 %r7, %r6, %r5;
+; O0-NEXT: st.param.b32 [func_retval0], %r7;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_xor_computed(
+; O3: {
+; O3-NEXT: .reg .b32 %r<8>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: mov.b32 %r1, 0;
+; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
+; O3-NEXT: ld.param.b8 %r3, [test_xor_computed_param_0];
+; O3-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
+; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O3-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
+; O3-NEXT: xor.b32 %r7, %r6, %r5;
+; O3-NEXT: st.param.b32 [func_retval0], %r7;
+; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
%r = xor <4 x i8> %ins.1, %ins.0
@@ -641,37 +1014,63 @@ define <4 x i8> @test_xor_imm_1(<4 x i8> %a) #0 {
}
define <4 x i8> @test_and(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_and(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_and_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_and_param_0];
-; CHECK-NEXT: and.b32 %r3, %r1, %r2;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_and(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_and_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_and_param_0];
+; O0-NEXT: and.b32 %r3, %r1, %r2;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_and(
+; O3: {
+; O3-NEXT: .reg .b32 %r<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_and_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_and_param_1];
+; O3-NEXT: and.b32 %r3, %r1, %r2;
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%r = and <4 x i8> %a, %b
ret <4 x i8> %r
}
define <4 x i8> @test_and_computed(i8 %a) {
-; CHECK-LABEL: test_and_computed(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<8>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b8 %rs1, [test_and_computed_param_0];
-; CHECK-NEXT: mov.b32 %r1, 0;
-; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
-; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
-; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; CHECK-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
-; CHECK-NEXT: and.b32 %r7, %r6, %r5;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_and_computed(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<2>;
+; O0-NEXT: .reg .b32 %r<8>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b8 %rs1, [test_and_computed_param_0];
+; O0-NEXT: mov.b32 %r1, 0;
+; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
+; O0-NEXT: cvt.u32.u16 %r3, %rs1;
+; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
+; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
+; O0-NEXT: and.b32 %r7, %r6, %r5;
+; O0-NEXT: st.param.b32 [func_retval0], %r7;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_and_computed(
+; O3: {
+; O3-NEXT: .reg .b32 %r<8>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: mov.b32 %r1, 0;
+; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
+; O3-NEXT: ld.param.b8 %r3, [test_and_computed_param_0];
+; O3-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
+; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O3-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
+; O3-NEXT: and.b32 %r7, %r6, %r5;
+; O3-NEXT: st.param.b32 [func_retval0], %r7;
+; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
%r = and <4 x i8> %ins.1, %ins.0
@@ -707,76 +1106,132 @@ define <4 x i8> @test_and_imm_1(<4 x i8> %a) #0 {
}
define void @test_ldst_v2i8(ptr %a, ptr %b) {
-; CHECK-LABEL: test_ldst_v2i8(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0];
-; CHECK-NEXT: ld.b32 %r1, [%rd1];
-; CHECK-NEXT: st.b32 [%rd2], %r1;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_ldst_v2i8(
+; O0: {
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-NEXT: .reg .b64 %rd<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0];
+; O0-NEXT: ld.b32 %r1, [%rd1];
+; O0-NEXT: st.b32 [%rd2], %r1;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_ldst_v2i8(
+; O3: {
+; O3-NEXT: .reg .b32 %r<2>;
+; O3-NEXT: .reg .b64 %rd<3>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0];
+; O3-NEXT: ld.b32 %r1, [%rd1];
+; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1];
+; O3-NEXT: st.b32 [%rd2], %r1;
+; O3-NEXT: ret;
%t1 = load <4 x i8>, ptr %a
store <4 x i8> %t1, ptr %b, align 16
ret void
}
define void @test_ldst_v3i8(ptr %a, ptr %b) {
-; CHECK-LABEL: test_ldst_v3i8(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0];
-; CHECK-NEXT: ld.b32 %r1, [%rd1];
-; CHECK-NEXT: st.b16 [%rd2], %r1;
-; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;
-; CHECK-NEXT: st.b8 [%rd2+2], %r2;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_ldst_v3i8(
+; O0: {
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-NEXT: .reg .b64 %rd<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0];
+; O0-NEXT: ld.b32 %r1, [%rd1];
+; O0-NEXT: st.b16 [%rd2], %r1;
+; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;
+; O0-NEXT: st.b8 [%rd2+2], %r2;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_ldst_v3i8(
+; O3: {
+; O3-NEXT: .reg .b32 %r<3>;
+; O3-NEXT: .reg .b64 %rd<3>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0];
+; O3-NEXT: ld.b32 %r1, [%rd1];
+; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1];
+; O3-NEXT: st.b16 [%rd2], %r1;
+; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;
+; O3-NEXT: st.b8 [%rd2+2], %r2;
+; O3-NEXT: ret;
%t1 = load <3 x i8>, ptr %a
store <3 x i8> %t1, ptr %b, align 16
ret void
}
define void @test_ldst_v4i8(ptr %a, ptr %b) {
-; CHECK-LABEL: test_ldst_v4i8(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<2>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0];
-; CHECK-NEXT: ld.b32 %r1, [%rd1];
-; CHECK-NEXT: st.b32 [%rd2], %r1;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_ldst_v4i8(
+; O0: {
+; O0-NEXT: .reg .b32 %r<2>;
+; O0-NEXT: .reg .b64 %rd<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0];
+; O0-NEXT: ld.b32 %r1, [%rd1];
+; O0-NEXT: st.b32 [%rd2], %r1;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_ldst_v4i8(
+; O3: {
+; O3-NEXT: .reg .b32 %r<2>;
+; O3-NEXT: .reg .b64 %rd<3>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0];
+; O3-NEXT: ld.b32 %r1, [%rd1];
+; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1];
+; O3-NEXT: st.b32 [%rd2], %r1;
+; O3-NEXT: ret;
%t1 = load <4 x i8>, ptr %a
store <4 x i8> %t1, ptr %b, align 16
ret void
}
define void @test_ldst_v4i8_unaligned(ptr %a, ptr %b) {
-; CHECK-LABEL: test_ldst_v4i8_unaligned(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0];
-; CHECK-NEXT: ld.b8 %r1, [%rd1];
-; CHECK-NEXT: ld.b8 %r2, [%rd1+1];
-; CHECK-NEXT: ld.b8 %r3, [%rd1+2];
-; CHECK-NEXT: ld.b8 %r4, [%rd1+3];
-; CHECK-NEXT: st.b8 [%rd2+3], %r4;
-; CHECK-NEXT: st.b8 [%rd2+2], %r3;
-; CHECK-NEXT: st.b8 [%rd2+1], %r2;
-; CHECK-NEXT: st.b8 [%rd2], %r1;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_ldst_v4i8_unaligned(
+; O0: {
+; O0-NEXT: .reg .b32 %r<5>;
+; O0-NEXT: .reg .b64 %rd<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0];
+; O0-NEXT: ld.b8 %r1, [%rd1];
+; O0-NEXT: ld.b8 %r2, [%rd1+1];
+; O0-NEXT: ld.b8 %r3, [%rd1+2];
+; O0-NEXT: ld.b8 %r4, [%rd1+3];
+; O0-NEXT: st.b8 [%rd2+3], %r4;
+; O0-NEXT: st.b8 [%rd2+2], %r3;
+; O0-NEXT: st.b8 [%rd2+1], %r2;
+; O0-NEXT: st.b8 [%rd2], %r1;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_ldst_v4i8_unaligned(
+; O3: {
+; O3-NEXT: .reg .b32 %r<5>;
+; O3-NEXT: .reg .b64 %rd<3>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0];
+; O3-NEXT: ld.b8 %r1, [%rd1+1];
+; O3-NEXT: ld.b8 %r2, [%rd1];
+; O3-NEXT: ld.b8 %r3, [%rd1+3];
+; O3-NEXT: ld.b8 %r4, [%rd1+2];
+; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1];
+; O3-NEXT: st.b8 [%rd2+2], %r4;
+; O3-NEXT: st.b8 [%rd2+3], %r3;
+; O3-NEXT: st.b8 [%rd2], %r2;
+; O3-NEXT: st.b8 [%rd2+1], %r1;
+; O3-NEXT: ret;
%t1 = load <4 x i8>, ptr %a, align 1
store <4 x i8> %t1, ptr %b, align 1
ret void
@@ -784,17 +1239,29 @@ define void @test_ldst_v4i8_unaligned(ptr %a, ptr %b) {
define void @test_ldst_v8i8(ptr %a, ptr %b) {
-; CHECK-LABEL: test_ldst_v8i8(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<3>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0];
-; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];
-; CHECK-NEXT: st.v2.b32 [%rd2], {%r1, %r2};
-; CHECK-NEXT: ret;
+; O0-LABEL: test_ldst_v8i8(
+; O0: {
+; O0-NEXT: .reg .b32 %r<3>;
+; O0-NEXT: .reg .b64 %rd<3>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0];
+; O0-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];
+; O0-NEXT: st.v2.b32 [%rd2], {%r1, %r2};
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_ldst_v8i8(
+; O3: {
+; O3-NEXT: .reg .b32 %r<3>;
+; O3-NEXT: .reg .b64 %rd<3>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0];
+; O3-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];
+; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1];
+; O3-NEXT: st.v2.b32 [%rd2], {%r1, %r2};
+; O3-NEXT: ret;
%t1 = load <8 x i8>, ptr %a
store <8 x i8> %t1, ptr %b, align 16
ret void
@@ -803,168 +1270,310 @@ define void @test_ldst_v8i8(ptr %a, ptr %b) {
declare <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) #0
define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_call(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_call_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_call_param_0];
-; CHECK-NEXT: { // callseq 0, 0
-; CHECK-NEXT: .param .align 4 .b8 param0[4];
-; CHECK-NEXT: st.param.b32 [param0], %r1;
-; CHECK-NEXT: .param .align 4 .b8 param1[4];
-; CHECK-NEXT: st.param.b32 [param1], %r2;
-; CHECK-NEXT: .param .align 4 .b8 retval0[4];
-; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);
-; CHECK-NEXT: ld.param.b32 %r3, [retval0];
-; CHECK-NEXT: } // callseq 0
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_call(
+; O0: {
+; O0-NEXT: .reg .b32 %r<5>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_call_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_call_param_0];
+; O0-NEXT: { // callseq 0, 0
+; O0-NEXT: .param .align 4 .b8 param0[4];
+; O0-NEXT: st.param.b32 [param0], %r1;
+; O0-NEXT: .param .align 4 .b8 param1[4];
+; O0-NEXT: st.param.b32 [param1], %r2;
+; O0-NEXT: .param .align 4 .b8 retval0[4];
+; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
+; O0-NEXT: ld.param.b32 %r3, [retval0];
+; O0-NEXT: } // callseq 0
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_call(
+; O3: {
+; O3-NEXT: .reg .b32 %r<5>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_call_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_call_param_1];
+; O3-NEXT: { // callseq 0, 0
+; O3-NEXT: .param .align 4 .b8 param0[4];
+; O3-NEXT: st.param.b32 [param0], %r1;
+; O3-NEXT: .param .align 4 .b8 param1[4];
+; O3-NEXT: st.param.b32 [param1], %r2;
+; O3-NEXT: .param .align 4 .b8 retval0[4];
+; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
+; O3-NEXT: ld.param.b32 %r3, [retval0];
+; O3-NEXT: } // callseq 0
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%r = call <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b)
ret <4 x i8> %r
}
define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_call_flipped(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];
-; CHECK-NEXT: { // callseq 1, 0
-; CHECK-NEXT: .param .align 4 .b8 param0[4];
-; CHECK-NEXT: st.param.b32 [param0], %r2;
-; CHECK-NEXT: .param .align 4 .b8 param1[4];
-; CHECK-NEXT: st.param.b32 [param1], %r1;
-; CHECK-NEXT: .param .align 4 .b8 retval0[4];
-; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);
-; CHECK-NEXT: ld.param.b32 %r3, [retval0];
-; CHECK-NEXT: } // callseq 1
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_call_flipped(
+; O0: {
+; O0-NEXT: .reg .b32 %r<5>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];
+; O0-NEXT: { // callseq 1, 0
+; O0-NEXT: .param .align 4 .b8 param0[4];
+; O0-NEXT: st.param.b32 [param0], %r2;
+; O0-NEXT: .param .align 4 .b8 param1[4];
+; O0-NEXT: st.param.b32 [param1], %r1;
+; O0-NEXT: .param .align 4 .b8 retval0[4];
+; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
+; O0-NEXT: ld.param.b32 %r3, [retval0];
+; O0-NEXT: } // callseq 1
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_call_flipped(
+; O3: {
+; O3-NEXT: .reg .b32 %r<5>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];
+; O3-NEXT: { // callseq 1, 0
+; O3-NEXT: .param .align 4 .b8 param0[4];
+; O3-NEXT: st.param.b32 [param0], %r2;
+; O3-NEXT: .param .align 4 .b8 param1[4];
+; O3-NEXT: st.param.b32 [param1], %r1;
+; O3-NEXT: .param .align 4 .b8 retval0[4];
+; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
+; O3-NEXT: ld.param.b32 %r3, [retval0];
+; O3-NEXT: } // callseq 1
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%r = call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a)
ret <4 x i8> %r
}
define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_tailcall_flipped(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];
-; CHECK-NEXT: { // callseq 2, 0
-; CHECK-NEXT: .param .align 4 .b8 param0[4];
-; CHECK-NEXT: st.param.b32 [param0], %r2;
-; CHECK-NEXT: .param .align 4 .b8 param1[4];
-; CHECK-NEXT: st.param.b32 [param1], %r1;
-; CHECK-NEXT: .param .align 4 .b8 retval0[4];
-; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);
-; CHECK-NEXT: ld.param.b32 %r3, [retval0];
-; CHECK-NEXT: } // callseq 2
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_tailcall_flipped(
+; O0: {
+; O0-NEXT: .reg .b32 %r<5>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];
+; O0-NEXT: { // callseq 2, 0
+; O0-NEXT: .param .align 4 .b8 param0[4];
+; O0-NEXT: st.param.b32 [param0], %r2;
+; O0-NEXT: .param .align 4 .b8 param1[4];
+; O0-NEXT: st.param.b32 [param1], %r1;
+; O0-NEXT: .param .align 4 .b8 retval0[4];
+; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);
+; O0-NEXT: ld.param.b32 %r3, [retval0];
+; O0-NEXT: } // callseq 2
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_tailcall_flipped(
+; O3: {
+; O3-NEXT: .reg .b32 %r<5>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];
+; O3-NEXT: { // callseq 2, 0
+; O3-NEXT: .param .align 4 .b8 param0[4];
+; O3-NEXT: st.param.b32 [param0], %r2;
+; O3-NEXT: .param .align 4 .b8 param1[4];
+; O3-NEXT: st.param.b32 [param1], %r1;
+; O3-NEXT: .param .align 4 .b8 retval0[4];
+; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);
+; O3-NEXT: ld.param.b32 %r3, [retval0];
+; O3-NEXT: } // callseq 2
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%r = tail call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a)
ret <4 x i8> %r
}
define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 {
-; CHECK-LABEL: test_select(
-; CHECK: {
-; CHECK-NEXT: .reg .pred %p<2>;
-; CHECK-NEXT: .reg .b16 %rs<3>;
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b8 %rs1, [test_select_param_2];
-; CHECK-NEXT: and.b16 %rs2, %rs1, 1;
-; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
-; CHECK-NEXT: ld.param.b32 %r2, [test_select_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_select_param_0];
-; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_select(
+; O0: {
+; O0-NEXT: .reg .pred %p<2>;
+; O0-NEXT: .reg .b16 %rs<3>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b8 %rs1, [test_select_param_2];
+; O0-NEXT: and.b16 %rs2, %rs1, 1;
+; O0-NEXT: setp.ne.b16 %p1, %rs2, 0;
+; O0-NEXT: ld.param.b32 %r2, [test_select_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_select_param_0];
+; O0-NEXT: selp.b32 %r3, %r1, %r2, %p1;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_select(
+; O3: {
+; O3-NEXT: .reg .pred %p<2>;
+; O3-NEXT: .reg .b16 %rs<3>;
+; O3-NEXT: .reg .b32 %r<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b8 %rs1, [test_select_param_2];
+; O3-NEXT: and.b16 %rs2, %rs1, 1;
+; O3-NEXT: setp.ne.b16 %p1, %rs2, 0;
+; O3-NEXT: ld.param.b32 %r1, [test_select_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_select_param_1];
+; O3-NEXT: selp.b32 %r3, %r1, %r2, %p1;
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%r = select i1 %c, <4 x i8> %a, <4 x i8> %b
ret <4 x i8> %r
}
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<28>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];
-; CHECK-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];
-; CHECK-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];
-; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r6, %r3, 0, 0x7770U;
-; CHECK-NEXT: setp.ne.b32 %p1, %r6, %r5;
-; CHECK-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r8, %r3, 0, 0x7771U;
-; CHECK-NEXT: setp.ne.b32 %p2, %r8, %r7;
-; CHECK-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r10, %r3, 0, 0x7772U;
-; CHECK-NEXT: setp.ne.b32 %p3, %r10, %r9;
-; CHECK-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;
-; CHECK-NEXT: setp.ne.b32 %p4, %r12, %r11;
-; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;
-; CHECK-NEXT: selp.b32 %r15, %r14, %r13, %p4;
-; CHECK-NEXT: prmt.b32 %r16, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;
-; CHECK-NEXT: selp.b32 %r18, %r17, %r16, %p3;
-; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r20, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U;
-; CHECK-NEXT: selp.b32 %r22, %r21, %r20, %p2;
-; CHECK-NEXT: prmt.b32 %r23, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U;
-; CHECK-NEXT: selp.b32 %r25, %r24, %r23, %p1;
-; CHECK-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r27;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_select_cc(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<28>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];
+; O0-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];
+; O0-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];
+; O0-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r6, %r3, 0, 0x7770U;
+; O0-NEXT: setp.ne.b32 %p1, %r6, %r5;
+; O0-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r8, %r3, 0, 0x7771U;
+; O0-NEXT: setp.ne.b32 %p2, %r8, %r7;
+; O0-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r10, %r3, 0, 0x7772U;
+; O0-NEXT: setp.ne.b32 %p3, %r10, %r9;
+; O0-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;
+; O0-NEXT: setp.ne.b32 %p4, %r12, %r11;
+; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;
+; O0-NEXT: selp.b32 %r15, %r14, %r13, %p4;
+; O0-NEXT: prmt.b32 %r16, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;
+; O0-NEXT: selp.b32 %r18, %r17, %r16, %p3;
+; O0-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U;
+; O0-NEXT: prmt.b32 %r20, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U;
+; O0-NEXT: selp.b32 %r22, %r21, %r20, %p2;
+; O0-NEXT: prmt.b32 %r23, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U;
+; O0-NEXT: selp.b32 %r25, %r24, %r23, %p1;
+; O0-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U;
+; O0-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r27;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_select_cc(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<28>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_select_cc_param_3];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: ld.param.b32 %r4, [test_select_cc_param_2];
+; O3-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U;
+; O3-NEXT: setp.ne.b32 %p1, %r5, %r3;
+; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U;
+; O3-NEXT: setp.ne.b32 %p2, %r7, %r6;
+; O3-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U;
+; O3-NEXT: setp.ne.b32 %p3, %r9, %r8;
+; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U;
+; O3-NEXT: setp.ne.b32 %p4, %r11, %r10;
+; O3-NEXT: ld.param.b32 %r12, [test_select_cc_param_1];
+; O3-NEXT: prmt.b32 %r13, %r12, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;
+; O3-NEXT: selp.b32 %r15, %r14, %r13, %p4;
+; O3-NEXT: prmt.b32 %r16, %r12, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;
+; O3-NEXT: selp.b32 %r18, %r17, %r16, %p3;
+; O3-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U;
+; O3-NEXT: prmt.b32 %r20, %r12, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U;
+; O3-NEXT: selp.b32 %r22, %r21, %r20, %p2;
+; O3-NEXT: prmt.b32 %r23, %r12, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U;
+; O3-NEXT: selp.b32 %r25, %r24, %r23, %p1;
+; O3-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U;
+; O3-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r27;
+; O3-NEXT: ret;
%cc = icmp ne <4 x i8> %c, %d
%r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b
ret <4 x i8> %r
}
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-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1];
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0];
-; CHECK-NEXT: ld.param.b32 %r10, [test_select_cc_i32_i8_param_3];
-; CHECK-NEXT: ld.param.b32 %r9, [test_select_cc_i32_i8_param_2];
-; CHECK-NEXT: prmt.b32 %r11, %r10, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r12, %r9, 0, 0x7770U;
-; CHECK-NEXT: setp.ne.b32 %p1, %r12, %r11;
-; CHECK-NEXT: prmt.b32 %r13, %r10, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r14, %r9, 0, 0x7771U;
-; CHECK-NEXT: setp.ne.b32 %p2, %r14, %r13;
-; CHECK-NEXT: prmt.b32 %r15, %r10, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r16, %r9, 0, 0x7772U;
-; CHECK-NEXT: setp.ne.b32 %p3, %r16, %r15;
-; CHECK-NEXT: prmt.b32 %r17, %r10, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r18, %r9, 0, 0x7773U;
-; CHECK-NEXT: setp.ne.b32 %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], {%r22, %r21, %r20, %r19};
-; CHECK-NEXT: ret;
+; O0-LABEL: test_select_cc_i32_i8(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<23>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1];
+; O0-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0];
+; O0-NEXT: ld.param.b32 %r10, [test_select_cc_i32_i8_param_3];
+; O0-NEXT: ld.param.b32 %r9, [test_select_cc_i32_i8_param_2];
+; O0-NEXT: prmt.b32 %r11, %r10, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r12, %r9, 0, 0x7770U;
+; O0-NEXT: setp.ne.b32 %p1, %r12, %r11;
+; O0-NEXT: prmt.b32 %r13, %r10, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r14, %r9, 0, 0x7771U;
+; O0-NEXT: setp.ne.b32 %p2, %r14, %r13;
+; O0-NEXT: prmt.b32 %r15, %r10, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r16, %r9, 0, 0x7772U;
+; O0-NEXT: setp.ne.b32 %p3, %r16, %r15;
+; O0-NEXT: prmt.b32 %r17, %r10, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r18, %r9, 0, 0x7773U;
+; O0-NEXT: setp.ne.b32 %p4, %r18, %r17;
+; O0-NEXT: selp.b32 %r19, %r4, %r8, %p4;
+; O0-NEXT: selp.b32 %r20, %r3, %r7, %p3;
+; O0-NEXT: selp.b32 %r21, %r2, %r6, %p2;
+; O0-NEXT: selp.b32 %r22, %r1, %r5, %p1;
+; O0-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19};
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_select_cc_i32_i8(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<23>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0];
+; O3-NEXT: ld.param.b32 %r5, [test_select_cc_i32_i8_param_3];
+; O3-NEXT: prmt.b32 %r6, %r5, 0, 0x7770U;
+; O3-NEXT: ld.param.b32 %r7, [test_select_cc_i32_i8_param_2];
+; O3-NEXT: prmt.b32 %r8, %r7, 0, 0x7770U;
+; O3-NEXT: setp.ne.b32 %p1, %r8, %r6;
+; O3-NEXT: prmt.b32 %r9, %r5, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r10, %r7, 0, 0x7771U;
+; O3-NEXT: setp.ne.b32 %p2, %r10, %r9;
+; O3-NEXT: prmt.b32 %r11, %r5, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r12, %r7, 0, 0x7772U;
+; O3-NEXT: setp.ne.b32 %p3, %r12, %r11;
+; O3-NEXT: prmt.b32 %r13, %r5, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r14, %r7, 0, 0x7773U;
+; O3-NEXT: setp.ne.b32 %p4, %r14, %r13;
+; O3-NEXT: ld.param.v4.b32 {%r15, %r16, %r17, %r18}, [test_select_cc_i32_i8_param_1];
+; O3-NEXT: selp.b32 %r19, %r4, %r18, %p4;
+; O3-NEXT: selp.b32 %r20, %r3, %r17, %p3;
+; O3-NEXT: selp.b32 %r21, %r2, %r16, %p2;
+; O3-NEXT: selp.b32 %r22, %r1, %r15, %p1;
+; O3-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19};
+; O3-NEXT: ret;
<4 x i8> %c, <4 x i8> %d) #0 {
%cc = icmp ne <4 x i8> %c, %d
%r = select <4 x i1> %cc, <4 x i32> %a, <4 x i32> %b
@@ -972,37 +1581,69 @@ define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b,
}
define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b,
-; CHECK-LABEL: test_select_cc_i8_i32(
-; CHECK: {
-; CHECK-NEXT: .reg .pred %p<5>;
-; CHECK-NEXT: .reg .b32 %r<26>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r7, %r8, %r9, %r10}, [test_select_cc_i8_i32_param_3];
-; CHECK-NEXT: ld.param.v4.b32 {%r3, %r4, %r5, %r6}, [test_select_cc_i8_i32_param_2];
-; CHECK-NEXT: ld.param.b32 %r2, [test_select_cc_i8_i32_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0];
-; CHECK-NEXT: setp.ne.b32 %p1, %r3, %r7;
-; CHECK-NEXT: setp.ne.b32 %p2, %r4, %r8;
-; CHECK-NEXT: setp.ne.b32 %p3, %r5, %r9;
-; CHECK-NEXT: setp.ne.b32 %p4, %r6, %r10;
-; CHECK-NEXT: prmt.b32 %r11, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U;
-; CHECK-NEXT: selp.b32 %r13, %r12, %r11, %p4;
-; CHECK-NEXT: prmt.b32 %r14, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U;
-; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p3;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r18, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;
-; CHECK-NEXT: selp.b32 %r20, %r19, %r18, %p2;
-; CHECK-NEXT: prmt.b32 %r21, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;
-; CHECK-NEXT: selp.b32 %r23, %r22, %r21, %p1;
-; CHECK-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r25;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_select_cc_i8_i32(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<26>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.v4.b32 {%r7, %r8, %r9, %r10}, [test_select_cc_i8_i32_param_3];
+; O0-NEXT: ld.param.v4.b32 {%r3, %r4, %r5, %r6}, [test_select_cc_i8_i32_param_2];
+; O0-NEXT: ld.param.b32 %r2, [test_select_cc_i8_i32_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0];
+; O0-NEXT: setp.ne.b32 %p1, %r3, %r7;
+; O0-NEXT: setp.ne.b32 %p2, %r4, %r8;
+; O0-NEXT: setp.ne.b32 %p3, %r5, %r9;
+; O0-NEXT: setp.ne.b32 %p4, %r6, %r10;
+; O0-NEXT: prmt.b32 %r11, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U;
+; O0-NEXT: selp.b32 %r13, %r12, %r11, %p4;
+; O0-NEXT: prmt.b32 %r14, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U;
+; O0-NEXT: selp.b32 %r16, %r15, %r14, %p3;
+; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U;
+; O0-NEXT: prmt.b32 %r18, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;
+; O0-NEXT: selp.b32 %r20, %r19, %r18, %p2;
+; O0-NEXT: prmt.b32 %r21, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;
+; O0-NEXT: selp.b32 %r23, %r22, %r21, %p1;
+; O0-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U;
+; O0-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r25;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_select_cc_i8_i32(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<26>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0];
+; O3-NEXT: ld.param.v4.b32 {%r2, %r3, %r4, %r5}, [test_select_cc_i8_i32_param_2];
+; O3-NEXT: ld.param.v4.b32 {%r6, %r7, %r8, %r9}, [test_select_cc_i8_i32_param_3];
+; O3-NEXT: setp.ne.b32 %p1, %r2, %r6;
+; O3-NEXT: setp.ne.b32 %p2, %r3, %r7;
+; O3-NEXT: setp.ne.b32 %p3, %r4, %r8;
+; O3-NEXT: setp.ne.b32 %p4, %r5, %r9;
+; O3-NEXT: ld.param.b32 %r10, [test_select_cc_i8_i32_param_1];
+; O3-NEXT: prmt.b32 %r11, %r10, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U;
+; O3-NEXT: selp.b32 %r13, %r12, %r11, %p4;
+; O3-NEXT: prmt.b32 %r14, %r10, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U;
+; O3-NEXT: selp.b32 %r16, %r15, %r14, %p3;
+; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U;
+; O3-NEXT: prmt.b32 %r18, %r10, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;
+; O3-NEXT: selp.b32 %r20, %r19, %r18, %p2;
+; O3-NEXT: prmt.b32 %r21, %r10, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;
+; O3-NEXT: selp.b32 %r23, %r22, %r21, %p1;
+; O3-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U;
+; O3-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r25;
+; O3-NEXT: ret;
<4 x i32> %c, <4 x i32> %d) #0 {
%cc = icmp ne <4 x i32> %c, %d
%r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b
@@ -1027,23 +1668,41 @@ define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 {
}
define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 {
-; CHECK-LABEL: test_trunc_2xi64(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<8>;
-; CHECK-NEXT: .reg .b64 %rd<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16];
-; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];
-; CHECK-NEXT: cvt.u32.u64 %r1, %rd4;
-; CHECK-NEXT: cvt.u32.u64 %r2, %rd3;
-; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;
-; CHECK-NEXT: cvt.u32.u64 %r4, %rd2;
-; CHECK-NEXT: cvt.u32.u64 %r5, %rd1;
-; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_trunc_2xi64(
+; O0: {
+; O0-NEXT: .reg .b32 %r<8>;
+; O0-NEXT: .reg .b64 %rd<5>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16];
+; O0-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];
+; O0-NEXT: cvt.u32.u64 %r1, %rd4;
+; O0-NEXT: cvt.u32.u64 %r2, %rd3;
+; O0-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;
+; O0-NEXT: cvt.u32.u64 %r4, %rd2;
+; O0-NEXT: cvt.u32.u64 %r5, %rd1;
+; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
+; O0-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r7;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_trunc_2xi64(
+; O3: {
+; O3-NEXT: .reg .b32 %r<8>;
+; O3-NEXT: .reg .b64 %rd<5>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];
+; O3-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16];
+; O3-NEXT: cvt.u32.u64 %r1, %rd4;
+; O3-NEXT: cvt.u32.u64 %r2, %rd3;
+; O3-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;
+; O3-NEXT: cvt.u32.u64 %r4, %rd2;
+; O3-NEXT: cvt.u32.u64 %r5, %rd1;
+; O3-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
+; O3-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r7;
+; O3-NEXT: ret;
%r = trunc <4 x i64> %a to <4 x i8>
ret <4 x i8> %r
}
@@ -1066,24 +1725,43 @@ define <4 x i32> @test_zext_2xi32(<4 x i8> %a) #0 {
}
define <4 x i64> @test_zext_2xi64(<4 x i8> %a) #0 {
-; CHECK-LABEL: test_zext_2xi64(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<6>;
-; CHECK-NEXT: .reg .b64 %rd<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0];
-; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;
-; CHECK-NEXT: cvt.u64.u32 %rd1, %r2;
-; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0x7772U;
-; CHECK-NEXT: cvt.u64.u32 %rd2, %r3;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7771U;
-; CHECK-NEXT: cvt.u64.u32 %rd3, %r4;
-; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;
-; CHECK-NEXT: cvt.u64.u32 %rd4, %r5;
-; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd3};
-; CHECK-NEXT: st.param.v2.b64 [func_retval0+16], {%rd2, %rd1};
-; CHECK-NEXT: ret;
+; O0-LABEL: test_zext_2xi64(
+; O0: {
+; O0-NEXT: .reg .b32 %r<6>;
+; O0-NEXT: .reg .b64 %rd<5>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0];
+; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;
+; O0-NEXT: cvt.u64.u32 %rd1, %r2;
+; O0-NEXT: prmt.b32 %r3, %r1, 0, 0x7772U;
+; O0-NEXT: cvt.u64.u32 %rd2, %r3;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7771U;
+; O0-NEXT: cvt.u64.u32 %rd3, %r4;
+; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;
+; O0-NEXT: cvt.u64.u32 %rd4, %r5;
+; O0-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd3};
+; O0-NEXT: st.param.v2.b64 [func_retval0+16], {%rd2, %rd1};
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_zext_2xi64(
+; O3: {
+; O3-NEXT: .reg .b32 %r<6>;
+; O3-NEXT: .reg .b64 %rd<5>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0];
+; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x7771U;
+; O3-NEXT: cvt.u64.u32 %rd1, %r2;
+; O3-NEXT: prmt.b32 %r3, %r1, 0, 0x7770U;
+; O3-NEXT: cvt.u64.u32 %rd2, %r3;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;
+; O3-NEXT: cvt.u64.u32 %rd3, %r4;
+; O3-NEXT: prmt.b32 %r5, %r1, 0, 0x7772U;
+; O3-NEXT: cvt.u64.u32 %rd4, %r5;
+; O3-NEXT: st.param.v2.b64 [func_retval0+16], {%rd4, %rd3};
+; O3-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
+; O3-NEXT: ret;
%r = zext <4 x i8> %a to <4 x i64>
ret <4 x i64> %r
}
@@ -1142,20 +1820,33 @@ define float @test_bitcast_4xi8_to_float(<4 x i8> %a) #0 {
define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 {
-; CHECK-LABEL: test_bitcast_4xi8_to_2xhalf(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<6>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0];
-; CHECK-NEXT: mov.b32 %r1, 6;
-; CHECK-NEXT: prmt.b32 %r2, %r1, 7, 0x3340U;
-; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
-; CHECK-NEXT: prmt.b32 %r4, %r3, 5, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_bitcast_4xi8_to_2xhalf(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<2>;
+; O0-NEXT: .reg .b32 %r<6>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0];
+; O0-NEXT: mov.b32 %r1, 6;
+; O0-NEXT: prmt.b32 %r2, %r1, 7, 0x3340U;
+; O0-NEXT: cvt.u32.u16 %r3, %rs1;
+; O0-NEXT: prmt.b32 %r4, %r3, 5, 0x3340U;
+; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r5;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_bitcast_4xi8_to_2xhalf(
+; O3: {
+; O3-NEXT: .reg .b32 %r<6>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: mov.b32 %r1, 6;
+; O3-NEXT: prmt.b32 %r2, %r1, 7, 0x3340U;
+; O3-NEXT: ld.param.b8 %r3, [test_bitcast_4xi8_to_2xhalf_param_0];
+; O3-NEXT: prmt.b32 %r4, %r3, 5, 0x3340U;
+; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r5;
+; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> undef, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
%ins.2 = insertelement <4 x i8> %ins.1, i8 6, i32 2
@@ -1166,153 +1857,277 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 {
define <4 x i8> @test_shufflevector(<4 x i8> %a) #0 {
-; CHECK-LABEL: test_shufflevector(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];
-; CHECK-NEXT: // implicit-def: %r3
-; CHECK-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_shufflevector(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];
+; O0-NEXT: // implicit-def: %r3
+; O0-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U;
+; O0-NEXT: st.param.b32 [func_retval0], %r2;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_shufflevector(
+; O3: {
+; O3-NEXT: .reg .b32 %r<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];
+; O3-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U;
+; O3-NEXT: st.param.b32 [func_retval0], %r2;
+; O3-NEXT: ret;
%s = shufflevector <4 x i8> %a, <4 x i8> undef, <4 x i32> <i32 3, i32 2, i32 1, i32 0>
ret <4 x i8> %s
}
define <4 x i8> @test_shufflevector_2(<4 x i8> %a, <4 x i8> %b) #0 {
-; CHECK-LABEL: test_shufflevector_2(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0];
-; CHECK-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_shufflevector_2(
+; O0: {
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0];
+; O0-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_shufflevector_2(
+; O3: {
+; O3-NEXT: .reg .b32 %r<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0];
+; O3-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1];
+; O3-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U;
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%s = shufflevector <4 x i8> %a, <4 x i8> %b, <4 x i32> <i32 7, i32 3, i32 5, i32 2>
ret <4 x i8> %s
}
define <4 x i8> @test_insertelement(<4 x i8> %a, i8 %x) #0 {
-; CHECK-LABEL: test_insertelement(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<2>;
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b8 %rs1, [test_insertelement_param_1];
-; CHECK-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];
-; CHECK-NEXT: cvt.u32.u16 %r2, %rs1;
-; CHECK-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_insertelement(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<2>;
+; O0-NEXT: .reg .b32 %r<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.b8 %rs1, [test_insertelement_param_1];
+; O0-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];
+; O0-NEXT: cvt.u32.u16 %r2, %rs1;
+; O0-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8;
+; O0-NEXT: st.param.b32 [func_retval0], %r3;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_insertelement(
+; O3: {
+; O3-NEXT: .reg .b32 %r<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];
+; O3-NEXT: ld.param.b8 %r2, [test_insertelement_param_1];
+; O3-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8;
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
+; O3-NEXT: ret;
%i = insertelement <4 x i8> %a, i8 %x, i64 1
ret <4 x i8> %i
}
define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 {
-; CHECK-LABEL: test_fptosi_4xhalf_to_4xi8(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<13>;
-; CHECK-NEXT: .reg .b32 %r<12>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0];
-; CHECK-NEXT: cvt.rzi.s16.f16 %rs5, %rs4;
-; CHECK-NEXT: cvt.rzi.s16.f16 %rs6, %rs3;
-; CHECK-NEXT: mov.b32 %r3, {%rs6, %rs5};
-; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3;
-; CHECK-NEXT: cvt.u32.u16 %r4, %rs8;
-; CHECK-NEXT: cvt.u32.u16 %r5, %rs7;
-; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
-; CHECK-NEXT: cvt.rzi.s16.f16 %rs9, %rs2;
-; CHECK-NEXT: cvt.rzi.s16.f16 %rs10, %rs1;
-; CHECK-NEXT: mov.b32 %r7, {%rs10, %rs9};
-; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r7;
-; CHECK-NEXT: cvt.u32.u16 %r8, %rs12;
-; CHECK-NEXT: cvt.u32.u16 %r9, %rs11;
-; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r11;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_fptosi_4xhalf_to_4xi8(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<13>;
+; O0-NEXT: .reg .b32 %r<12>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0];
+; O0-NEXT: cvt.rzi.s16.f16 %rs5, %rs4;
+; O0-NEXT: cvt.rzi.s16.f16 %rs6, %rs3;
+; O0-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; O0-NEXT: mov.b32 {%rs7, %rs8}, %r3;
+; O0-NEXT: cvt.u32.u16 %r4, %rs8;
+; O0-NEXT: cvt.u32.u16 %r5, %rs7;
+; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
+; O0-NEXT: cvt.rzi.s16.f16 %rs9, %rs2;
+; O0-NEXT: cvt.rzi.s16.f16 %rs10, %rs1;
+; O0-NEXT: mov.b32 %r7, {%rs10, %rs9};
+; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7;
+; O0-NEXT: cvt.u32.u16 %r8, %rs12;
+; O0-NEXT: cvt.u32.u16 %r9, %rs11;
+; O0-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U;
+; O0-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r11;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_fptosi_4xhalf_to_4xi8(
+; O3: {
+; O3-NEXT: .reg .b16 %rs<13>;
+; O3-NEXT: .reg .b32 %r<10>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0];
+; O3-NEXT: cvt.rzi.s16.f16 %rs5, %rs4;
+; O3-NEXT: cvt.rzi.s16.f16 %rs6, %rs3;
+; O3-NEXT: mov.b32 %r1, {%rs6, %rs5};
+; O3-NEXT: mov.b32 {%rs7, %rs8}, %r1;
+; O3-NEXT: cvt.u32.u16 %r2, %rs8;
+; O3-NEXT: cvt.u32.u16 %r3, %rs7;
+; O3-NEXT: prmt.b32 %r4, %r3, %r2, 0x3340U;
+; O3-NEXT: cvt.rzi.s16.f16 %rs9, %rs2;
+; O3-NEXT: cvt.rzi.s16.f16 %rs10, %rs1;
+; O3-NEXT: mov.b32 %r5, {%rs10, %rs9};
+; O3-NEXT: mov.b32 {%rs11, %rs12}, %r5;
+; O3-NEXT: cvt.u32.u16 %r6, %rs12;
+; O3-NEXT: cvt.u32.u16 %r7, %rs11;
+; O3-NEXT: prmt.b32 %r8, %r7, %r6, 0x3340U;
+; O3-NEXT: prmt.b32 %r9, %r8, %r4, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r9;
+; O3-NEXT: ret;
%r = fptosi <4 x half> %a to <4 x i8>
ret <4 x i8> %r
}
define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 {
-; CHECK-LABEL: test_fptoui_4xhalf_to_4xi8(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<13>;
-; CHECK-NEXT: .reg .b32 %r<12>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0];
-; CHECK-NEXT: cvt.rzi.u16.f16 %rs5, %rs4;
-; CHECK-NEXT: cvt.rzi.u16.f16 %rs6, %rs3;
-; CHECK-NEXT: mov.b32 %r3, {%rs6, %rs5};
-; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3;
-; CHECK-NEXT: cvt.u32.u16 %r4, %rs8;
-; CHECK-NEXT: cvt.u32.u16 %r5, %rs7;
-; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
-; CHECK-NEXT: cvt.rzi.u16.f16 %rs9, %rs2;
-; CHECK-NEXT: cvt.rzi.u16.f16 %rs10, %rs1;
-; CHECK-NEXT: mov.b32 %r7, {%rs10, %rs9};
-; CHECK-NEXT: mov.b32 {%rs11, %rs12}, %r7;
-; CHECK-NEXT: cvt.u32.u16 %r8, %rs12;
-; CHECK-NEXT: cvt.u32.u16 %r9, %rs11;
-; CHECK-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r11;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_fptoui_4xhalf_to_4xi8(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<13>;
+; O0-NEXT: .reg .b32 %r<12>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0:
+; O0-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0];
+; O0-NEXT: cvt.rzi.u16.f16 %rs5, %rs4;
+; O0-NEXT: cvt.rzi.u16.f16 %rs6, %rs3;
+; O0-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; O0-NEXT: mov.b32 {%rs7, %rs8}, %r3;
+; O0-NEXT: cvt.u32.u16 %r4, %rs8;
+; O0-NEXT: cvt.u32.u16 %r5, %rs7;
+; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
+; O0-NEXT: cvt.rzi.u16.f16 %rs9, %rs2;
+; O0-NEXT: cvt.rzi.u16.f16 %rs10, %rs1;
+; O0-NEXT: mov.b32 %r7, {%rs10, %rs9};
+; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7;
+; O0-NEXT: cvt.u32.u16 %r8, %rs12;
+; O0-NEXT: cvt.u32.u16 %r9, %rs11;
+; O0-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U;
+; O0-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U;
+; O0-NEXT: st.param.b32 [func_retval0], %r11;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_fptoui_4xhalf_to_4xi8(
+; O3: {
+; O3-NEXT: .reg .b16 %rs<13>;
+; O3-NEXT: .reg .b32 %r<10>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0:
+; O3-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0];
+; O3-NEXT: cvt.rzi.u16.f16 %rs5, %rs4;
+; O3-NEXT: cvt.rzi.u16.f16 %rs6, %rs3;
+; O3-NEXT: mov.b32 %r1, {%rs6, %rs5};
+; O3-NEXT: mov.b32 {%rs7, %rs8}, %r1;
+; O3-NEXT: cvt.u32.u16 %r2, %rs8;
+; O3-NEXT: cvt.u32.u16 %r3, %rs7;
+; O3-NEXT: prmt.b32 %r4, %r3, %r2, 0x3340U;
+; O3-NEXT: cvt.rzi.u16.f16 %rs9, %rs2;
+; O3-NEXT: cvt.rzi.u16.f16 %rs10, %rs1;
+; O3-NEXT: mov.b32 %r5, {%rs10, %rs9};
+; O3-NEXT: mov.b32 {%rs11, %rs12}, %r5;
+; O3-NEXT: cvt.u32.u16 %r6, %rs12;
+; O3-NEXT: cvt.u32.u16 %r7, %rs11;
+; O3-NEXT: prmt.b32 %r8, %r7, %r6, 0x3340U;
+; O3-NEXT: prmt.b32 %r9, %r8, %r4, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r9;
+; O3-NEXT: ret;
%r = fptoui <4 x half> %a to <4 x i8>
ret <4 x i8> %r
}
define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) {
-; CHECK-LABEL: test_srem_v4i8(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<13>;
-; CHECK-NEXT: .reg .b32 %r<18>;
-; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0: // %entry
-; CHECK-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2];
-; CHECK-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0];
-; CHECK-NEXT: ld.b32 %r1, [%rd1];
-; CHECK-NEXT: ld.b32 %r2, [%rd2];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U;
-; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U;
-; CHECK-NEXT: cvt.u16.u32 %rs2, %r4;
-; CHECK-NEXT: rem.s16 %rs3, %rs2, %rs1;
-; CHECK-NEXT: cvt.u32.u16 %r5, %rs3;
-; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U;
-; CHECK-NEXT: cvt.u16.u32 %rs4, %r6;
-; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U;
-; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
-; CHECK-NEXT: rem.s16 %rs6, %rs5, %rs4;
-; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
-; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x9991U;
-; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;
-; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x9991U;
-; CHECK-NEXT: cvt.u16.u32 %rs8, %r11;
-; CHECK-NEXT: rem.s16 %rs9, %rs8, %rs7;
-; CHECK-NEXT: cvt.u32.u16 %r12, %rs9;
-; CHECK-NEXT: prmt.b32 %r13, %r2, 0, 0x8880U;
-; CHECK-NEXT: cvt.u16.u32 %rs10, %r13;
-; CHECK-NEXT: prmt.b32 %r14, %r1, 0, 0x8880U;
-; CHECK-NEXT: cvt.u16.u32 %rs11, %r14;
-; CHECK-NEXT: rem.s16 %rs12, %rs11, %rs10;
-; CHECK-NEXT: cvt.u32.u16 %r15, %rs12;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
-; CHECK-NEXT: st.b32 [%rd3], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_srem_v4i8(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<13>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0: // %entry
+; O0-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2];
+; O0-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0];
+; O0-NEXT: ld.b32 %r1, [%rd1];
+; O0-NEXT: ld.b32 %r2, [%rd2];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U;
+; O0-NEXT: cvt.u16.u32 %rs1, %r3;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U;
+; O0-NEXT: cvt.u16.u32 %rs2, %r4;
+; O0-NEXT: rem.s16 %rs3, %rs2, %rs1;
+; O0-NEXT: cvt.u32.u16 %r5, %rs3;
+; O0-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U;
+; O0-NEXT: cvt.u16.u32 %rs4, %r6;
+; O0-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U;
+; O0-NEXT: cvt.u16.u32 %rs5, %r7;
+; O0-NEXT: rem.s16 %rs6, %rs5, %rs4;
+; O0-NEXT: cvt.u32.u16 %r8, %rs6;
+; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x9991U;
+; O0-NEXT: cvt.u16.u32 %rs7, %r10;
+; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x9991U;
+; O0-NEXT: cvt.u16.u32 %rs8, %r11;
+; O0-NEXT: rem.s16 %rs9, %rs8, %rs7;
+; O0-NEXT: cvt.u32.u16 %r12, %rs9;
+; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x8880U;
+; O0-NEXT: cvt.u16.u32 %rs10, %r13;
+; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x8880U;
+; O0-NEXT: cvt.u16.u32 %rs11, %r14;
+; O0-NEXT: rem.s16 %rs12, %rs11, %rs10;
+; O0-NEXT: cvt.u32.u16 %r15, %rs12;
+; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O0-NEXT: st.b32 [%rd3], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_srem_v4i8(
+; O3: {
+; O3-NEXT: .reg .b16 %rs<13>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-NEXT: .reg .b64 %rd<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0: // %entry
+; O3-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0];
+; O3-NEXT: ld.b32 %r1, [%rd1];
+; O3-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1];
+; O3-NEXT: ld.b32 %r2, [%rd2];
+; O3-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U;
+; O3-NEXT: cvt.u16.u32 %rs1, %r3;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U;
+; O3-NEXT: cvt.u16.u32 %rs2, %r4;
+; O3-NEXT: rem.s16 %rs3, %rs2, %rs1;
+; O3-NEXT: cvt.u32.u16 %r5, %rs3;
+; O3-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U;
+; O3-NEXT: cvt.u16.u32 %rs4, %r6;
+; O3-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U;
+; O3-NEXT: cvt.u16.u32 %rs5, %r7;
+; O3-NEXT: rem.s16 %rs6, %rs5, %rs4;
+; O3-NEXT: cvt.u32.u16 %r8, %rs6;
+; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x9991U;
+; O3-NEXT: cvt.u16.u32 %rs7, %r10;
+; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x9991U;
+; O3-NEXT: cvt.u16.u32 %rs8, %r11;
+; O3-NEXT: rem.s16 %rs9, %rs8, %rs7;
+; O3-NEXT: cvt.u32.u16 %r12, %rs9;
+; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x8880U;
+; O3-NEXT: cvt.u16.u32 %rs10, %r13;
+; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x8880U;
+; O3-NEXT: cvt.u16.u32 %rs11, %r14;
+; O3-NEXT: rem.s16 %rs12, %rs11, %rs10;
+; O3-NEXT: cvt.u32.u16 %r15, %rs12;
+; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;
+; O3-NEXT: st.b32 [%rd3], %r17;
+; O3-NEXT: ret;
entry:
%t57 = load <4 x i8>, ptr %a, align 4
%t59 = load <4 x i8>, ptr %b, align 4
@@ -1328,52 +2143,97 @@ entry:
;; Ideally we want to split it into element-wise ops, but legalizer can't handle
;; odd-sized vectors. TL;DR; don't use odd-sized vectors of v8.
define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) {
-; CHECK-LABEL: test_srem_v3i8(
-; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<20>;
-; CHECK-NEXT: .reg .b32 %r<14>;
-; CHECK-NEXT: .reg .b64 %rd<4>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0: // %entry
-; CHECK-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2];
-; CHECK-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0];
-; CHECK-NEXT: ld.b8 %rs1, [%rd1];
-; CHECK-NEXT: ld.b8 %rs2, [%rd1+1];
-; CHECK-NEXT: shl.b16 %rs3, %rs2, 8;
-; CHECK-NEXT: or.b16 %rs4, %rs3, %rs1;
-; CHECK-NEXT: cvt.u32.u16 %r1, %rs4;
-; CHECK-NEXT: ld.s8 %rs5, [%rd1+2];
-; CHECK-NEXT: ld.b8 %rs6, [%rd2];
-; CHECK-NEXT: ld.b8 %rs7, [%rd2+1];
-; CHECK-NEXT: shl.b16 %rs8, %rs7, 8;
-; CHECK-NEXT: or.b16 %rs9, %rs8, %rs6;
-; CHECK-NEXT: cvt.u32.u16 %r2, %rs9;
-; CHECK-NEXT: ld.s8 %rs10, [%rd2+2];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x9991U;
-; CHECK-NEXT: cvt.u16.u32 %rs11, %r3;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x9991U;
-; CHECK-NEXT: cvt.u16.u32 %rs12, %r4;
-; CHECK-NEXT: rem.s16 %rs13, %rs12, %rs11;
-; CHECK-NEXT: cvt.u32.u16 %r5, %rs13;
-; CHECK-NEXT: prmt.b32 %r6, %r2, 0, 0x8880U;
-; CHECK-NEXT: cvt.u16.u32 %rs14, %r6;
-; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x8880U;
-; CHECK-NEXT: cvt.u16.u32 %rs15, %r7;
-; CHECK-NEXT: rem.s16 %rs16, %rs15, %rs14;
-; CHECK-NEXT: cvt.u32.u16 %r8, %rs16;
-; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
-; CHECK-NEXT: // implicit-def: %r11
-; CHECK-NEXT: // implicit-def: %r12
-; CHECK-NEXT: prmt.b32 %r10, %r11, %r12, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U;
-; CHECK-NEXT: rem.s16 %rs17, %rs5, %rs10;
-; CHECK-NEXT: mov.b32 {%rs18, _}, %r13;
-; CHECK-NEXT: st.b8 [%rd3], %rs18;
-; CHECK-NEXT: shr.u16 %rs19, %rs18, 8;
-; CHECK-NEXT: st.b8 [%rd3+1], %rs19;
-; CHECK-NEXT: st.b8 [%rd3+2], %rs17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_srem_v3i8(
+; O0: {
+; O0-NEXT: .reg .b16 %rs<20>;
+; O0-NEXT: .reg .b32 %r<14>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0: // %entry
+; O0-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2];
+; O0-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0];
+; O0-NEXT: ld.b8 %rs1, [%rd1];
+; O0-NEXT: ld.b8 %rs2, [%rd1+1];
+; O0-NEXT: shl.b16 %rs3, %rs2, 8;
+; O0-NEXT: or.b16 %rs4, %rs3, %rs1;
+; O0-NEXT: cvt.u32.u16 %r1, %rs4;
+; O0-NEXT: ld.s8 %rs5, [%rd1+2];
+; O0-NEXT: ld.b8 %rs6, [%rd2];
+; O0-NEXT: ld.b8 %rs7, [%rd2+1];
+; O0-NEXT: shl.b16 %rs8, %rs7, 8;
+; O0-NEXT: or.b16 %rs9, %rs8, %rs6;
+; O0-NEXT: cvt.u32.u16 %r2, %rs9;
+; O0-NEXT: ld.s8 %rs10, [%rd2+2];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x9991U;
+; O0-NEXT: cvt.u16.u32 %rs11, %r3;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x9991U;
+; O0-NEXT: cvt.u16.u32 %rs12, %r4;
+; O0-NEXT: rem.s16 %rs13, %rs12, %rs11;
+; O0-NEXT: cvt.u32.u16 %r5, %rs13;
+; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x8880U;
+; O0-NEXT: cvt.u16.u32 %rs14, %r6;
+; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x8880U;
+; O0-NEXT: cvt.u16.u32 %rs15, %r7;
+; O0-NEXT: rem.s16 %rs16, %rs15, %rs14;
+; O0-NEXT: cvt.u32.u16 %r8, %rs16;
+; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O0-NEXT: // implicit-def: %r11
+; O0-NEXT: // implicit-def: %r12
+; O0-NEXT: prmt.b32 %r10, %r11, %r12, 0x3340U;
+; O0-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U;
+; O0-NEXT: rem.s16 %rs17, %rs5, %rs10;
+; O0-NEXT: mov.b32 {%rs18, _}, %r13;
+; O0-NEXT: st.b8 [%rd3], %rs18;
+; O0-NEXT: shr.u16 %rs19, %rs18, 8;
+; O0-NEXT: st.b8 [%rd3+1], %rs19;
+; O0-NEXT: st.b8 [%rd3+2], %rs17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_srem_v3i8(
+; O3: {
+; O3-NEXT: .reg .b16 %rs<20>;
+; O3-NEXT: .reg .b32 %r<14>;
+; O3-NEXT: .reg .b64 %rd<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0: // %entry
+; O3-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0];
+; O3-NEXT: ld.b8 %rs1, [%rd1];
+; O3-NEXT: ld.b8 %rs2, [%rd1+1];
+; O3-NEXT: shl.b16 %rs3, %rs2, 8;
+; O3-NEXT: or.b16 %rs4, %rs3, %rs1;
+; O3-NEXT: cvt.u32.u16 %r1, %rs4;
+; O3-NEXT: ld.s8 %rs5, [%rd1+2];
+; O3-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1];
+; O3-NEXT: ld.b8 %rs6, [%rd2];
+; O3-NEXT: ld.b8 %rs7, [%rd2+1];
+; O3-NEXT: shl.b16 %rs8, %rs7, 8;
+; O3-NEXT: or.b16 %rs9, %rs8, %rs6;
+; O3-NEXT: cvt.u32.u16 %r2, %rs9;
+; O3-NEXT: ld.s8 %rs10, [%rd2+2];
+; O3-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x9991U;
+; O3-NEXT: cvt.u16.u32 %rs11, %r3;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x9991U;
+; O3-NEXT: cvt.u16.u32 %rs12, %r4;
+; O3-NEXT: rem.s16 %rs13, %rs12, %rs11;
+; O3-NEXT: cvt.u32.u16 %r5, %rs13;
+; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x8880U;
+; O3-NEXT: cvt.u16.u32 %rs14, %r6;
+; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x8880U;
+; O3-NEXT: cvt.u16.u32 %rs15, %r7;
+; O3-NEXT: rem.s16 %rs16, %rs15, %rs14;
+; O3-NEXT: cvt.u32.u16 %r8, %rs16;
+; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;
+; O3-NEXT: prmt.b32 %r10, %r11, %r12, 0x3340U;
+; O3-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U;
+; O3-NEXT: rem.s16 %rs17, %rs5, %rs10;
+; O3-NEXT: st.b8 [%rd3+2], %rs17;
+; O3-NEXT: mov.b32 {%rs18, _}, %r13;
+; O3-NEXT: st.b8 [%rd3], %rs18;
+; O3-NEXT: shr.u16 %rs19, %rs18, 8;
+; O3-NEXT: st.b8 [%rd3+1], %rs19;
+; O3-NEXT: ret;
entry:
%t57 = load <3 x i8>, ptr %a, align 1
%t59 = load <3 x i8>, ptr %b, align 1
@@ -1383,39 +2243,73 @@ 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-EMPTY:
-; CHECK-NEXT: // %bb.0: // %entry
-; CHECK-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2];
-; CHECK-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1];
-; CHECK-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0];
-; CHECK-NEXT: ld.b32 %r1, [%rd1];
-; CHECK-NEXT: ld.b32 %r2, [%rd2];
-; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
-; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
-; CHECK-NEXT: setp.gt.u32 %p1, %r4, %r3;
-; CHECK-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
-; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
-; CHECK-NEXT: setp.gt.u32 %p2, %r6, %r5;
-; CHECK-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
-; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
-; CHECK-NEXT: setp.gt.u32 %p3, %r8, %r7;
-; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
-; CHECK-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
-; CHECK-NEXT: setp.gt.u32 %p4, %r10, %r9;
-; CHECK-NEXT: selp.b32 %r11, -1, 0, %p4;
-; CHECK-NEXT: selp.b32 %r12, -1, 0, %p3;
-; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
-; CHECK-NEXT: selp.b32 %r14, -1, 0, %p2;
-; CHECK-NEXT: selp.b32 %r15, -1, 0, %p1;
-; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
-; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
-; CHECK-NEXT: st.b32 [%rd3], %r17;
-; CHECK-NEXT: ret;
+; O0-LABEL: test_sext_v4i1_to_v4i8(
+; O0: {
+; O0-NEXT: .reg .pred %p<5>;
+; O0-NEXT: .reg .b32 %r<18>;
+; O0-NEXT: .reg .b64 %rd<4>;
+; O0-EMPTY:
+; O0-NEXT: // %bb.0: // %entry
+; O0-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2];
+; O0-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1];
+; O0-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0];
+; O0-NEXT: ld.b32 %r1, [%rd1];
+; O0-NEXT: ld.b32 %r2, [%rd2];
+; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O0-NEXT: setp.gt.u32 %p1, %r4, %r3;
+; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O0-NEXT: setp.gt.u32 %p2, %r6, %r5;
+; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O0-NEXT: setp.gt.u32 %p3, %r8, %r7;
+; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O0-NEXT: setp.gt.u32 %p4, %r10, %r9;
+; O0-NEXT: selp.b32 %r11, -1, 0, %p4;
+; O0-NEXT: selp.b32 %r12, -1, 0, %p3;
+; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O0-NEXT: selp.b32 %r14, -1, 0, %p2;
+; O0-NEXT: selp.b32 %r15, -1, 0, %p1;
+; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O0-NEXT: st.b32 [%rd3], %r17;
+; O0-NEXT: ret;
+;
+; O3-LABEL: test_sext_v4i1_to_v4i8(
+; O3: {
+; O3-NEXT: .reg .pred %p<5>;
+; O3-NEXT: .reg .b32 %r<18>;
+; O3-NEXT: .reg .b64 %rd<4>;
+; O3-EMPTY:
+; O3-NEXT: // %bb.0: // %entry
+; O3-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0];
+; O3-NEXT: ld.b32 %r1, [%rd1];
+; O3-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1];
+; O3-NEXT: ld.b32 %r2, [%rd2];
+; O3-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2];
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;
+; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;
+; O3-NEXT: setp.gt.u32 %p1, %r4, %r3;
+; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;
+; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;
+; O3-NEXT: setp.gt.u32 %p2, %r6, %r5;
+; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;
+; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;
+; O3-NEXT: setp.gt.u32 %p3, %r8, %r7;
+; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;
+; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;
+; O3-NEXT: setp.gt.u32 %p4, %r10, %r9;
+; O3-NEXT: selp.b32 %r11, -1, 0, %p4;
+; O3-NEXT: selp.b32 %r12, -1, 0, %p3;
+; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;
+; O3-NEXT: selp.b32 %r14, -1, 0, %p2;
+; O3-NEXT: selp.b32 %r15, -1, 0, %p1;
+; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;
+; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;
+; O3-NEXT: st.b32 [%rd3], %r17;
+; O3-NEXT: ret;
entry:
%t1 = load <4 x i8>, ptr %a, align 4
%t2 = load <4 x i8>, ptr %b, align 4
>From e509eef10a2f89a2c3f47c52a68fad74e3117798 Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Tue, 15 Jul 2025 16:35:55 +0000
Subject: [PATCH 2/3] [NVPTX] Add PRMT constant folding and cleanup usage of
PRMT node
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 236 +++++++++++++-----
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 23 +-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 18 --
llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 58 ++---
llvm/test/CodeGen/NVPTX/prmt-const-folding.ll | 171 +++++++++++++
5 files changed, 391 insertions(+), 115 deletions(-)
create mode 100644 llvm/test/CodeGen/NVPTX/prmt-const-folding.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index d017c658c53a3..d84e4fd0eb40a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -1048,9 +1048,12 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
MVT::v32i32, MVT::v64i32, MVT::v128i32},
Custom);
- setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::Other, Custom);
- // Enable custom lowering for the i128 bit operand with clusterlaunchcontrol
- setOperationAction(ISD::INTRINSIC_WO_CHAIN, MVT::i128, Custom);
+ // Enable custom lowering for the following:
+ // * MVT::i128 - clusterlaunchcontrol
+ // * MVT::i32 - prmt
+ // * MVT::Other - internal.addrspace.wrap
+ setOperationAction(ISD::INTRINSIC_WO_CHAIN, {MVT::i32, MVT::i128, MVT::Other},
+ Custom);
}
const char *NVPTXTargetLowering::getTargetNodeName(unsigned Opcode) const {
@@ -2060,6 +2063,13 @@ NVPTXTargetLowering::LowerCONCAT_VECTORS(SDValue Op, SelectionDAG &DAG) const {
return DAG.getBuildVector(Node->getValueType(0), dl, Ops);
}
+static SDValue getPRMT(SDValue A, SDValue B, SDValue Selector, SDLoc DL,
+ SelectionDAG &DAG,
+ unsigned Mode = NVPTX::PTXPrmtMode::NONE) {
+ return DAG.getNode(NVPTXISD::PRMT, DL, MVT::i32,
+ {A, B, Selector, DAG.getConstant(Mode, DL, MVT::i32)});
+}
+
SDValue NVPTXTargetLowering::LowerBITCAST(SDValue Op, SelectionDAG &DAG) const {
// Handle bitcasting from v2i8 without hitting the default promotion
// strategy which goes through stack memory.
@@ -2111,15 +2121,13 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op,
L = DAG.getAnyExtOrTrunc(L, DL, MVT::i32);
R = DAG.getAnyExtOrTrunc(R, DL, MVT::i32);
}
- return DAG.getNode(
- NVPTXISD::PRMT, DL, MVT::v4i8,
- {L, R, DAG.getConstant(SelectionValue, DL, MVT::i32),
- DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)});
+ return getPRMT(L, R, DAG.getConstant(SelectionValue, DL, MVT::i32), DL,
+ DAG);
};
auto PRMT__10 = GetPRMT(Op->getOperand(0), Op->getOperand(1), true, 0x3340);
auto PRMT__32 = GetPRMT(Op->getOperand(2), Op->getOperand(3), true, 0x3340);
auto PRMT3210 = GetPRMT(PRMT__10, PRMT__32, false, 0x5410);
- return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210);
+ return DAG.getBitcast(VT, PRMT3210);
}
// Get value or the Nth operand as an APInt(32). Undef values treated as 0.
@@ -2176,11 +2184,14 @@ SDValue NVPTXTargetLowering::LowerEXTRACT_VECTOR_ELT(SDValue Op,
SDValue Selector = DAG.getNode(ISD::OR, DL, MVT::i32,
DAG.getZExtOrTrunc(Index, DL, MVT::i32),
DAG.getConstant(0x7770, DL, MVT::i32));
- SDValue PRMT = DAG.getNode(
- NVPTXISD::PRMT, DL, MVT::i32,
- {DAG.getBitcast(MVT::i32, Vector), DAG.getConstant(0, DL, MVT::i32),
- Selector, DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)});
- return DAG.getAnyExtOrTrunc(PRMT, DL, Op->getValueType(0));
+ SDValue PRMT = getPRMT(DAG.getBitcast(MVT::i32, Vector),
+ DAG.getConstant(0, DL, MVT::i32), Selector, DL, DAG);
+ SDValue Ext = DAG.getAnyExtOrTrunc(PRMT, DL, Op->getValueType(0));
+ SDNodeFlags Flags;
+ Flags.setNoSignedWrap(Ext.getScalarValueSizeInBits() > 8);
+ Flags.setNoUnsignedWrap(Ext.getScalarValueSizeInBits() >= 8);
+ Ext->setFlags(Flags);
+ return Ext;
}
// Constant index will be matched by tablegen.
@@ -2242,9 +2253,10 @@ SDValue NVPTXTargetLowering::LowerVECTOR_SHUFFLE(SDValue Op,
}
SDLoc DL(Op);
- return DAG.getNode(NVPTXISD::PRMT, DL, MVT::v4i8, V1, V2,
- DAG.getConstant(Selector, DL, MVT::i32),
- DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32));
+ SDValue PRMT =
+ getPRMT(DAG.getBitcast(MVT::i32, V1), DAG.getBitcast(MVT::i32, V2),
+ DAG.getConstant(Selector, DL, MVT::i32), DL, DAG);
+ return DAG.getBitcast(Op.getValueType(), PRMT);
}
/// LowerShiftRightParts - Lower SRL_PARTS, SRA_PARTS, which
/// 1) returns two i32 values and take a 2 x i32 value to shift plus a shift
@@ -2729,10 +2741,46 @@ static SDValue LowerClusterLaunchControlQueryCancel(SDValue Op,
{TryCancelResponse0, TryCancelResponse1});
}
+static SDValue lowerPrmtIntrinsic(SDValue Op, SelectionDAG &DAG) {
+ const unsigned Mode = [&]() {
+ switch (Op->getConstantOperandVal(0)) {
+ case Intrinsic::nvvm_prmt:
+ return NVPTX::PTXPrmtMode::NONE;
+ case Intrinsic::nvvm_prmt_b4e:
+ return NVPTX::PTXPrmtMode::B4E;
+ case Intrinsic::nvvm_prmt_ecl:
+ return NVPTX::PTXPrmtMode::ECL;
+ case Intrinsic::nvvm_prmt_ecr:
+ return NVPTX::PTXPrmtMode::ECR;
+ case Intrinsic::nvvm_prmt_f4e:
+ return NVPTX::PTXPrmtMode::F4E;
+ case Intrinsic::nvvm_prmt_rc16:
+ return NVPTX::PTXPrmtMode::RC16;
+ case Intrinsic::nvvm_prmt_rc8:
+ return NVPTX::PTXPrmtMode::RC8;
+ default:
+ llvm_unreachable("unsupported/unhandled intrinsic");
+ }
+ }();
+ SDLoc DL(Op);
+ SDValue A = Op->getOperand(1);
+ SDValue B = Op.getNumOperands() == 4 ? Op.getOperand(2)
+ : DAG.getConstant(0, DL, MVT::i32);
+ SDValue Selector = (Op->op_end() - 1)->get();
+ return getPRMT(A, B, Selector, DL, DAG, Mode);
+}
static SDValue lowerIntrinsicWOChain(SDValue Op, SelectionDAG &DAG) {
switch (Op->getConstantOperandVal(0)) {
default:
return Op;
+ case Intrinsic::nvvm_prmt:
+ case Intrinsic::nvvm_prmt_b4e:
+ case Intrinsic::nvvm_prmt_ecl:
+ case Intrinsic::nvvm_prmt_ecr:
+ case Intrinsic::nvvm_prmt_f4e:
+ case Intrinsic::nvvm_prmt_rc16:
+ case Intrinsic::nvvm_prmt_rc8:
+ return lowerPrmtIntrinsic(Op, DAG);
case Intrinsic::nvvm_internal_addrspace_wrap:
return Op.getOperand(1);
case Intrinsic::nvvm_clusterlaunchcontrol_query_cancel_is_canceled:
@@ -5775,11 +5823,10 @@ PerformBUILD_VECTORCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
SDLoc DL(N);
auto &DAG = DCI.DAG;
- auto PRMT = DAG.getNode(
- NVPTXISD::PRMT, DL, MVT::v4i8,
- {Op0, Op1, DAG.getConstant((Op1Bytes << 8) | Op0Bytes, DL, MVT::i32),
- DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)});
- return DAG.getNode(ISD::BITCAST, DL, VT, PRMT);
+ auto PRMT = getPRMT(
+ DAG.getBitcast(MVT::i32, Op0), DAG.getBitcast(MVT::i32, Op1),
+ DAG.getConstant((Op1Bytes << 8) | Op0Bytes, DL, MVT::i32), DL, DAG);
+ return DAG.getBitcast(VT, PRMT);
}
static SDValue combineADDRSPACECAST(SDNode *N,
@@ -5797,47 +5844,116 @@ static SDValue combineADDRSPACECAST(SDNode *N,
return SDValue();
}
+static APInt getPRMTSelector(APInt Selector, unsigned Mode) {
+ if (Mode == NVPTX::PTXPrmtMode::NONE)
+ return Selector;
+
+ unsigned V = Selector.trunc(2).getZExtValue();
+
+ const auto GetSelector = [](unsigned S0, unsigned S1, unsigned S2,
+ unsigned S3) {
+ return APInt(32, S0 | (S1 << 4) | (S2 << 8) | (S3 << 12));
+ };
+
+ switch (Mode) {
+ case NVPTX::PTXPrmtMode::F4E:
+ return GetSelector(V, V + 1, V + 2, V + 3);
+ case NVPTX::PTXPrmtMode::B4E:
+ return GetSelector(V, (V - 1) & 7, (V - 2) & 7, (V - 3) & 7);
+ case NVPTX::PTXPrmtMode::RC8:
+ return GetSelector(V, V, V, V);
+ case NVPTX::PTXPrmtMode::ECL:
+ return GetSelector(V, std::max(V, 1U), std::max(V, 2U), 3U);
+ case NVPTX::PTXPrmtMode::ECR:
+ return GetSelector(0, std::min(V, 1U), std::min(V, 2U), V);
+ case NVPTX::PTXPrmtMode::RC16: {
+ unsigned V1 = (V & 1) << 1;
+ return GetSelector(V1, V1 + 1, V1, V1 + 1);
+ }
+ default:
+ llvm_unreachable("Invalid PRMT mode");
+ }
+}
+
+static APInt computePRMT(APInt A, APInt B, APInt Selector, unsigned Mode) {
+ // {b, a} = {{b7, b6, b5, b4}, {b3, b2, b1, b0}}
+ APInt BitField = B.concat(A);
+ APInt SelectorVal = getPRMTSelector(Selector, Mode);
+ APInt Result(32, 0);
+ for (unsigned I : llvm::seq(4U)) {
+ APInt Sel = SelectorVal.extractBits(4, I * 4);
+ unsigned Idx = Sel.getLoBits(3).getZExtValue();
+ unsigned Sign = Sel.getHiBits(1).getZExtValue();
+ APInt Byte = BitField.extractBits(8, Idx * 8);
+ if (Sign)
+ Byte = Byte.ashr(8);
+ Result.insertBits(Byte, I * 8);
+ }
+ return Result;
+}
+
+static SDValue combinePRMT(SDNode *N, TargetLowering::DAGCombinerInfo &DCI,
+ CodeGenOptLevel OptLevel) {
+ if (OptLevel == CodeGenOptLevel::None)
+ return SDValue();
+
+ // Constant fold PRMT
+ if (isa<ConstantSDNode>(N->getOperand(0)) &&
+ isa<ConstantSDNode>(N->getOperand(1)) &&
+ isa<ConstantSDNode>(N->getOperand(2)))
+ return DCI.DAG.getConstant(computePRMT(N->getConstantOperandAPInt(0),
+ N->getConstantOperandAPInt(1),
+ N->getConstantOperandAPInt(2),
+ N->getConstantOperandVal(3)),
+ SDLoc(N), N->getValueType(0));
+
+ return SDValue();
+}
+
SDValue NVPTXTargetLowering::PerformDAGCombine(SDNode *N,
DAGCombinerInfo &DCI) const {
CodeGenOptLevel OptLevel = getTargetMachine().getOptLevel();
switch (N->getOpcode()) {
- default: break;
- case ISD::ADD:
- return PerformADDCombine(N, DCI, OptLevel);
- case ISD::FADD:
- return PerformFADDCombine(N, DCI, OptLevel);
- case ISD::MUL:
- return PerformMULCombine(N, DCI, OptLevel);
- case ISD::SHL:
- return PerformSHLCombine(N, DCI, OptLevel);
- case ISD::AND:
- return PerformANDCombine(N, DCI);
- case ISD::UREM:
- case ISD::SREM:
- return PerformREMCombine(N, DCI, OptLevel);
- case ISD::SETCC:
- return PerformSETCCCombine(N, DCI, STI.getSmVersion());
- case ISD::LOAD:
- case NVPTXISD::LoadParamV2:
- case NVPTXISD::LoadV2:
- case NVPTXISD::LoadV4:
- return combineUnpackingMovIntoLoad(N, DCI);
- case NVPTXISD::StoreParam:
- case NVPTXISD::StoreParamV2:
- case NVPTXISD::StoreParamV4:
- return PerformStoreParamCombine(N, DCI);
- case ISD::STORE:
- case NVPTXISD::StoreV2:
- case NVPTXISD::StoreV4:
- return PerformStoreCombine(N, DCI);
- case ISD::EXTRACT_VECTOR_ELT:
- return PerformEXTRACTCombine(N, DCI);
- case ISD::VSELECT:
- return PerformVSELECTCombine(N, DCI);
- case ISD::BUILD_VECTOR:
- return PerformBUILD_VECTORCombine(N, DCI);
- case ISD::ADDRSPACECAST:
- return combineADDRSPACECAST(N, DCI);
+ default:
+ break;
+ case ISD::ADD:
+ return PerformADDCombine(N, DCI, OptLevel);
+ case ISD::FADD:
+ return PerformFADDCombine(N, DCI, OptLevel);
+ case ISD::MUL:
+ return PerformMULCombine(N, DCI, OptLevel);
+ case ISD::SHL:
+ return PerformSHLCombine(N, DCI, OptLevel);
+ case ISD::AND:
+ return PerformANDCombine(N, DCI);
+ case ISD::UREM:
+ case ISD::SREM:
+ return PerformREMCombine(N, DCI, OptLevel);
+ case ISD::SETCC:
+ return PerformSETCCCombine(N, DCI, STI.getSmVersion());
+ case ISD::LOAD:
+ case NVPTXISD::LoadParamV2:
+ case NVPTXISD::LoadV2:
+ case NVPTXISD::LoadV4:
+ return combineUnpackingMovIntoLoad(N, DCI);
+ case NVPTXISD::StoreParam:
+ case NVPTXISD::StoreParamV2:
+ case NVPTXISD::StoreParamV4:
+ return PerformStoreParamCombine(N, DCI);
+ case ISD::STORE:
+ case NVPTXISD::StoreV2:
+ case NVPTXISD::StoreV4:
+ return PerformStoreCombine(N, DCI);
+ case ISD::EXTRACT_VECTOR_ELT:
+ return PerformEXTRACTCombine(N, DCI);
+ case ISD::VSELECT:
+ return PerformVSELECTCombine(N, DCI);
+ case ISD::BUILD_VECTOR:
+ return PerformBUILD_VECTORCombine(N, DCI);
+ case ISD::ADDRSPACECAST:
+ return combineADDRSPACECAST(N, DCI);
+ case NVPTXISD::PRMT:
+ return combinePRMT(N, DCI, OptLevel);
}
return SDValue();
}
@@ -6387,7 +6503,7 @@ static void computeKnownBitsForPRMT(const SDValue Op, KnownBits &Known,
ConstantSDNode *Selector = dyn_cast<ConstantSDNode>(Op.getOperand(2));
unsigned Mode = Op.getConstantOperandVal(3);
- if (Mode != NVPTX::PTXPrmtMode::NONE || !Selector)
+ if (!Selector)
return;
KnownBits AKnown = DAG.computeKnownBits(A, Depth);
@@ -6396,7 +6512,7 @@ static void computeKnownBitsForPRMT(const SDValue Op, KnownBits &Known,
// {b, a} = {{b7, b6, b5, b4}, {b3, b2, b1, b0}}
KnownBits BitField = BKnown.concat(AKnown);
- APInt SelectorVal = Selector->getAPIntValue();
+ APInt SelectorVal = getPRMTSelector(Selector->getAPIntValue(), Mode);
for (unsigned I : llvm::seq(std::min(4U, Known.getBitWidth() / 8))) {
APInt Sel = SelectorVal.extractBits(4, I * 4);
unsigned Idx = Sel.getLoBits(3).getZExtValue();
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 4eef6c939720c..a5bb83dfadb84 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1453,18 +1453,33 @@ let hasSideEffects = false in {
(ins PrmtMode:$mode),
"prmt.b32$mode",
[(set i32:$d, (prmt i32:$a, i32:$b, imm:$c, imm:$mode))]>;
+ def PRMT_B32rir
+ : BasicFlagsNVPTXInst<(outs B32:$d),
+ (ins B32:$a, i32imm:$b, B32:$c),
+ (ins PrmtMode:$mode),
+ "prmt.b32$mode",
+ [(set i32:$d, (prmt i32:$a, imm:$b, i32:$c, imm:$mode))]>;
def PRMT_B32rii
: BasicFlagsNVPTXInst<(outs B32:$d),
(ins B32:$a, i32imm:$b, Hexu32imm:$c),
(ins PrmtMode:$mode),
"prmt.b32$mode",
[(set i32:$d, (prmt i32:$a, imm:$b, imm:$c, imm:$mode))]>;
- def PRMT_B32rir
+ def PRMT_B32irr
: BasicFlagsNVPTXInst<(outs B32:$d),
- (ins B32:$a, i32imm:$b, B32:$c),
- (ins PrmtMode:$mode),
+ (ins i32imm:$a, B32:$b, B32:$c), (ins PrmtMode:$mode),
+ "prmt.b32$mode",
+ [(set i32:$d, (prmt imm:$a, i32:$b, i32:$c, imm:$mode))]>;
+ def PRMT_B32iri
+ : BasicFlagsNVPTXInst<(outs B32:$d),
+ (ins i32imm:$a, B32:$b, Hexu32imm:$c), (ins PrmtMode:$mode),
+ "prmt.b32$mode",
+ [(set i32:$d, (prmt imm:$a, i32:$b, imm:$c, imm:$mode))]>;
+ def PRMT_B32iir
+ : BasicFlagsNVPTXInst<(outs B32:$d),
+ (ins i32imm:$a, i32imm:$b, B32:$c), (ins PrmtMode:$mode),
"prmt.b32$mode",
- [(set i32:$d, (prmt i32:$a, imm:$b, i32:$c, imm:$mode))]>;
+ [(set i32:$d, (prmt imm:$a, imm:$b, i32:$c, imm:$mode))]>;
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index bad4c3c4c5f3a..70150bdfc8d16 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -1047,24 +1047,6 @@ class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
// MISC
//
-class PRMT3Pat<Intrinsic prmt_intrinsic, PatLeaf prmt_mode>
- : Pat<(prmt_intrinsic i32:$a, i32:$b, i32:$c),
- (PRMT_B32rrr $a, $b, $c, prmt_mode)>;
-
-class PRMT2Pat<Intrinsic prmt_intrinsic, PatLeaf prmt_mode>
- : Pat<(prmt_intrinsic i32:$a, i32:$c),
- (PRMT_B32rir $a, (i32 0), $c, prmt_mode)>;
-
-def : PRMT3Pat<int_nvvm_prmt, PrmtNONE>;
-def : PRMT3Pat<int_nvvm_prmt_f4e, PrmtF4E>;
-def : PRMT3Pat<int_nvvm_prmt_b4e, PrmtB4E>;
-
-def : PRMT2Pat<int_nvvm_prmt_rc8, PrmtRC8>;
-def : PRMT2Pat<int_nvvm_prmt_ecl, PrmtECL>;
-def : PRMT2Pat<int_nvvm_prmt_ecr, PrmtECR>;
-def : PRMT2Pat<int_nvvm_prmt_rc16, PrmtRC16>;
-
-
def INT_NVVM_NANOSLEEP_I : BasicNVPTXInst<(outs), (ins i32imm:$i), "nanosleep.u32",
[(int_nvvm_nanosleep imm:$i)]>,
Requires<[hasPTX<63>, hasSM<70>]>;
diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index 728636573a42d..cbc9f700b1f01 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -875,17 +875,15 @@ define <4 x i8> @test_or_computed(i8 %a) {
;
; O3-LABEL: test_or_computed(
; O3: {
-; O3-NEXT: .reg .b32 %r<8>;
+; O3-NEXT: .reg .b32 %r<6>;
; O3-EMPTY:
; O3-NEXT: // %bb.0:
-; O3-NEXT: mov.b32 %r1, 0;
+; O3-NEXT: ld.param.b8 %r1, [test_or_computed_param_0];
; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
-; O3-NEXT: ld.param.b8 %r3, [test_or_computed_param_0];
-; O3-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
-; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; O3-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
-; O3-NEXT: or.b32 %r7, %r6, %r5;
-; O3-NEXT: st.param.b32 [func_retval0], %r7;
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U;
+; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8;
+; O3-NEXT: or.b32 %r5, %r4, %r3;
+; O3-NEXT: st.param.b32 [func_retval0], %r5;
; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
@@ -967,17 +965,15 @@ define <4 x i8> @test_xor_computed(i8 %a) {
;
; O3-LABEL: test_xor_computed(
; O3: {
-; O3-NEXT: .reg .b32 %r<8>;
+; O3-NEXT: .reg .b32 %r<6>;
; O3-EMPTY:
; O3-NEXT: // %bb.0:
-; O3-NEXT: mov.b32 %r1, 0;
+; O3-NEXT: ld.param.b8 %r1, [test_xor_computed_param_0];
; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
-; O3-NEXT: ld.param.b8 %r3, [test_xor_computed_param_0];
-; O3-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
-; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; O3-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
-; O3-NEXT: xor.b32 %r7, %r6, %r5;
-; O3-NEXT: st.param.b32 [func_retval0], %r7;
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U;
+; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8;
+; O3-NEXT: xor.b32 %r5, %r4, %r3;
+; O3-NEXT: st.param.b32 [func_retval0], %r5;
; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
@@ -1059,17 +1055,15 @@ define <4 x i8> @test_and_computed(i8 %a) {
;
; O3-LABEL: test_and_computed(
; O3: {
-; O3-NEXT: .reg .b32 %r<8>;
+; O3-NEXT: .reg .b32 %r<6>;
; O3-EMPTY:
; O3-NEXT: // %bb.0:
-; O3-NEXT: mov.b32 %r1, 0;
+; O3-NEXT: ld.param.b8 %r1, [test_and_computed_param_0];
; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;
-; O3-NEXT: ld.param.b8 %r3, [test_and_computed_param_0];
-; O3-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;
-; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; O3-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;
-; O3-NEXT: and.b32 %r7, %r6, %r5;
-; O3-NEXT: st.param.b32 [func_retval0], %r7;
+; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U;
+; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8;
+; O3-NEXT: and.b32 %r5, %r4, %r3;
+; O3-NEXT: st.param.b32 [func_retval0], %r5;
; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
@@ -1837,15 +1831,13 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 {
;
; O3-LABEL: test_bitcast_4xi8_to_2xhalf(
; O3: {
-; O3-NEXT: .reg .b32 %r<6>;
+; O3-NEXT: .reg .b32 %r<4>;
; O3-EMPTY:
; O3-NEXT: // %bb.0:
-; O3-NEXT: mov.b32 %r1, 6;
-; O3-NEXT: prmt.b32 %r2, %r1, 7, 0x3340U;
-; O3-NEXT: ld.param.b8 %r3, [test_bitcast_4xi8_to_2xhalf_param_0];
-; O3-NEXT: prmt.b32 %r4, %r3, 5, 0x3340U;
-; O3-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;
-; O3-NEXT: st.param.b32 [func_retval0], %r5;
+; O3-NEXT: ld.param.b8 %r1, [test_bitcast_4xi8_to_2xhalf_param_0];
+; O3-NEXT: prmt.b32 %r2, %r1, 5, 0x3340U;
+; O3-NEXT: prmt.b32 %r3, %r2, 1798, 0x5410U;
+; O3-NEXT: st.param.b32 [func_retval0], %r3;
; O3-NEXT: ret;
%ins.0 = insertelement <4 x i8> undef, i8 %a, i32 0
%ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1
@@ -2183,7 +2175,7 @@ define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) {
; O0-NEXT: prmt.b32 %r10, %r11, %r12, 0x3340U;
; O0-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U;
; O0-NEXT: rem.s16 %rs17, %rs5, %rs10;
-; O0-NEXT: mov.b32 {%rs18, _}, %r13;
+; O0-NEXT: cvt.u16.u32 %rs18, %r13;
; O0-NEXT: st.b8 [%rd3], %rs18;
; O0-NEXT: shr.u16 %rs19, %rs18, 8;
; O0-NEXT: st.b8 [%rd3+1], %rs19;
@@ -2229,7 +2221,7 @@ define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) {
; O3-NEXT: prmt.b32 %r13, %r9, %r10, 0x5410U;
; O3-NEXT: rem.s16 %rs17, %rs5, %rs10;
; O3-NEXT: st.b8 [%rd3+2], %rs17;
-; O3-NEXT: mov.b32 {%rs18, _}, %r13;
+; O3-NEXT: cvt.u16.u32 %rs18, %r13;
; O3-NEXT: st.b8 [%rd3], %rs18;
; O3-NEXT: shr.u16 %rs19, %rs18, 8;
; O3-NEXT: st.b8 [%rd3+1], %rs19;
diff --git a/llvm/test/CodeGen/NVPTX/prmt-const-folding.ll b/llvm/test/CodeGen/NVPTX/prmt-const-folding.ll
new file mode 100644
index 0000000000000..7afead127c84f
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/prmt-const-folding.ll
@@ -0,0 +1,171 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -verify-machineinstrs -mcpu=sm_50 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -verify-machineinstrs -mcpu=sm_50 | %ptxas-verify %}
+
+target triple = "nvptx64-nvidia-cuda"
+
+ at g = global i32 0
+
+define void @test_prmt_f4e() {
+; CHECK-LABEL: test_prmt_f4e(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.volatile.global.b32 [g], 50462976;
+; CHECK-NEXT: st.volatile.global.b32 [g], 67305985;
+; CHECK-NEXT: st.volatile.global.b32 [g], 84148994;
+; CHECK-NEXT: st.volatile.global.b32 [g], 100992003;
+; CHECK-NEXT: st.volatile.global.b32 [g], 50462976;
+; CHECK-NEXT: ret;
+ %v1 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x0)
+ store volatile i32 %v1, ptr @g
+ %v2 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x1)
+ store volatile i32 %v2, ptr @g
+ %v3 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x2)
+ store volatile i32 %v3, ptr @g
+ %v4 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x3)
+ store volatile i32 %v4, ptr @g
+ %v5 = call i32 @llvm.nvvm.prmt.f4e(i32 u0x03020100, i32 u0x07060504, i32 u0x4)
+ store volatile i32 %v5, ptr @g
+ ret void
+}
+
+define void @test_prmt_b4e() {
+; CHECK-LABEL: test_prmt_b4e(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.volatile.global.b32 [g], 84281088;
+; CHECK-NEXT: st.volatile.global.b32 [g], 101122049;
+; CHECK-NEXT: st.volatile.global.b32 [g], 117440770;
+; CHECK-NEXT: st.volatile.global.b32 [g], 66051;
+; CHECK-NEXT: st.volatile.global.b32 [g], 84281088;
+; CHECK-NEXT: ret;
+ %v1 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x0)
+ store volatile i32 %v1, ptr @g
+ %v2 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x1)
+ store volatile i32 %v2, ptr @g
+ %v3 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x2)
+ store volatile i32 %v3, ptr @g
+ %v4 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x3)
+ store volatile i32 %v4, ptr @g
+ %v5 = call i32 @llvm.nvvm.prmt.b4e(i32 u0x03020100, i32 u0x07060504, i32 u0x4)
+ store volatile i32 %v5, ptr @g
+ ret void
+}
+
+define void @test_prmt_ecl() {
+; CHECK-LABEL: test_prmt_ecl(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.volatile.global.b32 [g], 50462976;
+; CHECK-NEXT: st.volatile.global.b32 [g], 50462977;
+; CHECK-NEXT: st.volatile.global.b32 [g], 50463234;
+; CHECK-NEXT: st.volatile.global.b32 [g], 50529027;
+; CHECK-NEXT: ret;
+ %v1 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x0)
+ store volatile i32 %v1, ptr @g
+ %v2 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x1)
+ store volatile i32 %v2, ptr @g
+ %v3 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x2)
+ store volatile i32 %v3, ptr @g
+ %v4 = call i32 @llvm.nvvm.prmt.ecl(i32 u0x03020100, i32 u0x3)
+ store volatile i32 %v4, ptr @g
+ ret void
+}
+
+define void @test_prmt_ecr() {
+; CHECK-LABEL: test_prmt_ecr(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.volatile.global.b32 [g], 0;
+; CHECK-NEXT: st.volatile.global.b32 [g], 16843008;
+; CHECK-NEXT: st.volatile.global.b32 [g], 33685760;
+; CHECK-NEXT: st.volatile.global.b32 [g], 50462976;
+; CHECK-NEXT: ret;
+ %v1 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x0)
+ store volatile i32 %v1, ptr @g
+ %v2 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x1)
+ store volatile i32 %v2, ptr @g
+ %v3 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x2)
+ store volatile i32 %v3, ptr @g
+ %v4 = call i32 @llvm.nvvm.prmt.ecr(i32 u0x03020100, i32 u0x3)
+ store volatile i32 %v4, ptr @g
+ ret void
+}
+
+define void @test_prmt_rc8() {
+; CHECK-LABEL: test_prmt_rc8(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.volatile.global.b32 [g], 0;
+; CHECK-NEXT: st.volatile.global.b32 [g], 16843009;
+; CHECK-NEXT: st.volatile.global.b32 [g], 33686018;
+; CHECK-NEXT: st.volatile.global.b32 [g], 50529027;
+; CHECK-NEXT: st.volatile.global.b32 [g], 0;
+; CHECK-NEXT: ret;
+ %v1 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x0)
+ store volatile i32 %v1, ptr @g
+ %v2 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x1)
+ store volatile i32 %v2, ptr @g
+ %v3 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x2)
+ store volatile i32 %v3, ptr @g
+ %v4 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x3)
+ store volatile i32 %v4, ptr @g
+ %v5 = call i32 @llvm.nvvm.prmt.rc8(i32 u0x03020100, i32 u0x4)
+ store volatile i32 %v5, ptr @g
+ ret void
+}
+
+define void @test_prmt_rc16() {
+; CHECK-LABEL: test_prmt_rc16(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.volatile.global.b32 [g], 16777472;
+; CHECK-NEXT: st.volatile.global.b32 [g], 50463490;
+; CHECK-NEXT: st.volatile.global.b32 [g], 16777472;
+; CHECK-NEXT: ret;
+ %v1 = call i32 @llvm.nvvm.prmt.rc16(i32 u0x03020100, i32 u0x0)
+ store volatile i32 %v1, ptr @g
+ %v2 = call i32 @llvm.nvvm.prmt.rc16(i32 u0x03020100, i32 u0x1)
+ store volatile i32 %v2, ptr @g
+ %v3 = call i32 @llvm.nvvm.prmt.rc16(i32 u0x03020100, i32 u0x2)
+ store volatile i32 %v3, ptr @g
+ ret void
+}
+
+
+define void @test_prmt_basic() {
+; CHECK-LABEL: test_prmt_basic(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: st.volatile.global.b32 [g], 66051;
+; CHECK-NEXT: st.volatile.global.b32 [g], 117507841;
+; CHECK-NEXT: st.volatile.global.b32 [g], 1146447479;
+; CHECK-NEXT: st.volatile.global.b32 [g], 0;
+; CHECK-NEXT: st.volatile.global.b32 [g], -16711936;
+; CHECK-NEXT: ret;
+ %v1 = call i32 @llvm.nvvm.prmt(i32 u0x03020100, i32 u0x07060504, i32 u0x0123)
+ store volatile i32 %v1, ptr @g
+ %v2 = call i32 @llvm.nvvm.prmt(i32 u0x03020100, i32 u0x07060504, i32 u0x7171)
+ store volatile i32 %v2, ptr @g
+ %v3 = call i32 @llvm.nvvm.prmt(i32 u0x33221100, i32 u0x77665544, i32 u0x4567)
+ store volatile i32 %v3, ptr @g
+ %v4 = call i32 @llvm.nvvm.prmt(i32 u0x33221100, i32 u0x77665544, i32 u0xBA98)
+ store volatile i32 %v4, ptr @g
+ %v5 = call i32 @llvm.nvvm.prmt(i32 u0xF322F100, i32 u0x77665544, i32 u0xBA98)
+ store volatile i32 %v5, ptr @g
+ ret void
+}
>From 81a2e32ce1f84b210c8117710f3b83a4276be4cf Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Thu, 17 Jul 2025 16:50:01 +0000
Subject: [PATCH 3/3] address comments
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 64 ++++++++++++---------
1 file changed, 36 insertions(+), 28 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index d84e4fd0eb40a..7aa06f9079b09 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -2070,6 +2070,12 @@ static SDValue getPRMT(SDValue A, SDValue B, SDValue Selector, SDLoc DL,
{A, B, Selector, DAG.getConstant(Mode, DL, MVT::i32)});
}
+static SDValue getPRMT(SDValue A, SDValue B, uint64_t Selector, SDLoc DL,
+ SelectionDAG &DAG,
+ unsigned Mode = NVPTX::PTXPrmtMode::NONE) {
+ return getPRMT(A, B, DAG.getConstant(Selector, DL, MVT::i32), DL, DAG, Mode);
+}
+
SDValue NVPTXTargetLowering::LowerBITCAST(SDValue Op, SelectionDAG &DAG) const {
// Handle bitcasting from v2i8 without hitting the default promotion
// strategy which goes through stack memory.
@@ -2121,8 +2127,7 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op,
L = DAG.getAnyExtOrTrunc(L, DL, MVT::i32);
R = DAG.getAnyExtOrTrunc(R, DL, MVT::i32);
}
- return getPRMT(L, R, DAG.getConstant(SelectionValue, DL, MVT::i32), DL,
- DAG);
+ return getPRMT(L, R, SelectionValue, DL, DAG);
};
auto PRMT__10 = GetPRMT(Op->getOperand(0), Op->getOperand(1), true, 0x3340);
auto PRMT__32 = GetPRMT(Op->getOperand(2), Op->getOperand(3), true, 0x3340);
@@ -2253,9 +2258,8 @@ SDValue NVPTXTargetLowering::LowerVECTOR_SHUFFLE(SDValue Op,
}
SDLoc DL(Op);
- SDValue PRMT =
- getPRMT(DAG.getBitcast(MVT::i32, V1), DAG.getBitcast(MVT::i32, V2),
- DAG.getConstant(Selector, DL, MVT::i32), DL, DAG);
+ SDValue PRMT = getPRMT(DAG.getBitcast(MVT::i32, V1),
+ DAG.getBitcast(MVT::i32, V2), Selector, DL, DAG);
return DAG.getBitcast(Op.getValueType(), PRMT);
}
/// LowerShiftRightParts - Lower SRL_PARTS, SRA_PARTS, which
@@ -5823,9 +5827,9 @@ PerformBUILD_VECTORCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
SDLoc DL(N);
auto &DAG = DCI.DAG;
- auto PRMT = getPRMT(
- DAG.getBitcast(MVT::i32, Op0), DAG.getBitcast(MVT::i32, Op1),
- DAG.getConstant((Op1Bytes << 8) | Op0Bytes, DL, MVT::i32), DL, DAG);
+ auto PRMT =
+ getPRMT(DAG.getBitcast(MVT::i32, Op0), DAG.getBitcast(MVT::i32, Op1),
+ (Op1Bytes << 8) | Op0Bytes, DL, DAG);
return DAG.getBitcast(VT, PRMT);
}
@@ -5844,11 +5848,15 @@ static SDValue combineADDRSPACECAST(SDNode *N,
return SDValue();
}
-static APInt getPRMTSelector(APInt Selector, unsigned Mode) {
+// Given a constant selector value and a prmt mode, return the selector value
+// normalized to the generic prmt mode. See the PTX ISA documentation for more
+// details:
+// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-prmt
+static APInt getPRMTSelector(const APInt &Selector, unsigned Mode) {
if (Mode == NVPTX::PTXPrmtMode::NONE)
return Selector;
- unsigned V = Selector.trunc(2).getZExtValue();
+ const unsigned V = Selector.trunc(2).getZExtValue();
const auto GetSelector = [](unsigned S0, unsigned S1, unsigned S2,
unsigned S3) {
@@ -5918,24 +5926,32 @@ SDValue NVPTXTargetLowering::PerformDAGCombine(SDNode *N,
break;
case ISD::ADD:
return PerformADDCombine(N, DCI, OptLevel);
+ case ISD::ADDRSPACECAST:
+ return combineADDRSPACECAST(N, DCI);
+ case ISD::AND:
+ return PerformANDCombine(N, DCI);
+ case ISD::BUILD_VECTOR:
+ return PerformBUILD_VECTORCombine(N, DCI);
+ case ISD::EXTRACT_VECTOR_ELT:
+ return PerformEXTRACTCombine(N, DCI);
case ISD::FADD:
return PerformFADDCombine(N, DCI, OptLevel);
+ case ISD::LOAD:
+ case NVPTXISD::LoadParamV2:
+ case NVPTXISD::LoadV2:
+ case NVPTXISD::LoadV4:
+ return combineUnpackingMovIntoLoad(N, DCI);
case ISD::MUL:
return PerformMULCombine(N, DCI, OptLevel);
+ case NVPTXISD::PRMT:
+ return combinePRMT(N, DCI, OptLevel);
+ case ISD::SETCC:
+ return PerformSETCCCombine(N, DCI, STI.getSmVersion());
case ISD::SHL:
return PerformSHLCombine(N, DCI, OptLevel);
- case ISD::AND:
- return PerformANDCombine(N, DCI);
- case ISD::UREM:
case ISD::SREM:
+ case ISD::UREM:
return PerformREMCombine(N, DCI, OptLevel);
- case ISD::SETCC:
- return PerformSETCCCombine(N, DCI, STI.getSmVersion());
- case ISD::LOAD:
- case NVPTXISD::LoadParamV2:
- case NVPTXISD::LoadV2:
- case NVPTXISD::LoadV4:
- return combineUnpackingMovIntoLoad(N, DCI);
case NVPTXISD::StoreParam:
case NVPTXISD::StoreParamV2:
case NVPTXISD::StoreParamV4:
@@ -5944,16 +5960,8 @@ SDValue NVPTXTargetLowering::PerformDAGCombine(SDNode *N,
case NVPTXISD::StoreV2:
case NVPTXISD::StoreV4:
return PerformStoreCombine(N, DCI);
- case ISD::EXTRACT_VECTOR_ELT:
- return PerformEXTRACTCombine(N, DCI);
case ISD::VSELECT:
return PerformVSELECTCombine(N, DCI);
- case ISD::BUILD_VECTOR:
- return PerformBUILD_VECTORCombine(N, DCI);
- case ISD::ADDRSPACECAST:
- return combineADDRSPACECAST(N, DCI);
- case NVPTXISD::PRMT:
- return combinePRMT(N, DCI, OptLevel);
}
return SDValue();
}
More information about the llvm-commits
mailing list