2385 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 32; ## Support i16x2 instructions3; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \4; RUN: | FileCheck %s --check-prefixes=CHECK,O05; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -verify-machineinstrs \6; RUN: | FileCheck %s --check-prefixes=CHECK,O37; RUN: %if ptxas %{ \8; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \9; RUN: | %ptxas-verify -arch=sm_90 \10; RUN: %}11; RUN: %if ptxas %{ \12; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -verify-machineinstrs \13; RUN: | %ptxas-verify -arch=sm_90 \14; RUN: %}15 16target triple = "nvptx64-nvidia-cuda"17target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"18 19define <4 x i8> @test_ret_const() #0 {20; CHECK-LABEL: test_ret_const(21; CHECK: {22; CHECK-EMPTY:23; CHECK-EMPTY:24; CHECK-NEXT: // %bb.0:25; CHECK-NEXT: st.param.b32 [func_retval0], -66911489;26; CHECK-NEXT: ret;27 ret <4 x i8> <i8 -1, i8 2, i8 3, i8 -4>28}29 30define i8 @test_extract_0(<4 x i8> %a) #0 {31; CHECK-LABEL: test_extract_0(32; CHECK: {33; CHECK-NEXT: .reg .b32 %r<3>;34; CHECK-EMPTY:35; CHECK-NEXT: // %bb.0:36; CHECK-NEXT: ld.param.b32 %r1, [test_extract_0_param_0];37; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7770U;38; CHECK-NEXT: st.param.b32 [func_retval0], %r2;39; CHECK-NEXT: ret;40 %e = extractelement <4 x i8> %a, i32 041 ret i8 %e42}43 44define i8 @test_extract_1(<4 x i8> %a) #0 {45; CHECK-LABEL: test_extract_1(46; CHECK: {47; CHECK-NEXT: .reg .b32 %r<3>;48; CHECK-EMPTY:49; CHECK-NEXT: // %bb.0:50; CHECK-NEXT: ld.param.b32 %r1, [test_extract_1_param_0];51; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7771U;52; CHECK-NEXT: st.param.b32 [func_retval0], %r2;53; CHECK-NEXT: ret;54 %e = extractelement <4 x i8> %a, i32 155 ret i8 %e56}57 58define i8 @test_extract_2(<4 x i8> %a) #0 {59; CHECK-LABEL: test_extract_2(60; CHECK: {61; CHECK-NEXT: .reg .b32 %r<3>;62; CHECK-EMPTY:63; CHECK-NEXT: // %bb.0:64; CHECK-NEXT: ld.param.b32 %r1, [test_extract_2_param_0];65; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;66; CHECK-NEXT: st.param.b32 [func_retval0], %r2;67; CHECK-NEXT: ret;68 %e = extractelement <4 x i8> %a, i32 269 ret i8 %e70}71 72define i8 @test_extract_3(<4 x i8> %a) #0 {73; CHECK-LABEL: test_extract_3(74; CHECK: {75; CHECK-NEXT: .reg .b32 %r<3>;76; CHECK-EMPTY:77; CHECK-NEXT: // %bb.0:78; CHECK-NEXT: ld.param.b32 %r1, [test_extract_3_param_0];79; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;80; CHECK-NEXT: st.param.b32 [func_retval0], %r2;81; CHECK-NEXT: ret;82 %e = extractelement <4 x i8> %a, i32 383 ret i8 %e84}85 86define i8 @test_extract_i(<4 x i8> %a, i64 %idx) #0 {87; O0-LABEL: test_extract_i(88; O0: {89; O0-NEXT: .reg .b32 %r<5>;90; O0-NEXT: .reg .b64 %rd<2>;91; O0-EMPTY:92; O0-NEXT: // %bb.0:93; O0-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];94; O0-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];95; O0-NEXT: cvt.u32.u64 %r2, %rd1;96; O0-NEXT: or.b32 %r3, %r2, 30576;97; O0-NEXT: prmt.b32 %r4, %r1, 0, %r3;98; O0-NEXT: st.param.b32 [func_retval0], %r4;99; O0-NEXT: ret;100;101; O3-LABEL: test_extract_i(102; O3: {103; O3-NEXT: .reg .b32 %r<5>;104; O3-EMPTY:105; O3-NEXT: // %bb.0:106; O3-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];107; O3-NEXT: ld.param.b32 %r2, [test_extract_i_param_1];108; O3-NEXT: or.b32 %r3, %r2, 30576;109; O3-NEXT: prmt.b32 %r4, %r1, 0, %r3;110; O3-NEXT: st.param.b32 [func_retval0], %r4;111; O3-NEXT: ret;112 %e = extractelement <4 x i8> %a, i64 %idx113 ret i8 %e114}115 116define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 {117; O0-LABEL: test_add(118; O0: {119; O0-NEXT: .reg .b16 %rs<13>;120; O0-NEXT: .reg .b32 %r<18>;121; O0-EMPTY:122; O0-NEXT: // %bb.0:123; O0-NEXT: ld.param.b32 %r2, [test_add_param_1];124; O0-NEXT: ld.param.b32 %r1, [test_add_param_0];125; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;126; O0-NEXT: cvt.u16.u32 %rs1, %r3;127; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;128; O0-NEXT: cvt.u16.u32 %rs2, %r4;129; O0-NEXT: add.s16 %rs3, %rs2, %rs1;130; O0-NEXT: cvt.u32.u16 %r5, %rs3;131; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;132; O0-NEXT: cvt.u16.u32 %rs4, %r6;133; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;134; O0-NEXT: cvt.u16.u32 %rs5, %r7;135; O0-NEXT: add.s16 %rs6, %rs5, %rs4;136; O0-NEXT: cvt.u32.u16 %r8, %rs6;137; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;138; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;139; O0-NEXT: cvt.u16.u32 %rs7, %r10;140; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;141; O0-NEXT: cvt.u16.u32 %rs8, %r11;142; O0-NEXT: add.s16 %rs9, %rs8, %rs7;143; O0-NEXT: cvt.u32.u16 %r12, %rs9;144; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;145; O0-NEXT: cvt.u16.u32 %rs10, %r13;146; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;147; O0-NEXT: cvt.u16.u32 %rs11, %r14;148; O0-NEXT: add.s16 %rs12, %rs11, %rs10;149; O0-NEXT: cvt.u32.u16 %r15, %rs12;150; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;151; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;152; O0-NEXT: st.param.b32 [func_retval0], %r17;153; O0-NEXT: ret;154;155; O3-LABEL: test_add(156; O3: {157; O3-NEXT: .reg .b16 %rs<13>;158; O3-NEXT: .reg .b32 %r<18>;159; O3-EMPTY:160; O3-NEXT: // %bb.0:161; O3-NEXT: ld.param.b32 %r1, [test_add_param_0];162; O3-NEXT: ld.param.b32 %r2, [test_add_param_1];163; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;164; O3-NEXT: cvt.u16.u32 %rs1, %r3;165; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;166; O3-NEXT: cvt.u16.u32 %rs2, %r4;167; O3-NEXT: add.s16 %rs3, %rs2, %rs1;168; O3-NEXT: cvt.u32.u16 %r5, %rs3;169; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;170; O3-NEXT: cvt.u16.u32 %rs4, %r6;171; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;172; O3-NEXT: cvt.u16.u32 %rs5, %r7;173; O3-NEXT: add.s16 %rs6, %rs5, %rs4;174; O3-NEXT: cvt.u32.u16 %r8, %rs6;175; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;176; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;177; O3-NEXT: cvt.u16.u32 %rs7, %r10;178; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;179; O3-NEXT: cvt.u16.u32 %rs8, %r11;180; O3-NEXT: add.s16 %rs9, %rs8, %rs7;181; O3-NEXT: cvt.u32.u16 %r12, %rs9;182; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;183; O3-NEXT: cvt.u16.u32 %rs10, %r13;184; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;185; O3-NEXT: cvt.u16.u32 %rs11, %r14;186; O3-NEXT: add.s16 %rs12, %rs11, %rs10;187; O3-NEXT: cvt.u32.u16 %r15, %rs12;188; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;189; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;190; O3-NEXT: st.param.b32 [func_retval0], %r17;191; O3-NEXT: ret;192 %r = add <4 x i8> %a, %b193 ret <4 x i8> %r194}195 196define <4 x i8> @test_add_imm_0(<4 x i8> %a) #0 {197; CHECK-LABEL: test_add_imm_0(198; CHECK: {199; CHECK-NEXT: .reg .b16 %rs<9>;200; CHECK-NEXT: .reg .b32 %r<13>;201; CHECK-EMPTY:202; CHECK-NEXT: // %bb.0:203; CHECK-NEXT: ld.param.b32 %r1, [test_add_imm_0_param_0];204; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;205; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;206; CHECK-NEXT: add.s16 %rs2, %rs1, 4;207; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;208; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;209; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;210; CHECK-NEXT: add.s16 %rs4, %rs3, 3;211; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;212; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;213; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;214; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;215; CHECK-NEXT: add.s16 %rs6, %rs5, 2;216; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;217; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;218; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;219; CHECK-NEXT: add.s16 %rs8, %rs7, 1;220; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;221; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;222; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;223; CHECK-NEXT: st.param.b32 [func_retval0], %r12;224; CHECK-NEXT: ret;225 %r = add <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a226 ret <4 x i8> %r227}228 229define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 {230; CHECK-LABEL: test_add_imm_1(231; CHECK: {232; CHECK-NEXT: .reg .b16 %rs<9>;233; CHECK-NEXT: .reg .b32 %r<13>;234; CHECK-EMPTY:235; CHECK-NEXT: // %bb.0:236; CHECK-NEXT: ld.param.b32 %r1, [test_add_imm_1_param_0];237; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;238; CHECK-NEXT: cvt.u16.u32 %rs1, %r2;239; CHECK-NEXT: add.s16 %rs2, %rs1, 4;240; CHECK-NEXT: cvt.u32.u16 %r3, %rs2;241; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;242; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;243; CHECK-NEXT: add.s16 %rs4, %rs3, 3;244; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;245; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 0x3340U;246; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;247; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;248; CHECK-NEXT: add.s16 %rs6, %rs5, 2;249; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;250; CHECK-NEXT: prmt.b32 %r9, %r1, 0, 0x7770U;251; CHECK-NEXT: cvt.u16.u32 %rs7, %r9;252; CHECK-NEXT: add.s16 %rs8, %rs7, 1;253; CHECK-NEXT: cvt.u32.u16 %r10, %rs8;254; CHECK-NEXT: prmt.b32 %r11, %r10, %r8, 0x3340U;255; CHECK-NEXT: prmt.b32 %r12, %r11, %r6, 0x5410U;256; CHECK-NEXT: st.param.b32 [func_retval0], %r12;257; CHECK-NEXT: ret;258 %r = add <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4>259 ret <4 x i8> %r260}261 262define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 {263; O0-LABEL: test_sub(264; O0: {265; O0-NEXT: .reg .b16 %rs<13>;266; O0-NEXT: .reg .b32 %r<18>;267; O0-EMPTY:268; O0-NEXT: // %bb.0:269; O0-NEXT: ld.param.b32 %r2, [test_sub_param_1];270; O0-NEXT: ld.param.b32 %r1, [test_sub_param_0];271; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;272; O0-NEXT: cvt.u16.u32 %rs1, %r3;273; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;274; O0-NEXT: cvt.u16.u32 %rs2, %r4;275; O0-NEXT: sub.s16 %rs3, %rs2, %rs1;276; O0-NEXT: cvt.u32.u16 %r5, %rs3;277; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;278; O0-NEXT: cvt.u16.u32 %rs4, %r6;279; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;280; O0-NEXT: cvt.u16.u32 %rs5, %r7;281; O0-NEXT: sub.s16 %rs6, %rs5, %rs4;282; O0-NEXT: cvt.u32.u16 %r8, %rs6;283; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;284; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;285; O0-NEXT: cvt.u16.u32 %rs7, %r10;286; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;287; O0-NEXT: cvt.u16.u32 %rs8, %r11;288; O0-NEXT: sub.s16 %rs9, %rs8, %rs7;289; O0-NEXT: cvt.u32.u16 %r12, %rs9;290; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;291; O0-NEXT: cvt.u16.u32 %rs10, %r13;292; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;293; O0-NEXT: cvt.u16.u32 %rs11, %r14;294; O0-NEXT: sub.s16 %rs12, %rs11, %rs10;295; O0-NEXT: cvt.u32.u16 %r15, %rs12;296; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;297; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;298; O0-NEXT: st.param.b32 [func_retval0], %r17;299; O0-NEXT: ret;300;301; O3-LABEL: test_sub(302; O3: {303; O3-NEXT: .reg .b16 %rs<13>;304; O3-NEXT: .reg .b32 %r<18>;305; O3-EMPTY:306; O3-NEXT: // %bb.0:307; O3-NEXT: ld.param.b32 %r1, [test_sub_param_0];308; O3-NEXT: ld.param.b32 %r2, [test_sub_param_1];309; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;310; O3-NEXT: cvt.u16.u32 %rs1, %r3;311; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;312; O3-NEXT: cvt.u16.u32 %rs2, %r4;313; O3-NEXT: sub.s16 %rs3, %rs2, %rs1;314; O3-NEXT: cvt.u32.u16 %r5, %rs3;315; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;316; O3-NEXT: cvt.u16.u32 %rs4, %r6;317; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;318; O3-NEXT: cvt.u16.u32 %rs5, %r7;319; O3-NEXT: sub.s16 %rs6, %rs5, %rs4;320; O3-NEXT: cvt.u32.u16 %r8, %rs6;321; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;322; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;323; O3-NEXT: cvt.u16.u32 %rs7, %r10;324; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;325; O3-NEXT: cvt.u16.u32 %rs8, %r11;326; O3-NEXT: sub.s16 %rs9, %rs8, %rs7;327; O3-NEXT: cvt.u32.u16 %r12, %rs9;328; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;329; O3-NEXT: cvt.u16.u32 %rs10, %r13;330; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;331; O3-NEXT: cvt.u16.u32 %rs11, %r14;332; O3-NEXT: sub.s16 %rs12, %rs11, %rs10;333; O3-NEXT: cvt.u32.u16 %r15, %rs12;334; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;335; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;336; O3-NEXT: st.param.b32 [func_retval0], %r17;337; O3-NEXT: ret;338 %r = sub <4 x i8> %a, %b339 ret <4 x i8> %r340}341 342define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 {343; O0-LABEL: test_smax(344; O0: {345; O0-NEXT: .reg .pred %p<5>;346; O0-NEXT: .reg .b32 %r<26>;347; O0-EMPTY:348; O0-NEXT: // %bb.0:349; O0-NEXT: ld.param.b32 %r2, [test_smax_param_1];350; O0-NEXT: ld.param.b32 %r1, [test_smax_param_0];351; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x8880U;352; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x8880U;353; O0-NEXT: setp.gt.s32 %p1, %r4, %r3;354; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x9991U;355; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x9991U;356; O0-NEXT: setp.gt.s32 %p2, %r6, %r5;357; O0-NEXT: prmt.b32 %r7, %r2, 0, 0xaaa2U;358; O0-NEXT: prmt.b32 %r8, %r1, 0, 0xaaa2U;359; O0-NEXT: setp.gt.s32 %p3, %r8, %r7;360; O0-NEXT: prmt.b32 %r9, %r2, 0, 0xbbb3U;361; O0-NEXT: prmt.b32 %r10, %r1, 0, 0xbbb3U;362; O0-NEXT: setp.gt.s32 %p4, %r10, %r9;363; O0-NEXT: prmt.b32 %r11, %r2, 0, 0x7770U;364; O0-NEXT: prmt.b32 %r12, %r2, 0, 0x7771U;365; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7772U;366; O0-NEXT: prmt.b32 %r14, %r2, 0, 0x7773U;367; O0-NEXT: prmt.b32 %r15, %r1, 0, 0x7773U;368; O0-NEXT: selp.b32 %r16, %r15, %r14, %p4;369; O0-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;370; O0-NEXT: selp.b32 %r18, %r17, %r13, %p3;371; O0-NEXT: prmt.b32 %r19, %r18, %r16, 0x3340U;372; O0-NEXT: prmt.b32 %r20, %r1, 0, 0x7771U;373; O0-NEXT: selp.b32 %r21, %r20, %r12, %p2;374; O0-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;375; O0-NEXT: selp.b32 %r23, %r22, %r11, %p1;376; O0-NEXT: prmt.b32 %r24, %r23, %r21, 0x3340U;377; O0-NEXT: prmt.b32 %r25, %r24, %r19, 0x5410U;378; O0-NEXT: st.param.b32 [func_retval0], %r25;379; O0-NEXT: ret;380;381; O3-LABEL: test_smax(382; O3: {383; O3-NEXT: .reg .pred %p<5>;384; O3-NEXT: .reg .b32 %r<26>;385; O3-EMPTY:386; O3-NEXT: // %bb.0:387; O3-NEXT: ld.param.b32 %r1, [test_smax_param_0];388; O3-NEXT: ld.param.b32 %r2, [test_smax_param_1];389; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x8880U;390; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x8880U;391; O3-NEXT: setp.gt.s32 %p1, %r4, %r3;392; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x9991U;393; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x9991U;394; O3-NEXT: setp.gt.s32 %p2, %r6, %r5;395; O3-NEXT: prmt.b32 %r7, %r2, 0, 0xaaa2U;396; O3-NEXT: prmt.b32 %r8, %r1, 0, 0xaaa2U;397; O3-NEXT: setp.gt.s32 %p3, %r8, %r7;398; O3-NEXT: prmt.b32 %r9, %r2, 0, 0xbbb3U;399; O3-NEXT: prmt.b32 %r10, %r1, 0, 0xbbb3U;400; O3-NEXT: setp.gt.s32 %p4, %r10, %r9;401; O3-NEXT: prmt.b32 %r11, %r2, 0, 0x7770U;402; O3-NEXT: prmt.b32 %r12, %r2, 0, 0x7771U;403; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7772U;404; O3-NEXT: prmt.b32 %r14, %r2, 0, 0x7773U;405; O3-NEXT: prmt.b32 %r15, %r1, 0, 0x7773U;406; O3-NEXT: selp.b32 %r16, %r15, %r14, %p4;407; O3-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;408; O3-NEXT: selp.b32 %r18, %r17, %r13, %p3;409; O3-NEXT: prmt.b32 %r19, %r18, %r16, 0x3340U;410; O3-NEXT: prmt.b32 %r20, %r1, 0, 0x7771U;411; O3-NEXT: selp.b32 %r21, %r20, %r12, %p2;412; O3-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;413; O3-NEXT: selp.b32 %r23, %r22, %r11, %p1;414; O3-NEXT: prmt.b32 %r24, %r23, %r21, 0x3340U;415; O3-NEXT: prmt.b32 %r25, %r24, %r19, 0x5410U;416; O3-NEXT: st.param.b32 [func_retval0], %r25;417; O3-NEXT: ret;418 %cmp = icmp sgt <4 x i8> %a, %b419 %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b420 ret <4 x i8> %r421}422 423define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 {424; O0-LABEL: test_umax(425; O0: {426; O0-NEXT: .reg .pred %p<5>;427; O0-NEXT: .reg .b32 %r<18>;428; O0-EMPTY:429; O0-NEXT: // %bb.0:430; O0-NEXT: ld.param.b32 %r2, [test_umax_param_1];431; O0-NEXT: ld.param.b32 %r1, [test_umax_param_0];432; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;433; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;434; O0-NEXT: setp.gt.u32 %p1, %r4, %r3;435; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;436; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;437; O0-NEXT: setp.gt.u32 %p2, %r6, %r5;438; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;439; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;440; O0-NEXT: setp.gt.u32 %p3, %r8, %r7;441; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;442; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;443; O0-NEXT: setp.gt.u32 %p4, %r10, %r9;444; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4;445; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3;446; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;447; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2;448; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1;449; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;450; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;451; O0-NEXT: st.param.b32 [func_retval0], %r17;452; O0-NEXT: ret;453;454; O3-LABEL: test_umax(455; O3: {456; O3-NEXT: .reg .pred %p<5>;457; O3-NEXT: .reg .b32 %r<18>;458; O3-EMPTY:459; O3-NEXT: // %bb.0:460; O3-NEXT: ld.param.b32 %r1, [test_umax_param_0];461; O3-NEXT: ld.param.b32 %r2, [test_umax_param_1];462; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;463; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;464; O3-NEXT: setp.gt.u32 %p1, %r4, %r3;465; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;466; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;467; O3-NEXT: setp.gt.u32 %p2, %r6, %r5;468; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;469; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;470; O3-NEXT: setp.gt.u32 %p3, %r8, %r7;471; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;472; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;473; O3-NEXT: setp.gt.u32 %p4, %r10, %r9;474; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4;475; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3;476; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;477; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2;478; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1;479; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;480; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;481; O3-NEXT: st.param.b32 [func_retval0], %r17;482; O3-NEXT: ret;483 %cmp = icmp ugt <4 x i8> %a, %b484 %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b485 ret <4 x i8> %r486}487 488define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 {489; O0-LABEL: test_smin(490; O0: {491; O0-NEXT: .reg .pred %p<5>;492; O0-NEXT: .reg .b32 %r<26>;493; O0-EMPTY:494; O0-NEXT: // %bb.0:495; O0-NEXT: ld.param.b32 %r2, [test_smin_param_1];496; O0-NEXT: ld.param.b32 %r1, [test_smin_param_0];497; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x8880U;498; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x8880U;499; O0-NEXT: setp.le.s32 %p1, %r4, %r3;500; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x9991U;501; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x9991U;502; O0-NEXT: setp.le.s32 %p2, %r6, %r5;503; O0-NEXT: prmt.b32 %r7, %r2, 0, 0xaaa2U;504; O0-NEXT: prmt.b32 %r8, %r1, 0, 0xaaa2U;505; O0-NEXT: setp.le.s32 %p3, %r8, %r7;506; O0-NEXT: prmt.b32 %r9, %r2, 0, 0xbbb3U;507; O0-NEXT: prmt.b32 %r10, %r1, 0, 0xbbb3U;508; O0-NEXT: setp.le.s32 %p4, %r10, %r9;509; O0-NEXT: prmt.b32 %r11, %r2, 0, 0x7770U;510; O0-NEXT: prmt.b32 %r12, %r2, 0, 0x7771U;511; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7772U;512; O0-NEXT: prmt.b32 %r14, %r2, 0, 0x7773U;513; O0-NEXT: prmt.b32 %r15, %r1, 0, 0x7773U;514; O0-NEXT: selp.b32 %r16, %r15, %r14, %p4;515; O0-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;516; O0-NEXT: selp.b32 %r18, %r17, %r13, %p3;517; O0-NEXT: prmt.b32 %r19, %r18, %r16, 0x3340U;518; O0-NEXT: prmt.b32 %r20, %r1, 0, 0x7771U;519; O0-NEXT: selp.b32 %r21, %r20, %r12, %p2;520; O0-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;521; O0-NEXT: selp.b32 %r23, %r22, %r11, %p1;522; O0-NEXT: prmt.b32 %r24, %r23, %r21, 0x3340U;523; O0-NEXT: prmt.b32 %r25, %r24, %r19, 0x5410U;524; O0-NEXT: st.param.b32 [func_retval0], %r25;525; O0-NEXT: ret;526;527; O3-LABEL: test_smin(528; O3: {529; O3-NEXT: .reg .pred %p<5>;530; O3-NEXT: .reg .b32 %r<26>;531; O3-EMPTY:532; O3-NEXT: // %bb.0:533; O3-NEXT: ld.param.b32 %r1, [test_smin_param_0];534; O3-NEXT: ld.param.b32 %r2, [test_smin_param_1];535; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x8880U;536; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x8880U;537; O3-NEXT: setp.le.s32 %p1, %r4, %r3;538; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x9991U;539; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x9991U;540; O3-NEXT: setp.le.s32 %p2, %r6, %r5;541; O3-NEXT: prmt.b32 %r7, %r2, 0, 0xaaa2U;542; O3-NEXT: prmt.b32 %r8, %r1, 0, 0xaaa2U;543; O3-NEXT: setp.le.s32 %p3, %r8, %r7;544; O3-NEXT: prmt.b32 %r9, %r2, 0, 0xbbb3U;545; O3-NEXT: prmt.b32 %r10, %r1, 0, 0xbbb3U;546; O3-NEXT: setp.le.s32 %p4, %r10, %r9;547; O3-NEXT: prmt.b32 %r11, %r2, 0, 0x7770U;548; O3-NEXT: prmt.b32 %r12, %r2, 0, 0x7771U;549; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7772U;550; O3-NEXT: prmt.b32 %r14, %r2, 0, 0x7773U;551; O3-NEXT: prmt.b32 %r15, %r1, 0, 0x7773U;552; O3-NEXT: selp.b32 %r16, %r15, %r14, %p4;553; O3-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;554; O3-NEXT: selp.b32 %r18, %r17, %r13, %p3;555; O3-NEXT: prmt.b32 %r19, %r18, %r16, 0x3340U;556; O3-NEXT: prmt.b32 %r20, %r1, 0, 0x7771U;557; O3-NEXT: selp.b32 %r21, %r20, %r12, %p2;558; O3-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;559; O3-NEXT: selp.b32 %r23, %r22, %r11, %p1;560; O3-NEXT: prmt.b32 %r24, %r23, %r21, 0x3340U;561; O3-NEXT: prmt.b32 %r25, %r24, %r19, 0x5410U;562; O3-NEXT: st.param.b32 [func_retval0], %r25;563; O3-NEXT: ret;564 %cmp = icmp sle <4 x i8> %a, %b565 %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b566 ret <4 x i8> %r567}568 569define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 {570; O0-LABEL: test_umin(571; O0: {572; O0-NEXT: .reg .pred %p<5>;573; O0-NEXT: .reg .b32 %r<18>;574; O0-EMPTY:575; O0-NEXT: // %bb.0:576; O0-NEXT: ld.param.b32 %r2, [test_umin_param_1];577; O0-NEXT: ld.param.b32 %r1, [test_umin_param_0];578; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;579; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;580; O0-NEXT: setp.le.u32 %p1, %r4, %r3;581; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;582; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;583; O0-NEXT: setp.le.u32 %p2, %r6, %r5;584; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;585; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;586; O0-NEXT: setp.le.u32 %p3, %r8, %r7;587; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;588; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;589; O0-NEXT: setp.le.u32 %p4, %r10, %r9;590; O0-NEXT: selp.b32 %r11, %r10, %r9, %p4;591; O0-NEXT: selp.b32 %r12, %r8, %r7, %p3;592; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;593; O0-NEXT: selp.b32 %r14, %r6, %r5, %p2;594; O0-NEXT: selp.b32 %r15, %r4, %r3, %p1;595; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;596; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;597; O0-NEXT: st.param.b32 [func_retval0], %r17;598; O0-NEXT: ret;599;600; O3-LABEL: test_umin(601; O3: {602; O3-NEXT: .reg .pred %p<5>;603; O3-NEXT: .reg .b32 %r<18>;604; O3-EMPTY:605; O3-NEXT: // %bb.0:606; O3-NEXT: ld.param.b32 %r1, [test_umin_param_0];607; O3-NEXT: ld.param.b32 %r2, [test_umin_param_1];608; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;609; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;610; O3-NEXT: setp.le.u32 %p1, %r4, %r3;611; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;612; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;613; O3-NEXT: setp.le.u32 %p2, %r6, %r5;614; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;615; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;616; O3-NEXT: setp.le.u32 %p3, %r8, %r7;617; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;618; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;619; O3-NEXT: setp.le.u32 %p4, %r10, %r9;620; O3-NEXT: selp.b32 %r11, %r10, %r9, %p4;621; O3-NEXT: selp.b32 %r12, %r8, %r7, %p3;622; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;623; O3-NEXT: selp.b32 %r14, %r6, %r5, %p2;624; O3-NEXT: selp.b32 %r15, %r4, %r3, %p1;625; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;626; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;627; O3-NEXT: st.param.b32 [func_retval0], %r17;628; O3-NEXT: ret;629 %cmp = icmp ule <4 x i8> %a, %b630 %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %b631 ret <4 x i8> %r632}633 634define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 {635; O0-LABEL: test_eq(636; O0: {637; O0-NEXT: .reg .pred %p<5>;638; O0-NEXT: .reg .b32 %r<23>;639; O0-EMPTY:640; O0-NEXT: // %bb.0:641; O0-NEXT: ld.param.b32 %r3, [test_eq_param_2];642; O0-NEXT: ld.param.b32 %r2, [test_eq_param_1];643; O0-NEXT: ld.param.b32 %r1, [test_eq_param_0];644; O0-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U;645; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;646; O0-NEXT: setp.eq.b32 %p1, %r5, %r4;647; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;648; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;649; O0-NEXT: setp.eq.b32 %p2, %r7, %r6;650; O0-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;651; O0-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U;652; O0-NEXT: setp.eq.b32 %p3, %r9, %r8;653; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;654; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U;655; O0-NEXT: setp.eq.b32 %p4, %r11, %r10;656; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;657; O0-NEXT: selp.b32 %r13, %r11, %r12, %p4;658; O0-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U;659; O0-NEXT: selp.b32 %r15, %r9, %r14, %p3;660; O0-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;661; O0-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U;662; O0-NEXT: selp.b32 %r18, %r7, %r17, %p2;663; O0-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U;664; O0-NEXT: selp.b32 %r20, %r5, %r19, %p1;665; O0-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;666; O0-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;667; O0-NEXT: st.param.b32 [func_retval0], %r22;668; O0-NEXT: ret;669;670; O3-LABEL: test_eq(671; O3: {672; O3-NEXT: .reg .pred %p<5>;673; O3-NEXT: .reg .b32 %r<23>;674; O3-EMPTY:675; O3-NEXT: // %bb.0:676; O3-NEXT: ld.param.b32 %r1, [test_eq_param_0];677; O3-NEXT: ld.param.b32 %r2, [test_eq_param_1];678; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;679; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;680; O3-NEXT: setp.eq.b32 %p1, %r4, %r3;681; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;682; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;683; O3-NEXT: setp.eq.b32 %p2, %r6, %r5;684; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;685; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;686; O3-NEXT: setp.eq.b32 %p3, %r8, %r7;687; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;688; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;689; O3-NEXT: setp.eq.b32 %p4, %r10, %r9;690; O3-NEXT: ld.param.b32 %r11, [test_eq_param_2];691; O3-NEXT: prmt.b32 %r12, %r11, 0, 0x7773U;692; O3-NEXT: selp.b32 %r13, %r10, %r12, %p4;693; O3-NEXT: prmt.b32 %r14, %r11, 0, 0x7772U;694; O3-NEXT: selp.b32 %r15, %r8, %r14, %p3;695; O3-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;696; O3-NEXT: prmt.b32 %r17, %r11, 0, 0x7771U;697; O3-NEXT: selp.b32 %r18, %r6, %r17, %p2;698; O3-NEXT: prmt.b32 %r19, %r11, 0, 0x7770U;699; O3-NEXT: selp.b32 %r20, %r4, %r19, %p1;700; O3-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;701; O3-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;702; O3-NEXT: st.param.b32 [func_retval0], %r22;703; O3-NEXT: ret;704 %cmp = icmp eq <4 x i8> %a, %b705 %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c706 ret <4 x i8> %r707}708 709define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 {710; O0-LABEL: test_ne(711; O0: {712; O0-NEXT: .reg .pred %p<5>;713; O0-NEXT: .reg .b32 %r<23>;714; O0-EMPTY:715; O0-NEXT: // %bb.0:716; O0-NEXT: ld.param.b32 %r3, [test_ne_param_2];717; O0-NEXT: ld.param.b32 %r2, [test_ne_param_1];718; O0-NEXT: ld.param.b32 %r1, [test_ne_param_0];719; O0-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U;720; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;721; O0-NEXT: setp.ne.b32 %p1, %r5, %r4;722; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;723; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7771U;724; O0-NEXT: setp.ne.b32 %p2, %r7, %r6;725; O0-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;726; O0-NEXT: prmt.b32 %r9, %r1, 0, 0x7772U;727; O0-NEXT: setp.ne.b32 %p3, %r9, %r8;728; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;729; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7773U;730; O0-NEXT: setp.ne.b32 %p4, %r11, %r10;731; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;732; O0-NEXT: selp.b32 %r13, %r11, %r12, %p4;733; O0-NEXT: prmt.b32 %r14, %r3, 0, 0x7772U;734; O0-NEXT: selp.b32 %r15, %r9, %r14, %p3;735; O0-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;736; O0-NEXT: prmt.b32 %r17, %r3, 0, 0x7771U;737; O0-NEXT: selp.b32 %r18, %r7, %r17, %p2;738; O0-NEXT: prmt.b32 %r19, %r3, 0, 0x7770U;739; O0-NEXT: selp.b32 %r20, %r5, %r19, %p1;740; O0-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;741; O0-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;742; O0-NEXT: st.param.b32 [func_retval0], %r22;743; O0-NEXT: ret;744;745; O3-LABEL: test_ne(746; O3: {747; O3-NEXT: .reg .pred %p<5>;748; O3-NEXT: .reg .b32 %r<23>;749; O3-EMPTY:750; O3-NEXT: // %bb.0:751; O3-NEXT: ld.param.b32 %r1, [test_ne_param_0];752; O3-NEXT: ld.param.b32 %r2, [test_ne_param_1];753; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;754; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;755; O3-NEXT: setp.ne.b32 %p1, %r4, %r3;756; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;757; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;758; O3-NEXT: setp.ne.b32 %p2, %r6, %r5;759; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;760; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;761; O3-NEXT: setp.ne.b32 %p3, %r8, %r7;762; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;763; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;764; O3-NEXT: setp.ne.b32 %p4, %r10, %r9;765; O3-NEXT: ld.param.b32 %r11, [test_ne_param_2];766; O3-NEXT: prmt.b32 %r12, %r11, 0, 0x7773U;767; O3-NEXT: selp.b32 %r13, %r10, %r12, %p4;768; O3-NEXT: prmt.b32 %r14, %r11, 0, 0x7772U;769; O3-NEXT: selp.b32 %r15, %r8, %r14, %p3;770; O3-NEXT: prmt.b32 %r16, %r15, %r13, 0x3340U;771; O3-NEXT: prmt.b32 %r17, %r11, 0, 0x7771U;772; O3-NEXT: selp.b32 %r18, %r6, %r17, %p2;773; O3-NEXT: prmt.b32 %r19, %r11, 0, 0x7770U;774; O3-NEXT: selp.b32 %r20, %r4, %r19, %p1;775; O3-NEXT: prmt.b32 %r21, %r20, %r18, 0x3340U;776; O3-NEXT: prmt.b32 %r22, %r21, %r16, 0x5410U;777; O3-NEXT: st.param.b32 [func_retval0], %r22;778; O3-NEXT: ret;779 %cmp = icmp ne <4 x i8> %a, %b780 %r = select <4 x i1> %cmp, <4 x i8> %a, <4 x i8> %c781 ret <4 x i8> %r782}783 784define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 {785; O0-LABEL: test_mul(786; O0: {787; O0-NEXT: .reg .b16 %rs<13>;788; O0-NEXT: .reg .b32 %r<18>;789; O0-EMPTY:790; O0-NEXT: // %bb.0:791; O0-NEXT: ld.param.b32 %r2, [test_mul_param_1];792; O0-NEXT: ld.param.b32 %r1, [test_mul_param_0];793; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;794; O0-NEXT: cvt.u16.u32 %rs1, %r3;795; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;796; O0-NEXT: cvt.u16.u32 %rs2, %r4;797; O0-NEXT: mul.lo.s16 %rs3, %rs2, %rs1;798; O0-NEXT: cvt.u32.u16 %r5, %rs3;799; O0-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;800; O0-NEXT: cvt.u16.u32 %rs4, %r6;801; O0-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;802; O0-NEXT: cvt.u16.u32 %rs5, %r7;803; O0-NEXT: mul.lo.s16 %rs6, %rs5, %rs4;804; O0-NEXT: cvt.u32.u16 %r8, %rs6;805; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;806; O0-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;807; O0-NEXT: cvt.u16.u32 %rs7, %r10;808; O0-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;809; O0-NEXT: cvt.u16.u32 %rs8, %r11;810; O0-NEXT: mul.lo.s16 %rs9, %rs8, %rs7;811; O0-NEXT: cvt.u32.u16 %r12, %rs9;812; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;813; O0-NEXT: cvt.u16.u32 %rs10, %r13;814; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;815; O0-NEXT: cvt.u16.u32 %rs11, %r14;816; O0-NEXT: mul.lo.s16 %rs12, %rs11, %rs10;817; O0-NEXT: cvt.u32.u16 %r15, %rs12;818; O0-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;819; O0-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;820; O0-NEXT: st.param.b32 [func_retval0], %r17;821; O0-NEXT: ret;822;823; O3-LABEL: test_mul(824; O3: {825; O3-NEXT: .reg .b16 %rs<13>;826; O3-NEXT: .reg .b32 %r<18>;827; O3-EMPTY:828; O3-NEXT: // %bb.0:829; O3-NEXT: ld.param.b32 %r1, [test_mul_param_0];830; O3-NEXT: ld.param.b32 %r2, [test_mul_param_1];831; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7773U;832; O3-NEXT: cvt.u16.u32 %rs1, %r3;833; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;834; O3-NEXT: cvt.u16.u32 %rs2, %r4;835; O3-NEXT: mul.lo.s16 %rs3, %rs2, %rs1;836; O3-NEXT: cvt.u32.u16 %r5, %rs3;837; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7772U;838; O3-NEXT: cvt.u16.u32 %rs4, %r6;839; O3-NEXT: prmt.b32 %r7, %r1, 0, 0x7772U;840; O3-NEXT: cvt.u16.u32 %rs5, %r7;841; O3-NEXT: mul.lo.s16 %rs6, %rs5, %rs4;842; O3-NEXT: cvt.u32.u16 %r8, %rs6;843; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;844; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7771U;845; O3-NEXT: cvt.u16.u32 %rs7, %r10;846; O3-NEXT: prmt.b32 %r11, %r1, 0, 0x7771U;847; O3-NEXT: cvt.u16.u32 %rs8, %r11;848; O3-NEXT: mul.lo.s16 %rs9, %rs8, %rs7;849; O3-NEXT: cvt.u32.u16 %r12, %rs9;850; O3-NEXT: prmt.b32 %r13, %r2, 0, 0x7770U;851; O3-NEXT: cvt.u16.u32 %rs10, %r13;852; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7770U;853; O3-NEXT: cvt.u16.u32 %rs11, %r14;854; O3-NEXT: mul.lo.s16 %rs12, %rs11, %rs10;855; O3-NEXT: cvt.u32.u16 %r15, %rs12;856; O3-NEXT: prmt.b32 %r16, %r15, %r12, 0x3340U;857; O3-NEXT: prmt.b32 %r17, %r16, %r9, 0x5410U;858; O3-NEXT: st.param.b32 [func_retval0], %r17;859; O3-NEXT: ret;860 %r = mul <4 x i8> %a, %b861 ret <4 x i8> %r862}863 864define <4 x i8> @test_or(<4 x i8> %a, <4 x i8> %b) #0 {865; O0-LABEL: test_or(866; O0: {867; O0-NEXT: .reg .b32 %r<4>;868; O0-EMPTY:869; O0-NEXT: // %bb.0:870; O0-NEXT: ld.param.b32 %r2, [test_or_param_1];871; O0-NEXT: ld.param.b32 %r1, [test_or_param_0];872; O0-NEXT: or.b32 %r3, %r1, %r2;873; O0-NEXT: st.param.b32 [func_retval0], %r3;874; O0-NEXT: ret;875;876; O3-LABEL: test_or(877; O3: {878; O3-NEXT: .reg .b32 %r<4>;879; O3-EMPTY:880; O3-NEXT: // %bb.0:881; O3-NEXT: ld.param.b32 %r1, [test_or_param_0];882; O3-NEXT: ld.param.b32 %r2, [test_or_param_1];883; O3-NEXT: or.b32 %r3, %r1, %r2;884; O3-NEXT: st.param.b32 [func_retval0], %r3;885; O3-NEXT: ret;886 %r = or <4 x i8> %a, %b887 ret <4 x i8> %r888}889 890define <4 x i8> @test_or_computed(i8 %a) {891; O0-LABEL: test_or_computed(892; O0: {893; O0-NEXT: .reg .b16 %rs<2>;894; O0-NEXT: .reg .b32 %r<8>;895; O0-EMPTY:896; O0-NEXT: // %bb.0:897; O0-NEXT: ld.param.b8 %rs1, [test_or_computed_param_0];898; O0-NEXT: mov.b32 %r1, 0;899; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;900; O0-NEXT: cvt.u32.u16 %r3, %rs1;901; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;902; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;903; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;904; O0-NEXT: or.b32 %r7, %r6, %r5;905; O0-NEXT: st.param.b32 [func_retval0], %r7;906; O0-NEXT: ret;907;908; O3-LABEL: test_or_computed(909; O3: {910; O3-NEXT: .reg .b32 %r<6>;911; O3-EMPTY:912; O3-NEXT: // %bb.0:913; O3-NEXT: ld.param.b8 %r1, [test_or_computed_param_0];914; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;915; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U;916; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8;917; O3-NEXT: or.b32 %r5, %r4, %r3;918; O3-NEXT: st.param.b32 [func_retval0], %r5;919; O3-NEXT: ret;920 %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 0921 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 1922 %r = or <4 x i8> %ins.1, %ins.0923 ret <4 x i8> %r924}925 926define <4 x i8> @test_or_imm_0(<4 x i8> %a) #0 {927; CHECK-LABEL: test_or_imm_0(928; CHECK: {929; CHECK-NEXT: .reg .b32 %r<3>;930; CHECK-EMPTY:931; CHECK-NEXT: // %bb.0:932; CHECK-NEXT: ld.param.b32 %r1, [test_or_imm_0_param_0];933; CHECK-NEXT: or.b32 %r2, %r1, 67305985;934; CHECK-NEXT: st.param.b32 [func_retval0], %r2;935; CHECK-NEXT: ret;936 %r = or <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a937 ret <4 x i8> %r938}939 940define <4 x i8> @test_or_imm_1(<4 x i8> %a) #0 {941; CHECK-LABEL: test_or_imm_1(942; CHECK: {943; CHECK-NEXT: .reg .b32 %r<3>;944; CHECK-EMPTY:945; CHECK-NEXT: // %bb.0:946; CHECK-NEXT: ld.param.b32 %r1, [test_or_imm_1_param_0];947; CHECK-NEXT: or.b32 %r2, %r1, 67305985;948; CHECK-NEXT: st.param.b32 [func_retval0], %r2;949; CHECK-NEXT: ret;950 %r = or <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4>951 ret <4 x i8> %r952}953 954define <4 x i8> @test_xor(<4 x i8> %a, <4 x i8> %b) #0 {955; O0-LABEL: test_xor(956; O0: {957; O0-NEXT: .reg .b32 %r<4>;958; O0-EMPTY:959; O0-NEXT: // %bb.0:960; O0-NEXT: ld.param.b32 %r2, [test_xor_param_1];961; O0-NEXT: ld.param.b32 %r1, [test_xor_param_0];962; O0-NEXT: xor.b32 %r3, %r1, %r2;963; O0-NEXT: st.param.b32 [func_retval0], %r3;964; O0-NEXT: ret;965;966; O3-LABEL: test_xor(967; O3: {968; O3-NEXT: .reg .b32 %r<4>;969; O3-EMPTY:970; O3-NEXT: // %bb.0:971; O3-NEXT: ld.param.b32 %r1, [test_xor_param_0];972; O3-NEXT: ld.param.b32 %r2, [test_xor_param_1];973; O3-NEXT: xor.b32 %r3, %r1, %r2;974; O3-NEXT: st.param.b32 [func_retval0], %r3;975; O3-NEXT: ret;976 %r = xor <4 x i8> %a, %b977 ret <4 x i8> %r978}979 980define <4 x i8> @test_xor_computed(i8 %a) {981; O0-LABEL: test_xor_computed(982; O0: {983; O0-NEXT: .reg .b16 %rs<2>;984; O0-NEXT: .reg .b32 %r<8>;985; O0-EMPTY:986; O0-NEXT: // %bb.0:987; O0-NEXT: ld.param.b8 %rs1, [test_xor_computed_param_0];988; O0-NEXT: mov.b32 %r1, 0;989; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;990; O0-NEXT: cvt.u32.u16 %r3, %rs1;991; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;992; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;993; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;994; O0-NEXT: xor.b32 %r7, %r6, %r5;995; O0-NEXT: st.param.b32 [func_retval0], %r7;996; O0-NEXT: ret;997;998; O3-LABEL: test_xor_computed(999; O3: {1000; O3-NEXT: .reg .b32 %r<6>;1001; O3-EMPTY:1002; O3-NEXT: // %bb.0:1003; O3-NEXT: ld.param.b8 %r1, [test_xor_computed_param_0];1004; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;1005; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U;1006; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8;1007; O3-NEXT: xor.b32 %r5, %r4, %r3;1008; O3-NEXT: st.param.b32 [func_retval0], %r5;1009; O3-NEXT: ret;1010 %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 01011 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 11012 %r = xor <4 x i8> %ins.1, %ins.01013 ret <4 x i8> %r1014}1015 1016define <4 x i8> @test_xor_imm_0(<4 x i8> %a) #0 {1017; CHECK-LABEL: test_xor_imm_0(1018; CHECK: {1019; CHECK-NEXT: .reg .b32 %r<3>;1020; CHECK-EMPTY:1021; CHECK-NEXT: // %bb.0:1022; CHECK-NEXT: ld.param.b32 %r1, [test_xor_imm_0_param_0];1023; CHECK-NEXT: xor.b32 %r2, %r1, 67305985;1024; CHECK-NEXT: st.param.b32 [func_retval0], %r2;1025; CHECK-NEXT: ret;1026 %r = xor <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a1027 ret <4 x i8> %r1028}1029 1030define <4 x i8> @test_xor_imm_1(<4 x i8> %a) #0 {1031; CHECK-LABEL: test_xor_imm_1(1032; CHECK: {1033; CHECK-NEXT: .reg .b32 %r<3>;1034; CHECK-EMPTY:1035; CHECK-NEXT: // %bb.0:1036; CHECK-NEXT: ld.param.b32 %r1, [test_xor_imm_1_param_0];1037; CHECK-NEXT: xor.b32 %r2, %r1, 67305985;1038; CHECK-NEXT: st.param.b32 [func_retval0], %r2;1039; CHECK-NEXT: ret;1040 %r = xor <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4>1041 ret <4 x i8> %r1042}1043 1044define <4 x i8> @test_and(<4 x i8> %a, <4 x i8> %b) #0 {1045; O0-LABEL: test_and(1046; O0: {1047; O0-NEXT: .reg .b32 %r<4>;1048; O0-EMPTY:1049; O0-NEXT: // %bb.0:1050; O0-NEXT: ld.param.b32 %r2, [test_and_param_1];1051; O0-NEXT: ld.param.b32 %r1, [test_and_param_0];1052; O0-NEXT: and.b32 %r3, %r1, %r2;1053; O0-NEXT: st.param.b32 [func_retval0], %r3;1054; O0-NEXT: ret;1055;1056; O3-LABEL: test_and(1057; O3: {1058; O3-NEXT: .reg .b32 %r<4>;1059; O3-EMPTY:1060; O3-NEXT: // %bb.0:1061; O3-NEXT: ld.param.b32 %r1, [test_and_param_0];1062; O3-NEXT: ld.param.b32 %r2, [test_and_param_1];1063; O3-NEXT: and.b32 %r3, %r1, %r2;1064; O3-NEXT: st.param.b32 [func_retval0], %r3;1065; O3-NEXT: ret;1066 %r = and <4 x i8> %a, %b1067 ret <4 x i8> %r1068}1069 1070define <4 x i8> @test_and_computed(i8 %a) {1071; O0-LABEL: test_and_computed(1072; O0: {1073; O0-NEXT: .reg .b16 %rs<2>;1074; O0-NEXT: .reg .b32 %r<8>;1075; O0-EMPTY:1076; O0-NEXT: // %bb.0:1077; O0-NEXT: ld.param.b8 %rs1, [test_and_computed_param_0];1078; O0-NEXT: mov.b32 %r1, 0;1079; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;1080; O0-NEXT: cvt.u32.u16 %r3, %rs1;1081; O0-NEXT: prmt.b32 %r4, %r3, 0, 0x3340U;1082; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;1083; O0-NEXT: bfi.b32 %r6, 5, %r5, 8, 8;1084; O0-NEXT: and.b32 %r7, %r6, %r5;1085; O0-NEXT: st.param.b32 [func_retval0], %r7;1086; O0-NEXT: ret;1087;1088; O3-LABEL: test_and_computed(1089; O3: {1090; O3-NEXT: .reg .b32 %r<6>;1091; O3-EMPTY:1092; O3-NEXT: // %bb.0:1093; O3-NEXT: ld.param.b8 %r1, [test_and_computed_param_0];1094; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x3340U;1095; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x5410U;1096; O3-NEXT: bfi.b32 %r4, 5, %r3, 8, 8;1097; O3-NEXT: and.b32 %r5, %r4, %r3;1098; O3-NEXT: st.param.b32 [func_retval0], %r5;1099; O3-NEXT: ret;1100 %ins.0 = insertelement <4 x i8> zeroinitializer, i8 %a, i32 01101 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 11102 %r = and <4 x i8> %ins.1, %ins.01103 ret <4 x i8> %r1104}1105 1106define <4 x i8> @test_and_imm_0(<4 x i8> %a) #0 {1107; CHECK-LABEL: test_and_imm_0(1108; CHECK: {1109; CHECK-NEXT: .reg .b32 %r<3>;1110; CHECK-EMPTY:1111; CHECK-NEXT: // %bb.0:1112; CHECK-NEXT: ld.param.b32 %r1, [test_and_imm_0_param_0];1113; CHECK-NEXT: and.b32 %r2, %r1, 67305985;1114; CHECK-NEXT: st.param.b32 [func_retval0], %r2;1115; CHECK-NEXT: ret;1116 %r = and <4 x i8> <i8 1, i8 2, i8 3, i8 4>, %a1117 ret <4 x i8> %r1118}1119 1120define <4 x i8> @test_and_imm_1(<4 x i8> %a) #0 {1121; CHECK-LABEL: test_and_imm_1(1122; CHECK: {1123; CHECK-NEXT: .reg .b32 %r<3>;1124; CHECK-EMPTY:1125; CHECK-NEXT: // %bb.0:1126; CHECK-NEXT: ld.param.b32 %r1, [test_and_imm_1_param_0];1127; CHECK-NEXT: and.b32 %r2, %r1, 67305985;1128; CHECK-NEXT: st.param.b32 [func_retval0], %r2;1129; CHECK-NEXT: ret;1130 %r = and <4 x i8> %a, <i8 1, i8 2, i8 3, i8 4>1131 ret <4 x i8> %r1132}1133 1134define void @test_ldst_v2i8(ptr %a, ptr %b) {1135; O0-LABEL: test_ldst_v2i8(1136; O0: {1137; O0-NEXT: .reg .b32 %r<2>;1138; O0-NEXT: .reg .b64 %rd<3>;1139; O0-EMPTY:1140; O0-NEXT: // %bb.0:1141; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1];1142; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0];1143; O0-NEXT: ld.b32 %r1, [%rd1];1144; O0-NEXT: st.b32 [%rd2], %r1;1145; O0-NEXT: ret;1146;1147; O3-LABEL: test_ldst_v2i8(1148; O3: {1149; O3-NEXT: .reg .b32 %r<2>;1150; O3-NEXT: .reg .b64 %rd<3>;1151; O3-EMPTY:1152; O3-NEXT: // %bb.0:1153; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v2i8_param_0];1154; O3-NEXT: ld.b32 %r1, [%rd1];1155; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v2i8_param_1];1156; O3-NEXT: st.b32 [%rd2], %r1;1157; O3-NEXT: ret;1158 %t1 = load <4 x i8>, ptr %a1159 store <4 x i8> %t1, ptr %b, align 161160 ret void1161}1162 1163define void @test_ldst_v3i8(ptr %a, ptr %b) {1164; O0-LABEL: test_ldst_v3i8(1165; O0: {1166; O0-NEXT: .reg .b32 %r<3>;1167; O0-NEXT: .reg .b64 %rd<3>;1168; O0-EMPTY:1169; O0-NEXT: // %bb.0:1170; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1];1171; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0];1172; O0-NEXT: ld.b32 %r1, [%rd1];1173; O0-NEXT: st.b16 [%rd2], %r1;1174; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;1175; O0-NEXT: st.b8 [%rd2+2], %r2;1176; O0-NEXT: ret;1177;1178; O3-LABEL: test_ldst_v3i8(1179; O3: {1180; O3-NEXT: .reg .b32 %r<3>;1181; O3-NEXT: .reg .b64 %rd<3>;1182; O3-EMPTY:1183; O3-NEXT: // %bb.0:1184; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v3i8_param_0];1185; O3-NEXT: ld.b32 %r1, [%rd1];1186; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v3i8_param_1];1187; O3-NEXT: st.b16 [%rd2], %r1;1188; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;1189; O3-NEXT: st.b8 [%rd2+2], %r2;1190; O3-NEXT: ret;1191 %t1 = load <3 x i8>, ptr %a1192 store <3 x i8> %t1, ptr %b, align 161193 ret void1194}1195 1196define void @test_ldst_v4i8(ptr %a, ptr %b) {1197; O0-LABEL: test_ldst_v4i8(1198; O0: {1199; O0-NEXT: .reg .b32 %r<2>;1200; O0-NEXT: .reg .b64 %rd<3>;1201; O0-EMPTY:1202; O0-NEXT: // %bb.0:1203; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1];1204; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0];1205; O0-NEXT: ld.b32 %r1, [%rd1];1206; O0-NEXT: st.b32 [%rd2], %r1;1207; O0-NEXT: ret;1208;1209; O3-LABEL: test_ldst_v4i8(1210; O3: {1211; O3-NEXT: .reg .b32 %r<2>;1212; O3-NEXT: .reg .b64 %rd<3>;1213; O3-EMPTY:1214; O3-NEXT: // %bb.0:1215; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_param_0];1216; O3-NEXT: ld.b32 %r1, [%rd1];1217; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_param_1];1218; O3-NEXT: st.b32 [%rd2], %r1;1219; O3-NEXT: ret;1220 %t1 = load <4 x i8>, ptr %a1221 store <4 x i8> %t1, ptr %b, align 161222 ret void1223}1224 1225define void @test_ldst_v4i8_unaligned(ptr %a, ptr %b) {1226; O0-LABEL: test_ldst_v4i8_unaligned(1227; O0: {1228; O0-NEXT: .reg .b32 %r<5>;1229; O0-NEXT: .reg .b64 %rd<3>;1230; O0-EMPTY:1231; O0-NEXT: // %bb.0:1232; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1];1233; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0];1234; O0-NEXT: ld.b8 %r1, [%rd1];1235; O0-NEXT: ld.b8 %r2, [%rd1+1];1236; O0-NEXT: ld.b8 %r3, [%rd1+2];1237; O0-NEXT: ld.b8 %r4, [%rd1+3];1238; O0-NEXT: st.b8 [%rd2+3], %r4;1239; O0-NEXT: st.b8 [%rd2+2], %r3;1240; O0-NEXT: st.b8 [%rd2+1], %r2;1241; O0-NEXT: st.b8 [%rd2], %r1;1242; O0-NEXT: ret;1243;1244; O3-LABEL: test_ldst_v4i8_unaligned(1245; O3: {1246; O3-NEXT: .reg .b32 %r<5>;1247; O3-NEXT: .reg .b64 %rd<3>;1248; O3-EMPTY:1249; O3-NEXT: // %bb.0:1250; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v4i8_unaligned_param_0];1251; O3-NEXT: ld.b8 %r1, [%rd1+1];1252; O3-NEXT: ld.b8 %r2, [%rd1];1253; O3-NEXT: ld.b8 %r3, [%rd1+3];1254; O3-NEXT: ld.b8 %r4, [%rd1+2];1255; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v4i8_unaligned_param_1];1256; O3-NEXT: st.b8 [%rd2+2], %r4;1257; O3-NEXT: st.b8 [%rd2+3], %r3;1258; O3-NEXT: st.b8 [%rd2], %r2;1259; O3-NEXT: st.b8 [%rd2+1], %r1;1260; O3-NEXT: ret;1261 %t1 = load <4 x i8>, ptr %a, align 11262 store <4 x i8> %t1, ptr %b, align 11263 ret void1264}1265 1266 1267define void @test_ldst_v8i8(ptr %a, ptr %b) {1268; O0-LABEL: test_ldst_v8i8(1269; O0: {1270; O0-NEXT: .reg .b32 %r<3>;1271; O0-NEXT: .reg .b64 %rd<3>;1272; O0-EMPTY:1273; O0-NEXT: // %bb.0:1274; O0-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1];1275; O0-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0];1276; O0-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];1277; O0-NEXT: st.v2.b32 [%rd2], {%r1, %r2};1278; O0-NEXT: ret;1279;1280; O3-LABEL: test_ldst_v8i8(1281; O3: {1282; O3-NEXT: .reg .b32 %r<3>;1283; O3-NEXT: .reg .b64 %rd<3>;1284; O3-EMPTY:1285; O3-NEXT: // %bb.0:1286; O3-NEXT: ld.param.b64 %rd1, [test_ldst_v8i8_param_0];1287; O3-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];1288; O3-NEXT: ld.param.b64 %rd2, [test_ldst_v8i8_param_1];1289; O3-NEXT: st.v2.b32 [%rd2], {%r1, %r2};1290; O3-NEXT: ret;1291 %t1 = load <8 x i8>, ptr %a1292 store <8 x i8> %t1, ptr %b, align 161293 ret void1294}1295 1296declare <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b) #01297 1298define <4 x i8> @test_call(<4 x i8> %a, <4 x i8> %b) #0 {1299; O0-LABEL: test_call(1300; O0: {1301; O0-NEXT: .reg .b32 %r<4>;1302; O0-EMPTY:1303; O0-NEXT: // %bb.0:1304; O0-NEXT: ld.param.b32 %r2, [test_call_param_1];1305; O0-NEXT: ld.param.b32 %r1, [test_call_param_0];1306; O0-NEXT: { // callseq 0, 01307; O0-NEXT: .param .align 4 .b8 param0[4];1308; O0-NEXT: .param .align 4 .b8 param1[4];1309; O0-NEXT: .param .align 4 .b8 retval0[4];1310; O0-NEXT: st.param.b32 [param1], %r2;1311; O0-NEXT: st.param.b32 [param0], %r1;1312; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);1313; O0-NEXT: ld.param.b32 %r3, [retval0];1314; O0-NEXT: } // callseq 01315; O0-NEXT: st.param.b32 [func_retval0], %r3;1316; O0-NEXT: ret;1317;1318; O3-LABEL: test_call(1319; O3: {1320; O3-NEXT: .reg .b32 %r<4>;1321; O3-EMPTY:1322; O3-NEXT: // %bb.0:1323; O3-NEXT: ld.param.b32 %r1, [test_call_param_0];1324; O3-NEXT: { // callseq 0, 01325; O3-NEXT: .param .align 4 .b8 param0[4];1326; O3-NEXT: .param .align 4 .b8 param1[4];1327; O3-NEXT: .param .align 4 .b8 retval0[4];1328; O3-NEXT: ld.param.b32 %r2, [test_call_param_1];1329; O3-NEXT: st.param.b32 [param1], %r2;1330; O3-NEXT: st.param.b32 [param0], %r1;1331; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);1332; O3-NEXT: ld.param.b32 %r3, [retval0];1333; O3-NEXT: } // callseq 01334; O3-NEXT: st.param.b32 [func_retval0], %r3;1335; O3-NEXT: ret;1336 %r = call <4 x i8> @test_callee(<4 x i8> %a, <4 x i8> %b)1337 ret <4 x i8> %r1338}1339 1340define <4 x i8> @test_call_flipped(<4 x i8> %a, <4 x i8> %b) #0 {1341; O0-LABEL: test_call_flipped(1342; O0: {1343; O0-NEXT: .reg .b32 %r<4>;1344; O0-EMPTY:1345; O0-NEXT: // %bb.0:1346; O0-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];1347; O0-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];1348; O0-NEXT: { // callseq 1, 01349; O0-NEXT: .param .align 4 .b8 param0[4];1350; O0-NEXT: .param .align 4 .b8 param1[4];1351; O0-NEXT: .param .align 4 .b8 retval0[4];1352; O0-NEXT: st.param.b32 [param1], %r1;1353; O0-NEXT: st.param.b32 [param0], %r2;1354; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);1355; O0-NEXT: ld.param.b32 %r3, [retval0];1356; O0-NEXT: } // callseq 11357; O0-NEXT: st.param.b32 [func_retval0], %r3;1358; O0-NEXT: ret;1359;1360; O3-LABEL: test_call_flipped(1361; O3: {1362; O3-NEXT: .reg .b32 %r<4>;1363; O3-EMPTY:1364; O3-NEXT: // %bb.0:1365; O3-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];1366; O3-NEXT: { // callseq 1, 01367; O3-NEXT: .param .align 4 .b8 param0[4];1368; O3-NEXT: .param .align 4 .b8 param1[4];1369; O3-NEXT: .param .align 4 .b8 retval0[4];1370; O3-NEXT: st.param.b32 [param1], %r1;1371; O3-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];1372; O3-NEXT: st.param.b32 [param0], %r2;1373; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);1374; O3-NEXT: ld.param.b32 %r3, [retval0];1375; O3-NEXT: } // callseq 11376; O3-NEXT: st.param.b32 [func_retval0], %r3;1377; O3-NEXT: ret;1378 %r = call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a)1379 ret <4 x i8> %r1380}1381 1382define <4 x i8> @test_tailcall_flipped(<4 x i8> %a, <4 x i8> %b) #0 {1383; O0-LABEL: test_tailcall_flipped(1384; O0: {1385; O0-NEXT: .reg .b32 %r<4>;1386; O0-EMPTY:1387; O0-NEXT: // %bb.0:1388; O0-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];1389; O0-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];1390; O0-NEXT: { // callseq 2, 01391; O0-NEXT: .param .align 4 .b8 param0[4];1392; O0-NEXT: .param .align 4 .b8 param1[4];1393; O0-NEXT: .param .align 4 .b8 retval0[4];1394; O0-NEXT: st.param.b32 [param1], %r1;1395; O0-NEXT: st.param.b32 [param0], %r2;1396; O0-NEXT: call.uni (retval0), test_callee, (param0, param1);1397; O0-NEXT: ld.param.b32 %r3, [retval0];1398; O0-NEXT: } // callseq 21399; O0-NEXT: st.param.b32 [func_retval0], %r3;1400; O0-NEXT: ret;1401;1402; O3-LABEL: test_tailcall_flipped(1403; O3: {1404; O3-NEXT: .reg .b32 %r<4>;1405; O3-EMPTY:1406; O3-NEXT: // %bb.0:1407; O3-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];1408; O3-NEXT: { // callseq 2, 01409; O3-NEXT: .param .align 4 .b8 param0[4];1410; O3-NEXT: .param .align 4 .b8 param1[4];1411; O3-NEXT: .param .align 4 .b8 retval0[4];1412; O3-NEXT: st.param.b32 [param1], %r1;1413; O3-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];1414; O3-NEXT: st.param.b32 [param0], %r2;1415; O3-NEXT: call.uni (retval0), test_callee, (param0, param1);1416; O3-NEXT: ld.param.b32 %r3, [retval0];1417; O3-NEXT: } // callseq 21418; O3-NEXT: st.param.b32 [func_retval0], %r3;1419; O3-NEXT: ret;1420 %r = tail call <4 x i8> @test_callee(<4 x i8> %b, <4 x i8> %a)1421 ret <4 x i8> %r1422}1423 1424define <4 x i8> @test_select(<4 x i8> %a, <4 x i8> %b, i1 zeroext %c) #0 {1425; O0-LABEL: test_select(1426; O0: {1427; O0-NEXT: .reg .pred %p<2>;1428; O0-NEXT: .reg .b16 %rs<3>;1429; O0-NEXT: .reg .b32 %r<4>;1430; O0-EMPTY:1431; O0-NEXT: // %bb.0:1432; O0-NEXT: ld.param.b8 %rs1, [test_select_param_2];1433; O0-NEXT: and.b16 %rs2, %rs1, 1;1434; O0-NEXT: setp.ne.b16 %p1, %rs2, 0;1435; O0-NEXT: ld.param.b32 %r2, [test_select_param_1];1436; O0-NEXT: ld.param.b32 %r1, [test_select_param_0];1437; O0-NEXT: selp.b32 %r3, %r1, %r2, %p1;1438; O0-NEXT: st.param.b32 [func_retval0], %r3;1439; O0-NEXT: ret;1440;1441; O3-LABEL: test_select(1442; O3: {1443; O3-NEXT: .reg .pred %p<2>;1444; O3-NEXT: .reg .b16 %rs<3>;1445; O3-NEXT: .reg .b32 %r<2>;1446; O3-NEXT: .reg .b64 %rd<4>;1447; O3-EMPTY:1448; O3-NEXT: // %bb.0:1449; O3-NEXT: ld.param.b8 %rs1, [test_select_param_2];1450; O3-NEXT: and.b16 %rs2, %rs1, 1;1451; O3-NEXT: setp.ne.b16 %p1, %rs2, 0;1452; O3-NEXT: mov.b64 %rd1, test_select_param_1;1453; O3-NEXT: mov.b64 %rd2, test_select_param_0;1454; O3-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;1455; O3-NEXT: ld.param.b32 %r1, [%rd3];1456; O3-NEXT: st.param.b32 [func_retval0], %r1;1457; O3-NEXT: ret;1458 %r = select i1 %c, <4 x i8> %a, <4 x i8> %b1459 ret <4 x i8> %r1460}1461 1462define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8> %d) #0 {1463; O0-LABEL: test_select_cc(1464; O0: {1465; O0-NEXT: .reg .pred %p<5>;1466; O0-NEXT: .reg .b32 %r<28>;1467; O0-EMPTY:1468; O0-NEXT: // %bb.0:1469; O0-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];1470; O0-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];1471; O0-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];1472; O0-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];1473; O0-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U;1474; O0-NEXT: prmt.b32 %r6, %r3, 0, 0x7770U;1475; O0-NEXT: setp.ne.b32 %p1, %r6, %r5;1476; O0-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U;1477; O0-NEXT: prmt.b32 %r8, %r3, 0, 0x7771U;1478; O0-NEXT: setp.ne.b32 %p2, %r8, %r7;1479; O0-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U;1480; O0-NEXT: prmt.b32 %r10, %r3, 0, 0x7772U;1481; O0-NEXT: setp.ne.b32 %p3, %r10, %r9;1482; O0-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U;1483; O0-NEXT: prmt.b32 %r12, %r3, 0, 0x7773U;1484; O0-NEXT: setp.ne.b32 %p4, %r12, %r11;1485; O0-NEXT: prmt.b32 %r13, %r2, 0, 0x7773U;1486; O0-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;1487; O0-NEXT: selp.b32 %r15, %r14, %r13, %p4;1488; O0-NEXT: prmt.b32 %r16, %r2, 0, 0x7772U;1489; O0-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;1490; O0-NEXT: selp.b32 %r18, %r17, %r16, %p3;1491; O0-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U;1492; O0-NEXT: prmt.b32 %r20, %r2, 0, 0x7771U;1493; O0-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U;1494; O0-NEXT: selp.b32 %r22, %r21, %r20, %p2;1495; O0-NEXT: prmt.b32 %r23, %r2, 0, 0x7770U;1496; O0-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U;1497; O0-NEXT: selp.b32 %r25, %r24, %r23, %p1;1498; O0-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U;1499; O0-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U;1500; O0-NEXT: st.param.b32 [func_retval0], %r27;1501; O0-NEXT: ret;1502;1503; O3-LABEL: test_select_cc(1504; O3: {1505; O3-NEXT: .reg .pred %p<5>;1506; O3-NEXT: .reg .b32 %r<28>;1507; O3-EMPTY:1508; O3-NEXT: // %bb.0:1509; O3-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];1510; O3-NEXT: ld.param.b32 %r2, [test_select_cc_param_3];1511; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;1512; O3-NEXT: ld.param.b32 %r4, [test_select_cc_param_2];1513; O3-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U;1514; O3-NEXT: setp.ne.b32 %p1, %r5, %r3;1515; O3-NEXT: prmt.b32 %r6, %r2, 0, 0x7771U;1516; O3-NEXT: prmt.b32 %r7, %r4, 0, 0x7771U;1517; O3-NEXT: setp.ne.b32 %p2, %r7, %r6;1518; O3-NEXT: prmt.b32 %r8, %r2, 0, 0x7772U;1519; O3-NEXT: prmt.b32 %r9, %r4, 0, 0x7772U;1520; O3-NEXT: setp.ne.b32 %p3, %r9, %r8;1521; O3-NEXT: prmt.b32 %r10, %r2, 0, 0x7773U;1522; O3-NEXT: prmt.b32 %r11, %r4, 0, 0x7773U;1523; O3-NEXT: setp.ne.b32 %p4, %r11, %r10;1524; O3-NEXT: ld.param.b32 %r12, [test_select_cc_param_1];1525; O3-NEXT: prmt.b32 %r13, %r12, 0, 0x7773U;1526; O3-NEXT: prmt.b32 %r14, %r1, 0, 0x7773U;1527; O3-NEXT: selp.b32 %r15, %r14, %r13, %p4;1528; O3-NEXT: prmt.b32 %r16, %r12, 0, 0x7772U;1529; O3-NEXT: prmt.b32 %r17, %r1, 0, 0x7772U;1530; O3-NEXT: selp.b32 %r18, %r17, %r16, %p3;1531; O3-NEXT: prmt.b32 %r19, %r18, %r15, 0x3340U;1532; O3-NEXT: prmt.b32 %r20, %r12, 0, 0x7771U;1533; O3-NEXT: prmt.b32 %r21, %r1, 0, 0x7771U;1534; O3-NEXT: selp.b32 %r22, %r21, %r20, %p2;1535; O3-NEXT: prmt.b32 %r23, %r12, 0, 0x7770U;1536; O3-NEXT: prmt.b32 %r24, %r1, 0, 0x7770U;1537; O3-NEXT: selp.b32 %r25, %r24, %r23, %p1;1538; O3-NEXT: prmt.b32 %r26, %r25, %r22, 0x3340U;1539; O3-NEXT: prmt.b32 %r27, %r26, %r19, 0x5410U;1540; O3-NEXT: st.param.b32 [func_retval0], %r27;1541; O3-NEXT: ret;1542 %cc = icmp ne <4 x i8> %c, %d1543 %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b1544 ret <4 x i8> %r1545}1546 1547define <4 x i32> @test_select_cc_i32_i8(<4 x i32> %a, <4 x i32> %b,1548; O0-LABEL: test_select_cc_i32_i8(1549; O0: {1550; O0-NEXT: .reg .pred %p<5>;1551; O0-NEXT: .reg .b32 %r<23>;1552; O0-EMPTY:1553; O0-NEXT: // %bb.0:1554; O0-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_select_cc_i32_i8_param_1];1555; O0-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0];1556; O0-NEXT: ld.param.b32 %r10, [test_select_cc_i32_i8_param_3];1557; O0-NEXT: ld.param.b32 %r9, [test_select_cc_i32_i8_param_2];1558; O0-NEXT: prmt.b32 %r11, %r10, 0, 0x7770U;1559; O0-NEXT: prmt.b32 %r12, %r9, 0, 0x7770U;1560; O0-NEXT: setp.ne.b32 %p1, %r12, %r11;1561; O0-NEXT: prmt.b32 %r13, %r10, 0, 0x7771U;1562; O0-NEXT: prmt.b32 %r14, %r9, 0, 0x7771U;1563; O0-NEXT: setp.ne.b32 %p2, %r14, %r13;1564; O0-NEXT: prmt.b32 %r15, %r10, 0, 0x7772U;1565; O0-NEXT: prmt.b32 %r16, %r9, 0, 0x7772U;1566; O0-NEXT: setp.ne.b32 %p3, %r16, %r15;1567; O0-NEXT: prmt.b32 %r17, %r10, 0, 0x7773U;1568; O0-NEXT: prmt.b32 %r18, %r9, 0, 0x7773U;1569; O0-NEXT: setp.ne.b32 %p4, %r18, %r17;1570; O0-NEXT: selp.b32 %r19, %r4, %r8, %p4;1571; O0-NEXT: selp.b32 %r20, %r3, %r7, %p3;1572; O0-NEXT: selp.b32 %r21, %r2, %r6, %p2;1573; O0-NEXT: selp.b32 %r22, %r1, %r5, %p1;1574; O0-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19};1575; O0-NEXT: ret;1576;1577; O3-LABEL: test_select_cc_i32_i8(1578; O3: {1579; O3-NEXT: .reg .pred %p<5>;1580; O3-NEXT: .reg .b32 %r<23>;1581; O3-EMPTY:1582; O3-NEXT: // %bb.0:1583; O3-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_select_cc_i32_i8_param_0];1584; O3-NEXT: ld.param.b32 %r5, [test_select_cc_i32_i8_param_3];1585; O3-NEXT: prmt.b32 %r6, %r5, 0, 0x7770U;1586; O3-NEXT: ld.param.b32 %r7, [test_select_cc_i32_i8_param_2];1587; O3-NEXT: prmt.b32 %r8, %r7, 0, 0x7770U;1588; O3-NEXT: setp.ne.b32 %p1, %r8, %r6;1589; O3-NEXT: prmt.b32 %r9, %r5, 0, 0x7771U;1590; O3-NEXT: prmt.b32 %r10, %r7, 0, 0x7771U;1591; O3-NEXT: setp.ne.b32 %p2, %r10, %r9;1592; O3-NEXT: prmt.b32 %r11, %r5, 0, 0x7772U;1593; O3-NEXT: prmt.b32 %r12, %r7, 0, 0x7772U;1594; O3-NEXT: setp.ne.b32 %p3, %r12, %r11;1595; O3-NEXT: prmt.b32 %r13, %r5, 0, 0x7773U;1596; O3-NEXT: prmt.b32 %r14, %r7, 0, 0x7773U;1597; O3-NEXT: setp.ne.b32 %p4, %r14, %r13;1598; O3-NEXT: ld.param.v4.b32 {%r15, %r16, %r17, %r18}, [test_select_cc_i32_i8_param_1];1599; O3-NEXT: selp.b32 %r19, %r4, %r18, %p4;1600; O3-NEXT: selp.b32 %r20, %r3, %r17, %p3;1601; O3-NEXT: selp.b32 %r21, %r2, %r16, %p2;1602; O3-NEXT: selp.b32 %r22, %r1, %r15, %p1;1603; O3-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r21, %r20, %r19};1604; O3-NEXT: ret;1605 <4 x i8> %c, <4 x i8> %d) #0 {1606 %cc = icmp ne <4 x i8> %c, %d1607 %r = select <4 x i1> %cc, <4 x i32> %a, <4 x i32> %b1608 ret <4 x i32> %r1609}1610 1611define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b,1612; O0-LABEL: test_select_cc_i8_i32(1613; O0: {1614; O0-NEXT: .reg .pred %p<5>;1615; O0-NEXT: .reg .b32 %r<26>;1616; O0-EMPTY:1617; O0-NEXT: // %bb.0:1618; O0-NEXT: ld.param.v4.b32 {%r7, %r8, %r9, %r10}, [test_select_cc_i8_i32_param_3];1619; O0-NEXT: ld.param.v4.b32 {%r3, %r4, %r5, %r6}, [test_select_cc_i8_i32_param_2];1620; O0-NEXT: ld.param.b32 %r2, [test_select_cc_i8_i32_param_1];1621; O0-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0];1622; O0-NEXT: setp.ne.b32 %p1, %r3, %r7;1623; O0-NEXT: setp.ne.b32 %p2, %r4, %r8;1624; O0-NEXT: setp.ne.b32 %p3, %r5, %r9;1625; O0-NEXT: setp.ne.b32 %p4, %r6, %r10;1626; O0-NEXT: prmt.b32 %r11, %r2, 0, 0x7773U;1627; O0-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U;1628; O0-NEXT: selp.b32 %r13, %r12, %r11, %p4;1629; O0-NEXT: prmt.b32 %r14, %r2, 0, 0x7772U;1630; O0-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U;1631; O0-NEXT: selp.b32 %r16, %r15, %r14, %p3;1632; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U;1633; O0-NEXT: prmt.b32 %r18, %r2, 0, 0x7771U;1634; O0-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;1635; O0-NEXT: selp.b32 %r20, %r19, %r18, %p2;1636; O0-NEXT: prmt.b32 %r21, %r2, 0, 0x7770U;1637; O0-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;1638; O0-NEXT: selp.b32 %r23, %r22, %r21, %p1;1639; O0-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U;1640; O0-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U;1641; O0-NEXT: st.param.b32 [func_retval0], %r25;1642; O0-NEXT: ret;1643;1644; O3-LABEL: test_select_cc_i8_i32(1645; O3: {1646; O3-NEXT: .reg .pred %p<5>;1647; O3-NEXT: .reg .b32 %r<26>;1648; O3-EMPTY:1649; O3-NEXT: // %bb.0:1650; O3-NEXT: ld.param.b32 %r1, [test_select_cc_i8_i32_param_0];1651; O3-NEXT: ld.param.v4.b32 {%r2, %r3, %r4, %r5}, [test_select_cc_i8_i32_param_2];1652; O3-NEXT: ld.param.v4.b32 {%r6, %r7, %r8, %r9}, [test_select_cc_i8_i32_param_3];1653; O3-NEXT: setp.ne.b32 %p1, %r2, %r6;1654; O3-NEXT: setp.ne.b32 %p2, %r3, %r7;1655; O3-NEXT: setp.ne.b32 %p3, %r4, %r8;1656; O3-NEXT: setp.ne.b32 %p4, %r5, %r9;1657; O3-NEXT: ld.param.b32 %r10, [test_select_cc_i8_i32_param_1];1658; O3-NEXT: prmt.b32 %r11, %r10, 0, 0x7773U;1659; O3-NEXT: prmt.b32 %r12, %r1, 0, 0x7773U;1660; O3-NEXT: selp.b32 %r13, %r12, %r11, %p4;1661; O3-NEXT: prmt.b32 %r14, %r10, 0, 0x7772U;1662; O3-NEXT: prmt.b32 %r15, %r1, 0, 0x7772U;1663; O3-NEXT: selp.b32 %r16, %r15, %r14, %p3;1664; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x3340U;1665; O3-NEXT: prmt.b32 %r18, %r10, 0, 0x7771U;1666; O3-NEXT: prmt.b32 %r19, %r1, 0, 0x7771U;1667; O3-NEXT: selp.b32 %r20, %r19, %r18, %p2;1668; O3-NEXT: prmt.b32 %r21, %r10, 0, 0x7770U;1669; O3-NEXT: prmt.b32 %r22, %r1, 0, 0x7770U;1670; O3-NEXT: selp.b32 %r23, %r22, %r21, %p1;1671; O3-NEXT: prmt.b32 %r24, %r23, %r20, 0x3340U;1672; O3-NEXT: prmt.b32 %r25, %r24, %r17, 0x5410U;1673; O3-NEXT: st.param.b32 [func_retval0], %r25;1674; O3-NEXT: ret;1675 <4 x i32> %c, <4 x i32> %d) #0 {1676 %cc = icmp ne <4 x i32> %c, %d1677 %r = select <4 x i1> %cc, <4 x i8> %a, <4 x i8> %b1678 ret <4 x i8> %r1679}1680 1681 1682define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 {1683; CHECK-LABEL: test_trunc_2xi32(1684; CHECK: {1685; CHECK-NEXT: .reg .b32 %r<8>;1686; CHECK-EMPTY:1687; CHECK-NEXT: // %bb.0:1688; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_trunc_2xi32_param_0];1689; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 0x3340U;1690; CHECK-NEXT: prmt.b32 %r6, %r1, %r2, 0x3340U;1691; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 0x5410U;1692; CHECK-NEXT: st.param.b32 [func_retval0], %r7;1693; CHECK-NEXT: ret;1694 %r = trunc <4 x i32> %a to <4 x i8>1695 ret <4 x i8> %r1696}1697 1698define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 {1699; O0-LABEL: test_trunc_2xi64(1700; O0: {1701; O0-NEXT: .reg .b32 %r<8>;1702; O0-NEXT: .reg .b64 %rd<5>;1703; O0-EMPTY:1704; O0-NEXT: // %bb.0:1705; O0-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16];1706; O0-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];1707; O0-NEXT: cvt.u32.u64 %r1, %rd4;1708; O0-NEXT: cvt.u32.u64 %r2, %rd3;1709; O0-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;1710; O0-NEXT: cvt.u32.u64 %r4, %rd2;1711; O0-NEXT: cvt.u32.u64 %r5, %rd1;1712; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;1713; O0-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U;1714; O0-NEXT: st.param.b32 [func_retval0], %r7;1715; O0-NEXT: ret;1716;1717; O3-LABEL: test_trunc_2xi64(1718; O3: {1719; O3-NEXT: .reg .b32 %r<8>;1720; O3-NEXT: .reg .b64 %rd<5>;1721; O3-EMPTY:1722; O3-NEXT: // %bb.0:1723; O3-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];1724; O3-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_trunc_2xi64_param_0+16];1725; O3-NEXT: cvt.u32.u64 %r1, %rd4;1726; O3-NEXT: cvt.u32.u64 %r2, %rd3;1727; O3-NEXT: prmt.b32 %r3, %r2, %r1, 0x3340U;1728; O3-NEXT: cvt.u32.u64 %r4, %rd2;1729; O3-NEXT: cvt.u32.u64 %r5, %rd1;1730; O3-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;1731; O3-NEXT: prmt.b32 %r7, %r6, %r3, 0x5410U;1732; O3-NEXT: st.param.b32 [func_retval0], %r7;1733; O3-NEXT: ret;1734 %r = trunc <4 x i64> %a to <4 x i8>1735 ret <4 x i8> %r1736}1737 1738define <4 x i32> @test_zext_2xi32(<4 x i8> %a) #0 {1739; CHECK-LABEL: test_zext_2xi32(1740; CHECK: {1741; CHECK-NEXT: .reg .b32 %r<6>;1742; CHECK-EMPTY:1743; CHECK-NEXT: // %bb.0:1744; CHECK-NEXT: ld.param.b32 %r1, [test_zext_2xi32_param_0];1745; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;1746; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0x7772U;1747; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7771U;1748; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;1749; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r5, %r4, %r3, %r2};1750; CHECK-NEXT: ret;1751 %r = zext <4 x i8> %a to <4 x i32>1752 ret <4 x i32> %r1753}1754 1755define <4 x i64> @test_zext_2xi64(<4 x i8> %a) #0 {1756; O0-LABEL: test_zext_2xi64(1757; O0: {1758; O0-NEXT: .reg .b32 %r<6>;1759; O0-NEXT: .reg .b64 %rd<5>;1760; O0-EMPTY:1761; O0-NEXT: // %bb.0:1762; O0-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0];1763; O0-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;1764; O0-NEXT: cvt.u64.u32 %rd1, %r2;1765; O0-NEXT: prmt.b32 %r3, %r1, 0, 0x7772U;1766; O0-NEXT: cvt.u64.u32 %rd2, %r3;1767; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7771U;1768; O0-NEXT: cvt.u64.u32 %rd3, %r4;1769; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x7770U;1770; O0-NEXT: cvt.u64.u32 %rd4, %r5;1771; O0-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd3};1772; O0-NEXT: st.param.v2.b64 [func_retval0+16], {%rd2, %rd1};1773; O0-NEXT: ret;1774;1775; O3-LABEL: test_zext_2xi64(1776; O3: {1777; O3-NEXT: .reg .b32 %r<6>;1778; O3-NEXT: .reg .b64 %rd<5>;1779; O3-EMPTY:1780; O3-NEXT: // %bb.0:1781; O3-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0];1782; O3-NEXT: prmt.b32 %r2, %r1, 0, 0x7771U;1783; O3-NEXT: cvt.u64.u32 %rd1, %r2;1784; O3-NEXT: prmt.b32 %r3, %r1, 0, 0x7770U;1785; O3-NEXT: cvt.u64.u32 %rd2, %r3;1786; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7773U;1787; O3-NEXT: cvt.u64.u32 %rd3, %r4;1788; O3-NEXT: prmt.b32 %r5, %r1, 0, 0x7772U;1789; O3-NEXT: cvt.u64.u32 %rd4, %r5;1790; O3-NEXT: st.param.v2.b64 [func_retval0+16], {%rd4, %rd3};1791; O3-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};1792; O3-NEXT: ret;1793 %r = zext <4 x i8> %a to <4 x i64>1794 ret <4 x i64> %r1795}1796 1797define <4 x i8> @test_bitcast_i32_to_4xi8(i32 %a) #0 {1798; CHECK-LABEL: test_bitcast_i32_to_4xi8(1799; CHECK: {1800; CHECK-NEXT: .reg .b32 %r<2>;1801; CHECK-EMPTY:1802; CHECK-NEXT: // %bb.0:1803; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_i32_to_4xi8_param_0];1804; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1805; CHECK-NEXT: ret;1806 %r = bitcast i32 %a to <4 x i8>1807 ret <4 x i8> %r1808}1809 1810define <4 x i8> @test_bitcast_float_to_4xi8(float %a) #0 {1811; CHECK-LABEL: test_bitcast_float_to_4xi8(1812; CHECK: {1813; CHECK-NEXT: .reg .b32 %r<2>;1814; CHECK-EMPTY:1815; CHECK-NEXT: // %bb.0:1816; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_float_to_4xi8_param_0];1817; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1818; CHECK-NEXT: ret;1819 %r = bitcast float %a to <4 x i8>1820 ret <4 x i8> %r1821}1822 1823define i32 @test_bitcast_4xi8_to_i32(<4 x i8> %a) #0 {1824; CHECK-LABEL: test_bitcast_4xi8_to_i32(1825; CHECK: {1826; CHECK-NEXT: .reg .b32 %r<2>;1827; CHECK-EMPTY:1828; CHECK-NEXT: // %bb.0:1829; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_4xi8_to_i32_param_0];1830; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1831; CHECK-NEXT: ret;1832 %r = bitcast <4 x i8> %a to i321833 ret i32 %r1834}1835 1836define float @test_bitcast_4xi8_to_float(<4 x i8> %a) #0 {1837; CHECK-LABEL: test_bitcast_4xi8_to_float(1838; CHECK: {1839; CHECK-NEXT: .reg .b32 %r<2>;1840; CHECK-EMPTY:1841; CHECK-NEXT: // %bb.0:1842; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_4xi8_to_float_param_0];1843; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1844; CHECK-NEXT: ret;1845 %r = bitcast <4 x i8> %a to float1846 ret float %r1847}1848 1849 1850define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 {1851; O0-LABEL: test_bitcast_4xi8_to_2xhalf(1852; O0: {1853; O0-NEXT: .reg .b16 %rs<2>;1854; O0-NEXT: .reg .b32 %r<6>;1855; O0-EMPTY:1856; O0-NEXT: // %bb.0:1857; O0-NEXT: ld.param.b8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0];1858; O0-NEXT: mov.b32 %r1, 6;1859; O0-NEXT: prmt.b32 %r2, %r1, 7, 0x3340U;1860; O0-NEXT: cvt.u32.u16 %r3, %rs1;1861; O0-NEXT: prmt.b32 %r4, %r3, 5, 0x3340U;1862; O0-NEXT: prmt.b32 %r5, %r4, %r2, 0x5410U;1863; O0-NEXT: st.param.b32 [func_retval0], %r5;1864; O0-NEXT: ret;1865;1866; O3-LABEL: test_bitcast_4xi8_to_2xhalf(1867; O3: {1868; O3-NEXT: .reg .b32 %r<4>;1869; O3-EMPTY:1870; O3-NEXT: // %bb.0:1871; O3-NEXT: ld.param.b8 %r1, [test_bitcast_4xi8_to_2xhalf_param_0];1872; O3-NEXT: prmt.b32 %r2, %r1, 5, 0x3340U;1873; O3-NEXT: prmt.b32 %r3, %r2, 1798, 0x5410U;1874; O3-NEXT: st.param.b32 [func_retval0], %r3;1875; O3-NEXT: ret;1876 %ins.0 = insertelement <4 x i8> undef, i8 %a, i32 01877 %ins.1 = insertelement <4 x i8> %ins.0, i8 5, i32 11878 %ins.2 = insertelement <4 x i8> %ins.1, i8 6, i32 21879 %ins.3 = insertelement <4 x i8> %ins.2, i8 7, i32 31880 %r = bitcast <4 x i8> %ins.3 to <2 x half>1881 ret <2 x half> %r1882}1883 1884 1885define <4 x i8> @test_shufflevector(<4 x i8> %a) #0 {1886; O0-LABEL: test_shufflevector(1887; O0: {1888; O0-NEXT: .reg .b32 %r<4>;1889; O0-EMPTY:1890; O0-NEXT: // %bb.0:1891; O0-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];1892; O0-NEXT: // implicit-def: %r31893; O0-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U;1894; O0-NEXT: st.param.b32 [func_retval0], %r2;1895; O0-NEXT: ret;1896;1897; O3-LABEL: test_shufflevector(1898; O3: {1899; O3-NEXT: .reg .b32 %r<4>;1900; O3-EMPTY:1901; O3-NEXT: // %bb.0:1902; O3-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];1903; O3-NEXT: prmt.b32 %r2, %r1, %r3, 0x123U;1904; O3-NEXT: st.param.b32 [func_retval0], %r2;1905; O3-NEXT: ret;1906 %s = shufflevector <4 x i8> %a, <4 x i8> undef, <4 x i32> <i32 3, i32 2, i32 1, i32 0>1907 ret <4 x i8> %s1908}1909 1910define <4 x i8> @test_shufflevector_2(<4 x i8> %a, <4 x i8> %b) #0 {1911; O0-LABEL: test_shufflevector_2(1912; O0: {1913; O0-NEXT: .reg .b32 %r<4>;1914; O0-EMPTY:1915; O0-NEXT: // %bb.0:1916; O0-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1];1917; O0-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0];1918; O0-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U;1919; O0-NEXT: st.param.b32 [func_retval0], %r3;1920; O0-NEXT: ret;1921;1922; O3-LABEL: test_shufflevector_2(1923; O3: {1924; O3-NEXT: .reg .b32 %r<4>;1925; O3-EMPTY:1926; O3-NEXT: // %bb.0:1927; O3-NEXT: ld.param.b32 %r1, [test_shufflevector_2_param_0];1928; O3-NEXT: ld.param.b32 %r2, [test_shufflevector_2_param_1];1929; O3-NEXT: prmt.b32 %r3, %r1, %r2, 0x2537U;1930; O3-NEXT: st.param.b32 [func_retval0], %r3;1931; O3-NEXT: ret;1932 %s = shufflevector <4 x i8> %a, <4 x i8> %b, <4 x i32> <i32 7, i32 3, i32 5, i32 2>1933 ret <4 x i8> %s1934}1935 1936 1937define <4 x i8> @test_insertelement(<4 x i8> %a, i8 %x) #0 {1938; O0-LABEL: test_insertelement(1939; O0: {1940; O0-NEXT: .reg .b16 %rs<2>;1941; O0-NEXT: .reg .b32 %r<4>;1942; O0-EMPTY:1943; O0-NEXT: // %bb.0:1944; O0-NEXT: ld.param.b8 %rs1, [test_insertelement_param_1];1945; O0-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];1946; O0-NEXT: cvt.u32.u16 %r2, %rs1;1947; O0-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8;1948; O0-NEXT: st.param.b32 [func_retval0], %r3;1949; O0-NEXT: ret;1950;1951; O3-LABEL: test_insertelement(1952; O3: {1953; O3-NEXT: .reg .b32 %r<4>;1954; O3-EMPTY:1955; O3-NEXT: // %bb.0:1956; O3-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];1957; O3-NEXT: ld.param.b8 %r2, [test_insertelement_param_1];1958; O3-NEXT: bfi.b32 %r3, %r2, %r1, 8, 8;1959; O3-NEXT: st.param.b32 [func_retval0], %r3;1960; O3-NEXT: ret;1961 %i = insertelement <4 x i8> %a, i8 %x, i64 11962 ret <4 x i8> %i1963}1964 1965define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 {1966; O0-LABEL: test_fptosi_4xhalf_to_4xi8(1967; O0: {1968; O0-NEXT: .reg .b16 %rs<13>;1969; O0-NEXT: .reg .b32 %r<12>;1970; O0-EMPTY:1971; O0-NEXT: // %bb.0:1972; O0-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_4xhalf_to_4xi8_param_0];1973; O0-NEXT: mov.b32 {%rs1, %rs2}, %r2;1974; O0-NEXT: cvt.rzi.s16.f16 %rs3, %rs2;1975; O0-NEXT: cvt.rzi.s16.f16 %rs4, %rs1;1976; O0-NEXT: mov.b32 %r3, {%rs4, %rs3};1977; O0-NEXT: mov.b32 {%rs5, %rs6}, %r3;1978; O0-NEXT: cvt.u32.u16 %r4, %rs6;1979; O0-NEXT: cvt.u32.u16 %r5, %rs5;1980; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;1981; O0-NEXT: mov.b32 {%rs7, %rs8}, %r1;1982; O0-NEXT: cvt.rzi.s16.f16 %rs9, %rs8;1983; O0-NEXT: cvt.rzi.s16.f16 %rs10, %rs7;1984; O0-NEXT: mov.b32 %r7, {%rs10, %rs9};1985; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7;1986; O0-NEXT: cvt.u32.u16 %r8, %rs12;1987; O0-NEXT: cvt.u32.u16 %r9, %rs11;1988; O0-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U;1989; O0-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U;1990; O0-NEXT: st.param.b32 [func_retval0], %r11;1991; O0-NEXT: ret;1992;1993; O3-LABEL: test_fptosi_4xhalf_to_4xi8(1994; O3: {1995; O3-NEXT: .reg .b16 %rs<13>;1996; O3-NEXT: .reg .b32 %r<10>;1997; O3-EMPTY:1998; O3-NEXT: // %bb.0:1999; O3-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0];2000; O3-NEXT: cvt.rzi.s16.f16 %rs5, %rs4;2001; O3-NEXT: cvt.rzi.s16.f16 %rs6, %rs3;2002; O3-NEXT: mov.b32 %r1, {%rs6, %rs5};2003; O3-NEXT: mov.b32 {%rs7, %rs8}, %r1;2004; O3-NEXT: cvt.u32.u16 %r2, %rs8;2005; O3-NEXT: cvt.u32.u16 %r3, %rs7;2006; O3-NEXT: prmt.b32 %r4, %r3, %r2, 0x3340U;2007; O3-NEXT: cvt.rzi.s16.f16 %rs9, %rs2;2008; O3-NEXT: cvt.rzi.s16.f16 %rs10, %rs1;2009; O3-NEXT: mov.b32 %r5, {%rs10, %rs9};2010; O3-NEXT: mov.b32 {%rs11, %rs12}, %r5;2011; O3-NEXT: cvt.u32.u16 %r6, %rs12;2012; O3-NEXT: cvt.u32.u16 %r7, %rs11;2013; O3-NEXT: prmt.b32 %r8, %r7, %r6, 0x3340U;2014; O3-NEXT: prmt.b32 %r9, %r8, %r4, 0x5410U;2015; O3-NEXT: st.param.b32 [func_retval0], %r9;2016; O3-NEXT: ret;2017 %r = fptosi <4 x half> %a to <4 x i8>2018 ret <4 x i8> %r2019}2020 2021define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 {2022; O0-LABEL: test_fptoui_4xhalf_to_4xi8(2023; O0: {2024; O0-NEXT: .reg .b16 %rs<13>;2025; O0-NEXT: .reg .b32 %r<12>;2026; O0-EMPTY:2027; O0-NEXT: // %bb.0:2028; O0-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_4xhalf_to_4xi8_param_0];2029; O0-NEXT: mov.b32 {%rs1, %rs2}, %r2;2030; O0-NEXT: cvt.rzi.u16.f16 %rs3, %rs2;2031; O0-NEXT: cvt.rzi.u16.f16 %rs4, %rs1;2032; O0-NEXT: mov.b32 %r3, {%rs4, %rs3};2033; O0-NEXT: mov.b32 {%rs5, %rs6}, %r3;2034; O0-NEXT: cvt.u32.u16 %r4, %rs6;2035; O0-NEXT: cvt.u32.u16 %r5, %rs5;2036; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;2037; O0-NEXT: mov.b32 {%rs7, %rs8}, %r1;2038; O0-NEXT: cvt.rzi.u16.f16 %rs9, %rs8;2039; O0-NEXT: cvt.rzi.u16.f16 %rs10, %rs7;2040; O0-NEXT: mov.b32 %r7, {%rs10, %rs9};2041; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7;2042; O0-NEXT: cvt.u32.u16 %r8, %rs12;2043; O0-NEXT: cvt.u32.u16 %r9, %rs11;2044; O0-NEXT: prmt.b32 %r10, %r9, %r8, 0x3340U;2045; O0-NEXT: prmt.b32 %r11, %r10, %r6, 0x5410U;2046; O0-NEXT: st.param.b32 [func_retval0], %r11;2047; O0-NEXT: ret;2048;2049; O3-LABEL: test_fptoui_4xhalf_to_4xi8(2050; O3: {2051; O3-NEXT: .reg .b16 %rs<13>;2052; O3-NEXT: .reg .b32 %r<10>;2053; O3-EMPTY:2054; O3-NEXT: // %bb.0:2055; O3-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0];2056; O3-NEXT: cvt.rzi.u16.f16 %rs5, %rs4;2057; O3-NEXT: cvt.rzi.u16.f16 %rs6, %rs3;2058; O3-NEXT: mov.b32 %r1, {%rs6, %rs5};2059; O3-NEXT: mov.b32 {%rs7, %rs8}, %r1;2060; O3-NEXT: cvt.u32.u16 %r2, %rs8;2061; O3-NEXT: cvt.u32.u16 %r3, %rs7;2062; O3-NEXT: prmt.b32 %r4, %r3, %r2, 0x3340U;2063; O3-NEXT: cvt.rzi.u16.f16 %rs9, %rs2;2064; O3-NEXT: cvt.rzi.u16.f16 %rs10, %rs1;2065; O3-NEXT: mov.b32 %r5, {%rs10, %rs9};2066; O3-NEXT: mov.b32 {%rs11, %rs12}, %r5;2067; O3-NEXT: cvt.u32.u16 %r6, %rs12;2068; O3-NEXT: cvt.u32.u16 %r7, %rs11;2069; O3-NEXT: prmt.b32 %r8, %r7, %r6, 0x3340U;2070; O3-NEXT: prmt.b32 %r9, %r8, %r4, 0x5410U;2071; O3-NEXT: st.param.b32 [func_retval0], %r9;2072; O3-NEXT: ret;2073 %r = fptoui <4 x half> %a to <4 x i8>2074 ret <4 x i8> %r2075}2076 2077define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) {2078; O0-LABEL: test_srem_v4i8(2079; O0: {2080; O0-NEXT: .reg .b16 %rs<13>;2081; O0-NEXT: .reg .b32 %r<16>;2082; O0-NEXT: .reg .b64 %rd<4>;2083; O0-EMPTY:2084; O0-NEXT: // %bb.0: // %entry2085; O0-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2];2086; O0-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1];2087; O0-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0];2088; O0-NEXT: ld.b32 %r1, [%rd1];2089; O0-NEXT: ld.b32 %r2, [%rd2];2090; O0-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U;2091; O0-NEXT: cvt.u16.u32 %rs1, %r3;2092; O0-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U;2093; O0-NEXT: cvt.u16.u32 %rs2, %r4;2094; O0-NEXT: rem.s16 %rs3, %rs2, %rs1;2095; O0-NEXT: cvt.u32.u16 %r5, %rs3;2096; O0-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U;2097; O0-NEXT: cvt.u16.u32 %rs4, %r6;2098; O0-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U;2099; O0-NEXT: cvt.u16.u32 %rs5, %r7;2100; O0-NEXT: rem.s16 %rs6, %rs5, %rs4;2101; O0-NEXT: cvt.u32.u16 %r8, %rs6;2102; O0-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;2103; O0-NEXT: cvt.s8.s32 %rs7, %r2;2104; O0-NEXT: cvt.s8.s32 %rs8, %r1;2105; O0-NEXT: rem.s16 %rs9, %rs8, %rs7;2106; O0-NEXT: cvt.u32.u16 %r10, %rs9;2107; O0-NEXT: prmt.b32 %r11, %r2, 0, 0x9991U;2108; O0-NEXT: cvt.u16.u32 %rs10, %r11;2109; O0-NEXT: prmt.b32 %r12, %r1, 0, 0x9991U;2110; O0-NEXT: cvt.u16.u32 %rs11, %r12;2111; O0-NEXT: rem.s16 %rs12, %rs11, %rs10;2112; O0-NEXT: cvt.u32.u16 %r13, %rs12;2113; O0-NEXT: prmt.b32 %r14, %r10, %r13, 0x3340U;2114; O0-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;2115; O0-NEXT: st.b32 [%rd3], %r15;2116; O0-NEXT: ret;2117;2118; O3-LABEL: test_srem_v4i8(2119; O3: {2120; O3-NEXT: .reg .b16 %rs<13>;2121; O3-NEXT: .reg .b32 %r<16>;2122; O3-NEXT: .reg .b64 %rd<4>;2123; O3-EMPTY:2124; O3-NEXT: // %bb.0: // %entry2125; O3-NEXT: ld.param.b64 %rd1, [test_srem_v4i8_param_0];2126; O3-NEXT: ld.b32 %r1, [%rd1];2127; O3-NEXT: ld.param.b64 %rd2, [test_srem_v4i8_param_1];2128; O3-NEXT: ld.b32 %r2, [%rd2];2129; O3-NEXT: ld.param.b64 %rd3, [test_srem_v4i8_param_2];2130; O3-NEXT: prmt.b32 %r3, %r2, 0, 0xbbb3U;2131; O3-NEXT: cvt.u16.u32 %rs1, %r3;2132; O3-NEXT: prmt.b32 %r4, %r1, 0, 0xbbb3U;2133; O3-NEXT: cvt.u16.u32 %rs2, %r4;2134; O3-NEXT: rem.s16 %rs3, %rs2, %rs1;2135; O3-NEXT: cvt.u32.u16 %r5, %rs3;2136; O3-NEXT: prmt.b32 %r6, %r2, 0, 0xaaa2U;2137; O3-NEXT: cvt.u16.u32 %rs4, %r6;2138; O3-NEXT: prmt.b32 %r7, %r1, 0, 0xaaa2U;2139; O3-NEXT: cvt.u16.u32 %rs5, %r7;2140; O3-NEXT: rem.s16 %rs6, %rs5, %rs4;2141; O3-NEXT: cvt.u32.u16 %r8, %rs6;2142; O3-NEXT: prmt.b32 %r9, %r8, %r5, 0x3340U;2143; O3-NEXT: cvt.s8.s32 %rs7, %r2;2144; O3-NEXT: cvt.s8.s32 %rs8, %r1;2145; O3-NEXT: rem.s16 %rs9, %rs8, %rs7;2146; O3-NEXT: cvt.u32.u16 %r10, %rs9;2147; O3-NEXT: prmt.b32 %r11, %r2, 0, 0x9991U;2148; O3-NEXT: cvt.u16.u32 %rs10, %r11;2149; O3-NEXT: prmt.b32 %r12, %r1, 0, 0x9991U;2150; O3-NEXT: cvt.u16.u32 %rs11, %r12;2151; O3-NEXT: rem.s16 %rs12, %rs11, %rs10;2152; O3-NEXT: cvt.u32.u16 %r13, %rs12;2153; O3-NEXT: prmt.b32 %r14, %r10, %r13, 0x3340U;2154; O3-NEXT: prmt.b32 %r15, %r14, %r9, 0x5410U;2155; O3-NEXT: st.b32 [%rd3], %r15;2156; O3-NEXT: ret;2157entry:2158 %t57 = load <4 x i8>, ptr %a, align 42159 %t59 = load <4 x i8>, ptr %b, align 42160 %x = srem <4 x i8> %t57, %t592161 store <4 x i8> %x, ptr %c, align 42162 ret void2163}2164 2165;; v3i8 lowering, especially for unaligned loads is terrible. We end up doing2166;; tons of pointless scalar_to_vector/bitcast/extract_elt on v2i16/v4i8, which2167;; is further complicated by LLVM trying to use i16 as an intermediate type,2168;; because we don't have i8 registers. It's a mess.2169;; Ideally we want to split it into element-wise ops, but legalizer can't handle2170;; odd-sized vectors. TL;DR; don't use odd-sized vectors of v8.2171define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) {2172; O0-LABEL: test_srem_v3i8(2173; O0: {2174; O0-NEXT: .reg .b16 %rs<20>;2175; O0-NEXT: .reg .b32 %r<8>;2176; O0-NEXT: .reg .b64 %rd<4>;2177; O0-EMPTY:2178; O0-NEXT: // %bb.0: // %entry2179; O0-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2];2180; O0-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1];2181; O0-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0];2182; O0-NEXT: ld.b8 %rs1, [%rd1];2183; O0-NEXT: ld.b8 %rs2, [%rd1+1];2184; O0-NEXT: shl.b16 %rs3, %rs2, 8;2185; O0-NEXT: or.b16 %rs4, %rs3, %rs1;2186; O0-NEXT: cvt.u32.u16 %r1, %rs4;2187; O0-NEXT: ld.s8 %rs5, [%rd1+2];2188; O0-NEXT: ld.b8 %rs6, [%rd2];2189; O0-NEXT: ld.b8 %rs7, [%rd2+1];2190; O0-NEXT: shl.b16 %rs8, %rs7, 8;2191; O0-NEXT: or.b16 %rs9, %rs8, %rs6;2192; O0-NEXT: cvt.u32.u16 %r2, %rs9;2193; O0-NEXT: ld.s8 %rs10, [%rd2+2];2194; O0-NEXT: cvt.s16.s8 %rs11, %rs9;2195; O0-NEXT: cvt.s16.s8 %rs12, %rs4;2196; O0-NEXT: rem.s16 %rs13, %rs12, %rs11;2197; O0-NEXT: cvt.u32.u16 %r3, %rs13;2198; O0-NEXT: prmt.b32 %r4, %r2, 0, 0x9991U;2199; O0-NEXT: cvt.u16.u32 %rs14, %r4;2200; O0-NEXT: prmt.b32 %r5, %r1, 0, 0x9991U;2201; O0-NEXT: cvt.u16.u32 %rs15, %r5;2202; O0-NEXT: rem.s16 %rs16, %rs15, %rs14;2203; O0-NEXT: cvt.u32.u16 %r6, %rs16;2204; O0-NEXT: prmt.b32 %r7, %r3, %r6, 0x3340U;2205; O0-NEXT: rem.s16 %rs17, %rs5, %rs10;2206; O0-NEXT: cvt.u16.u32 %rs18, %r7;2207; O0-NEXT: st.b8 [%rd3], %rs18;2208; O0-NEXT: shr.u16 %rs19, %rs18, 8;2209; O0-NEXT: st.b8 [%rd3+1], %rs19;2210; O0-NEXT: st.b8 [%rd3+2], %rs17;2211; O0-NEXT: ret;2212;2213; O3-LABEL: test_srem_v3i8(2214; O3: {2215; O3-NEXT: .reg .b16 %rs<20>;2216; O3-NEXT: .reg .b32 %r<8>;2217; O3-NEXT: .reg .b64 %rd<4>;2218; O3-EMPTY:2219; O3-NEXT: // %bb.0: // %entry2220; O3-NEXT: ld.param.b64 %rd1, [test_srem_v3i8_param_0];2221; O3-NEXT: ld.b8 %rs1, [%rd1];2222; O3-NEXT: ld.b8 %rs2, [%rd1+1];2223; O3-NEXT: shl.b16 %rs3, %rs2, 8;2224; O3-NEXT: or.b16 %rs4, %rs3, %rs1;2225; O3-NEXT: cvt.u32.u16 %r1, %rs4;2226; O3-NEXT: ld.s8 %rs5, [%rd1+2];2227; O3-NEXT: ld.param.b64 %rd2, [test_srem_v3i8_param_1];2228; O3-NEXT: ld.b8 %rs6, [%rd2];2229; O3-NEXT: ld.b8 %rs7, [%rd2+1];2230; O3-NEXT: shl.b16 %rs8, %rs7, 8;2231; O3-NEXT: or.b16 %rs9, %rs8, %rs6;2232; O3-NEXT: cvt.u32.u16 %r2, %rs9;2233; O3-NEXT: ld.s8 %rs10, [%rd2+2];2234; O3-NEXT: ld.param.b64 %rd3, [test_srem_v3i8_param_2];2235; O3-NEXT: cvt.s16.s8 %rs11, %rs9;2236; O3-NEXT: cvt.s16.s8 %rs12, %rs4;2237; O3-NEXT: rem.s16 %rs13, %rs12, %rs11;2238; O3-NEXT: cvt.u32.u16 %r3, %rs13;2239; O3-NEXT: prmt.b32 %r4, %r2, 0, 0x9991U;2240; O3-NEXT: cvt.u16.u32 %rs14, %r4;2241; O3-NEXT: prmt.b32 %r5, %r1, 0, 0x9991U;2242; O3-NEXT: cvt.u16.u32 %rs15, %r5;2243; O3-NEXT: rem.s16 %rs16, %rs15, %rs14;2244; O3-NEXT: cvt.u32.u16 %r6, %rs16;2245; O3-NEXT: prmt.b32 %r7, %r3, %r6, 0x3340U;2246; O3-NEXT: rem.s16 %rs17, %rs5, %rs10;2247; O3-NEXT: st.b8 [%rd3+2], %rs17;2248; O3-NEXT: cvt.u16.u32 %rs18, %r7;2249; O3-NEXT: st.b8 [%rd3], %rs18;2250; O3-NEXT: shr.u16 %rs19, %rs18, 8;2251; O3-NEXT: st.b8 [%rd3+1], %rs19;2252; O3-NEXT: ret;2253entry:2254 %t57 = load <3 x i8>, ptr %a, align 12255 %t59 = load <3 x i8>, ptr %b, align 12256 %x = srem <3 x i8> %t57, %t592257 store <3 x i8> %x, ptr %c, align 12258 ret void2259}2260 2261define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) {2262; O0-LABEL: test_sext_v4i1_to_v4i8(2263; O0: {2264; O0-NEXT: .reg .pred %p<5>;2265; O0-NEXT: .reg .b32 %r<18>;2266; O0-NEXT: .reg .b64 %rd<4>;2267; O0-EMPTY:2268; O0-NEXT: // %bb.0: // %entry2269; O0-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2];2270; O0-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1];2271; O0-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0];2272; O0-NEXT: ld.b32 %r1, [%rd1];2273; O0-NEXT: ld.b32 %r2, [%rd2];2274; O0-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;2275; O0-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;2276; O0-NEXT: setp.gt.u32 %p1, %r4, %r3;2277; O0-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;2278; O0-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;2279; O0-NEXT: setp.gt.u32 %p2, %r6, %r5;2280; O0-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;2281; O0-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;2282; O0-NEXT: setp.gt.u32 %p3, %r8, %r7;2283; O0-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;2284; O0-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;2285; O0-NEXT: setp.gt.u32 %p4, %r10, %r9;2286; O0-NEXT: selp.b32 %r11, -1, 0, %p4;2287; O0-NEXT: selp.b32 %r12, -1, 0, %p3;2288; O0-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;2289; O0-NEXT: selp.b32 %r14, -1, 0, %p2;2290; O0-NEXT: selp.b32 %r15, -1, 0, %p1;2291; O0-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;2292; O0-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;2293; O0-NEXT: st.b32 [%rd3], %r17;2294; O0-NEXT: ret;2295;2296; O3-LABEL: test_sext_v4i1_to_v4i8(2297; O3: {2298; O3-NEXT: .reg .pred %p<5>;2299; O3-NEXT: .reg .b32 %r<18>;2300; O3-NEXT: .reg .b64 %rd<4>;2301; O3-EMPTY:2302; O3-NEXT: // %bb.0: // %entry2303; O3-NEXT: ld.param.b64 %rd1, [test_sext_v4i1_to_v4i8_param_0];2304; O3-NEXT: ld.b32 %r1, [%rd1];2305; O3-NEXT: ld.param.b64 %rd2, [test_sext_v4i1_to_v4i8_param_1];2306; O3-NEXT: ld.b32 %r2, [%rd2];2307; O3-NEXT: ld.param.b64 %rd3, [test_sext_v4i1_to_v4i8_param_2];2308; O3-NEXT: prmt.b32 %r3, %r2, 0, 0x7770U;2309; O3-NEXT: prmt.b32 %r4, %r1, 0, 0x7770U;2310; O3-NEXT: setp.gt.u32 %p1, %r4, %r3;2311; O3-NEXT: prmt.b32 %r5, %r2, 0, 0x7771U;2312; O3-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;2313; O3-NEXT: setp.gt.u32 %p2, %r6, %r5;2314; O3-NEXT: prmt.b32 %r7, %r2, 0, 0x7772U;2315; O3-NEXT: prmt.b32 %r8, %r1, 0, 0x7772U;2316; O3-NEXT: setp.gt.u32 %p3, %r8, %r7;2317; O3-NEXT: prmt.b32 %r9, %r2, 0, 0x7773U;2318; O3-NEXT: prmt.b32 %r10, %r1, 0, 0x7773U;2319; O3-NEXT: setp.gt.u32 %p4, %r10, %r9;2320; O3-NEXT: selp.b32 %r11, -1, 0, %p4;2321; O3-NEXT: selp.b32 %r12, -1, 0, %p3;2322; O3-NEXT: prmt.b32 %r13, %r12, %r11, 0x3340U;2323; O3-NEXT: selp.b32 %r14, -1, 0, %p2;2324; O3-NEXT: selp.b32 %r15, -1, 0, %p1;2325; O3-NEXT: prmt.b32 %r16, %r15, %r14, 0x3340U;2326; O3-NEXT: prmt.b32 %r17, %r16, %r13, 0x5410U;2327; O3-NEXT: st.b32 [%rd3], %r17;2328; O3-NEXT: ret;2329entry:2330 %t1 = load <4 x i8>, ptr %a, align 42331 %t2 = load <4 x i8>, ptr %b, align 42332 %t5 = icmp ugt <4 x i8> %t1, %t22333 %t6 = sext <4 x i1> %t5 to <4 x i8>2334 store <4 x i8> %t6, ptr %c, align 42335 ret void2336}2337 2338define <4 x float> @test_uitofp_v4i8(<4 x i8> %a) {2339; CHECK-LABEL: test_uitofp_v4i8(2340; CHECK: {2341; CHECK-NEXT: .reg .b32 %r<10>;2342; CHECK-EMPTY:2343; CHECK-NEXT: // %bb.0:2344; CHECK-NEXT: ld.param.b32 %r1, [test_uitofp_v4i8_param_0];2345; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7773U;2346; CHECK-NEXT: cvt.rn.f32.u32 %r3, %r2;2347; CHECK-NEXT: prmt.b32 %r4, %r1, 0, 0x7772U;2348; CHECK-NEXT: cvt.rn.f32.u32 %r5, %r4;2349; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7771U;2350; CHECK-NEXT: cvt.rn.f32.u32 %r7, %r6;2351; CHECK-NEXT: prmt.b32 %r8, %r1, 0, 0x7770U;2352; CHECK-NEXT: cvt.rn.f32.u32 %r9, %r8;2353; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r9, %r7, %r5, %r3};2354; CHECK-NEXT: ret;2355 %r = uitofp <4 x i8> %a to <4 x float>2356 ret <4 x float> %r2357}2358 2359define <4 x float> @test_sitofp_v4i8(<4 x i8> %a) {2360; CHECK-LABEL: test_sitofp_v4i8(2361; CHECK: {2362; CHECK-NEXT: .reg .b16 %rs<5>;2363; CHECK-NEXT: .reg .b32 %r<9>;2364; CHECK-EMPTY:2365; CHECK-NEXT: // %bb.0:2366; CHECK-NEXT: ld.param.b32 %r1, [test_sitofp_v4i8_param_0];2367; CHECK-NEXT: cvt.s8.s32 %rs1, %r1;2368; CHECK-NEXT: cvt.rn.f32.s16 %r2, %rs1;2369; CHECK-NEXT: prmt.b32 %r3, %r1, 0, 0xbbb3U;2370; CHECK-NEXT: cvt.u16.u32 %rs2, %r3;2371; CHECK-NEXT: cvt.rn.f32.s16 %r4, %rs2;2372; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0xaaa2U;2373; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;2374; CHECK-NEXT: cvt.rn.f32.s16 %r6, %rs3;2375; CHECK-NEXT: prmt.b32 %r7, %r1, 0, 0x9991U;2376; CHECK-NEXT: cvt.u16.u32 %rs4, %r7;2377; CHECK-NEXT: cvt.rn.f32.s16 %r8, %rs4;2378; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r2, %r8, %r6, %r4};2379; CHECK-NEXT: ret;2380 %r = sitofp <4 x i8> %a to <4 x float>2381 ret <4 x float> %r2382}2383 2384attributes #0 = { nounwind }2385