1775 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 32; RUN: llc < %s -mcpu=sm_80 -O0 -disable-post-ra -frame-pointer=all \3; RUN: -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-NOI32X2 %s4; RUN: %if ptxas-sm_80 %{ \5; RUN: llc < %s -mcpu=sm_80 -O0 -disable-post-ra -frame-pointer=all \6; RUN: -verify-machineinstrs | %ptxas-verify -arch=sm_80 \7; RUN: %}8; RUN: llc < %s -mcpu=sm_100 -O0 -disable-post-ra -frame-pointer=all \9; RUN: -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-I32X2 %s10; RUN: %if ptxas-sm_100 %{ \11; RUN: llc < %s -mcpu=sm_100 -O0 -disable-post-ra -frame-pointer=all \12; RUN: -verify-machineinstrs | %ptxas-verify -arch=sm_100 \13; RUN: %}14 15target triple = "nvptx64-nvidia-cuda"16target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"17 18define <2 x i32> @test_ret_const() #0 {19; CHECK-LABEL: test_ret_const(20; CHECK: {21; CHECK-EMPTY:22; CHECK-EMPTY:23; CHECK-NEXT: // %bb.0:24; CHECK-NEXT: st.param.v2.b32 [func_retval0], {-1, 2};25; CHECK-NEXT: ret;26 ret <2 x i32> <i32 -1, i32 2>27}28 29define i32 @test_extract_0(<2 x i32> %a) #0 {30; CHECK-NOI32X2-LABEL: test_extract_0(31; CHECK-NOI32X2: {32; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;33; CHECK-NOI32X2-EMPTY:34; CHECK-NOI32X2-NEXT: // %bb.0:35; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_0_param_0];36; CHECK-NOI32X2-NEXT: st.param.b32 [func_retval0], %r1;37; CHECK-NOI32X2-NEXT: ret;38;39; CHECK-I32X2-LABEL: test_extract_0(40; CHECK-I32X2: {41; CHECK-I32X2-NEXT: .reg .b32 %r<2>;42; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;43; CHECK-I32X2-EMPTY:44; CHECK-I32X2-NEXT: // %bb.0:45; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_extract_0_param_0];46; CHECK-I32X2-NEXT: mov.b64 {%r1, _}, %rd1;47; CHECK-I32X2-NEXT: st.param.b32 [func_retval0], %r1;48; CHECK-I32X2-NEXT: ret;49 %e = extractelement <2 x i32> %a, i32 050 ret i32 %e51}52 53define i32 @test_extract_1(<2 x i32> %a) #0 {54; CHECK-NOI32X2-LABEL: test_extract_1(55; CHECK-NOI32X2: {56; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;57; CHECK-NOI32X2-EMPTY:58; CHECK-NOI32X2-NEXT: // %bb.0:59; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_1_param_0];60; CHECK-NOI32X2-NEXT: st.param.b32 [func_retval0], %r2;61; CHECK-NOI32X2-NEXT: ret;62;63; CHECK-I32X2-LABEL: test_extract_1(64; CHECK-I32X2: {65; CHECK-I32X2-NEXT: .reg .b32 %r<2>;66; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;67; CHECK-I32X2-EMPTY:68; CHECK-I32X2-NEXT: // %bb.0:69; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_extract_1_param_0];70; CHECK-I32X2-NEXT: mov.b64 {_, %r1}, %rd1;71; CHECK-I32X2-NEXT: st.param.b32 [func_retval0], %r1;72; CHECK-I32X2-NEXT: ret;73 %e = extractelement <2 x i32> %a, i32 174 ret i32 %e75}76 77define i32 @test_extract_i(<2 x i32> %a, i64 %idx) #0 {78; CHECK-NOI32X2-LABEL: test_extract_i(79; CHECK-NOI32X2: {80; CHECK-NOI32X2-NEXT: .local .align 8 .b8 __local_depot3[8];81; CHECK-NOI32X2-NEXT: .reg .b64 %SP;82; CHECK-NOI32X2-NEXT: .reg .b64 %SPL;83; CHECK-NOI32X2-NEXT: .reg .b32 %r<4>;84; CHECK-NOI32X2-NEXT: .reg .b64 %rd<6>;85; CHECK-NOI32X2-EMPTY:86; CHECK-NOI32X2-NEXT: // %bb.0:87; CHECK-NOI32X2-NEXT: mov.b64 %SPL, __local_depot3;88; CHECK-NOI32X2-NEXT: cvta.local.u64 %SP, %SPL;89; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_i_param_0];90; CHECK-NOI32X2-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];91; CHECK-NOI32X2-NEXT: st.v2.b32 [%SP], {%r1, %r2};92; CHECK-NOI32X2-NEXT: and.b64 %rd2, %rd1, 1;93; CHECK-NOI32X2-NEXT: shl.b64 %rd3, %rd2, 2;94; CHECK-NOI32X2-NEXT: add.u64 %rd4, %SP, 0;95; CHECK-NOI32X2-NEXT: or.b64 %rd5, %rd4, %rd3;96; CHECK-NOI32X2-NEXT: ld.b32 %r3, [%rd5];97; CHECK-NOI32X2-NEXT: st.param.b32 [func_retval0], %r3;98; CHECK-NOI32X2-NEXT: ret;99;100; CHECK-I32X2-LABEL: test_extract_i(101; CHECK-I32X2: {102; CHECK-I32X2-NEXT: .reg .pred %p<2>;103; CHECK-I32X2-NEXT: .reg .b32 %r<4>;104; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;105; CHECK-I32X2-EMPTY:106; CHECK-I32X2-NEXT: // %bb.0:107; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_extract_i_param_1];108; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_extract_i_param_0];109; CHECK-I32X2-NEXT: setp.eq.b64 %p1, %rd2, 0;110; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;111; CHECK-I32X2-NEXT: selp.b32 %r3, %r1, %r2, %p1;112; CHECK-I32X2-NEXT: st.param.b32 [func_retval0], %r3;113; CHECK-I32X2-NEXT: ret;114 %e = extractelement <2 x i32> %a, i64 %idx115 ret i32 %e116}117 118define <2 x i32> @test_add(<2 x i32> %a, <2 x i32> %b) #0 {119; CHECK-NOI32X2-LABEL: test_add(120; CHECK-NOI32X2: {121; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;122; CHECK-NOI32X2-EMPTY:123; CHECK-NOI32X2-NEXT: // %bb.0:124; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_add_param_1];125; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_add_param_0];126; CHECK-NOI32X2-NEXT: add.s32 %r5, %r2, %r4;127; CHECK-NOI32X2-NEXT: add.s32 %r6, %r1, %r3;128; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};129; CHECK-NOI32X2-NEXT: ret;130;131; CHECK-I32X2-LABEL: test_add(132; CHECK-I32X2: {133; CHECK-I32X2-NEXT: .reg .b32 %r<7>;134; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;135; CHECK-I32X2-EMPTY:136; CHECK-I32X2-NEXT: // %bb.0:137; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_add_param_1];138; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_add_param_0];139; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;140; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;141; CHECK-I32X2-NEXT: add.s32 %r5, %r4, %r2;142; CHECK-I32X2-NEXT: add.s32 %r6, %r3, %r1;143; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};144; CHECK-I32X2-NEXT: ret;145 %r = add <2 x i32> %a, %b146 ret <2 x i32> %r147}148 149define <2 x i32> @test_add_imm_0(<2 x i32> %a) #0 {150; CHECK-NOI32X2-LABEL: test_add_imm_0(151; CHECK-NOI32X2: {152; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;153; CHECK-NOI32X2-EMPTY:154; CHECK-NOI32X2-NEXT: // %bb.0:155; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_add_imm_0_param_0];156; CHECK-NOI32X2-NEXT: add.s32 %r3, %r2, 2;157; CHECK-NOI32X2-NEXT: add.s32 %r4, %r1, 1;158; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};159; CHECK-NOI32X2-NEXT: ret;160;161; CHECK-I32X2-LABEL: test_add_imm_0(162; CHECK-I32X2: {163; CHECK-I32X2-NEXT: .reg .b32 %r<5>;164; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;165; CHECK-I32X2-EMPTY:166; CHECK-I32X2-NEXT: // %bb.0:167; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_add_imm_0_param_0];168; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;169; CHECK-I32X2-NEXT: add.s32 %r3, %r2, 2;170; CHECK-I32X2-NEXT: add.s32 %r4, %r1, 1;171; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};172; CHECK-I32X2-NEXT: ret;173 %r = add <2 x i32> <i32 1, i32 2>, %a174 ret <2 x i32> %r175}176 177define <2 x i32> @test_add_imm_1(<2 x i32> %a) #0 {178; CHECK-NOI32X2-LABEL: test_add_imm_1(179; CHECK-NOI32X2: {180; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;181; CHECK-NOI32X2-EMPTY:182; CHECK-NOI32X2-NEXT: // %bb.0:183; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_add_imm_1_param_0];184; CHECK-NOI32X2-NEXT: add.s32 %r3, %r2, 2;185; CHECK-NOI32X2-NEXT: add.s32 %r4, %r1, 1;186; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};187; CHECK-NOI32X2-NEXT: ret;188;189; CHECK-I32X2-LABEL: test_add_imm_1(190; CHECK-I32X2: {191; CHECK-I32X2-NEXT: .reg .b32 %r<5>;192; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;193; CHECK-I32X2-EMPTY:194; CHECK-I32X2-NEXT: // %bb.0:195; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_add_imm_1_param_0];196; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;197; CHECK-I32X2-NEXT: add.s32 %r3, %r2, 2;198; CHECK-I32X2-NEXT: add.s32 %r4, %r1, 1;199; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};200; CHECK-I32X2-NEXT: ret;201 %r = add <2 x i32> %a, <i32 1, i32 2>202 ret <2 x i32> %r203}204 205define <2 x i32> @test_sub(<2 x i32> %a, <2 x i32> %b) #0 {206; CHECK-NOI32X2-LABEL: test_sub(207; CHECK-NOI32X2: {208; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;209; CHECK-NOI32X2-EMPTY:210; CHECK-NOI32X2-NEXT: // %bb.0:211; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_sub_param_1];212; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sub_param_0];213; CHECK-NOI32X2-NEXT: sub.s32 %r5, %r2, %r4;214; CHECK-NOI32X2-NEXT: sub.s32 %r6, %r1, %r3;215; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};216; CHECK-NOI32X2-NEXT: ret;217;218; CHECK-I32X2-LABEL: test_sub(219; CHECK-I32X2: {220; CHECK-I32X2-NEXT: .reg .b32 %r<7>;221; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;222; CHECK-I32X2-EMPTY:223; CHECK-I32X2-NEXT: // %bb.0:224; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_sub_param_1];225; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_sub_param_0];226; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;227; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;228; CHECK-I32X2-NEXT: sub.s32 %r5, %r4, %r2;229; CHECK-I32X2-NEXT: sub.s32 %r6, %r3, %r1;230; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};231; CHECK-I32X2-NEXT: ret;232 %r = sub <2 x i32> %a, %b233 ret <2 x i32> %r234}235 236define <2 x i32> @test_smax(<2 x i32> %a, <2 x i32> %b) #0 {237; CHECK-NOI32X2-LABEL: test_smax(238; CHECK-NOI32X2: {239; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;240; CHECK-NOI32X2-EMPTY:241; CHECK-NOI32X2-NEXT: // %bb.0:242; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_smax_param_1];243; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_smax_param_0];244; CHECK-NOI32X2-NEXT: max.s32 %r5, %r2, %r4;245; CHECK-NOI32X2-NEXT: max.s32 %r6, %r1, %r3;246; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};247; CHECK-NOI32X2-NEXT: ret;248;249; CHECK-I32X2-LABEL: test_smax(250; CHECK-I32X2: {251; CHECK-I32X2-NEXT: .reg .b32 %r<7>;252; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;253; CHECK-I32X2-EMPTY:254; CHECK-I32X2-NEXT: // %bb.0:255; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_smax_param_1];256; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_smax_param_0];257; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;258; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;259; CHECK-I32X2-NEXT: max.s32 %r5, %r4, %r2;260; CHECK-I32X2-NEXT: max.s32 %r6, %r3, %r1;261; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};262; CHECK-I32X2-NEXT: ret;263 %cmp = icmp sgt <2 x i32> %a, %b264 %r = select <2 x i1> %cmp, <2 x i32> %a, <2 x i32> %b265 ret <2 x i32> %r266}267 268define <2 x i32> @test_umax(<2 x i32> %a, <2 x i32> %b) #0 {269; CHECK-NOI32X2-LABEL: test_umax(270; CHECK-NOI32X2: {271; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;272; CHECK-NOI32X2-EMPTY:273; CHECK-NOI32X2-NEXT: // %bb.0:274; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_umax_param_1];275; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_umax_param_0];276; CHECK-NOI32X2-NEXT: max.u32 %r5, %r2, %r4;277; CHECK-NOI32X2-NEXT: max.u32 %r6, %r1, %r3;278; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};279; CHECK-NOI32X2-NEXT: ret;280;281; CHECK-I32X2-LABEL: test_umax(282; CHECK-I32X2: {283; CHECK-I32X2-NEXT: .reg .b32 %r<7>;284; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;285; CHECK-I32X2-EMPTY:286; CHECK-I32X2-NEXT: // %bb.0:287; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_umax_param_1];288; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_umax_param_0];289; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;290; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;291; CHECK-I32X2-NEXT: max.u32 %r5, %r4, %r2;292; CHECK-I32X2-NEXT: max.u32 %r6, %r3, %r1;293; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};294; CHECK-I32X2-NEXT: ret;295 %cmp = icmp ugt <2 x i32> %a, %b296 %r = select <2 x i1> %cmp, <2 x i32> %a, <2 x i32> %b297 ret <2 x i32> %r298}299 300define <2 x i32> @test_smin(<2 x i32> %a, <2 x i32> %b) #0 {301; CHECK-NOI32X2-LABEL: test_smin(302; CHECK-NOI32X2: {303; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;304; CHECK-NOI32X2-EMPTY:305; CHECK-NOI32X2-NEXT: // %bb.0:306; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_smin_param_1];307; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_smin_param_0];308; CHECK-NOI32X2-NEXT: min.s32 %r5, %r2, %r4;309; CHECK-NOI32X2-NEXT: min.s32 %r6, %r1, %r3;310; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};311; CHECK-NOI32X2-NEXT: ret;312;313; CHECK-I32X2-LABEL: test_smin(314; CHECK-I32X2: {315; CHECK-I32X2-NEXT: .reg .b32 %r<7>;316; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;317; CHECK-I32X2-EMPTY:318; CHECK-I32X2-NEXT: // %bb.0:319; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_smin_param_1];320; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_smin_param_0];321; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;322; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;323; CHECK-I32X2-NEXT: min.s32 %r5, %r4, %r2;324; CHECK-I32X2-NEXT: min.s32 %r6, %r3, %r1;325; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};326; CHECK-I32X2-NEXT: ret;327 %cmp = icmp sle <2 x i32> %a, %b328 %r = select <2 x i1> %cmp, <2 x i32> %a, <2 x i32> %b329 ret <2 x i32> %r330}331 332define <2 x i32> @test_umin(<2 x i32> %a, <2 x i32> %b) #0 {333; CHECK-NOI32X2-LABEL: test_umin(334; CHECK-NOI32X2: {335; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;336; CHECK-NOI32X2-EMPTY:337; CHECK-NOI32X2-NEXT: // %bb.0:338; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_umin_param_1];339; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_umin_param_0];340; CHECK-NOI32X2-NEXT: min.u32 %r5, %r2, %r4;341; CHECK-NOI32X2-NEXT: min.u32 %r6, %r1, %r3;342; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};343; CHECK-NOI32X2-NEXT: ret;344;345; CHECK-I32X2-LABEL: test_umin(346; CHECK-I32X2: {347; CHECK-I32X2-NEXT: .reg .b32 %r<7>;348; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;349; CHECK-I32X2-EMPTY:350; CHECK-I32X2-NEXT: // %bb.0:351; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_umin_param_1];352; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_umin_param_0];353; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;354; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;355; CHECK-I32X2-NEXT: min.u32 %r5, %r4, %r2;356; CHECK-I32X2-NEXT: min.u32 %r6, %r3, %r1;357; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};358; CHECK-I32X2-NEXT: ret;359 %cmp = icmp ule <2 x i32> %a, %b360 %r = select <2 x i1> %cmp, <2 x i32> %a, <2 x i32> %b361 ret <2 x i32> %r362}363 364define <2 x i32> @test_eq(<2 x i32> %a, <2 x i32> %b, <2 x i32> %c) #0 {365; CHECK-NOI32X2-LABEL: test_eq(366; CHECK-NOI32X2: {367; CHECK-NOI32X2-NEXT: .reg .pred %p<3>;368; CHECK-NOI32X2-NEXT: .reg .b32 %r<9>;369; CHECK-NOI32X2-EMPTY:370; CHECK-NOI32X2-NEXT: // %bb.0:371; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_eq_param_2];372; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_eq_param_1];373; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_eq_param_0];374; CHECK-NOI32X2-NEXT: setp.eq.b32 %p1, %r1, %r3;375; CHECK-NOI32X2-NEXT: setp.eq.b32 %p2, %r2, %r4;376; CHECK-NOI32X2-NEXT: selp.b32 %r7, %r2, %r6, %p2;377; CHECK-NOI32X2-NEXT: selp.b32 %r8, %r1, %r5, %p1;378; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};379; CHECK-NOI32X2-NEXT: ret;380;381; CHECK-I32X2-LABEL: test_eq(382; CHECK-I32X2: {383; CHECK-I32X2-NEXT: .reg .pred %p<3>;384; CHECK-I32X2-NEXT: .reg .b32 %r<9>;385; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;386; CHECK-I32X2-EMPTY:387; CHECK-I32X2-NEXT: // %bb.0:388; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [test_eq_param_2];389; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_eq_param_1];390; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_eq_param_0];391; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;392; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;393; CHECK-I32X2-NEXT: setp.eq.b32 %p1, %r3, %r1;394; CHECK-I32X2-NEXT: setp.eq.b32 %p2, %r4, %r2;395; CHECK-I32X2-NEXT: mov.b64 {%r5, %r6}, %rd3;396; CHECK-I32X2-NEXT: selp.b32 %r7, %r4, %r6, %p2;397; CHECK-I32X2-NEXT: selp.b32 %r8, %r3, %r5, %p1;398; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};399; CHECK-I32X2-NEXT: ret;400 %cmp = icmp eq <2 x i32> %a, %b401 %r = select <2 x i1> %cmp, <2 x i32> %a, <2 x i32> %c402 ret <2 x i32> %r403}404 405define <2 x i32> @test_ne(<2 x i32> %a, <2 x i32> %b, <2 x i32> %c) #0 {406; CHECK-NOI32X2-LABEL: test_ne(407; CHECK-NOI32X2: {408; CHECK-NOI32X2-NEXT: .reg .pred %p<3>;409; CHECK-NOI32X2-NEXT: .reg .b32 %r<9>;410; CHECK-NOI32X2-EMPTY:411; CHECK-NOI32X2-NEXT: // %bb.0:412; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_ne_param_2];413; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_ne_param_1];414; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_ne_param_0];415; CHECK-NOI32X2-NEXT: setp.ne.b32 %p1, %r1, %r3;416; CHECK-NOI32X2-NEXT: setp.ne.b32 %p2, %r2, %r4;417; CHECK-NOI32X2-NEXT: selp.b32 %r7, %r2, %r6, %p2;418; CHECK-NOI32X2-NEXT: selp.b32 %r8, %r1, %r5, %p1;419; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};420; CHECK-NOI32X2-NEXT: ret;421;422; CHECK-I32X2-LABEL: test_ne(423; CHECK-I32X2: {424; CHECK-I32X2-NEXT: .reg .pred %p<3>;425; CHECK-I32X2-NEXT: .reg .b32 %r<9>;426; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;427; CHECK-I32X2-EMPTY:428; CHECK-I32X2-NEXT: // %bb.0:429; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [test_ne_param_2];430; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_ne_param_1];431; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_ne_param_0];432; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;433; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;434; CHECK-I32X2-NEXT: setp.ne.b32 %p1, %r3, %r1;435; CHECK-I32X2-NEXT: setp.ne.b32 %p2, %r4, %r2;436; CHECK-I32X2-NEXT: mov.b64 {%r5, %r6}, %rd3;437; CHECK-I32X2-NEXT: selp.b32 %r7, %r4, %r6, %p2;438; CHECK-I32X2-NEXT: selp.b32 %r8, %r3, %r5, %p1;439; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};440; CHECK-I32X2-NEXT: ret;441 %cmp = icmp ne <2 x i32> %a, %b442 %r = select <2 x i1> %cmp, <2 x i32> %a, <2 x i32> %c443 ret <2 x i32> %r444}445 446define <2 x i32> @test_mul(<2 x i32> %a, <2 x i32> %b) #0 {447; CHECK-NOI32X2-LABEL: test_mul(448; CHECK-NOI32X2: {449; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;450; CHECK-NOI32X2-EMPTY:451; CHECK-NOI32X2-NEXT: // %bb.0:452; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_mul_param_1];453; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_mul_param_0];454; CHECK-NOI32X2-NEXT: mul.lo.s32 %r5, %r2, %r4;455; CHECK-NOI32X2-NEXT: mul.lo.s32 %r6, %r1, %r3;456; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};457; CHECK-NOI32X2-NEXT: ret;458;459; CHECK-I32X2-LABEL: test_mul(460; CHECK-I32X2: {461; CHECK-I32X2-NEXT: .reg .b32 %r<7>;462; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;463; CHECK-I32X2-EMPTY:464; CHECK-I32X2-NEXT: // %bb.0:465; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_mul_param_1];466; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_mul_param_0];467; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;468; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;469; CHECK-I32X2-NEXT: mul.lo.s32 %r5, %r4, %r2;470; CHECK-I32X2-NEXT: mul.lo.s32 %r6, %r3, %r1;471; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};472; CHECK-I32X2-NEXT: ret;473 %r = mul <2 x i32> %a, %b474 ret <2 x i32> %r475}476 477define <2 x i32> @test_or(<2 x i32> %a, <2 x i32> %b) #0 {478; CHECK-NOI32X2-LABEL: test_or(479; CHECK-NOI32X2: {480; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;481; CHECK-NOI32X2-EMPTY:482; CHECK-NOI32X2-NEXT: // %bb.0:483; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_or_param_1];484; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_or_param_0];485; CHECK-NOI32X2-NEXT: or.b32 %r5, %r2, %r4;486; CHECK-NOI32X2-NEXT: or.b32 %r6, %r1, %r3;487; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};488; CHECK-NOI32X2-NEXT: ret;489;490; CHECK-I32X2-LABEL: test_or(491; CHECK-I32X2: {492; CHECK-I32X2-NEXT: .reg .b32 %r<7>;493; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;494; CHECK-I32X2-EMPTY:495; CHECK-I32X2-NEXT: // %bb.0:496; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_or_param_1];497; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_or_param_0];498; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;499; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;500; CHECK-I32X2-NEXT: or.b32 %r5, %r4, %r2;501; CHECK-I32X2-NEXT: or.b32 %r6, %r3, %r1;502; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};503; CHECK-I32X2-NEXT: ret;504 %r = or <2 x i32> %a, %b505 ret <2 x i32> %r506}507 508define <2 x i32> @test_or_computed(i32 %a) {509; CHECK-LABEL: test_or_computed(510; CHECK: {511; CHECK-NEXT: .reg .b32 %r<2>;512; CHECK-EMPTY:513; CHECK-NEXT: // %bb.0:514; CHECK-NEXT: ld.param.b32 %r1, [test_or_computed_param_0];515; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, 5};516; CHECK-NEXT: ret;517 %ins.0 = insertelement <2 x i32> zeroinitializer, i32 %a, i32 0518 %ins.1 = insertelement <2 x i32> %ins.0, i32 5, i32 1519 %r = or <2 x i32> %ins.1, %ins.0520 ret <2 x i32> %r521}522 523define <2 x i32> @test_or_imm_0(<2 x i32> %a) #0 {524; CHECK-NOI32X2-LABEL: test_or_imm_0(525; CHECK-NOI32X2: {526; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;527; CHECK-NOI32X2-EMPTY:528; CHECK-NOI32X2-NEXT: // %bb.0:529; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_or_imm_0_param_0];530; CHECK-NOI32X2-NEXT: or.b32 %r3, %r2, 2;531; CHECK-NOI32X2-NEXT: or.b32 %r4, %r1, 1;532; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};533; CHECK-NOI32X2-NEXT: ret;534;535; CHECK-I32X2-LABEL: test_or_imm_0(536; CHECK-I32X2: {537; CHECK-I32X2-NEXT: .reg .b32 %r<5>;538; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;539; CHECK-I32X2-EMPTY:540; CHECK-I32X2-NEXT: // %bb.0:541; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_or_imm_0_param_0];542; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;543; CHECK-I32X2-NEXT: or.b32 %r3, %r2, 2;544; CHECK-I32X2-NEXT: or.b32 %r4, %r1, 1;545; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};546; CHECK-I32X2-NEXT: ret;547 %r = or <2 x i32> <i32 1, i32 2>, %a548 ret <2 x i32> %r549}550 551define <2 x i32> @test_or_imm_1(<2 x i32> %a) #0 {552; CHECK-NOI32X2-LABEL: test_or_imm_1(553; CHECK-NOI32X2: {554; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;555; CHECK-NOI32X2-EMPTY:556; CHECK-NOI32X2-NEXT: // %bb.0:557; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_or_imm_1_param_0];558; CHECK-NOI32X2-NEXT: or.b32 %r3, %r2, 2;559; CHECK-NOI32X2-NEXT: or.b32 %r4, %r1, 1;560; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};561; CHECK-NOI32X2-NEXT: ret;562;563; CHECK-I32X2-LABEL: test_or_imm_1(564; CHECK-I32X2: {565; CHECK-I32X2-NEXT: .reg .b32 %r<5>;566; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;567; CHECK-I32X2-EMPTY:568; CHECK-I32X2-NEXT: // %bb.0:569; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_or_imm_1_param_0];570; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;571; CHECK-I32X2-NEXT: or.b32 %r3, %r2, 2;572; CHECK-I32X2-NEXT: or.b32 %r4, %r1, 1;573; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};574; CHECK-I32X2-NEXT: ret;575 %r = or <2 x i32> %a, <i32 1, i32 2>576 ret <2 x i32> %r577}578 579define <2 x i32> @test_xor(<2 x i32> %a, <2 x i32> %b) #0 {580; CHECK-NOI32X2-LABEL: test_xor(581; CHECK-NOI32X2: {582; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;583; CHECK-NOI32X2-EMPTY:584; CHECK-NOI32X2-NEXT: // %bb.0:585; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_xor_param_1];586; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_xor_param_0];587; CHECK-NOI32X2-NEXT: xor.b32 %r5, %r2, %r4;588; CHECK-NOI32X2-NEXT: xor.b32 %r6, %r1, %r3;589; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};590; CHECK-NOI32X2-NEXT: ret;591;592; CHECK-I32X2-LABEL: test_xor(593; CHECK-I32X2: {594; CHECK-I32X2-NEXT: .reg .b32 %r<7>;595; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;596; CHECK-I32X2-EMPTY:597; CHECK-I32X2-NEXT: // %bb.0:598; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_xor_param_1];599; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_xor_param_0];600; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;601; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;602; CHECK-I32X2-NEXT: xor.b32 %r5, %r4, %r2;603; CHECK-I32X2-NEXT: xor.b32 %r6, %r3, %r1;604; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};605; CHECK-I32X2-NEXT: ret;606 %r = xor <2 x i32> %a, %b607 ret <2 x i32> %r608}609 610define <2 x i32> @test_xor_computed(i32 %a) {611; CHECK-LABEL: test_xor_computed(612; CHECK: {613; CHECK-NEXT: .reg .b32 %r<2>;614; CHECK-EMPTY:615; CHECK-NEXT: // %bb.0:616; CHECK-NEXT: ld.param.b32 %r1, [test_xor_computed_param_0];617; CHECK-NEXT: st.param.v2.b32 [func_retval0], {0, 5};618; CHECK-NEXT: ret;619 %ins.0 = insertelement <2 x i32> zeroinitializer, i32 %a, i32 0620 %ins.1 = insertelement <2 x i32> %ins.0, i32 5, i32 1621 %r = xor <2 x i32> %ins.1, %ins.0622 ret <2 x i32> %r623}624 625define <2 x i32> @test_xor_imm_0(<2 x i32> %a) #0 {626; CHECK-NOI32X2-LABEL: test_xor_imm_0(627; CHECK-NOI32X2: {628; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;629; CHECK-NOI32X2-EMPTY:630; CHECK-NOI32X2-NEXT: // %bb.0:631; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_xor_imm_0_param_0];632; CHECK-NOI32X2-NEXT: xor.b32 %r3, %r2, 2;633; CHECK-NOI32X2-NEXT: xor.b32 %r4, %r1, 1;634; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};635; CHECK-NOI32X2-NEXT: ret;636;637; CHECK-I32X2-LABEL: test_xor_imm_0(638; CHECK-I32X2: {639; CHECK-I32X2-NEXT: .reg .b32 %r<5>;640; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;641; CHECK-I32X2-EMPTY:642; CHECK-I32X2-NEXT: // %bb.0:643; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_xor_imm_0_param_0];644; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;645; CHECK-I32X2-NEXT: xor.b32 %r3, %r2, 2;646; CHECK-I32X2-NEXT: xor.b32 %r4, %r1, 1;647; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};648; CHECK-I32X2-NEXT: ret;649 %r = xor <2 x i32> <i32 1, i32 2>, %a650 ret <2 x i32> %r651}652 653define <2 x i32> @test_xor_imm_1(<2 x i32> %a) #0 {654; CHECK-NOI32X2-LABEL: test_xor_imm_1(655; CHECK-NOI32X2: {656; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;657; CHECK-NOI32X2-EMPTY:658; CHECK-NOI32X2-NEXT: // %bb.0:659; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_xor_imm_1_param_0];660; CHECK-NOI32X2-NEXT: xor.b32 %r3, %r2, 2;661; CHECK-NOI32X2-NEXT: xor.b32 %r4, %r1, 1;662; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};663; CHECK-NOI32X2-NEXT: ret;664;665; CHECK-I32X2-LABEL: test_xor_imm_1(666; CHECK-I32X2: {667; CHECK-I32X2-NEXT: .reg .b32 %r<5>;668; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;669; CHECK-I32X2-EMPTY:670; CHECK-I32X2-NEXT: // %bb.0:671; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_xor_imm_1_param_0];672; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;673; CHECK-I32X2-NEXT: xor.b32 %r3, %r2, 2;674; CHECK-I32X2-NEXT: xor.b32 %r4, %r1, 1;675; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};676; CHECK-I32X2-NEXT: ret;677 %r = xor <2 x i32> %a, <i32 1, i32 2>678 ret <2 x i32> %r679}680 681define <2 x i32> @test_and(<2 x i32> %a, <2 x i32> %b) #0 {682; CHECK-NOI32X2-LABEL: test_and(683; CHECK-NOI32X2: {684; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;685; CHECK-NOI32X2-EMPTY:686; CHECK-NOI32X2-NEXT: // %bb.0:687; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_and_param_1];688; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_and_param_0];689; CHECK-NOI32X2-NEXT: and.b32 %r5, %r2, %r4;690; CHECK-NOI32X2-NEXT: and.b32 %r6, %r1, %r3;691; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};692; CHECK-NOI32X2-NEXT: ret;693;694; CHECK-I32X2-LABEL: test_and(695; CHECK-I32X2: {696; CHECK-I32X2-NEXT: .reg .b32 %r<7>;697; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;698; CHECK-I32X2-EMPTY:699; CHECK-I32X2-NEXT: // %bb.0:700; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_and_param_1];701; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_and_param_0];702; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;703; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;704; CHECK-I32X2-NEXT: and.b32 %r5, %r4, %r2;705; CHECK-I32X2-NEXT: and.b32 %r6, %r3, %r1;706; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};707; CHECK-I32X2-NEXT: ret;708 %r = and <2 x i32> %a, %b709 ret <2 x i32> %r710}711 712define <2 x i32> @test_and_computed(i32 %a) {713; CHECK-LABEL: test_and_computed(714; CHECK: {715; CHECK-NEXT: .reg .b32 %r<2>;716; CHECK-EMPTY:717; CHECK-NEXT: // %bb.0:718; CHECK-NEXT: ld.param.b32 %r1, [test_and_computed_param_0];719; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, 0};720; CHECK-NEXT: ret;721 %ins.0 = insertelement <2 x i32> zeroinitializer, i32 %a, i32 0722 %ins.1 = insertelement <2 x i32> %ins.0, i32 5, i32 1723 %r = and <2 x i32> %ins.1, %ins.0724 ret <2 x i32> %r725}726 727define <2 x i32> @test_and_imm_0(<2 x i32> %a) #0 {728; CHECK-NOI32X2-LABEL: test_and_imm_0(729; CHECK-NOI32X2: {730; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;731; CHECK-NOI32X2-EMPTY:732; CHECK-NOI32X2-NEXT: // %bb.0:733; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_and_imm_0_param_0];734; CHECK-NOI32X2-NEXT: and.b32 %r3, %r2, 2;735; CHECK-NOI32X2-NEXT: and.b32 %r4, %r1, 1;736; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};737; CHECK-NOI32X2-NEXT: ret;738;739; CHECK-I32X2-LABEL: test_and_imm_0(740; CHECK-I32X2: {741; CHECK-I32X2-NEXT: .reg .b32 %r<5>;742; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;743; CHECK-I32X2-EMPTY:744; CHECK-I32X2-NEXT: // %bb.0:745; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_and_imm_0_param_0];746; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;747; CHECK-I32X2-NEXT: and.b32 %r3, %r2, 2;748; CHECK-I32X2-NEXT: and.b32 %r4, %r1, 1;749; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};750; CHECK-I32X2-NEXT: ret;751 %r = and <2 x i32> <i32 1, i32 2>, %a752 ret <2 x i32> %r753}754 755define <2 x i32> @test_and_imm_1(<2 x i32> %a) #0 {756; CHECK-NOI32X2-LABEL: test_and_imm_1(757; CHECK-NOI32X2: {758; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;759; CHECK-NOI32X2-EMPTY:760; CHECK-NOI32X2-NEXT: // %bb.0:761; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_and_imm_1_param_0];762; CHECK-NOI32X2-NEXT: and.b32 %r3, %r2, 2;763; CHECK-NOI32X2-NEXT: and.b32 %r4, %r1, 1;764; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};765; CHECK-NOI32X2-NEXT: ret;766;767; CHECK-I32X2-LABEL: test_and_imm_1(768; CHECK-I32X2: {769; CHECK-I32X2-NEXT: .reg .b32 %r<5>;770; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;771; CHECK-I32X2-EMPTY:772; CHECK-I32X2-NEXT: // %bb.0:773; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_and_imm_1_param_0];774; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;775; CHECK-I32X2-NEXT: and.b32 %r3, %r2, 2;776; CHECK-I32X2-NEXT: and.b32 %r4, %r1, 1;777; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};778; CHECK-I32X2-NEXT: ret;779 %r = and <2 x i32> %a, <i32 1, i32 2>780 ret <2 x i32> %r781}782 783define void @test_ldst_v2i32(ptr %a, ptr %b) {784; CHECK-NOI32X2-LABEL: test_ldst_v2i32(785; CHECK-NOI32X2: {786; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;787; CHECK-NOI32X2-NEXT: .reg .b64 %rd<3>;788; CHECK-NOI32X2-EMPTY:789; CHECK-NOI32X2-NEXT: // %bb.0:790; CHECK-NOI32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v2i32_param_1];791; CHECK-NOI32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v2i32_param_0];792; CHECK-NOI32X2-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];793; CHECK-NOI32X2-NEXT: st.v2.b32 [%rd2], {%r1, %r2};794; CHECK-NOI32X2-NEXT: ret;795;796; CHECK-I32X2-LABEL: test_ldst_v2i32(797; CHECK-I32X2: {798; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;799; CHECK-I32X2-EMPTY:800; CHECK-I32X2-NEXT: // %bb.0:801; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v2i32_param_1];802; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v2i32_param_0];803; CHECK-I32X2-NEXT: ld.b64 %rd3, [%rd1];804; CHECK-I32X2-NEXT: st.b64 [%rd2], %rd3;805; CHECK-I32X2-NEXT: ret;806 %t1 = load <2 x i32>, ptr %a807 store <2 x i32> %t1, ptr %b, align 16808 ret void809}810 811define void @test_ldst_v3i32(ptr %a, ptr %b) {812; CHECK-LABEL: test_ldst_v3i32(813; CHECK: {814; CHECK-NEXT: .reg .b32 %r<2>;815; CHECK-NEXT: .reg .b64 %rd<4>;816; CHECK-EMPTY:817; CHECK-NEXT: // %bb.0:818; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v3i32_param_1];819; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v3i32_param_0];820; CHECK-NEXT: ld.b64 %rd3, [%rd1];821; CHECK-NEXT: ld.b32 %r1, [%rd1+8];822; CHECK-NEXT: st.b32 [%rd2+8], %r1;823; CHECK-NEXT: st.b64 [%rd2], %rd3;824; CHECK-NEXT: ret;825 %t1 = load <3 x i32>, ptr %a826 store <3 x i32> %t1, ptr %b, align 16827 ret void828}829 830define void @test_ldst_v4i32(ptr %a, ptr %b) {831; CHECK-NOI32X2-LABEL: test_ldst_v4i32(832; CHECK-NOI32X2: {833; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;834; CHECK-NOI32X2-NEXT: .reg .b64 %rd<3>;835; CHECK-NOI32X2-EMPTY:836; CHECK-NOI32X2-NEXT: // %bb.0:837; CHECK-NOI32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v4i32_param_1];838; CHECK-NOI32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v4i32_param_0];839; CHECK-NOI32X2-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];840; CHECK-NOI32X2-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};841; CHECK-NOI32X2-NEXT: ret;842;843; CHECK-I32X2-LABEL: test_ldst_v4i32(844; CHECK-I32X2: {845; CHECK-I32X2-NEXT: .reg .b64 %rd<5>;846; CHECK-I32X2-EMPTY:847; CHECK-I32X2-NEXT: // %bb.0:848; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v4i32_param_1];849; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v4i32_param_0];850; CHECK-I32X2-NEXT: ld.v2.b64 {%rd3, %rd4}, [%rd1];851; CHECK-I32X2-NEXT: st.v2.b64 [%rd2], {%rd3, %rd4};852; CHECK-I32X2-NEXT: ret;853 %t1 = load <4 x i32>, ptr %a854 store <4 x i32> %t1, ptr %b, align 16855 ret void856}857 858define void @test_ldst_v2i32_unaligned(ptr %a, ptr %b) {859; CHECK-NOI32X2-LABEL: test_ldst_v2i32_unaligned(860; CHECK-NOI32X2: {861; CHECK-NOI32X2-NEXT: .reg .b32 %r<13>;862; CHECK-NOI32X2-NEXT: .reg .b64 %rd<3>;863; CHECK-NOI32X2-EMPTY:864; CHECK-NOI32X2-NEXT: // %bb.0:865; CHECK-NOI32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v2i32_unaligned_param_1];866; CHECK-NOI32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v2i32_unaligned_param_0];867; CHECK-NOI32X2-NEXT: ld.b8 %r1, [%rd1+2];868; CHECK-NOI32X2-NEXT: shl.b32 %r2, %r1, 16;869; CHECK-NOI32X2-NEXT: ld.b8 %r3, [%rd1+3];870; CHECK-NOI32X2-NEXT: shl.b32 %r4, %r3, 24;871; CHECK-NOI32X2-NEXT: or.b32 %r5, %r4, %r2;872; CHECK-NOI32X2-NEXT: ld.b8 %r6, [%rd1];873; CHECK-NOI32X2-NEXT: ld.b8 %r7, [%rd1+1];874; CHECK-NOI32X2-NEXT: ld.b8 %r8, [%rd1+4];875; CHECK-NOI32X2-NEXT: ld.b8 %r9, [%rd1+5];876; CHECK-NOI32X2-NEXT: ld.b8 %r10, [%rd1+6];877; CHECK-NOI32X2-NEXT: ld.b8 %r11, [%rd1+7];878; CHECK-NOI32X2-NEXT: st.b8 [%rd2+7], %r11;879; CHECK-NOI32X2-NEXT: st.b8 [%rd2+6], %r10;880; CHECK-NOI32X2-NEXT: st.b8 [%rd2+5], %r9;881; CHECK-NOI32X2-NEXT: st.b8 [%rd2+4], %r8;882; CHECK-NOI32X2-NEXT: st.b8 [%rd2+1], %r7;883; CHECK-NOI32X2-NEXT: st.b8 [%rd2], %r6;884; CHECK-NOI32X2-NEXT: st.b8 [%rd2+3], %r3;885; CHECK-NOI32X2-NEXT: shr.u32 %r12, %r5, 16;886; CHECK-NOI32X2-NEXT: st.b8 [%rd2+2], %r12;887; CHECK-NOI32X2-NEXT: ret;888;889; CHECK-I32X2-LABEL: test_ldst_v2i32_unaligned(890; CHECK-I32X2: {891; CHECK-I32X2-NEXT: .reg .b64 %rd<28>;892; CHECK-I32X2-EMPTY:893; CHECK-I32X2-NEXT: // %bb.0:894; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v2i32_unaligned_param_1];895; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v2i32_unaligned_param_0];896; CHECK-I32X2-NEXT: ld.b8 %rd3, [%rd1];897; CHECK-I32X2-NEXT: ld.b8 %rd4, [%rd1+1];898; CHECK-I32X2-NEXT: shl.b64 %rd5, %rd4, 8;899; CHECK-I32X2-NEXT: or.b64 %rd6, %rd5, %rd3;900; CHECK-I32X2-NEXT: ld.b8 %rd7, [%rd1+2];901; CHECK-I32X2-NEXT: shl.b64 %rd8, %rd7, 16;902; CHECK-I32X2-NEXT: ld.b8 %rd9, [%rd1+3];903; CHECK-I32X2-NEXT: shl.b64 %rd10, %rd9, 24;904; CHECK-I32X2-NEXT: or.b64 %rd11, %rd10, %rd8;905; CHECK-I32X2-NEXT: or.b64 %rd12, %rd11, %rd6;906; CHECK-I32X2-NEXT: ld.b8 %rd13, [%rd1+4];907; CHECK-I32X2-NEXT: ld.b8 %rd14, [%rd1+5];908; CHECK-I32X2-NEXT: shl.b64 %rd15, %rd14, 8;909; CHECK-I32X2-NEXT: or.b64 %rd16, %rd15, %rd13;910; CHECK-I32X2-NEXT: ld.b8 %rd17, [%rd1+6];911; CHECK-I32X2-NEXT: shl.b64 %rd18, %rd17, 16;912; CHECK-I32X2-NEXT: ld.b8 %rd19, [%rd1+7];913; CHECK-I32X2-NEXT: shl.b64 %rd20, %rd19, 24;914; CHECK-I32X2-NEXT: or.b64 %rd21, %rd20, %rd18;915; CHECK-I32X2-NEXT: or.b64 %rd22, %rd21, %rd16;916; CHECK-I32X2-NEXT: shl.b64 %rd23, %rd22, 32;917; CHECK-I32X2-NEXT: or.b64 %rd24, %rd23, %rd12;918; CHECK-I32X2-NEXT: st.b8 [%rd2+6], %rd17;919; CHECK-I32X2-NEXT: shr.u64 %rd25, %rd24, 56;920; CHECK-I32X2-NEXT: st.b8 [%rd2+7], %rd25;921; CHECK-I32X2-NEXT: st.b8 [%rd2+4], %rd13;922; CHECK-I32X2-NEXT: shr.u64 %rd26, %rd24, 40;923; CHECK-I32X2-NEXT: st.b8 [%rd2+5], %rd26;924; CHECK-I32X2-NEXT: st.b8 [%rd2+1], %rd4;925; CHECK-I32X2-NEXT: st.b8 [%rd2], %rd3;926; CHECK-I32X2-NEXT: st.b8 [%rd2+3], %rd9;927; CHECK-I32X2-NEXT: shr.u64 %rd27, %rd24, 16;928; CHECK-I32X2-NEXT: st.b8 [%rd2+2], %rd27;929; CHECK-I32X2-NEXT: ret;930 %t1 = load <2 x i32>, ptr %a, align 1931 store <2 x i32> %t1, ptr %b, align 1932 ret void933}934 935declare <2 x i32> @test_callee(<2 x i32> %a, <2 x i32> %b) #0936 937define <2 x i32> @test_call(<2 x i32> %a, <2 x i32> %b) #0 {938; CHECK-NOI32X2-LABEL: test_call(939; CHECK-NOI32X2: {940; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;941; CHECK-NOI32X2-EMPTY:942; CHECK-NOI32X2-NEXT: // %bb.0:943; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_call_param_1];944; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_call_param_0];945; CHECK-NOI32X2-NEXT: { // callseq 0, 0946; CHECK-NOI32X2-NEXT: .param .align 8 .b8 param0[8];947; CHECK-NOI32X2-NEXT: .param .align 8 .b8 param1[8];948; CHECK-NOI32X2-NEXT: .param .align 8 .b8 retval0[8];949; CHECK-NOI32X2-NEXT: st.param.v2.b32 [param1], {%r3, %r4};950; CHECK-NOI32X2-NEXT: st.param.v2.b32 [param0], {%r1, %r2};951; CHECK-NOI32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);952; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [retval0];953; CHECK-NOI32X2-NEXT: } // callseq 0954; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r5, %r6};955; CHECK-NOI32X2-NEXT: ret;956;957; CHECK-I32X2-LABEL: test_call(958; CHECK-I32X2: {959; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;960; CHECK-I32X2-EMPTY:961; CHECK-I32X2-NEXT: // %bb.0:962; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_call_param_1];963; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_call_param_0];964; CHECK-I32X2-NEXT: { // callseq 0, 0965; CHECK-I32X2-NEXT: .param .align 8 .b8 param0[8];966; CHECK-I32X2-NEXT: .param .align 8 .b8 param1[8];967; CHECK-I32X2-NEXT: .param .align 8 .b8 retval0[8];968; CHECK-I32X2-NEXT: st.param.b64 [param1], %rd2;969; CHECK-I32X2-NEXT: st.param.b64 [param0], %rd1;970; CHECK-I32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);971; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [retval0];972; CHECK-I32X2-NEXT: } // callseq 0973; CHECK-I32X2-NEXT: st.param.b64 [func_retval0], %rd3;974; CHECK-I32X2-NEXT: ret;975 %r = call <2 x i32> @test_callee(<2 x i32> %a, <2 x i32> %b)976 ret <2 x i32> %r977}978 979define <2 x i32> @test_call_flipped(<2 x i32> %a, <2 x i32> %b) #0 {980; CHECK-NOI32X2-LABEL: test_call_flipped(981; CHECK-NOI32X2: {982; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;983; CHECK-NOI32X2-EMPTY:984; CHECK-NOI32X2-NEXT: // %bb.0:985; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_call_flipped_param_1];986; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_call_flipped_param_0];987; CHECK-NOI32X2-NEXT: { // callseq 1, 0988; CHECK-NOI32X2-NEXT: .param .align 8 .b8 param0[8];989; CHECK-NOI32X2-NEXT: .param .align 8 .b8 param1[8];990; CHECK-NOI32X2-NEXT: .param .align 8 .b8 retval0[8];991; CHECK-NOI32X2-NEXT: st.param.v2.b32 [param1], {%r1, %r2};992; CHECK-NOI32X2-NEXT: st.param.v2.b32 [param0], {%r3, %r4};993; CHECK-NOI32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);994; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [retval0];995; CHECK-NOI32X2-NEXT: } // callseq 1996; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r5, %r6};997; CHECK-NOI32X2-NEXT: ret;998;999; CHECK-I32X2-LABEL: test_call_flipped(1000; CHECK-I32X2: {1001; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;1002; CHECK-I32X2-EMPTY:1003; CHECK-I32X2-NEXT: // %bb.0:1004; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_call_flipped_param_1];1005; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_call_flipped_param_0];1006; CHECK-I32X2-NEXT: { // callseq 1, 01007; CHECK-I32X2-NEXT: .param .align 8 .b8 param0[8];1008; CHECK-I32X2-NEXT: .param .align 8 .b8 param1[8];1009; CHECK-I32X2-NEXT: .param .align 8 .b8 retval0[8];1010; CHECK-I32X2-NEXT: st.param.b64 [param1], %rd1;1011; CHECK-I32X2-NEXT: st.param.b64 [param0], %rd2;1012; CHECK-I32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);1013; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [retval0];1014; CHECK-I32X2-NEXT: } // callseq 11015; CHECK-I32X2-NEXT: st.param.b64 [func_retval0], %rd3;1016; CHECK-I32X2-NEXT: ret;1017 %r = call <2 x i32> @test_callee(<2 x i32> %b, <2 x i32> %a)1018 ret <2 x i32> %r1019}1020 1021define <2 x i32> @test_tailcall_flipped(<2 x i32> %a, <2 x i32> %b) #0 {1022; CHECK-NOI32X2-LABEL: test_tailcall_flipped(1023; CHECK-NOI32X2: {1024; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;1025; CHECK-NOI32X2-EMPTY:1026; CHECK-NOI32X2-NEXT: // %bb.0:1027; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_tailcall_flipped_param_1];1028; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_tailcall_flipped_param_0];1029; CHECK-NOI32X2-NEXT: { // callseq 2, 01030; CHECK-NOI32X2-NEXT: .param .align 8 .b8 param0[8];1031; CHECK-NOI32X2-NEXT: .param .align 8 .b8 param1[8];1032; CHECK-NOI32X2-NEXT: .param .align 8 .b8 retval0[8];1033; CHECK-NOI32X2-NEXT: st.param.v2.b32 [param1], {%r1, %r2};1034; CHECK-NOI32X2-NEXT: st.param.v2.b32 [param0], {%r3, %r4};1035; CHECK-NOI32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);1036; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [retval0];1037; CHECK-NOI32X2-NEXT: } // callseq 21038; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r5, %r6};1039; CHECK-NOI32X2-NEXT: ret;1040;1041; CHECK-I32X2-LABEL: test_tailcall_flipped(1042; CHECK-I32X2: {1043; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;1044; CHECK-I32X2-EMPTY:1045; CHECK-I32X2-NEXT: // %bb.0:1046; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_tailcall_flipped_param_1];1047; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_tailcall_flipped_param_0];1048; CHECK-I32X2-NEXT: { // callseq 2, 01049; CHECK-I32X2-NEXT: .param .align 8 .b8 param0[8];1050; CHECK-I32X2-NEXT: .param .align 8 .b8 param1[8];1051; CHECK-I32X2-NEXT: .param .align 8 .b8 retval0[8];1052; CHECK-I32X2-NEXT: st.param.b64 [param1], %rd1;1053; CHECK-I32X2-NEXT: st.param.b64 [param0], %rd2;1054; CHECK-I32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);1055; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [retval0];1056; CHECK-I32X2-NEXT: } // callseq 21057; CHECK-I32X2-NEXT: st.param.b64 [func_retval0], %rd3;1058; CHECK-I32X2-NEXT: ret;1059 %r = tail call <2 x i32> @test_callee(<2 x i32> %b, <2 x i32> %a)1060 ret <2 x i32> %r1061}1062 1063define <2 x i32> @test_select(<2 x i32> %a, <2 x i32> %b, i1 zeroext %c) #0 {1064; CHECK-NOI32X2-LABEL: test_select(1065; CHECK-NOI32X2: {1066; CHECK-NOI32X2-NEXT: .reg .pred %p<2>;1067; CHECK-NOI32X2-NEXT: .reg .b16 %rs<3>;1068; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;1069; CHECK-NOI32X2-EMPTY:1070; CHECK-NOI32X2-NEXT: // %bb.0:1071; CHECK-NOI32X2-NEXT: ld.param.b8 %rs1, [test_select_param_2];1072; CHECK-NOI32X2-NEXT: and.b16 %rs2, %rs1, 1;1073; CHECK-NOI32X2-NEXT: setp.ne.b16 %p1, %rs2, 0;1074; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_param_1];1075; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_param_0];1076; CHECK-NOI32X2-NEXT: selp.b32 %r5, %r2, %r4, %p1;1077; CHECK-NOI32X2-NEXT: selp.b32 %r6, %r1, %r3, %p1;1078; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};1079; CHECK-NOI32X2-NEXT: ret;1080;1081; CHECK-I32X2-LABEL: test_select(1082; CHECK-I32X2: {1083; CHECK-I32X2-NEXT: .reg .pred %p<2>;1084; CHECK-I32X2-NEXT: .reg .b16 %rs<3>;1085; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;1086; CHECK-I32X2-EMPTY:1087; CHECK-I32X2-NEXT: // %bb.0:1088; CHECK-I32X2-NEXT: ld.param.b8 %rs1, [test_select_param_2];1089; CHECK-I32X2-NEXT: and.b16 %rs2, %rs1, 1;1090; CHECK-I32X2-NEXT: setp.ne.b16 %p1, %rs2, 0;1091; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_select_param_1];1092; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_select_param_0];1093; CHECK-I32X2-NEXT: selp.b64 %rd3, %rd1, %rd2, %p1;1094; CHECK-I32X2-NEXT: st.param.b64 [func_retval0], %rd3;1095; CHECK-I32X2-NEXT: ret;1096 %r = select i1 %c, <2 x i32> %a, <2 x i32> %b1097 ret <2 x i32> %r1098}1099 1100define <2 x i32> @test_select_cc(<2 x i32> %a, <2 x i32> %b, <2 x i32> %c, <2 x i32> %d) #0 {1101; CHECK-NOI32X2-LABEL: test_select_cc(1102; CHECK-NOI32X2: {1103; CHECK-NOI32X2-NEXT: .reg .pred %p<3>;1104; CHECK-NOI32X2-NEXT: .reg .b32 %r<11>;1105; CHECK-NOI32X2-EMPTY:1106; CHECK-NOI32X2-NEXT: // %bb.0:1107; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r7, %r8}, [test_select_cc_param_3];1108; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_param_2];1109; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_param_1];1110; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_param_0];1111; CHECK-NOI32X2-NEXT: setp.ne.b32 %p1, %r5, %r7;1112; CHECK-NOI32X2-NEXT: setp.ne.b32 %p2, %r6, %r8;1113; CHECK-NOI32X2-NEXT: selp.b32 %r9, %r2, %r4, %p2;1114; CHECK-NOI32X2-NEXT: selp.b32 %r10, %r1, %r3, %p1;1115; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r10, %r9};1116; CHECK-NOI32X2-NEXT: ret;1117;1118; CHECK-I32X2-LABEL: test_select_cc(1119; CHECK-I32X2: {1120; CHECK-I32X2-NEXT: .reg .pred %p<3>;1121; CHECK-I32X2-NEXT: .reg .b32 %r<11>;1122; CHECK-I32X2-NEXT: .reg .b64 %rd<5>;1123; CHECK-I32X2-EMPTY:1124; CHECK-I32X2-NEXT: // %bb.0:1125; CHECK-I32X2-NEXT: ld.param.b64 %rd4, [test_select_cc_param_3];1126; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [test_select_cc_param_2];1127; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_select_cc_param_1];1128; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_select_cc_param_0];1129; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd4;1130; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd3;1131; CHECK-I32X2-NEXT: setp.ne.b32 %p1, %r3, %r1;1132; CHECK-I32X2-NEXT: setp.ne.b32 %p2, %r4, %r2;1133; CHECK-I32X2-NEXT: mov.b64 {%r5, %r6}, %rd2;1134; CHECK-I32X2-NEXT: mov.b64 {%r7, %r8}, %rd1;1135; CHECK-I32X2-NEXT: selp.b32 %r9, %r8, %r6, %p2;1136; CHECK-I32X2-NEXT: selp.b32 %r10, %r7, %r5, %p1;1137; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r10, %r9};1138; CHECK-I32X2-NEXT: ret;1139 %cc = icmp ne <2 x i32> %c, %d1140 %r = select <2 x i1> %cc, <2 x i32> %a, <2 x i32> %b1141 ret <2 x i32> %r1142}1143 1144define <2 x i16> @test_trunc_2xi32_to_2xi16(<2 x i32> %a) #0 {1145; CHECK-NOI32X2-LABEL: test_trunc_2xi32_to_2xi16(1146; CHECK-NOI32X2: {1147; CHECK-NOI32X2-NEXT: .reg .b32 %r<4>;1148; CHECK-NOI32X2-EMPTY:1149; CHECK-NOI32X2-NEXT: // %bb.0:1150; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_2xi32_to_2xi16_param_0];1151; CHECK-NOI32X2-NEXT: prmt.b32 %r3, %r1, %r2, 0x5410U;1152; CHECK-NOI32X2-NEXT: st.param.b32 [func_retval0], %r3;1153; CHECK-NOI32X2-NEXT: ret;1154;1155; CHECK-I32X2-LABEL: test_trunc_2xi32_to_2xi16(1156; CHECK-I32X2: {1157; CHECK-I32X2-NEXT: .reg .b32 %r<4>;1158; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1159; CHECK-I32X2-EMPTY:1160; CHECK-I32X2-NEXT: // %bb.0:1161; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_trunc_2xi32_to_2xi16_param_0];1162; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1163; CHECK-I32X2-NEXT: prmt.b32 %r3, %r1, %r2, 0x5410U;1164; CHECK-I32X2-NEXT: st.param.b32 [func_retval0], %r3;1165; CHECK-I32X2-NEXT: ret;1166 %r = trunc <2 x i32> %a to <2 x i16>1167 ret <2 x i16> %r1168}1169 1170define <2 x i8> @test_trunc_2xi32_to_2xi8(<2 x i32> %a) #0 {1171; CHECK-NOI32X2-LABEL: test_trunc_2xi32_to_2xi8(1172; CHECK-NOI32X2: {1173; CHECK-NOI32X2-NEXT: .reg .b16 %rs<3>;1174; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;1175; CHECK-NOI32X2-EMPTY:1176; CHECK-NOI32X2-NEXT: // %bb.0:1177; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_2xi32_to_2xi8_param_0];1178; CHECK-NOI32X2-NEXT: cvt.u16.u32 %rs1, %r2;1179; CHECK-NOI32X2-NEXT: cvt.u16.u32 %rs2, %r1;1180; CHECK-NOI32X2-NEXT: st.param.v2.b8 [func_retval0], {%rs2, %rs1};1181; CHECK-NOI32X2-NEXT: ret;1182;1183; CHECK-I32X2-LABEL: test_trunc_2xi32_to_2xi8(1184; CHECK-I32X2: {1185; CHECK-I32X2-NEXT: .reg .b16 %rs<3>;1186; CHECK-I32X2-NEXT: .reg .b32 %r<3>;1187; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1188; CHECK-I32X2-EMPTY:1189; CHECK-I32X2-NEXT: // %bb.0:1190; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_trunc_2xi32_to_2xi8_param_0];1191; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1192; CHECK-I32X2-NEXT: cvt.u16.u32 %rs1, %r2;1193; CHECK-I32X2-NEXT: cvt.u16.u32 %rs2, %r1;1194; CHECK-I32X2-NEXT: st.param.v2.b8 [func_retval0], {%rs2, %rs1};1195; CHECK-I32X2-NEXT: ret;1196 %r = trunc <2 x i32> %a to <2 x i8>1197 ret <2 x i8> %r1198}1199 1200define <2 x i1> @test_trunc_2xi32_to_2xi1(<2 x i32> %a) #0 {1201; CHECK-NOI32X2-LABEL: test_trunc_2xi32_to_2xi1(1202; CHECK-NOI32X2: {1203; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;1204; CHECK-NOI32X2-EMPTY:1205; CHECK-NOI32X2-NEXT: // %bb.0:1206; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_2xi32_to_2xi1_param_0];1207; CHECK-NOI32X2-NEXT: st.param.b8 [func_retval0], %r1;1208; CHECK-NOI32X2-NEXT: st.param.b8 [func_retval0+1], %r2;1209; CHECK-NOI32X2-NEXT: ret;1210;1211; CHECK-I32X2-LABEL: test_trunc_2xi32_to_2xi1(1212; CHECK-I32X2: {1213; CHECK-I32X2-NEXT: .reg .b32 %r<3>;1214; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1215; CHECK-I32X2-EMPTY:1216; CHECK-I32X2-NEXT: // %bb.0:1217; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_trunc_2xi32_to_2xi1_param_0];1218; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1219; CHECK-I32X2-NEXT: st.param.b8 [func_retval0], %r1;1220; CHECK-I32X2-NEXT: st.param.b8 [func_retval0+1], %r2;1221; CHECK-I32X2-NEXT: ret;1222 %r = trunc <2 x i32> %a to <2 x i1>1223 ret <2 x i1> %r1224}1225 1226define <2 x i32> @test_trunc_2xi64(<2 x i64> %a) #0 {1227; CHECK-LABEL: test_trunc_2xi64(1228; CHECK: {1229; CHECK-NEXT: .reg .b32 %r<3>;1230; CHECK-NEXT: .reg .b64 %rd<3>;1231; CHECK-EMPTY:1232; CHECK-NEXT: // %bb.0:1233; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];1234; CHECK-NEXT: cvt.u32.u64 %r1, %rd2;1235; CHECK-NEXT: cvt.u32.u64 %r2, %rd1;1236; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};1237; CHECK-NEXT: ret;1238 %r = trunc <2 x i64> %a to <2 x i32>1239 ret <2 x i32> %r1240}1241 1242define <2 x i32> @test_zext_2xi16_to_2xi32(<2 x i16> %a) #0 {1243; CHECK-LABEL: test_zext_2xi16_to_2xi32(1244; CHECK: {1245; CHECK-NEXT: .reg .b16 %rs<3>;1246; CHECK-NEXT: .reg .b32 %r<4>;1247; CHECK-EMPTY:1248; CHECK-NEXT: // %bb.0:1249; CHECK-NEXT: ld.param.b32 %r1, [test_zext_2xi16_to_2xi32_param_0];1250; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1251; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;1252; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;1253; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};1254; CHECK-NEXT: ret;1255 %r = zext <2 x i16> %a to <2 x i32>1256 ret <2 x i32> %r1257}1258 1259define <2 x i32> @test_zext_2xi8_to_2xi32(<2 x i8> %a) #0 {1260; CHECK-LABEL: test_zext_2xi8_to_2xi32(1261; CHECK: {1262; CHECK-NEXT: .reg .b16 %rs<3>;1263; CHECK-NEXT: .reg .b32 %r<4>;1264; CHECK-EMPTY:1265; CHECK-NEXT: // %bb.0:1266; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_zext_2xi8_to_2xi32_param_0];1267; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};1268; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;1269; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;1270; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};1271; CHECK-NEXT: ret;1272 %r = zext <2 x i8> %a to <2 x i32>1273 ret <2 x i32> %r1274}1275 1276define <2 x i32> @test_zext_2xi1_to_2xi32(<2 x i1> %a) #0 {1277; CHECK-LABEL: test_zext_2xi1_to_2xi32(1278; CHECK: {1279; CHECK-NEXT: .reg .pred %p<3>;1280; CHECK-NEXT: .reg .b16 %rs<5>;1281; CHECK-NEXT: .reg .b32 %r<5>;1282; CHECK-EMPTY:1283; CHECK-NEXT: // %bb.0:1284; CHECK-NEXT: ld.param.b8 %rs1, [test_zext_2xi1_to_2xi32_param_0+1];1285; CHECK-NEXT: and.b16 %rs2, %rs1, 1;1286; CHECK-NEXT: setp.ne.b16 %p2, %rs2, 0;1287; CHECK-NEXT: ld.param.b8 %rs3, [test_zext_2xi1_to_2xi32_param_0];1288; CHECK-NEXT: and.b16 %rs4, %rs3, 1;1289; CHECK-NEXT: setp.ne.b16 %p1, %rs4, 0;1290; CHECK-NEXT: cvt.u32.u16 %r1, %rs1;1291; CHECK-NEXT: and.b32 %r2, %r1, 1;1292; CHECK-NEXT: cvt.u32.u16 %r3, %rs3;1293; CHECK-NEXT: and.b32 %r4, %r3, 1;1294; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r2};1295; CHECK-NEXT: ret;1296 %r = zext <2 x i1> %a to <2 x i32>1297 ret <2 x i32> %r1298}1299 1300define <2 x i64> @test_zext_2xi64(<2 x i32> %a) #0 {1301; CHECK-NOI32X2-LABEL: test_zext_2xi64(1302; CHECK-NOI32X2: {1303; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;1304; CHECK-NOI32X2-NEXT: .reg .b64 %rd<3>;1305; CHECK-NOI32X2-EMPTY:1306; CHECK-NOI32X2-NEXT: // %bb.0:1307; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_zext_2xi64_param_0];1308; CHECK-NOI32X2-NEXT: cvt.u64.u32 %rd1, %r2;1309; CHECK-NOI32X2-NEXT: cvt.u64.u32 %rd2, %r1;1310; CHECK-NOI32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};1311; CHECK-NOI32X2-NEXT: ret;1312;1313; CHECK-I32X2-LABEL: test_zext_2xi64(1314; CHECK-I32X2: {1315; CHECK-I32X2-NEXT: .reg .b32 %r<3>;1316; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;1317; CHECK-I32X2-EMPTY:1318; CHECK-I32X2-NEXT: // %bb.0:1319; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_zext_2xi64_param_0];1320; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1321; CHECK-I32X2-NEXT: cvt.u64.u32 %rd2, %r2;1322; CHECK-I32X2-NEXT: cvt.u64.u32 %rd3, %r1;1323; CHECK-I32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};1324; CHECK-I32X2-NEXT: ret;1325 %r = zext <2 x i32> %a to <2 x i64>1326 ret <2 x i64> %r1327}1328 1329define <2 x i32> @test_bitcast_i64_to_2xi32(i64 %a) #0 {1330; CHECK-LABEL: test_bitcast_i64_to_2xi32(1331; CHECK: {1332; CHECK-NEXT: .reg .b64 %rd<2>;1333; CHECK-EMPTY:1334; CHECK-NEXT: // %bb.0:1335; CHECK-NEXT: ld.param.b64 %rd1, [test_bitcast_i64_to_2xi32_param_0];1336; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;1337; CHECK-NEXT: ret;1338 %r = bitcast i64 %a to <2 x i32>1339 ret <2 x i32> %r1340}1341 1342define <2 x i32> @test_bitcast_double_to_2xi32(double %a) #0 {1343; CHECK-LABEL: test_bitcast_double_to_2xi32(1344; CHECK: {1345; CHECK-NEXT: .reg .b64 %rd<2>;1346; CHECK-EMPTY:1347; CHECK-NEXT: // %bb.0:1348; CHECK-NEXT: ld.param.b64 %rd1, [test_bitcast_double_to_2xi32_param_0];1349; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;1350; CHECK-NEXT: ret;1351 %r = bitcast double %a to <2 x i32>1352 ret <2 x i32> %r1353}1354 1355define i64 @test_bitcast_2xi32_to_i64(<2 x i32> %a) #0 {1356; CHECK-NOI32X2-LABEL: test_bitcast_2xi32_to_i64(1357; CHECK-NOI32X2: {1358; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;1359; CHECK-NOI32X2-EMPTY:1360; CHECK-NOI32X2-NEXT: // %bb.0:1361; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_bitcast_2xi32_to_i64_param_0];1362; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};1363; CHECK-NOI32X2-NEXT: ret;1364;1365; CHECK-I32X2-LABEL: test_bitcast_2xi32_to_i64(1366; CHECK-I32X2: {1367; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1368; CHECK-I32X2-EMPTY:1369; CHECK-I32X2-NEXT: // %bb.0:1370; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_bitcast_2xi32_to_i64_param_0];1371; CHECK-I32X2-NEXT: st.param.b64 [func_retval0], %rd1;1372; CHECK-I32X2-NEXT: ret;1373 %r = bitcast <2 x i32> %a to i641374 ret i64 %r1375}1376 1377define double @test_bitcast_2xi32_to_double(<2 x i32> %a) #0 {1378; CHECK-NOI32X2-LABEL: test_bitcast_2xi32_to_double(1379; CHECK-NOI32X2: {1380; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;1381; CHECK-NOI32X2-EMPTY:1382; CHECK-NOI32X2-NEXT: // %bb.0:1383; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_bitcast_2xi32_to_double_param_0];1384; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};1385; CHECK-NOI32X2-NEXT: ret;1386;1387; CHECK-I32X2-LABEL: test_bitcast_2xi32_to_double(1388; CHECK-I32X2: {1389; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1390; CHECK-I32X2-EMPTY:1391; CHECK-I32X2-NEXT: // %bb.0:1392; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_bitcast_2xi32_to_double_param_0];1393; CHECK-I32X2-NEXT: st.param.b64 [func_retval0], %rd1;1394; CHECK-I32X2-NEXT: ret;1395 %r = bitcast <2 x i32> %a to double1396 ret double %r1397}1398 1399 1400define <4 x half> @test_bitcast_2xi32_to_4xhalf(i32 %a) #0 {1401; CHECK-LABEL: test_bitcast_2xi32_to_4xhalf(1402; CHECK: {1403; CHECK-NEXT: .reg .b32 %r<2>;1404; CHECK-EMPTY:1405; CHECK-NEXT: // %bb.0:1406; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_2xi32_to_4xhalf_param_0];1407; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r1, 5};1408; CHECK-NEXT: ret;1409 %ins.0 = insertelement <2 x i32> poison, i32 %a, i32 01410 %ins.1 = insertelement <2 x i32> %ins.0, i32 5, i32 11411 %r = bitcast <2 x i32> %ins.1 to <4 x half>1412 ret <4 x half> %r1413}1414 1415 1416define <2 x i32> @test_shufflevector(<2 x i32> %a) #0 {1417; CHECK-NOI32X2-LABEL: test_shufflevector(1418; CHECK-NOI32X2: {1419; CHECK-NOI32X2-NEXT: .reg .b32 %r<3>;1420; CHECK-NOI32X2-EMPTY:1421; CHECK-NOI32X2-NEXT: // %bb.0:1422; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_shufflevector_param_0];1423; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};1424; CHECK-NOI32X2-NEXT: ret;1425;1426; CHECK-I32X2-LABEL: test_shufflevector(1427; CHECK-I32X2: {1428; CHECK-I32X2-NEXT: .reg .b32 %r<3>;1429; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1430; CHECK-I32X2-EMPTY:1431; CHECK-I32X2-NEXT: // %bb.0:1432; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_shufflevector_param_0];1433; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1434; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};1435; CHECK-I32X2-NEXT: ret;1436 %s = shufflevector <2 x i32> %a, <2 x i32> poison, <2 x i32> <i32 1, i32 0>1437 ret <2 x i32> %s1438}1439 1440define <2 x i32> @test_shufflevector_2(<2 x i32> %a, <2 x i32> %b) #0 {1441; CHECK-NOI32X2-LABEL: test_shufflevector_2(1442; CHECK-NOI32X2: {1443; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;1444; CHECK-NOI32X2-EMPTY:1445; CHECK-NOI32X2-NEXT: // %bb.0:1446; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_shufflevector_2_param_1];1447; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_shufflevector_2_param_0];1448; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r4};1449; CHECK-NOI32X2-NEXT: ret;1450;1451; CHECK-I32X2-LABEL: test_shufflevector_2(1452; CHECK-I32X2: {1453; CHECK-I32X2-NEXT: .reg .b32 %r<3>;1454; CHECK-I32X2-NEXT: .reg .b64 %rd<3>;1455; CHECK-I32X2-EMPTY:1456; CHECK-I32X2-NEXT: // %bb.0:1457; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_shufflevector_2_param_1];1458; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_shufflevector_2_param_0];1459; CHECK-I32X2-NEXT: mov.b64 {_, %r1}, %rd2;1460; CHECK-I32X2-NEXT: mov.b64 {_, %r2}, %rd1;1461; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};1462; CHECK-I32X2-NEXT: ret;1463 %s = shufflevector <2 x i32> %a, <2 x i32> %b, <2 x i32> <i32 1, i32 3>1464 ret <2 x i32> %s1465}1466 1467 1468define <2 x i32> @test_insertelement(<2 x i32> %a, i32 %x) #0 {1469; CHECK-NOI32X2-LABEL: test_insertelement(1470; CHECK-NOI32X2: {1471; CHECK-NOI32X2-NEXT: .reg .b32 %r<4>;1472; CHECK-NOI32X2-EMPTY:1473; CHECK-NOI32X2-NEXT: // %bb.0:1474; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_insertelement_param_0];1475; CHECK-NOI32X2-NEXT: ld.param.b32 %r3, [test_insertelement_param_1];1476; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r3};1477; CHECK-NOI32X2-NEXT: ret;1478;1479; CHECK-I32X2-LABEL: test_insertelement(1480; CHECK-I32X2: {1481; CHECK-I32X2-NEXT: .reg .b32 %r<3>;1482; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1483; CHECK-I32X2-EMPTY:1484; CHECK-I32X2-NEXT: // %bb.0:1485; CHECK-I32X2-NEXT: ld.param.b32 %r1, [test_insertelement_param_1];1486; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_insertelement_param_0];1487; CHECK-I32X2-NEXT: mov.b64 {%r2, _}, %rd1;1488; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};1489; CHECK-I32X2-NEXT: ret;1490 %i = insertelement <2 x i32> %a, i32 %x, i64 11491 ret <2 x i32> %i1492}1493 1494define <2 x i32> @test_fptosi_2xhalf_to_2xi32(<2 x half> %a) #0 {1495; CHECK-LABEL: test_fptosi_2xhalf_to_2xi32(1496; CHECK: {1497; CHECK-NEXT: .reg .b16 %rs<3>;1498; CHECK-NEXT: .reg .b32 %r<4>;1499; CHECK-EMPTY:1500; CHECK-NEXT: // %bb.0:1501; CHECK-NEXT: ld.param.b32 %r1, [test_fptosi_2xhalf_to_2xi32_param_0];1502; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1503; CHECK-NEXT: cvt.rzi.s32.f16 %r2, %rs2;1504; CHECK-NEXT: cvt.rzi.s32.f16 %r3, %rs1;1505; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};1506; CHECK-NEXT: ret;1507 %r = fptosi <2 x half> %a to <2 x i32>1508 ret <2 x i32> %r1509}1510 1511define <2 x i32> @test_fptoui_2xhalf_to_2xi32(<2 x half> %a) #0 {1512; CHECK-LABEL: test_fptoui_2xhalf_to_2xi32(1513; CHECK: {1514; CHECK-NEXT: .reg .b16 %rs<3>;1515; CHECK-NEXT: .reg .b32 %r<4>;1516; CHECK-EMPTY:1517; CHECK-NEXT: // %bb.0:1518; CHECK-NEXT: ld.param.b32 %r1, [test_fptoui_2xhalf_to_2xi32_param_0];1519; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1520; CHECK-NEXT: cvt.rzi.u32.f16 %r2, %rs2;1521; CHECK-NEXT: cvt.rzi.u32.f16 %r3, %rs1;1522; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};1523; CHECK-NEXT: ret;1524 %r = fptoui <2 x half> %a to <2 x i32>1525 ret <2 x i32> %r1526}1527 1528define void @test_srem_v2i32(ptr %a, ptr %b, ptr %c) {1529; CHECK-LABEL: test_srem_v2i32(1530; CHECK: {1531; CHECK-NEXT: .reg .b32 %r<7>;1532; CHECK-NEXT: .reg .b64 %rd<4>;1533; CHECK-EMPTY:1534; CHECK-NEXT: // %bb.0: // %entry1535; CHECK-NEXT: ld.param.b64 %rd3, [test_srem_v2i32_param_2];1536; CHECK-NEXT: ld.param.b64 %rd2, [test_srem_v2i32_param_1];1537; CHECK-NEXT: ld.param.b64 %rd1, [test_srem_v2i32_param_0];1538; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];1539; CHECK-NEXT: ld.v2.b32 {%r3, %r4}, [%rd2];1540; CHECK-NEXT: rem.s32 %r5, %r2, %r4;1541; CHECK-NEXT: rem.s32 %r6, %r1, %r3;1542; CHECK-NEXT: st.v2.b32 [%rd3], {%r6, %r5};1543; CHECK-NEXT: ret;1544entry:1545 %t57 = load <2 x i32>, ptr %a, align 81546 %t59 = load <2 x i32>, ptr %b, align 81547 %x = srem <2 x i32> %t57, %t591548 store <2 x i32> %x, ptr %c, align 81549 ret void1550}1551 1552define void @test_srem_v3i32(ptr %a, ptr %b, ptr %c) {1553; CHECK-NOI32X2-LABEL: test_srem_v3i32(1554; CHECK-NOI32X2: {1555; CHECK-NOI32X2-NEXT: .reg .b32 %r<10>;1556; CHECK-NOI32X2-NEXT: .reg .b64 %rd<10>;1557; CHECK-NOI32X2-EMPTY:1558; CHECK-NOI32X2-NEXT: // %bb.0: // %entry1559; CHECK-NOI32X2-NEXT: ld.param.b64 %rd3, [test_srem_v3i32_param_2];1560; CHECK-NOI32X2-NEXT: ld.param.b64 %rd2, [test_srem_v3i32_param_1];1561; CHECK-NOI32X2-NEXT: ld.param.b64 %rd1, [test_srem_v3i32_param_0];1562; CHECK-NOI32X2-NEXT: ld.b32 %r1, [%rd1+8];1563; CHECK-NOI32X2-NEXT: ld.b64 %rd4, [%rd1];1564; CHECK-NOI32X2-NEXT: { .reg .b32 tmp; mov.b64 {tmp, %r2}, %rd4; }1565; CHECK-NOI32X2-NEXT: cvt.u32.u64 %r3, %rd4;1566; CHECK-NOI32X2-NEXT: ld.b32 %r4, [%rd2+8];1567; CHECK-NOI32X2-NEXT: ld.b64 %rd5, [%rd2];1568; CHECK-NOI32X2-NEXT: { .reg .b32 tmp; mov.b64 {tmp, %r5}, %rd5; }1569; CHECK-NOI32X2-NEXT: cvt.u32.u64 %r6, %rd5;1570; CHECK-NOI32X2-NEXT: rem.s32 %r7, %r3, %r6;1571; CHECK-NOI32X2-NEXT: cvt.u64.u32 %rd6, %r7;1572; CHECK-NOI32X2-NEXT: rem.s32 %r8, %r2, %r5;1573; CHECK-NOI32X2-NEXT: cvt.u64.u32 %rd7, %r8;1574; CHECK-NOI32X2-NEXT: shl.b64 %rd8, %rd7, 32;1575; CHECK-NOI32X2-NEXT: or.b64 %rd9, %rd6, %rd8;1576; CHECK-NOI32X2-NEXT: rem.s32 %r9, %r1, %r4;1577; CHECK-NOI32X2-NEXT: st.b32 [%rd3+8], %r9;1578; CHECK-NOI32X2-NEXT: st.b64 [%rd3], %rd9;1579; CHECK-NOI32X2-NEXT: ret;1580;1581; CHECK-I32X2-LABEL: test_srem_v3i32(1582; CHECK-I32X2: {1583; CHECK-I32X2-NEXT: .reg .b32 %r<10>;1584; CHECK-I32X2-NEXT: .reg .b64 %rd<4>;1585; CHECK-I32X2-EMPTY:1586; CHECK-I32X2-NEXT: // %bb.0: // %entry1587; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [test_srem_v3i32_param_2];1588; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_srem_v3i32_param_1];1589; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_srem_v3i32_param_0];1590; CHECK-I32X2-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];1591; CHECK-I32X2-NEXT: ld.b32 %r3, [%rd1+8];1592; CHECK-I32X2-NEXT: ld.v2.b32 {%r4, %r5}, [%rd2];1593; CHECK-I32X2-NEXT: ld.b32 %r6, [%rd2+8];1594; CHECK-I32X2-NEXT: rem.s32 %r7, %r3, %r6;1595; CHECK-I32X2-NEXT: rem.s32 %r8, %r2, %r5;1596; CHECK-I32X2-NEXT: rem.s32 %r9, %r1, %r4;1597; CHECK-I32X2-NEXT: st.v2.b32 [%rd3], {%r9, %r8};1598; CHECK-I32X2-NEXT: st.b32 [%rd3+8], %r7;1599; CHECK-I32X2-NEXT: ret;1600entry:1601 %t57 = load <3 x i32>, ptr %a, align 81602 %t59 = load <3 x i32>, ptr %b, align 81603 %x = srem <3 x i32> %t57, %t591604 store <3 x i32> %x, ptr %c, align 81605 ret void1606}1607 1608define void @test_sext_v2i1_to_v2i32(ptr %a, ptr %b, ptr %c) {1609; CHECK-NOI32X2-LABEL: test_sext_v2i1_to_v2i32(1610; CHECK-NOI32X2: {1611; CHECK-NOI32X2-NEXT: .reg .pred %p<3>;1612; CHECK-NOI32X2-NEXT: .reg .b32 %r<7>;1613; CHECK-NOI32X2-NEXT: .reg .b64 %rd<4>;1614; CHECK-NOI32X2-EMPTY:1615; CHECK-NOI32X2-NEXT: // %bb.0: // %entry1616; CHECK-NOI32X2-NEXT: ld.param.b64 %rd3, [test_sext_v2i1_to_v2i32_param_2];1617; CHECK-NOI32X2-NEXT: ld.param.b64 %rd2, [test_sext_v2i1_to_v2i32_param_1];1618; CHECK-NOI32X2-NEXT: ld.param.b64 %rd1, [test_sext_v2i1_to_v2i32_param_0];1619; CHECK-NOI32X2-NEXT: ld.b32 %r1, [%rd1];1620; CHECK-NOI32X2-NEXT: ld.b32 %r2, [%rd1+4];1621; CHECK-NOI32X2-NEXT: ld.b32 %r3, [%rd2];1622; CHECK-NOI32X2-NEXT: ld.b32 %r4, [%rd2+4];1623; CHECK-NOI32X2-NEXT: setp.gt.u32 %p1, %r2, %r4;1624; CHECK-NOI32X2-NEXT: setp.gt.u32 %p2, %r1, %r3;1625; CHECK-NOI32X2-NEXT: selp.b32 %r5, -1, 0, %p2;1626; CHECK-NOI32X2-NEXT: selp.b32 %r6, -1, 0, %p1;1627; CHECK-NOI32X2-NEXT: st.b32 [%rd3+4], %r6;1628; CHECK-NOI32X2-NEXT: st.b32 [%rd3], %r5;1629; CHECK-NOI32X2-NEXT: ret;1630;1631; CHECK-I32X2-LABEL: test_sext_v2i1_to_v2i32(1632; CHECK-I32X2: {1633; CHECK-I32X2-NEXT: .reg .pred %p<3>;1634; CHECK-I32X2-NEXT: .reg .b32 %r<7>;1635; CHECK-I32X2-NEXT: .reg .b64 %rd<14>;1636; CHECK-I32X2-EMPTY:1637; CHECK-I32X2-NEXT: // %bb.0: // %entry1638; CHECK-I32X2-NEXT: ld.param.b64 %rd3, [test_sext_v2i1_to_v2i32_param_2];1639; CHECK-I32X2-NEXT: ld.param.b64 %rd2, [test_sext_v2i1_to_v2i32_param_1];1640; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_sext_v2i1_to_v2i32_param_0];1641; CHECK-I32X2-NEXT: ld.b32 %rd4, [%rd1];1642; CHECK-I32X2-NEXT: ld.b32 %rd5, [%rd1+4];1643; CHECK-I32X2-NEXT: shl.b64 %rd6, %rd5, 32;1644; CHECK-I32X2-NEXT: or.b64 %rd7, %rd6, %rd4;1645; CHECK-I32X2-NEXT: ld.b32 %rd8, [%rd2];1646; CHECK-I32X2-NEXT: ld.b32 %rd9, [%rd2+4];1647; CHECK-I32X2-NEXT: shl.b64 %rd10, %rd9, 32;1648; CHECK-I32X2-NEXT: or.b64 %rd11, %rd10, %rd8;1649; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd11;1650; CHECK-I32X2-NEXT: mov.b64 {%r3, %r4}, %rd7;1651; CHECK-I32X2-NEXT: setp.gt.u32 %p1, %r3, %r1;1652; CHECK-I32X2-NEXT: setp.gt.u32 %p2, %r4, %r2;1653; CHECK-I32X2-NEXT: selp.b32 %r5, -1, 0, %p2;1654; CHECK-I32X2-NEXT: selp.b32 %r6, -1, 0, %p1;1655; CHECK-I32X2-NEXT: mov.b64 %rd12, {%r6, %r5};1656; CHECK-I32X2-NEXT: st.b32 [%rd3], %rd12;1657; CHECK-I32X2-NEXT: shr.u64 %rd13, %rd12, 32;1658; CHECK-I32X2-NEXT: st.b32 [%rd3+4], %rd13;1659; CHECK-I32X2-NEXT: ret;1660entry:1661 %t1 = load <2 x i32>, ptr %a, align 41662 %t2 = load <2 x i32>, ptr %b, align 41663 %t5 = icmp ugt <2 x i32> %t1, %t21664 %t6 = sext <2 x i1> %t5 to <2 x i32>1665 store <2 x i32> %t6, ptr %c, align 41666 ret void1667}1668 1669define <2 x i32> @test_sext_v2i8_to_v2i32 (<2 x i8> %a) {1670; CHECK-LABEL: test_sext_v2i8_to_v2i32(1671; CHECK: {1672; CHECK-NEXT: .reg .b16 %rs<3>;1673; CHECK-NEXT: .reg .b32 %r<6>;1674; CHECK-EMPTY:1675; CHECK-NEXT: // %bb.0:1676; CHECK-NEXT: ld.param.v2.b8 {%rs1, %rs2}, [test_sext_v2i8_to_v2i32_param_0];1677; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};1678; CHECK-NEXT: cvt.u32.u16 %r2, %rs2;1679; CHECK-NEXT: cvt.s32.s8 %r3, %r2;1680; CHECK-NEXT: cvt.u32.u16 %r4, %rs1;1681; CHECK-NEXT: cvt.s32.s8 %r5, %r4;1682; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r5, %r3};1683; CHECK-NEXT: ret;1684 %r = sext <2 x i8> %a to <2 x i32>1685 ret <2 x i32> %r1686}1687 1688define <2 x i32> @test_sext_v2i16_to_v2i32 (<2 x i16> %a) {1689; CHECK-NOI32X2-LABEL: test_sext_v2i16_to_v2i32(1690; CHECK-NOI32X2: {1691; CHECK-NOI32X2-NEXT: .reg .b16 %rs<2>;1692; CHECK-NOI32X2-NEXT: .reg .b32 %r<4>;1693; CHECK-NOI32X2-EMPTY:1694; CHECK-NOI32X2-NEXT: // %bb.0:1695; CHECK-NOI32X2-NEXT: ld.param.b32 %r1, [test_sext_v2i16_to_v2i32_param_0];1696; CHECK-NOI32X2-NEXT: cvt.s32.s16 %r2, %r1;1697; CHECK-NOI32X2-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }1698; CHECK-NOI32X2-NEXT: cvt.s32.s16 %r3, %rs1;1699; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r3};1700; CHECK-NOI32X2-NEXT: ret;1701;1702; CHECK-I32X2-LABEL: test_sext_v2i16_to_v2i32(1703; CHECK-I32X2: {1704; CHECK-I32X2-NEXT: .reg .b16 %rs<2>;1705; CHECK-I32X2-NEXT: .reg .b32 %r<4>;1706; CHECK-I32X2-EMPTY:1707; CHECK-I32X2-NEXT: // %bb.0:1708; CHECK-I32X2-NEXT: ld.param.b32 %r1, [test_sext_v2i16_to_v2i32_param_0];1709; CHECK-I32X2-NEXT: cvt.s32.s16 %r2, %r1;1710; CHECK-I32X2-NEXT: mov.b32 {_, %rs1}, %r1;1711; CHECK-I32X2-NEXT: cvt.s32.s16 %r3, %rs1;1712; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r3};1713; CHECK-I32X2-NEXT: ret;1714 %r = sext <2 x i16> %a to <2 x i32>1715 ret <2 x i32> %r1716}1717 1718define <2 x float> @test_uitofp_v2i32(<2 x i32> %a) {1719; CHECK-NOI32X2-LABEL: test_uitofp_v2i32(1720; CHECK-NOI32X2: {1721; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;1722; CHECK-NOI32X2-EMPTY:1723; CHECK-NOI32X2-NEXT: // %bb.0:1724; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_v2i32_param_0];1725; CHECK-NOI32X2-NEXT: cvt.rn.f32.u32 %r3, %r2;1726; CHECK-NOI32X2-NEXT: cvt.rn.f32.u32 %r4, %r1;1727; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1728; CHECK-NOI32X2-NEXT: ret;1729;1730; CHECK-I32X2-LABEL: test_uitofp_v2i32(1731; CHECK-I32X2: {1732; CHECK-I32X2-NEXT: .reg .b32 %r<5>;1733; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1734; CHECK-I32X2-EMPTY:1735; CHECK-I32X2-NEXT: // %bb.0:1736; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_uitofp_v2i32_param_0];1737; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1738; CHECK-I32X2-NEXT: cvt.rn.f32.u32 %r3, %r2;1739; CHECK-I32X2-NEXT: cvt.rn.f32.u32 %r4, %r1;1740; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1741; CHECK-I32X2-NEXT: ret;1742 %r = uitofp <2 x i32> %a to <2 x float>1743 ret <2 x float> %r1744}1745 1746define <2 x float> @test_sitofp_v2i32(<2 x i32> %a) {1747; CHECK-NOI32X2-LABEL: test_sitofp_v2i32(1748; CHECK-NOI32X2: {1749; CHECK-NOI32X2-NEXT: .reg .b32 %r<5>;1750; CHECK-NOI32X2-EMPTY:1751; CHECK-NOI32X2-NEXT: // %bb.0:1752; CHECK-NOI32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sitofp_v2i32_param_0];1753; CHECK-NOI32X2-NEXT: cvt.rn.f32.s32 %r3, %r2;1754; CHECK-NOI32X2-NEXT: cvt.rn.f32.s32 %r4, %r1;1755; CHECK-NOI32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1756; CHECK-NOI32X2-NEXT: ret;1757;1758; CHECK-I32X2-LABEL: test_sitofp_v2i32(1759; CHECK-I32X2: {1760; CHECK-I32X2-NEXT: .reg .b32 %r<5>;1761; CHECK-I32X2-NEXT: .reg .b64 %rd<2>;1762; CHECK-I32X2-EMPTY:1763; CHECK-I32X2-NEXT: // %bb.0:1764; CHECK-I32X2-NEXT: ld.param.b64 %rd1, [test_sitofp_v2i32_param_0];1765; CHECK-I32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1766; CHECK-I32X2-NEXT: cvt.rn.f32.s32 %r3, %r2;1767; CHECK-I32X2-NEXT: cvt.rn.f32.s32 %r4, %r1;1768; CHECK-I32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1769; CHECK-I32X2-NEXT: ret;1770 %r = sitofp <2 x i32> %a to <2 x float>1771 ret <2 x float> %r1772}1773 1774attributes #0 = { nounwind }1775