748 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py2; RUN: llc -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx70 < %s | FileCheck %s3 4; ============================================================================ ;5; 8-bit vector width6; ============================================================================ ;7 8define <1 x i8> @out_v1i8(<1 x i8> %x, <1 x i8> %y, <1 x i8> %mask) nounwind {9; CHECK-LABEL: out_v1i8(10; CHECK: {11; CHECK-NEXT: .reg .b16 %rs<7>;12; CHECK-EMPTY:13; CHECK-NEXT: // %bb.0:14; CHECK-NEXT: ld.param.b8 %rs1, [out_v1i8_param_0];15; CHECK-NEXT: ld.param.b8 %rs2, [out_v1i8_param_1];16; CHECK-NEXT: ld.param.b8 %rs3, [out_v1i8_param_2];17; CHECK-NEXT: xor.b16 %rs4, %rs1, %rs2;18; CHECK-NEXT: and.b16 %rs5, %rs4, %rs3;19; CHECK-NEXT: xor.b16 %rs6, %rs5, %rs2;20; CHECK-NEXT: st.param.b8 [func_retval0], %rs6;21; CHECK-NEXT: ret;22 %mx = and <1 x i8> %x, %mask23 %notmask = xor <1 x i8> %mask, <i8 -1>24 %my = and <1 x i8> %y, %notmask25 %r = or <1 x i8> %mx, %my26 ret <1 x i8> %r27}28 29; ============================================================================ ;30; 16-bit vector width31; ============================================================================ ;32 33define <1 x i16> @out_v1i16(<1 x i16> %x, <1 x i16> %y, <1 x i16> %mask) nounwind {34; CHECK-LABEL: out_v1i16(35; CHECK: {36; CHECK-NEXT: .reg .b16 %rs<7>;37; CHECK-EMPTY:38; CHECK-NEXT: // %bb.0:39; CHECK-NEXT: ld.param.b16 %rs1, [out_v1i16_param_0];40; CHECK-NEXT: ld.param.b16 %rs2, [out_v1i16_param_1];41; CHECK-NEXT: ld.param.b16 %rs3, [out_v1i16_param_2];42; CHECK-NEXT: xor.b16 %rs4, %rs1, %rs2;43; CHECK-NEXT: and.b16 %rs5, %rs4, %rs3;44; CHECK-NEXT: xor.b16 %rs6, %rs5, %rs2;45; CHECK-NEXT: st.param.b16 [func_retval0], %rs6;46; CHECK-NEXT: ret;47 %mx = and <1 x i16> %x, %mask48 %notmask = xor <1 x i16> %mask, <i16 -1>49 %my = and <1 x i16> %y, %notmask50 %r = or <1 x i16> %mx, %my51 ret <1 x i16> %r52}53 54; ============================================================================ ;55; 32-bit vector width56; ============================================================================ ;57 58define <4 x i8> @out_v4i8(<4 x i8> %x, <4 x i8> %y, <4 x i8> %mask) nounwind {59; CHECK-LABEL: out_v4i8(60; CHECK: {61; CHECK-NEXT: .reg .b32 %r<8>;62; CHECK-EMPTY:63; CHECK-NEXT: // %bb.0:64; CHECK-NEXT: ld.param.b32 %r1, [out_v4i8_param_0];65; CHECK-NEXT: ld.param.b32 %r2, [out_v4i8_param_2];66; CHECK-NEXT: and.b32 %r3, %r1, %r2;67; CHECK-NEXT: ld.param.b32 %r4, [out_v4i8_param_1];68; CHECK-NEXT: xor.b32 %r5, %r2, -1;69; CHECK-NEXT: and.b32 %r6, %r4, %r5;70; CHECK-NEXT: or.b32 %r7, %r3, %r6;71; CHECK-NEXT: st.param.b32 [func_retval0], %r7;72; CHECK-NEXT: ret;73 %mx = and <4 x i8> %x, %mask74 %notmask = xor <4 x i8> %mask, <i8 -1, i8 -1, i8 -1, i8 -1>75 %my = and <4 x i8> %y, %notmask76 %r = or <4 x i8> %mx, %my77 ret <4 x i8> %r78}79 80define <4 x i8> @out_v4i8_undef(<4 x i8> %x, <4 x i8> %y, <4 x i8> %mask) nounwind {81; CHECK-LABEL: out_v4i8_undef(82; CHECK: {83; CHECK-NEXT: .reg .b32 %r<8>;84; CHECK-EMPTY:85; CHECK-NEXT: // %bb.0:86; CHECK-NEXT: ld.param.b32 %r1, [out_v4i8_undef_param_0];87; CHECK-NEXT: ld.param.b32 %r2, [out_v4i8_undef_param_2];88; CHECK-NEXT: and.b32 %r3, %r1, %r2;89; CHECK-NEXT: ld.param.b32 %r4, [out_v4i8_undef_param_1];90; CHECK-NEXT: xor.b32 %r5, %r2, -16711681;91; CHECK-NEXT: and.b32 %r6, %r4, %r5;92; CHECK-NEXT: or.b32 %r7, %r3, %r6;93; CHECK-NEXT: st.param.b32 [func_retval0], %r7;94; CHECK-NEXT: ret;95 %mx = and <4 x i8> %x, %mask96 %notmask = xor <4 x i8> %mask, <i8 -1, i8 -1, i8 undef, i8 -1>97 %my = and <4 x i8> %y, %notmask98 %r = or <4 x i8> %mx, %my99 ret <4 x i8> %r100}101 102define <2 x i16> @out_v2i16(<2 x i16> %x, <2 x i16> %y, <2 x i16> %mask) nounwind {103; CHECK-LABEL: out_v2i16(104; CHECK: {105; CHECK-NEXT: .reg .b32 %r<8>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b32 %r1, [out_v2i16_param_0];109; CHECK-NEXT: ld.param.b32 %r2, [out_v2i16_param_2];110; CHECK-NEXT: and.b32 %r3, %r1, %r2;111; CHECK-NEXT: ld.param.b32 %r4, [out_v2i16_param_1];112; CHECK-NEXT: xor.b32 %r5, %r2, -1;113; CHECK-NEXT: and.b32 %r6, %r4, %r5;114; CHECK-NEXT: or.b32 %r7, %r3, %r6;115; CHECK-NEXT: st.param.b32 [func_retval0], %r7;116; CHECK-NEXT: ret;117 %mx = and <2 x i16> %x, %mask118 %notmask = xor <2 x i16> %mask, <i16 -1, i16 -1>119 %my = and <2 x i16> %y, %notmask120 %r = or <2 x i16> %mx, %my121 ret <2 x i16> %r122}123 124define <1 x i32> @out_v1i32(<1 x i32> %x, <1 x i32> %y, <1 x i32> %mask) nounwind {125; CHECK-LABEL: out_v1i32(126; CHECK: {127; CHECK-NEXT: .reg .b32 %r<7>;128; CHECK-EMPTY:129; CHECK-NEXT: // %bb.0:130; CHECK-NEXT: ld.param.b32 %r1, [out_v1i32_param_0];131; CHECK-NEXT: ld.param.b32 %r2, [out_v1i32_param_1];132; CHECK-NEXT: ld.param.b32 %r3, [out_v1i32_param_2];133; CHECK-NEXT: xor.b32 %r4, %r1, %r2;134; CHECK-NEXT: and.b32 %r5, %r4, %r3;135; CHECK-NEXT: xor.b32 %r6, %r5, %r2;136; CHECK-NEXT: st.param.b32 [func_retval0], %r6;137; CHECK-NEXT: ret;138 %mx = and <1 x i32> %x, %mask139 %notmask = xor <1 x i32> %mask, <i32 -1>140 %my = and <1 x i32> %y, %notmask141 %r = or <1 x i32> %mx, %my142 ret <1 x i32> %r143}144 145; ============================================================================ ;146; 64-bit vector width147; ============================================================================ ;148 149define <8 x i8> @out_v8i8(<8 x i8> %x, <8 x i8> %y, <8 x i8> %mask) nounwind {150; CHECK-LABEL: out_v8i8(151; CHECK: {152; CHECK-NEXT: .reg .b32 %r<15>;153; CHECK-EMPTY:154; CHECK-NEXT: // %bb.0:155; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [out_v8i8_param_0];156; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [out_v8i8_param_2];157; CHECK-NEXT: and.b32 %r5, %r1, %r3;158; CHECK-NEXT: and.b32 %r6, %r2, %r4;159; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [out_v8i8_param_1];160; CHECK-NEXT: xor.b32 %r9, %r4, -1;161; CHECK-NEXT: xor.b32 %r10, %r3, -1;162; CHECK-NEXT: and.b32 %r11, %r7, %r10;163; CHECK-NEXT: and.b32 %r12, %r8, %r9;164; CHECK-NEXT: or.b32 %r13, %r6, %r12;165; CHECK-NEXT: or.b32 %r14, %r5, %r11;166; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r13};167; CHECK-NEXT: ret;168 %mx = and <8 x i8> %x, %mask169 %notmask = xor <8 x i8> %mask, <i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1>170 %my = and <8 x i8> %y, %notmask171 %r = or <8 x i8> %mx, %my172 ret <8 x i8> %r173}174 175define <4 x i16> @out_v4i16(<4 x i16> %x, <4 x i16> %y, <4 x i16> %mask) nounwind {176; CHECK-LABEL: out_v4i16(177; CHECK: {178; CHECK-NEXT: .reg .b32 %r<15>;179; CHECK-EMPTY:180; CHECK-NEXT: // %bb.0:181; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [out_v4i16_param_0];182; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [out_v4i16_param_2];183; CHECK-NEXT: and.b32 %r5, %r1, %r3;184; CHECK-NEXT: and.b32 %r6, %r2, %r4;185; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [out_v4i16_param_1];186; CHECK-NEXT: xor.b32 %r9, %r4, -1;187; CHECK-NEXT: xor.b32 %r10, %r3, -1;188; CHECK-NEXT: and.b32 %r11, %r7, %r10;189; CHECK-NEXT: and.b32 %r12, %r8, %r9;190; CHECK-NEXT: or.b32 %r13, %r6, %r12;191; CHECK-NEXT: or.b32 %r14, %r5, %r11;192; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r13};193; CHECK-NEXT: ret;194 %mx = and <4 x i16> %x, %mask195 %notmask = xor <4 x i16> %mask, <i16 -1, i16 -1, i16 -1, i16 -1>196 %my = and <4 x i16> %y, %notmask197 %r = or <4 x i16> %mx, %my198 ret <4 x i16> %r199}200 201define <4 x i16> @out_v4i16_undef(<4 x i16> %x, <4 x i16> %y, <4 x i16> %mask) nounwind {202; CHECK-LABEL: out_v4i16_undef(203; CHECK: {204; CHECK-NEXT: .reg .b32 %r<15>;205; CHECK-EMPTY:206; CHECK-NEXT: // %bb.0:207; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [out_v4i16_undef_param_0];208; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [out_v4i16_undef_param_2];209; CHECK-NEXT: and.b32 %r5, %r1, %r3;210; CHECK-NEXT: and.b32 %r6, %r2, %r4;211; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [out_v4i16_undef_param_1];212; CHECK-NEXT: xor.b32 %r9, %r4, -65536;213; CHECK-NEXT: xor.b32 %r10, %r3, -1;214; CHECK-NEXT: and.b32 %r11, %r7, %r10;215; CHECK-NEXT: and.b32 %r12, %r8, %r9;216; CHECK-NEXT: or.b32 %r13, %r6, %r12;217; CHECK-NEXT: or.b32 %r14, %r5, %r11;218; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r13};219; CHECK-NEXT: ret;220 %mx = and <4 x i16> %x, %mask221 %notmask = xor <4 x i16> %mask, <i16 -1, i16 -1, i16 undef, i16 -1>222 %my = and <4 x i16> %y, %notmask223 %r = or <4 x i16> %mx, %my224 ret <4 x i16> %r225}226 227define <2 x i32> @out_v2i32(<2 x i32> %x, <2 x i32> %y, <2 x i32> %mask) nounwind {228; CHECK-LABEL: out_v2i32(229; CHECK: {230; CHECK-NEXT: .reg .b32 %r<13>;231; CHECK-EMPTY:232; CHECK-NEXT: // %bb.0:233; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [out_v2i32_param_0];234; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [out_v2i32_param_1];235; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [out_v2i32_param_2];236; CHECK-NEXT: xor.b32 %r7, %r2, %r4;237; CHECK-NEXT: and.b32 %r8, %r7, %r6;238; CHECK-NEXT: xor.b32 %r9, %r8, %r4;239; CHECK-NEXT: xor.b32 %r10, %r1, %r3;240; CHECK-NEXT: and.b32 %r11, %r10, %r5;241; CHECK-NEXT: xor.b32 %r12, %r11, %r3;242; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r12, %r9};243; CHECK-NEXT: ret;244 %mx = and <2 x i32> %x, %mask245 %notmask = xor <2 x i32> %mask, <i32 -1, i32 -1>246 %my = and <2 x i32> %y, %notmask247 %r = or <2 x i32> %mx, %my248 ret <2 x i32> %r249}250 251define <1 x i64> @out_v1i64(<1 x i64> %x, <1 x i64> %y, <1 x i64> %mask) nounwind {252; CHECK-LABEL: out_v1i64(253; CHECK: {254; CHECK-NEXT: .reg .b64 %rd<7>;255; CHECK-EMPTY:256; CHECK-NEXT: // %bb.0:257; CHECK-NEXT: ld.param.b64 %rd1, [out_v1i64_param_0];258; CHECK-NEXT: ld.param.b64 %rd2, [out_v1i64_param_1];259; CHECK-NEXT: ld.param.b64 %rd3, [out_v1i64_param_2];260; CHECK-NEXT: xor.b64 %rd4, %rd1, %rd2;261; CHECK-NEXT: and.b64 %rd5, %rd4, %rd3;262; CHECK-NEXT: xor.b64 %rd6, %rd5, %rd2;263; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;264; CHECK-NEXT: ret;265 %mx = and <1 x i64> %x, %mask266 %notmask = xor <1 x i64> %mask, <i64 -1>267 %my = and <1 x i64> %y, %notmask268 %r = or <1 x i64> %mx, %my269 ret <1 x i64> %r270}271 272; ============================================================================ ;273; 128-bit vector width274; ============================================================================ ;275 276define <16 x i8> @out_v16i8(<16 x i8> %x, <16 x i8> %y, <16 x i8> %mask) nounwind {277; CHECK-LABEL: out_v16i8(278; CHECK: {279; CHECK-NEXT: .reg .b32 %r<29>;280; CHECK-EMPTY:281; CHECK-NEXT: // %bb.0:282; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [out_v16i8_param_0];283; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [out_v16i8_param_2];284; CHECK-NEXT: and.b32 %r9, %r1, %r5;285; CHECK-NEXT: and.b32 %r10, %r2, %r6;286; CHECK-NEXT: and.b32 %r11, %r3, %r7;287; CHECK-NEXT: and.b32 %r12, %r4, %r8;288; CHECK-NEXT: ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [out_v16i8_param_1];289; CHECK-NEXT: xor.b32 %r17, %r8, -1;290; CHECK-NEXT: xor.b32 %r18, %r7, -1;291; CHECK-NEXT: xor.b32 %r19, %r6, -1;292; CHECK-NEXT: xor.b32 %r20, %r5, -1;293; CHECK-NEXT: and.b32 %r21, %r13, %r20;294; CHECK-NEXT: and.b32 %r22, %r14, %r19;295; CHECK-NEXT: and.b32 %r23, %r15, %r18;296; CHECK-NEXT: and.b32 %r24, %r16, %r17;297; CHECK-NEXT: or.b32 %r25, %r12, %r24;298; CHECK-NEXT: or.b32 %r26, %r11, %r23;299; CHECK-NEXT: or.b32 %r27, %r10, %r22;300; CHECK-NEXT: or.b32 %r28, %r9, %r21;301; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r28, %r27, %r26, %r25};302; CHECK-NEXT: ret;303 %mx = and <16 x i8> %x, %mask304 %notmask = xor <16 x i8> %mask, <i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1, i8 -1>305 %my = and <16 x i8> %y, %notmask306 %r = or <16 x i8> %mx, %my307 ret <16 x i8> %r308}309 310define <8 x i16> @out_v8i16(<8 x i16> %x, <8 x i16> %y, <8 x i16> %mask) nounwind {311; CHECK-LABEL: out_v8i16(312; CHECK: {313; CHECK-NEXT: .reg .b32 %r<29>;314; CHECK-EMPTY:315; CHECK-NEXT: // %bb.0:316; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [out_v8i16_param_0];317; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [out_v8i16_param_2];318; CHECK-NEXT: and.b32 %r9, %r1, %r5;319; CHECK-NEXT: and.b32 %r10, %r2, %r6;320; CHECK-NEXT: and.b32 %r11, %r3, %r7;321; CHECK-NEXT: and.b32 %r12, %r4, %r8;322; CHECK-NEXT: ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [out_v8i16_param_1];323; CHECK-NEXT: xor.b32 %r17, %r8, -1;324; CHECK-NEXT: xor.b32 %r18, %r7, -1;325; CHECK-NEXT: xor.b32 %r19, %r6, -1;326; CHECK-NEXT: xor.b32 %r20, %r5, -1;327; CHECK-NEXT: and.b32 %r21, %r13, %r20;328; CHECK-NEXT: and.b32 %r22, %r14, %r19;329; CHECK-NEXT: and.b32 %r23, %r15, %r18;330; CHECK-NEXT: and.b32 %r24, %r16, %r17;331; CHECK-NEXT: or.b32 %r25, %r12, %r24;332; CHECK-NEXT: or.b32 %r26, %r11, %r23;333; CHECK-NEXT: or.b32 %r27, %r10, %r22;334; CHECK-NEXT: or.b32 %r28, %r9, %r21;335; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r28, %r27, %r26, %r25};336; CHECK-NEXT: ret;337 %mx = and <8 x i16> %x, %mask338 %notmask = xor <8 x i16> %mask, <i16 -1, i16 -1, i16 -1, i16 -1, i16 -1, i16 -1, i16 -1, i16 -1>339 %my = and <8 x i16> %y, %notmask340 %r = or <8 x i16> %mx, %my341 ret <8 x i16> %r342}343 344define <4 x i32> @out_v4i32(<4 x i32> %x, <4 x i32> %y, <4 x i32> %mask) nounwind {345; CHECK-LABEL: out_v4i32(346; CHECK: {347; CHECK-NEXT: .reg .b32 %r<25>;348; CHECK-EMPTY:349; CHECK-NEXT: // %bb.0:350; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [out_v4i32_param_0];351; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [out_v4i32_param_1];352; CHECK-NEXT: ld.param.v4.b32 {%r9, %r10, %r11, %r12}, [out_v4i32_param_2];353; CHECK-NEXT: xor.b32 %r13, %r4, %r8;354; CHECK-NEXT: and.b32 %r14, %r13, %r12;355; CHECK-NEXT: xor.b32 %r15, %r14, %r8;356; CHECK-NEXT: xor.b32 %r16, %r3, %r7;357; CHECK-NEXT: and.b32 %r17, %r16, %r11;358; CHECK-NEXT: xor.b32 %r18, %r17, %r7;359; CHECK-NEXT: xor.b32 %r19, %r2, %r6;360; CHECK-NEXT: and.b32 %r20, %r19, %r10;361; CHECK-NEXT: xor.b32 %r21, %r20, %r6;362; CHECK-NEXT: xor.b32 %r22, %r1, %r5;363; CHECK-NEXT: and.b32 %r23, %r22, %r9;364; CHECK-NEXT: xor.b32 %r24, %r23, %r5;365; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r24, %r21, %r18, %r15};366; CHECK-NEXT: ret;367 %mx = and <4 x i32> %x, %mask368 %notmask = xor <4 x i32> %mask, <i32 -1, i32 -1, i32 -1, i32 -1>369 %my = and <4 x i32> %y, %notmask370 %r = or <4 x i32> %mx, %my371 ret <4 x i32> %r372}373 374define <4 x i32> @out_v4i32_undef(<4 x i32> %x, <4 x i32> %y, <4 x i32> %mask) nounwind {375; CHECK-LABEL: out_v4i32_undef(376; CHECK: {377; CHECK-NEXT: .reg .b32 %r<23>;378; CHECK-EMPTY:379; CHECK-NEXT: // %bb.0:380; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [out_v4i32_undef_param_0];381; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [out_v4i32_undef_param_2];382; CHECK-NEXT: and.b32 %r9, %r3, %r7;383; CHECK-NEXT: ld.param.v4.b32 {%r10, %r11, %r12, %r13}, [out_v4i32_undef_param_1];384; CHECK-NEXT: xor.b32 %r14, %r4, %r13;385; CHECK-NEXT: and.b32 %r15, %r14, %r8;386; CHECK-NEXT: xor.b32 %r16, %r15, %r13;387; CHECK-NEXT: xor.b32 %r17, %r2, %r11;388; CHECK-NEXT: and.b32 %r18, %r17, %r6;389; CHECK-NEXT: xor.b32 %r19, %r18, %r11;390; CHECK-NEXT: xor.b32 %r20, %r1, %r10;391; CHECK-NEXT: and.b32 %r21, %r20, %r5;392; CHECK-NEXT: xor.b32 %r22, %r21, %r10;393; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r22, %r19, %r9, %r16};394; CHECK-NEXT: ret;395 %mx = and <4 x i32> %x, %mask396 %notmask = xor <4 x i32> %mask, <i32 -1, i32 -1, i32 undef, i32 -1>397 %my = and <4 x i32> %y, %notmask398 %r = or <4 x i32> %mx, %my399 ret <4 x i32> %r400}401 402define <2 x i64> @out_v2i64(<2 x i64> %x, <2 x i64> %y, <2 x i64> %mask) nounwind {403; CHECK-LABEL: out_v2i64(404; CHECK: {405; CHECK-NEXT: .reg .b64 %rd<13>;406; CHECK-EMPTY:407; CHECK-NEXT: // %bb.0:408; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [out_v2i64_param_0];409; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [out_v2i64_param_1];410; CHECK-NEXT: ld.param.v2.b64 {%rd5, %rd6}, [out_v2i64_param_2];411; CHECK-NEXT: xor.b64 %rd7, %rd2, %rd4;412; CHECK-NEXT: and.b64 %rd8, %rd7, %rd6;413; CHECK-NEXT: xor.b64 %rd9, %rd8, %rd4;414; CHECK-NEXT: xor.b64 %rd10, %rd1, %rd3;415; CHECK-NEXT: and.b64 %rd11, %rd10, %rd5;416; CHECK-NEXT: xor.b64 %rd12, %rd11, %rd3;417; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd12, %rd9};418; CHECK-NEXT: ret;419 %mx = and <2 x i64> %x, %mask420 %notmask = xor <2 x i64> %mask, <i64 -1, i64 -1>421 %my = and <2 x i64> %y, %notmask422 %r = or <2 x i64> %mx, %my423 ret <2 x i64> %r424}425 426;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;427; Should be the same as the previous one.428;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;;429 430; ============================================================================ ;431; 8-bit vector width432; ============================================================================ ;433 434define <1 x i8> @in_v1i8(<1 x i8> %x, <1 x i8> %y, <1 x i8> %mask) nounwind {435; CHECK-LABEL: in_v1i8(436; CHECK: {437; CHECK-NEXT: .reg .b16 %rs<7>;438; CHECK-EMPTY:439; CHECK-NEXT: // %bb.0:440; CHECK-NEXT: ld.param.b8 %rs1, [in_v1i8_param_0];441; CHECK-NEXT: ld.param.b8 %rs2, [in_v1i8_param_1];442; CHECK-NEXT: xor.b16 %rs3, %rs1, %rs2;443; CHECK-NEXT: ld.param.b8 %rs4, [in_v1i8_param_2];444; CHECK-NEXT: and.b16 %rs5, %rs3, %rs4;445; CHECK-NEXT: xor.b16 %rs6, %rs5, %rs2;446; CHECK-NEXT: st.param.b8 [func_retval0], %rs6;447; CHECK-NEXT: ret;448 %n0 = xor <1 x i8> %x, %y449 %n1 = and <1 x i8> %n0, %mask450 %r = xor <1 x i8> %n1, %y451 ret <1 x i8> %r452}453 454; ============================================================================ ;455; 16-bit vector width456; ============================================================================ ;457 458define <1 x i16> @in_v1i16(<1 x i16> %x, <1 x i16> %y, <1 x i16> %mask) nounwind {459; CHECK-LABEL: in_v1i16(460; CHECK: {461; CHECK-NEXT: .reg .b16 %rs<7>;462; CHECK-EMPTY:463; CHECK-NEXT: // %bb.0:464; CHECK-NEXT: ld.param.b16 %rs1, [in_v1i16_param_0];465; CHECK-NEXT: ld.param.b16 %rs2, [in_v1i16_param_1];466; CHECK-NEXT: xor.b16 %rs3, %rs1, %rs2;467; CHECK-NEXT: ld.param.b16 %rs4, [in_v1i16_param_2];468; CHECK-NEXT: and.b16 %rs5, %rs3, %rs4;469; CHECK-NEXT: xor.b16 %rs6, %rs5, %rs2;470; CHECK-NEXT: st.param.b16 [func_retval0], %rs6;471; CHECK-NEXT: ret;472 %n0 = xor <1 x i16> %x, %y473 %n1 = and <1 x i16> %n0, %mask474 %r = xor <1 x i16> %n1, %y475 ret <1 x i16> %r476}477 478; ============================================================================ ;479; 32-bit vector width480; ============================================================================ ;481 482define <4 x i8> @in_v4i8(<4 x i8> %x, <4 x i8> %y, <4 x i8> %mask) nounwind {483; CHECK-LABEL: in_v4i8(484; CHECK: {485; CHECK-NEXT: .reg .b32 %r<7>;486; CHECK-EMPTY:487; CHECK-NEXT: // %bb.0:488; CHECK-NEXT: ld.param.b32 %r1, [in_v4i8_param_0];489; CHECK-NEXT: ld.param.b32 %r2, [in_v4i8_param_1];490; CHECK-NEXT: xor.b32 %r3, %r1, %r2;491; CHECK-NEXT: ld.param.b32 %r4, [in_v4i8_param_2];492; CHECK-NEXT: and.b32 %r5, %r3, %r4;493; CHECK-NEXT: xor.b32 %r6, %r5, %r2;494; CHECK-NEXT: st.param.b32 [func_retval0], %r6;495; CHECK-NEXT: ret;496 %n0 = xor <4 x i8> %x, %y497 %n1 = and <4 x i8> %n0, %mask498 %r = xor <4 x i8> %n1, %y499 ret <4 x i8> %r500}501 502define <2 x i16> @in_v2i16(<2 x i16> %x, <2 x i16> %y, <2 x i16> %mask) nounwind {503; CHECK-LABEL: in_v2i16(504; CHECK: {505; CHECK-NEXT: .reg .b32 %r<7>;506; CHECK-EMPTY:507; CHECK-NEXT: // %bb.0:508; CHECK-NEXT: ld.param.b32 %r1, [in_v2i16_param_0];509; CHECK-NEXT: ld.param.b32 %r2, [in_v2i16_param_1];510; CHECK-NEXT: xor.b32 %r3, %r1, %r2;511; CHECK-NEXT: ld.param.b32 %r4, [in_v2i16_param_2];512; CHECK-NEXT: and.b32 %r5, %r3, %r4;513; CHECK-NEXT: xor.b32 %r6, %r5, %r2;514; CHECK-NEXT: st.param.b32 [func_retval0], %r6;515; CHECK-NEXT: ret;516 %n0 = xor <2 x i16> %x, %y517 %n1 = and <2 x i16> %n0, %mask518 %r = xor <2 x i16> %n1, %y519 ret <2 x i16> %r520}521 522define <1 x i32> @in_v1i32(<1 x i32> %x, <1 x i32> %y, <1 x i32> %mask) nounwind {523; CHECK-LABEL: in_v1i32(524; CHECK: {525; CHECK-NEXT: .reg .b32 %r<7>;526; CHECK-EMPTY:527; CHECK-NEXT: // %bb.0:528; CHECK-NEXT: ld.param.b32 %r1, [in_v1i32_param_0];529; CHECK-NEXT: ld.param.b32 %r2, [in_v1i32_param_1];530; CHECK-NEXT: xor.b32 %r3, %r1, %r2;531; CHECK-NEXT: ld.param.b32 %r4, [in_v1i32_param_2];532; CHECK-NEXT: and.b32 %r5, %r3, %r4;533; CHECK-NEXT: xor.b32 %r6, %r5, %r2;534; CHECK-NEXT: st.param.b32 [func_retval0], %r6;535; CHECK-NEXT: ret;536 %n0 = xor <1 x i32> %x, %y537 %n1 = and <1 x i32> %n0, %mask538 %r = xor <1 x i32> %n1, %y539 ret <1 x i32> %r540}541 542; ============================================================================ ;543; 64-bit vector width544; ============================================================================ ;545 546define <8 x i8> @in_v8i8(<8 x i8> %x, <8 x i8> %y, <8 x i8> %mask) nounwind {547; CHECK-LABEL: in_v8i8(548; CHECK: {549; CHECK-NEXT: .reg .b32 %r<13>;550; CHECK-EMPTY:551; CHECK-NEXT: // %bb.0:552; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [in_v8i8_param_0];553; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [in_v8i8_param_1];554; CHECK-NEXT: xor.b32 %r5, %r2, %r4;555; CHECK-NEXT: xor.b32 %r6, %r1, %r3;556; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [in_v8i8_param_2];557; CHECK-NEXT: and.b32 %r9, %r6, %r7;558; CHECK-NEXT: and.b32 %r10, %r5, %r8;559; CHECK-NEXT: xor.b32 %r11, %r10, %r4;560; CHECK-NEXT: xor.b32 %r12, %r9, %r3;561; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r12, %r11};562; CHECK-NEXT: ret;563 %n0 = xor <8 x i8> %x, %y564 %n1 = and <8 x i8> %n0, %mask565 %r = xor <8 x i8> %n1, %y566 ret <8 x i8> %r567}568 569define <4 x i16> @in_v4i16(<4 x i16> %x, <4 x i16> %y, <4 x i16> %mask) nounwind {570; CHECK-LABEL: in_v4i16(571; CHECK: {572; CHECK-NEXT: .reg .b32 %r<13>;573; CHECK-EMPTY:574; CHECK-NEXT: // %bb.0:575; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [in_v4i16_param_0];576; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [in_v4i16_param_1];577; CHECK-NEXT: xor.b32 %r5, %r2, %r4;578; CHECK-NEXT: xor.b32 %r6, %r1, %r3;579; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [in_v4i16_param_2];580; CHECK-NEXT: and.b32 %r9, %r6, %r7;581; CHECK-NEXT: and.b32 %r10, %r5, %r8;582; CHECK-NEXT: xor.b32 %r11, %r10, %r4;583; CHECK-NEXT: xor.b32 %r12, %r9, %r3;584; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r12, %r11};585; CHECK-NEXT: ret;586 %n0 = xor <4 x i16> %x, %y587 %n1 = and <4 x i16> %n0, %mask588 %r = xor <4 x i16> %n1, %y589 ret <4 x i16> %r590}591 592define <2 x i32> @in_v2i32(<2 x i32> %x, <2 x i32> %y, <2 x i32> %mask) nounwind {593; CHECK-LABEL: in_v2i32(594; CHECK: {595; CHECK-NEXT: .reg .b32 %r<13>;596; CHECK-EMPTY:597; CHECK-NEXT: // %bb.0:598; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [in_v2i32_param_0];599; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [in_v2i32_param_1];600; CHECK-NEXT: xor.b32 %r5, %r2, %r4;601; CHECK-NEXT: xor.b32 %r6, %r1, %r3;602; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [in_v2i32_param_2];603; CHECK-NEXT: and.b32 %r9, %r6, %r7;604; CHECK-NEXT: and.b32 %r10, %r5, %r8;605; CHECK-NEXT: xor.b32 %r11, %r10, %r4;606; CHECK-NEXT: xor.b32 %r12, %r9, %r3;607; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r12, %r11};608; CHECK-NEXT: ret;609 %n0 = xor <2 x i32> %x, %y610 %n1 = and <2 x i32> %n0, %mask611 %r = xor <2 x i32> %n1, %y612 ret <2 x i32> %r613}614 615define <1 x i64> @in_v1i64(<1 x i64> %x, <1 x i64> %y, <1 x i64> %mask) nounwind {616; CHECK-LABEL: in_v1i64(617; CHECK: {618; CHECK-NEXT: .reg .b64 %rd<7>;619; CHECK-EMPTY:620; CHECK-NEXT: // %bb.0:621; CHECK-NEXT: ld.param.b64 %rd1, [in_v1i64_param_0];622; CHECK-NEXT: ld.param.b64 %rd2, [in_v1i64_param_1];623; CHECK-NEXT: xor.b64 %rd3, %rd1, %rd2;624; CHECK-NEXT: ld.param.b64 %rd4, [in_v1i64_param_2];625; CHECK-NEXT: and.b64 %rd5, %rd3, %rd4;626; CHECK-NEXT: xor.b64 %rd6, %rd5, %rd2;627; CHECK-NEXT: st.param.b64 [func_retval0], %rd6;628; CHECK-NEXT: ret;629 %n0 = xor <1 x i64> %x, %y630 %n1 = and <1 x i64> %n0, %mask631 %r = xor <1 x i64> %n1, %y632 ret <1 x i64> %r633}634 635; ============================================================================ ;636; 128-bit vector width637; ============================================================================ ;638 639define <16 x i8> @in_v16i8(<16 x i8> %x, <16 x i8> %y, <16 x i8> %mask) nounwind {640; CHECK-LABEL: in_v16i8(641; CHECK: {642; CHECK-NEXT: .reg .b32 %r<25>;643; CHECK-EMPTY:644; CHECK-NEXT: // %bb.0:645; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [in_v16i8_param_0];646; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [in_v16i8_param_1];647; CHECK-NEXT: xor.b32 %r9, %r4, %r8;648; CHECK-NEXT: xor.b32 %r10, %r3, %r7;649; CHECK-NEXT: xor.b32 %r11, %r2, %r6;650; CHECK-NEXT: xor.b32 %r12, %r1, %r5;651; CHECK-NEXT: ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [in_v16i8_param_2];652; CHECK-NEXT: and.b32 %r17, %r12, %r13;653; CHECK-NEXT: and.b32 %r18, %r11, %r14;654; CHECK-NEXT: and.b32 %r19, %r10, %r15;655; CHECK-NEXT: and.b32 %r20, %r9, %r16;656; CHECK-NEXT: xor.b32 %r21, %r20, %r8;657; CHECK-NEXT: xor.b32 %r22, %r19, %r7;658; CHECK-NEXT: xor.b32 %r23, %r18, %r6;659; CHECK-NEXT: xor.b32 %r24, %r17, %r5;660; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r24, %r23, %r22, %r21};661; CHECK-NEXT: ret;662 %n0 = xor <16 x i8> %x, %y663 %n1 = and <16 x i8> %n0, %mask664 %r = xor <16 x i8> %n1, %y665 ret <16 x i8> %r666}667 668define <8 x i16> @in_v8i16(<8 x i16> %x, <8 x i16> %y, <8 x i16> %mask) nounwind {669; CHECK-LABEL: in_v8i16(670; CHECK: {671; CHECK-NEXT: .reg .b32 %r<25>;672; CHECK-EMPTY:673; CHECK-NEXT: // %bb.0:674; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [in_v8i16_param_0];675; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [in_v8i16_param_1];676; CHECK-NEXT: xor.b32 %r9, %r4, %r8;677; CHECK-NEXT: xor.b32 %r10, %r3, %r7;678; CHECK-NEXT: xor.b32 %r11, %r2, %r6;679; CHECK-NEXT: xor.b32 %r12, %r1, %r5;680; CHECK-NEXT: ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [in_v8i16_param_2];681; CHECK-NEXT: and.b32 %r17, %r12, %r13;682; CHECK-NEXT: and.b32 %r18, %r11, %r14;683; CHECK-NEXT: and.b32 %r19, %r10, %r15;684; CHECK-NEXT: and.b32 %r20, %r9, %r16;685; CHECK-NEXT: xor.b32 %r21, %r20, %r8;686; CHECK-NEXT: xor.b32 %r22, %r19, %r7;687; CHECK-NEXT: xor.b32 %r23, %r18, %r6;688; CHECK-NEXT: xor.b32 %r24, %r17, %r5;689; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r24, %r23, %r22, %r21};690; CHECK-NEXT: ret;691 %n0 = xor <8 x i16> %x, %y692 %n1 = and <8 x i16> %n0, %mask693 %r = xor <8 x i16> %n1, %y694 ret <8 x i16> %r695}696 697define <4 x i32> @in_v4i32(<4 x i32> %x, <4 x i32> %y, <4 x i32> %mask) nounwind {698; CHECK-LABEL: in_v4i32(699; CHECK: {700; CHECK-NEXT: .reg .b32 %r<25>;701; CHECK-EMPTY:702; CHECK-NEXT: // %bb.0:703; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [in_v4i32_param_0];704; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [in_v4i32_param_1];705; CHECK-NEXT: xor.b32 %r9, %r4, %r8;706; CHECK-NEXT: xor.b32 %r10, %r3, %r7;707; CHECK-NEXT: xor.b32 %r11, %r2, %r6;708; CHECK-NEXT: xor.b32 %r12, %r1, %r5;709; CHECK-NEXT: ld.param.v4.b32 {%r13, %r14, %r15, %r16}, [in_v4i32_param_2];710; CHECK-NEXT: and.b32 %r17, %r12, %r13;711; CHECK-NEXT: and.b32 %r18, %r11, %r14;712; CHECK-NEXT: and.b32 %r19, %r10, %r15;713; CHECK-NEXT: and.b32 %r20, %r9, %r16;714; CHECK-NEXT: xor.b32 %r21, %r20, %r8;715; CHECK-NEXT: xor.b32 %r22, %r19, %r7;716; CHECK-NEXT: xor.b32 %r23, %r18, %r6;717; CHECK-NEXT: xor.b32 %r24, %r17, %r5;718; CHECK-NEXT: st.param.v4.b32 [func_retval0], {%r24, %r23, %r22, %r21};719; CHECK-NEXT: ret;720 %n0 = xor <4 x i32> %x, %y721 %n1 = and <4 x i32> %n0, %mask722 %r = xor <4 x i32> %n1, %y723 ret <4 x i32> %r724}725 726define <2 x i64> @in_v2i64(<2 x i64> %x, <2 x i64> %y, <2 x i64> %mask) nounwind {727; CHECK-LABEL: in_v2i64(728; CHECK: {729; CHECK-NEXT: .reg .b64 %rd<13>;730; CHECK-EMPTY:731; CHECK-NEXT: // %bb.0:732; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [in_v2i64_param_0];733; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [in_v2i64_param_1];734; CHECK-NEXT: xor.b64 %rd5, %rd2, %rd4;735; CHECK-NEXT: xor.b64 %rd6, %rd1, %rd3;736; CHECK-NEXT: ld.param.v2.b64 {%rd7, %rd8}, [in_v2i64_param_2];737; CHECK-NEXT: and.b64 %rd9, %rd6, %rd7;738; CHECK-NEXT: and.b64 %rd10, %rd5, %rd8;739; CHECK-NEXT: xor.b64 %rd11, %rd10, %rd4;740; CHECK-NEXT: xor.b64 %rd12, %rd9, %rd3;741; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd12, %rd11};742; CHECK-NEXT: ret;743 %n0 = xor <2 x i64> %x, %y744 %n1 = and <2 x i64> %n0, %mask745 %r = xor <2 x i64> %n1, %y746 ret <2 x i64> %r747}748