1053 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; ## Support i16x2 instructions3; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx80 \4; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \5; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,I16x2 %s6; RUN: %if ptxas-sm_90 %{ \7; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 \8; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \9; RUN: | %ptxas-verify -arch=sm_90 \10; RUN: %}11; ## No support for i16x2 instructions12; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \13; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \14; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes COMMON,NO-I16x2 %s15; RUN: %if ptxas-sm_53 %{ \16; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \17; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \18; RUN: | %ptxas-verify -arch=sm_53 \19; RUN: %}20 21target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"22 23define <2 x i16> @test_ret_const() #0 {24; COMMON-LABEL: test_ret_const(25; COMMON: {26; COMMON-EMPTY:27; COMMON-EMPTY:28; COMMON-NEXT: // %bb.0:29; COMMON-NEXT: st.param.b32 [func_retval0], 131073;30; COMMON-NEXT: ret;31 ret <2 x i16> <i16 1, i16 2>32}33 34define i16 @test_extract_0(<2 x i16> %a) #0 {35; I16x2-LABEL: test_extract_0(36; I16x2: {37; I16x2-NEXT: .reg .b16 %rs<2>;38; I16x2-NEXT: .reg .b32 %r<3>;39; I16x2-EMPTY:40; I16x2-NEXT: // %bb.0:41; I16x2-NEXT: ld.param.b32 %r1, [test_extract_0_param_0];42; I16x2-NEXT: mov.b32 {%rs1, _}, %r1;43; I16x2-NEXT: cvt.u32.u16 %r2, %rs1;44; I16x2-NEXT: st.param.b32 [func_retval0], %r2;45; I16x2-NEXT: ret;46;47; NO-I16x2-LABEL: test_extract_0(48; NO-I16x2: {49; NO-I16x2-NEXT: .reg .b16 %rs<2>;50; NO-I16x2-NEXT: .reg .b32 %r<3>;51; NO-I16x2-EMPTY:52; NO-I16x2-NEXT: // %bb.0:53; NO-I16x2-NEXT: ld.param.b32 %r1, [test_extract_0_param_0];54; NO-I16x2-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r1; }55; NO-I16x2-NEXT: cvt.u32.u16 %r2, %rs1;56; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r2;57; NO-I16x2-NEXT: ret;58 %e = extractelement <2 x i16> %a, i32 059 ret i16 %e60}61 62define i16 @test_extract_1(<2 x i16> %a) #0 {63; I16x2-LABEL: test_extract_1(64; I16x2: {65; I16x2-NEXT: .reg .b16 %rs<2>;66; I16x2-NEXT: .reg .b32 %r<3>;67; I16x2-EMPTY:68; I16x2-NEXT: // %bb.0:69; I16x2-NEXT: ld.param.b32 %r1, [test_extract_1_param_0];70; I16x2-NEXT: mov.b32 {_, %rs1}, %r1;71; I16x2-NEXT: cvt.u32.u16 %r2, %rs1;72; I16x2-NEXT: st.param.b32 [func_retval0], %r2;73; I16x2-NEXT: ret;74;75; NO-I16x2-LABEL: test_extract_1(76; NO-I16x2: {77; NO-I16x2-NEXT: .reg .b16 %rs<2>;78; NO-I16x2-NEXT: .reg .b32 %r<3>;79; NO-I16x2-EMPTY:80; NO-I16x2-NEXT: // %bb.0:81; NO-I16x2-NEXT: ld.param.b32 %r1, [test_extract_1_param_0];82; NO-I16x2-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }83; NO-I16x2-NEXT: cvt.u32.u16 %r2, %rs1;84; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r2;85; NO-I16x2-NEXT: ret;86 %e = extractelement <2 x i16> %a, i32 187 ret i16 %e88}89 90define i16 @test_extract_i(<2 x i16> %a, i64 %idx) #0 {91; COMMON-LABEL: test_extract_i(92; COMMON: {93; COMMON-NEXT: .reg .pred %p<2>;94; COMMON-NEXT: .reg .b16 %rs<4>;95; COMMON-NEXT: .reg .b32 %r<3>;96; COMMON-NEXT: .reg .b64 %rd<2>;97; COMMON-EMPTY:98; COMMON-NEXT: // %bb.0:99; COMMON-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];100; COMMON-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];101; COMMON-NEXT: setp.eq.b64 %p1, %rd1, 0;102; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;103; COMMON-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;104; COMMON-NEXT: cvt.u32.u16 %r2, %rs3;105; COMMON-NEXT: st.param.b32 [func_retval0], %r2;106; COMMON-NEXT: ret;107 %e = extractelement <2 x i16> %a, i64 %idx108 ret i16 %e109}110 111define <2 x i16> @test_add(<2 x i16> %a, <2 x i16> %b) #0 {112; I16x2-LABEL: test_add(113; I16x2: {114; I16x2-NEXT: .reg .b32 %r<4>;115; I16x2-EMPTY:116; I16x2-NEXT: // %bb.0:117; I16x2-NEXT: ld.param.b32 %r2, [test_add_param_1];118; I16x2-NEXT: ld.param.b32 %r1, [test_add_param_0];119; I16x2-NEXT: add.s16x2 %r3, %r1, %r2;120; I16x2-NEXT: st.param.b32 [func_retval0], %r3;121; I16x2-NEXT: ret;122;123; NO-I16x2-LABEL: test_add(124; NO-I16x2: {125; NO-I16x2-NEXT: .reg .b16 %rs<7>;126; NO-I16x2-NEXT: .reg .b32 %r<3>;127; NO-I16x2-EMPTY:128; NO-I16x2-NEXT: // %bb.0:129; NO-I16x2-NEXT: ld.param.b32 %r2, [test_add_param_1];130; NO-I16x2-NEXT: ld.param.b32 %r1, [test_add_param_0];131; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;132; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;133; NO-I16x2-NEXT: add.s16 %rs5, %rs4, %rs2;134; NO-I16x2-NEXT: add.s16 %rs6, %rs3, %rs1;135; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};136; NO-I16x2-NEXT: ret;137 %r = add <2 x i16> %a, %b138 ret <2 x i16> %r139}140 141; Check that we can lower add with immediate arguments.142define <2 x i16> @test_add_imm_0(<2 x i16> %a) #0 {143; I16x2-LABEL: test_add_imm_0(144; I16x2: {145; I16x2-NEXT: .reg .b32 %r<4>;146; I16x2-EMPTY:147; I16x2-NEXT: // %bb.0:148; I16x2-NEXT: ld.param.b32 %r1, [test_add_imm_0_param_0];149; I16x2-NEXT: mov.b32 %r2, 131073;150; I16x2-NEXT: add.s16x2 %r3, %r1, %r2;151; I16x2-NEXT: st.param.b32 [func_retval0], %r3;152; I16x2-NEXT: ret;153;154; NO-I16x2-LABEL: test_add_imm_0(155; NO-I16x2: {156; NO-I16x2-NEXT: .reg .b16 %rs<5>;157; NO-I16x2-NEXT: .reg .b32 %r<2>;158; NO-I16x2-EMPTY:159; NO-I16x2-NEXT: // %bb.0:160; NO-I16x2-NEXT: ld.param.b32 %r1, [test_add_imm_0_param_0];161; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r1;162; NO-I16x2-NEXT: add.s16 %rs3, %rs2, 2;163; NO-I16x2-NEXT: add.s16 %rs4, %rs1, 1;164; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};165; NO-I16x2-NEXT: ret;166 %r = add <2 x i16> <i16 1, i16 2>, %a167 ret <2 x i16> %r168}169 170define <2 x i16> @test_add_imm_1(<2 x i16> %a) #0 {171; I16x2-LABEL: test_add_imm_1(172; I16x2: {173; I16x2-NEXT: .reg .b32 %r<4>;174; I16x2-EMPTY:175; I16x2-NEXT: // %bb.0:176; I16x2-NEXT: ld.param.b32 %r1, [test_add_imm_1_param_0];177; I16x2-NEXT: mov.b32 %r2, 131073;178; I16x2-NEXT: add.s16x2 %r3, %r1, %r2;179; I16x2-NEXT: st.param.b32 [func_retval0], %r3;180; I16x2-NEXT: ret;181;182; NO-I16x2-LABEL: test_add_imm_1(183; NO-I16x2: {184; NO-I16x2-NEXT: .reg .b16 %rs<5>;185; NO-I16x2-NEXT: .reg .b32 %r<2>;186; NO-I16x2-EMPTY:187; NO-I16x2-NEXT: // %bb.0:188; NO-I16x2-NEXT: ld.param.b32 %r1, [test_add_imm_1_param_0];189; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r1;190; NO-I16x2-NEXT: add.s16 %rs3, %rs2, 2;191; NO-I16x2-NEXT: add.s16 %rs4, %rs1, 1;192; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};193; NO-I16x2-NEXT: ret;194 %r = add <2 x i16> %a, <i16 1, i16 2>195 ret <2 x i16> %r196}197 198define <2 x i16> @test_sub(<2 x i16> %a, <2 x i16> %b) #0 {199; COMMON-LABEL: test_sub(200; COMMON: {201; COMMON-NEXT: .reg .b16 %rs<7>;202; COMMON-NEXT: .reg .b32 %r<3>;203; COMMON-EMPTY:204; COMMON-NEXT: // %bb.0:205; COMMON-NEXT: ld.param.b32 %r2, [test_sub_param_1];206; COMMON-NEXT: ld.param.b32 %r1, [test_sub_param_0];207; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;208; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;209; COMMON-NEXT: sub.s16 %rs5, %rs4, %rs2;210; COMMON-NEXT: sub.s16 %rs6, %rs3, %rs1;211; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};212; COMMON-NEXT: ret;213 %r = sub <2 x i16> %a, %b214 ret <2 x i16> %r215}216 217define <2 x i16> @test_smax(<2 x i16> %a, <2 x i16> %b) #0 {218; I16x2-LABEL: test_smax(219; I16x2: {220; I16x2-NEXT: .reg .b32 %r<4>;221; I16x2-EMPTY:222; I16x2-NEXT: // %bb.0:223; I16x2-NEXT: ld.param.b32 %r2, [test_smax_param_1];224; I16x2-NEXT: ld.param.b32 %r1, [test_smax_param_0];225; I16x2-NEXT: max.s16x2 %r3, %r1, %r2;226; I16x2-NEXT: st.param.b32 [func_retval0], %r3;227; I16x2-NEXT: ret;228;229; NO-I16x2-LABEL: test_smax(230; NO-I16x2: {231; NO-I16x2-NEXT: .reg .b16 %rs<7>;232; NO-I16x2-NEXT: .reg .b32 %r<3>;233; NO-I16x2-EMPTY:234; NO-I16x2-NEXT: // %bb.0:235; NO-I16x2-NEXT: ld.param.b32 %r2, [test_smax_param_1];236; NO-I16x2-NEXT: ld.param.b32 %r1, [test_smax_param_0];237; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;238; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;239; NO-I16x2-NEXT: max.s16 %rs5, %rs4, %rs2;240; NO-I16x2-NEXT: max.s16 %rs6, %rs3, %rs1;241; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};242; NO-I16x2-NEXT: ret;243 %cmp = icmp sgt <2 x i16> %a, %b244 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b245 ret <2 x i16> %r246}247 248define <2 x i16> @test_umax(<2 x i16> %a, <2 x i16> %b) #0 {249; I16x2-LABEL: test_umax(250; I16x2: {251; I16x2-NEXT: .reg .b32 %r<4>;252; I16x2-EMPTY:253; I16x2-NEXT: // %bb.0:254; I16x2-NEXT: ld.param.b32 %r2, [test_umax_param_1];255; I16x2-NEXT: ld.param.b32 %r1, [test_umax_param_0];256; I16x2-NEXT: max.u16x2 %r3, %r1, %r2;257; I16x2-NEXT: st.param.b32 [func_retval0], %r3;258; I16x2-NEXT: ret;259;260; NO-I16x2-LABEL: test_umax(261; NO-I16x2: {262; NO-I16x2-NEXT: .reg .b16 %rs<7>;263; NO-I16x2-NEXT: .reg .b32 %r<3>;264; NO-I16x2-EMPTY:265; NO-I16x2-NEXT: // %bb.0:266; NO-I16x2-NEXT: ld.param.b32 %r2, [test_umax_param_1];267; NO-I16x2-NEXT: ld.param.b32 %r1, [test_umax_param_0];268; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;269; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;270; NO-I16x2-NEXT: max.u16 %rs5, %rs4, %rs2;271; NO-I16x2-NEXT: max.u16 %rs6, %rs3, %rs1;272; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};273; NO-I16x2-NEXT: ret;274 %cmp = icmp ugt <2 x i16> %a, %b275 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b276 ret <2 x i16> %r277}278 279define <2 x i16> @test_smin(<2 x i16> %a, <2 x i16> %b) #0 {280; I16x2-LABEL: test_smin(281; I16x2: {282; I16x2-NEXT: .reg .b32 %r<4>;283; I16x2-EMPTY:284; I16x2-NEXT: // %bb.0:285; I16x2-NEXT: ld.param.b32 %r2, [test_smin_param_1];286; I16x2-NEXT: ld.param.b32 %r1, [test_smin_param_0];287; I16x2-NEXT: min.s16x2 %r3, %r1, %r2;288; I16x2-NEXT: st.param.b32 [func_retval0], %r3;289; I16x2-NEXT: ret;290;291; NO-I16x2-LABEL: test_smin(292; NO-I16x2: {293; NO-I16x2-NEXT: .reg .b16 %rs<7>;294; NO-I16x2-NEXT: .reg .b32 %r<3>;295; NO-I16x2-EMPTY:296; NO-I16x2-NEXT: // %bb.0:297; NO-I16x2-NEXT: ld.param.b32 %r2, [test_smin_param_1];298; NO-I16x2-NEXT: ld.param.b32 %r1, [test_smin_param_0];299; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;300; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;301; NO-I16x2-NEXT: min.s16 %rs5, %rs4, %rs2;302; NO-I16x2-NEXT: min.s16 %rs6, %rs3, %rs1;303; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};304; NO-I16x2-NEXT: ret;305 %cmp = icmp sle <2 x i16> %a, %b306 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b307 ret <2 x i16> %r308}309 310define <2 x i16> @test_umin(<2 x i16> %a, <2 x i16> %b) #0 {311; I16x2-LABEL: test_umin(312; I16x2: {313; I16x2-NEXT: .reg .b32 %r<4>;314; I16x2-EMPTY:315; I16x2-NEXT: // %bb.0:316; I16x2-NEXT: ld.param.b32 %r2, [test_umin_param_1];317; I16x2-NEXT: ld.param.b32 %r1, [test_umin_param_0];318; I16x2-NEXT: min.u16x2 %r3, %r1, %r2;319; I16x2-NEXT: st.param.b32 [func_retval0], %r3;320; I16x2-NEXT: ret;321;322; NO-I16x2-LABEL: test_umin(323; NO-I16x2: {324; NO-I16x2-NEXT: .reg .b16 %rs<7>;325; NO-I16x2-NEXT: .reg .b32 %r<3>;326; NO-I16x2-EMPTY:327; NO-I16x2-NEXT: // %bb.0:328; NO-I16x2-NEXT: ld.param.b32 %r2, [test_umin_param_1];329; NO-I16x2-NEXT: ld.param.b32 %r1, [test_umin_param_0];330; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;331; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;332; NO-I16x2-NEXT: min.u16 %rs5, %rs4, %rs2;333; NO-I16x2-NEXT: min.u16 %rs6, %rs3, %rs1;334; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};335; NO-I16x2-NEXT: ret;336 %cmp = icmp ule <2 x i16> %a, %b337 %r = select <2 x i1> %cmp, <2 x i16> %a, <2 x i16> %b338 ret <2 x i16> %r339}340 341define <2 x i16> @test_mul(<2 x i16> %a, <2 x i16> %b) #0 {342; COMMON-LABEL: test_mul(343; COMMON: {344; COMMON-NEXT: .reg .b16 %rs<7>;345; COMMON-NEXT: .reg .b32 %r<3>;346; COMMON-EMPTY:347; COMMON-NEXT: // %bb.0:348; COMMON-NEXT: ld.param.b32 %r2, [test_mul_param_1];349; COMMON-NEXT: ld.param.b32 %r1, [test_mul_param_0];350; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;351; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;352; COMMON-NEXT: mul.lo.s16 %rs5, %rs4, %rs2;353; COMMON-NEXT: mul.lo.s16 %rs6, %rs3, %rs1;354; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};355; COMMON-NEXT: ret;356 %r = mul <2 x i16> %a, %b357 ret <2 x i16> %r358}359 360;; Logical ops are available on all GPUs as regular 32-bit logical ops361define <2 x i16> @test_or(<2 x i16> %a, <2 x i16> %b) #0 {362; COMMON-LABEL: test_or(363; COMMON: {364; COMMON-NEXT: .reg .b32 %r<4>;365; COMMON-EMPTY:366; COMMON-NEXT: // %bb.0:367; COMMON-NEXT: ld.param.b32 %r2, [test_or_param_1];368; COMMON-NEXT: ld.param.b32 %r1, [test_or_param_0];369; COMMON-NEXT: or.b32 %r3, %r1, %r2;370; COMMON-NEXT: st.param.b32 [func_retval0], %r3;371; COMMON-NEXT: ret;372 %r = or <2 x i16> %a, %b373 ret <2 x i16> %r374}375 376; Ops that operate on computed arguments go though a different lowering path.377; compared to the ones that operate on loaded data. So we test them separately.378define <2 x i16> @test_or_computed(i16 %a) {379; COMMON-LABEL: test_or_computed(380; COMMON: {381; COMMON-NEXT: .reg .b16 %rs<4>;382; COMMON-NEXT: .reg .b32 %r<4>;383; COMMON-EMPTY:384; COMMON-NEXT: // %bb.0:385; COMMON-NEXT: ld.param.b16 %rs1, [test_or_computed_param_0];386; COMMON-NEXT: mov.b16 %rs2, 0;387; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};388; COMMON-NEXT: mov.b16 %rs3, 5;389; COMMON-NEXT: mov.b32 %r2, {%rs1, %rs3};390; COMMON-NEXT: or.b32 %r3, %r2, %r1;391; COMMON-NEXT: st.param.b32 [func_retval0], %r3;392; COMMON-NEXT: ret;393 %ins.0 = insertelement <2 x i16> zeroinitializer, i16 %a, i32 0394 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1395 %r = or <2 x i16> %ins.1, %ins.0396 ret <2 x i16> %r397}398 399; Check that we can lower or with immediate arguments.400define <2 x i16> @test_or_imm_0(<2 x i16> %a) #0 {401; COMMON-LABEL: test_or_imm_0(402; COMMON: {403; COMMON-NEXT: .reg .b32 %r<3>;404; COMMON-EMPTY:405; COMMON-NEXT: // %bb.0:406; COMMON-NEXT: ld.param.b32 %r1, [test_or_imm_0_param_0];407; COMMON-NEXT: or.b32 %r2, %r1, 131073;408; COMMON-NEXT: st.param.b32 [func_retval0], %r2;409; COMMON-NEXT: ret;410 %r = or <2 x i16> <i16 1, i16 2>, %a411 ret <2 x i16> %r412}413 414define <2 x i16> @test_or_imm_1(<2 x i16> %a) #0 {415; COMMON-LABEL: test_or_imm_1(416; COMMON: {417; COMMON-NEXT: .reg .b32 %r<3>;418; COMMON-EMPTY:419; COMMON-NEXT: // %bb.0:420; COMMON-NEXT: ld.param.b32 %r1, [test_or_imm_1_param_0];421; COMMON-NEXT: or.b32 %r2, %r1, 131073;422; COMMON-NEXT: st.param.b32 [func_retval0], %r2;423; COMMON-NEXT: ret;424 %r = or <2 x i16> %a, <i16 1, i16 2>425 ret <2 x i16> %r426}427 428define <2 x i16> @test_xor(<2 x i16> %a, <2 x i16> %b) #0 {429; COMMON-LABEL: test_xor(430; COMMON: {431; COMMON-NEXT: .reg .b32 %r<4>;432; COMMON-EMPTY:433; COMMON-NEXT: // %bb.0:434; COMMON-NEXT: ld.param.b32 %r2, [test_xor_param_1];435; COMMON-NEXT: ld.param.b32 %r1, [test_xor_param_0];436; COMMON-NEXT: xor.b32 %r3, %r1, %r2;437; COMMON-NEXT: st.param.b32 [func_retval0], %r3;438; COMMON-NEXT: ret;439 %r = xor <2 x i16> %a, %b440 ret <2 x i16> %r441}442 443define <2 x i16> @test_xor_computed(i16 %a) {444; COMMON-LABEL: test_xor_computed(445; COMMON: {446; COMMON-NEXT: .reg .b16 %rs<4>;447; COMMON-NEXT: .reg .b32 %r<4>;448; COMMON-EMPTY:449; COMMON-NEXT: // %bb.0:450; COMMON-NEXT: ld.param.b16 %rs1, [test_xor_computed_param_0];451; COMMON-NEXT: mov.b16 %rs2, 0;452; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};453; COMMON-NEXT: mov.b16 %rs3, 5;454; COMMON-NEXT: mov.b32 %r2, {%rs1, %rs3};455; COMMON-NEXT: xor.b32 %r3, %r2, %r1;456; COMMON-NEXT: st.param.b32 [func_retval0], %r3;457; COMMON-NEXT: ret;458 %ins.0 = insertelement <2 x i16> zeroinitializer, i16 %a, i32 0459 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1460 %r = xor <2 x i16> %ins.1, %ins.0461 ret <2 x i16> %r462}463 464; Check that we can lower xor with immediate arguments.465define <2 x i16> @test_xor_imm_0(<2 x i16> %a) #0 {466; COMMON-LABEL: test_xor_imm_0(467; COMMON: {468; COMMON-NEXT: .reg .b32 %r<3>;469; COMMON-EMPTY:470; COMMON-NEXT: // %bb.0:471; COMMON-NEXT: ld.param.b32 %r1, [test_xor_imm_0_param_0];472; COMMON-NEXT: xor.b32 %r2, %r1, 131073;473; COMMON-NEXT: st.param.b32 [func_retval0], %r2;474; COMMON-NEXT: ret;475 %r = xor <2 x i16> <i16 1, i16 2>, %a476 ret <2 x i16> %r477}478 479define <2 x i16> @test_xor_imm_1(<2 x i16> %a) #0 {480; COMMON-LABEL: test_xor_imm_1(481; COMMON: {482; COMMON-NEXT: .reg .b32 %r<3>;483; COMMON-EMPTY:484; COMMON-NEXT: // %bb.0:485; COMMON-NEXT: ld.param.b32 %r1, [test_xor_imm_1_param_0];486; COMMON-NEXT: xor.b32 %r2, %r1, 131073;487; COMMON-NEXT: st.param.b32 [func_retval0], %r2;488; COMMON-NEXT: ret;489 %r = xor <2 x i16> %a, <i16 1, i16 2>490 ret <2 x i16> %r491}492 493define <2 x i16> @test_and(<2 x i16> %a, <2 x i16> %b) #0 {494; COMMON-LABEL: test_and(495; COMMON: {496; COMMON-NEXT: .reg .b32 %r<4>;497; COMMON-EMPTY:498; COMMON-NEXT: // %bb.0:499; COMMON-NEXT: ld.param.b32 %r2, [test_and_param_1];500; COMMON-NEXT: ld.param.b32 %r1, [test_and_param_0];501; COMMON-NEXT: and.b32 %r3, %r1, %r2;502; COMMON-NEXT: st.param.b32 [func_retval0], %r3;503; COMMON-NEXT: ret;504 %r = and <2 x i16> %a, %b505 ret <2 x i16> %r506}507 508; Ops that operate on computed arguments go though a different lowering path.509; compared to the ones that operate on loaded data. So we test them separately.510define <2 x i16> @test_and_computed(i16 %a) {511; COMMON-LABEL: test_and_computed(512; COMMON: {513; COMMON-NEXT: .reg .b16 %rs<4>;514; COMMON-NEXT: .reg .b32 %r<4>;515; COMMON-EMPTY:516; COMMON-NEXT: // %bb.0:517; COMMON-NEXT: ld.param.b16 %rs1, [test_and_computed_param_0];518; COMMON-NEXT: mov.b16 %rs2, 0;519; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};520; COMMON-NEXT: mov.b16 %rs3, 5;521; COMMON-NEXT: mov.b32 %r2, {%rs1, %rs3};522; COMMON-NEXT: and.b32 %r3, %r2, %r1;523; COMMON-NEXT: st.param.b32 [func_retval0], %r3;524; COMMON-NEXT: ret;525 %ins.0 = insertelement <2 x i16> zeroinitializer, i16 %a, i32 0526 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1527 %r = and <2 x i16> %ins.1, %ins.0528 ret <2 x i16> %r529}530 531; Check that we can lower and with immediate arguments.532define <2 x i16> @test_and_imm_0(<2 x i16> %a) #0 {533; COMMON-LABEL: test_and_imm_0(534; COMMON: {535; COMMON-NEXT: .reg .b32 %r<3>;536; COMMON-EMPTY:537; COMMON-NEXT: // %bb.0:538; COMMON-NEXT: ld.param.b32 %r1, [test_and_imm_0_param_0];539; COMMON-NEXT: and.b32 %r2, %r1, 131073;540; COMMON-NEXT: st.param.b32 [func_retval0], %r2;541; COMMON-NEXT: ret;542 %r = and <2 x i16> <i16 1, i16 2>, %a543 ret <2 x i16> %r544}545 546define <2 x i16> @test_and_imm_1(<2 x i16> %a) #0 {547; COMMON-LABEL: test_and_imm_1(548; COMMON: {549; COMMON-NEXT: .reg .b32 %r<3>;550; COMMON-EMPTY:551; COMMON-NEXT: // %bb.0:552; COMMON-NEXT: ld.param.b32 %r1, [test_and_imm_1_param_0];553; COMMON-NEXT: and.b32 %r2, %r1, 131073;554; COMMON-NEXT: st.param.b32 [func_retval0], %r2;555; COMMON-NEXT: ret;556 %r = and <2 x i16> %a, <i16 1, i16 2>557 ret <2 x i16> %r558}559 560define void @test_ldst_v2i16(ptr %a, ptr %b) {561; COMMON-LABEL: test_ldst_v2i16(562; COMMON: {563; COMMON-NEXT: .reg .b32 %r<2>;564; COMMON-NEXT: .reg .b64 %rd<3>;565; COMMON-EMPTY:566; COMMON-NEXT: // %bb.0:567; COMMON-NEXT: ld.param.b64 %rd2, [test_ldst_v2i16_param_1];568; COMMON-NEXT: ld.param.b64 %rd1, [test_ldst_v2i16_param_0];569; COMMON-NEXT: ld.b32 %r1, [%rd1];570; COMMON-NEXT: st.b32 [%rd2], %r1;571; COMMON-NEXT: ret;572 %t1 = load <2 x i16>, ptr %a573 store <2 x i16> %t1, ptr %b, align 16574 ret void575}576 577; -- v3 is inconvenient to capture as it's lowered as ld.b64 + fair578; number of bitshifting instructions that may change at llvm's whim.579; So we only verify that we only issue correct number of writes using580; correct offset, but not the values we write.581define void @test_ldst_v3i16(ptr %a, ptr %b) {582; COMMON-LABEL: test_ldst_v3i16(583; COMMON: {584; COMMON-NEXT: .reg .b64 %rd<5>;585; COMMON-EMPTY:586; COMMON-NEXT: // %bb.0:587; COMMON-NEXT: ld.param.b64 %rd2, [test_ldst_v3i16_param_1];588; COMMON-NEXT: ld.param.b64 %rd1, [test_ldst_v3i16_param_0];589; COMMON-NEXT: ld.b64 %rd3, [%rd1];590; COMMON-NEXT: shr.u64 %rd4, %rd3, 32;591; COMMON-NEXT: st.b32 [%rd2], %rd3;592; COMMON-NEXT: st.b16 [%rd2+4], %rd4;593; COMMON-NEXT: ret;594 %t1 = load <3 x i16>, ptr %a595 store <3 x i16> %t1, ptr %b, align 16596 ret void597}598 599define void @test_ldst_v4i16(ptr %a, ptr %b) {600; COMMON-LABEL: test_ldst_v4i16(601; COMMON: {602; COMMON-NEXT: .reg .b32 %r<3>;603; COMMON-NEXT: .reg .b64 %rd<3>;604; COMMON-EMPTY:605; COMMON-NEXT: // %bb.0:606; COMMON-NEXT: ld.param.b64 %rd2, [test_ldst_v4i16_param_1];607; COMMON-NEXT: ld.param.b64 %rd1, [test_ldst_v4i16_param_0];608; COMMON-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];609; COMMON-NEXT: st.v2.b32 [%rd2], {%r1, %r2};610; COMMON-NEXT: ret;611 %t1 = load <4 x i16>, ptr %a612 store <4 x i16> %t1, ptr %b, align 16613 ret void614}615 616define void @test_ldst_v8i16(ptr %a, ptr %b) {617; COMMON-LABEL: test_ldst_v8i16(618; COMMON: {619; COMMON-NEXT: .reg .b32 %r<5>;620; COMMON-NEXT: .reg .b64 %rd<3>;621; COMMON-EMPTY:622; COMMON-NEXT: // %bb.0:623; COMMON-NEXT: ld.param.b64 %rd2, [test_ldst_v8i16_param_1];624; COMMON-NEXT: ld.param.b64 %rd1, [test_ldst_v8i16_param_0];625; COMMON-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];626; COMMON-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};627; COMMON-NEXT: ret;628 %t1 = load <8 x i16>, ptr %a629 store <8 x i16> %t1, ptr %b, align 16630 ret void631}632 633declare <2 x i16> @test_callee(<2 x i16> %a, <2 x i16> %b) #0634 635define <2 x i16> @test_call(<2 x i16> %a, <2 x i16> %b) #0 {636; COMMON-LABEL: test_call(637; COMMON: {638; COMMON-NEXT: .reg .b32 %r<4>;639; COMMON-EMPTY:640; COMMON-NEXT: // %bb.0:641; COMMON-NEXT: ld.param.b32 %r2, [test_call_param_1];642; COMMON-NEXT: ld.param.b32 %r1, [test_call_param_0];643; COMMON-NEXT: { // callseq 0, 0644; COMMON-NEXT: .param .align 4 .b8 param0[4];645; COMMON-NEXT: .param .align 4 .b8 param1[4];646; COMMON-NEXT: .param .align 4 .b8 retval0[4];647; COMMON-NEXT: st.param.b32 [param1], %r2;648; COMMON-NEXT: st.param.b32 [param0], %r1;649; COMMON-NEXT: call.uni (retval0), test_callee, (param0, param1);650; COMMON-NEXT: ld.param.b32 %r3, [retval0];651; COMMON-NEXT: } // callseq 0652; COMMON-NEXT: st.param.b32 [func_retval0], %r3;653; COMMON-NEXT: ret;654 %r = call <2 x i16> @test_callee(<2 x i16> %a, <2 x i16> %b)655 ret <2 x i16> %r656}657 658define <2 x i16> @test_call_flipped(<2 x i16> %a, <2 x i16> %b) #0 {659; COMMON-LABEL: test_call_flipped(660; COMMON: {661; COMMON-NEXT: .reg .b32 %r<4>;662; COMMON-EMPTY:663; COMMON-NEXT: // %bb.0:664; COMMON-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];665; COMMON-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];666; COMMON-NEXT: { // callseq 1, 0667; COMMON-NEXT: .param .align 4 .b8 param0[4];668; COMMON-NEXT: .param .align 4 .b8 param1[4];669; COMMON-NEXT: .param .align 4 .b8 retval0[4];670; COMMON-NEXT: st.param.b32 [param1], %r1;671; COMMON-NEXT: st.param.b32 [param0], %r2;672; COMMON-NEXT: call.uni (retval0), test_callee, (param0, param1);673; COMMON-NEXT: ld.param.b32 %r3, [retval0];674; COMMON-NEXT: } // callseq 1675; COMMON-NEXT: st.param.b32 [func_retval0], %r3;676; COMMON-NEXT: ret;677 %r = call <2 x i16> @test_callee(<2 x i16> %b, <2 x i16> %a)678 ret <2 x i16> %r679}680 681define <2 x i16> @test_tailcall_flipped(<2 x i16> %a, <2 x i16> %b) #0 {682; COMMON-LABEL: test_tailcall_flipped(683; COMMON: {684; COMMON-NEXT: .reg .b32 %r<4>;685; COMMON-EMPTY:686; COMMON-NEXT: // %bb.0:687; COMMON-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];688; COMMON-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];689; COMMON-NEXT: { // callseq 2, 0690; COMMON-NEXT: .param .align 4 .b8 param0[4];691; COMMON-NEXT: .param .align 4 .b8 param1[4];692; COMMON-NEXT: .param .align 4 .b8 retval0[4];693; COMMON-NEXT: st.param.b32 [param1], %r1;694; COMMON-NEXT: st.param.b32 [param0], %r2;695; COMMON-NEXT: call.uni (retval0), test_callee, (param0, param1);696; COMMON-NEXT: ld.param.b32 %r3, [retval0];697; COMMON-NEXT: } // callseq 2698; COMMON-NEXT: st.param.b32 [func_retval0], %r3;699; COMMON-NEXT: ret;700 %r = tail call <2 x i16> @test_callee(<2 x i16> %b, <2 x i16> %a)701 ret <2 x i16> %r702}703 704define <2 x i16> @test_select(<2 x i16> %a, <2 x i16> %b, i1 zeroext %c) #0 {705; COMMON-LABEL: test_select(706; COMMON: {707; COMMON-NEXT: .reg .pred %p<2>;708; COMMON-NEXT: .reg .b16 %rs<3>;709; COMMON-NEXT: .reg .b32 %r<4>;710; COMMON-EMPTY:711; COMMON-NEXT: // %bb.0:712; COMMON-NEXT: ld.param.b8 %rs1, [test_select_param_2];713; COMMON-NEXT: and.b16 %rs2, %rs1, 1;714; COMMON-NEXT: setp.ne.b16 %p1, %rs2, 0;715; COMMON-NEXT: ld.param.b32 %r2, [test_select_param_1];716; COMMON-NEXT: ld.param.b32 %r1, [test_select_param_0];717; COMMON-NEXT: selp.b32 %r3, %r1, %r2, %p1;718; COMMON-NEXT: st.param.b32 [func_retval0], %r3;719; COMMON-NEXT: ret;720 %r = select i1 %c, <2 x i16> %a, <2 x i16> %b721 ret <2 x i16> %r722}723 724define <2 x i16> @test_select_cc(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> %d) #0 {725; COMMON-LABEL: test_select_cc(726; COMMON: {727; COMMON-NEXT: .reg .pred %p<3>;728; COMMON-NEXT: .reg .b16 %rs<11>;729; COMMON-NEXT: .reg .b32 %r<5>;730; COMMON-EMPTY:731; COMMON-NEXT: // %bb.0:732; COMMON-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];733; COMMON-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];734; COMMON-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];735; COMMON-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];736; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r4;737; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r3;738; COMMON-NEXT: setp.ne.b16 %p1, %rs3, %rs1;739; COMMON-NEXT: setp.ne.b16 %p2, %rs4, %rs2;740; COMMON-NEXT: mov.b32 {%rs5, %rs6}, %r2;741; COMMON-NEXT: mov.b32 {%rs7, %rs8}, %r1;742; COMMON-NEXT: selp.b16 %rs9, %rs8, %rs6, %p2;743; COMMON-NEXT: selp.b16 %rs10, %rs7, %rs5, %p1;744; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};745; COMMON-NEXT: ret;746 %cc = icmp ne <2 x i16> %c, %d747 %r = select <2 x i1> %cc, <2 x i16> %a, <2 x i16> %b748 ret <2 x i16> %r749}750 751define <2 x i32> @test_select_cc_i32_i16(<2 x i32> %a, <2 x i32> %b,752; COMMON-LABEL: test_select_cc_i32_i16(753; COMMON: {754; COMMON-NEXT: .reg .pred %p<3>;755; COMMON-NEXT: .reg .b16 %rs<5>;756; COMMON-NEXT: .reg .b32 %r<9>;757; COMMON-EMPTY:758; COMMON-NEXT: // %bb.0:759; COMMON-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_i32_i16_param_1];760; COMMON-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_i32_i16_param_0];761; COMMON-NEXT: ld.param.b32 %r6, [test_select_cc_i32_i16_param_3];762; COMMON-NEXT: ld.param.b32 %r5, [test_select_cc_i32_i16_param_2];763; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r6;764; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r5;765; COMMON-NEXT: setp.ne.b16 %p1, %rs3, %rs1;766; COMMON-NEXT: setp.ne.b16 %p2, %rs4, %rs2;767; COMMON-NEXT: selp.b32 %r7, %r2, %r4, %p2;768; COMMON-NEXT: selp.b32 %r8, %r1, %r3, %p1;769; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};770; COMMON-NEXT: ret;771 <2 x i16> %c, <2 x i16> %d) #0 {772 %cc = icmp ne <2 x i16> %c, %d773 %r = select <2 x i1> %cc, <2 x i32> %a, <2 x i32> %b774 ret <2 x i32> %r775}776 777define <2 x i16> @test_select_cc_i16_i32(<2 x i16> %a, <2 x i16> %b,778; COMMON-LABEL: test_select_cc_i16_i32(779; COMMON: {780; COMMON-NEXT: .reg .pred %p<3>;781; COMMON-NEXT: .reg .b16 %rs<7>;782; COMMON-NEXT: .reg .b32 %r<7>;783; COMMON-EMPTY:784; COMMON-NEXT: // %bb.0:785; COMMON-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_i16_i32_param_3];786; COMMON-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_i16_i32_param_2];787; COMMON-NEXT: ld.param.b32 %r2, [test_select_cc_i16_i32_param_1];788; COMMON-NEXT: ld.param.b32 %r1, [test_select_cc_i16_i32_param_0];789; COMMON-NEXT: setp.ne.b32 %p1, %r3, %r5;790; COMMON-NEXT: setp.ne.b32 %p2, %r4, %r6;791; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;792; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;793; COMMON-NEXT: selp.b16 %rs5, %rs4, %rs2, %p2;794; COMMON-NEXT: selp.b16 %rs6, %rs3, %rs1, %p1;795; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};796; COMMON-NEXT: ret;797 <2 x i32> %c, <2 x i32> %d) #0 {798 %cc = icmp ne <2 x i32> %c, %d799 %r = select <2 x i1> %cc, <2 x i16> %a, <2 x i16> %b800 ret <2 x i16> %r801}802 803 804define <2 x i16> @test_trunc_2xi32(<2 x i32> %a) #0 {805; COMMON-LABEL: test_trunc_2xi32(806; COMMON: {807; COMMON-NEXT: .reg .b32 %r<4>;808; COMMON-EMPTY:809; COMMON-NEXT: // %bb.0:810; COMMON-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_2xi32_param_0];811; COMMON-NEXT: prmt.b32 %r3, %r1, %r2, 0x5410U;812; COMMON-NEXT: st.param.b32 [func_retval0], %r3;813; COMMON-NEXT: ret;814 %r = trunc <2 x i32> %a to <2 x i16>815 ret <2 x i16> %r816}817 818define <2 x i16> @test_trunc_2xi32_muliple_use0(<2 x i32> %a, ptr %p) #0 {819; I16x2-LABEL: test_trunc_2xi32_muliple_use0(820; I16x2: {821; I16x2-NEXT: .reg .b32 %r<6>;822; I16x2-NEXT: .reg .b64 %rd<2>;823; I16x2-EMPTY:824; I16x2-NEXT: // %bb.0:825; I16x2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_2xi32_muliple_use0_param_0];826; I16x2-NEXT: ld.param.b64 %rd1, [test_trunc_2xi32_muliple_use0_param_1];827; I16x2-NEXT: prmt.b32 %r3, %r1, %r2, 0x5410U;828; I16x2-NEXT: mov.b32 %r4, 65537;829; I16x2-NEXT: add.s16x2 %r5, %r3, %r4;830; I16x2-NEXT: st.b32 [%rd1], %r5;831; I16x2-NEXT: st.param.b32 [func_retval0], %r3;832; I16x2-NEXT: ret;833;834; NO-I16x2-LABEL: test_trunc_2xi32_muliple_use0(835; NO-I16x2: {836; NO-I16x2-NEXT: .reg .b16 %rs<5>;837; NO-I16x2-NEXT: .reg .b32 %r<4>;838; NO-I16x2-NEXT: .reg .b64 %rd<2>;839; NO-I16x2-EMPTY:840; NO-I16x2-NEXT: // %bb.0:841; NO-I16x2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_2xi32_muliple_use0_param_0];842; NO-I16x2-NEXT: ld.param.b64 %rd1, [test_trunc_2xi32_muliple_use0_param_1];843; NO-I16x2-NEXT: cvt.u16.u32 %rs1, %r2;844; NO-I16x2-NEXT: cvt.u16.u32 %rs2, %r1;845; NO-I16x2-NEXT: mov.b32 %r3, {%rs2, %rs1};846; NO-I16x2-NEXT: add.s16 %rs3, %rs1, 1;847; NO-I16x2-NEXT: add.s16 %rs4, %rs2, 1;848; NO-I16x2-NEXT: st.v2.b16 [%rd1], {%rs4, %rs3};849; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r3;850; NO-I16x2-NEXT: ret;851 %r = trunc <2 x i32> %a to <2 x i16>852 ; Reuse the truncate - optimizing to PRMT when we don't have i16x2 vectors853 ; would increase register pressure854 %s = add <2 x i16> %r, splat (i16 1)855 store <2 x i16> %s, ptr %p856 ret <2 x i16> %r857}858 859define <2 x i16> @test_trunc_2xi32_muliple_use1(<2 x i32> %a, ptr %p) #0 {860; COMMON-LABEL: test_trunc_2xi32_muliple_use1(861; COMMON: {862; COMMON-NEXT: .reg .b32 %r<6>;863; COMMON-NEXT: .reg .b64 %rd<2>;864; COMMON-EMPTY:865; COMMON-NEXT: // %bb.0:866; COMMON-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_2xi32_muliple_use1_param_0];867; COMMON-NEXT: ld.param.b64 %rd1, [test_trunc_2xi32_muliple_use1_param_1];868; COMMON-NEXT: prmt.b32 %r3, %r1, %r2, 0x5410U;869; COMMON-NEXT: add.s32 %r4, %r2, 1;870; COMMON-NEXT: add.s32 %r5, %r1, 1;871; COMMON-NEXT: st.v2.b32 [%rd1], {%r5, %r4};872; COMMON-NEXT: st.param.b32 [func_retval0], %r3;873; COMMON-NEXT: ret;874 %r = trunc <2 x i32> %a to <2 x i16>875 ; Reuse the original value - optimizing to PRMT does not increase register876 ; pressure877 %s = add <2 x i32> %a, splat (i32 1)878 store <2 x i32> %s, ptr %p879 ret <2 x i16> %r880}881 882define <2 x i16> @test_trunc_2xi64(<2 x i64> %a) #0 {883; COMMON-LABEL: test_trunc_2xi64(884; COMMON: {885; COMMON-NEXT: .reg .b16 %rs<3>;886; COMMON-NEXT: .reg .b64 %rd<3>;887; COMMON-EMPTY:888; COMMON-NEXT: // %bb.0:889; COMMON-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];890; COMMON-NEXT: cvt.u16.u64 %rs1, %rd2;891; COMMON-NEXT: cvt.u16.u64 %rs2, %rd1;892; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};893; COMMON-NEXT: ret;894 %r = trunc <2 x i64> %a to <2 x i16>895 ret <2 x i16> %r896}897 898define <2 x i32> @test_zext_2xi32(<2 x i16> %a) #0 {899; COMMON-LABEL: test_zext_2xi32(900; COMMON: {901; COMMON-NEXT: .reg .b16 %rs<3>;902; COMMON-NEXT: .reg .b32 %r<4>;903; COMMON-EMPTY:904; COMMON-NEXT: // %bb.0:905; COMMON-NEXT: ld.param.b32 %r1, [test_zext_2xi32_param_0];906; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;907; COMMON-NEXT: cvt.u32.u16 %r2, %rs2;908; COMMON-NEXT: cvt.u32.u16 %r3, %rs1;909; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};910; COMMON-NEXT: ret;911 %r = zext <2 x i16> %a to <2 x i32>912 ret <2 x i32> %r913}914 915define <2 x i64> @test_zext_2xi64(<2 x i16> %a) #0 {916; COMMON-LABEL: test_zext_2xi64(917; COMMON: {918; COMMON-NEXT: .reg .b16 %rs<3>;919; COMMON-NEXT: .reg .b32 %r<2>;920; COMMON-NEXT: .reg .b64 %rd<3>;921; COMMON-EMPTY:922; COMMON-NEXT: // %bb.0:923; COMMON-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0];924; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;925; COMMON-NEXT: cvt.u64.u16 %rd1, %rs2;926; COMMON-NEXT: cvt.u64.u16 %rd2, %rs1;927; COMMON-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};928; COMMON-NEXT: ret;929 %r = zext <2 x i16> %a to <2 x i64>930 ret <2 x i64> %r931}932 933define <2 x i16> @test_bitcast_i32_to_2xi16(i32 %a) #0 {934; COMMON-LABEL: test_bitcast_i32_to_2xi16(935; COMMON: {936; COMMON-NEXT: .reg .b32 %r<2>;937; COMMON-EMPTY:938; COMMON-NEXT: // %bb.0:939; COMMON-NEXT: ld.param.b32 %r1, [test_bitcast_i32_to_2xi16_param_0];940; COMMON-NEXT: st.param.b32 [func_retval0], %r1;941; COMMON-NEXT: ret;942 %r = bitcast i32 %a to <2 x i16>943 ret <2 x i16> %r944}945 946define i32 @test_bitcast_2xi16_to_i32(<2 x i16> %a) #0 {947; COMMON-LABEL: test_bitcast_2xi16_to_i32(948; COMMON: {949; COMMON-NEXT: .reg .b32 %r<2>;950; COMMON-EMPTY:951; COMMON-NEXT: // %bb.0:952; COMMON-NEXT: ld.param.b32 %r1, [test_bitcast_2xi16_to_i32_param_0];953; COMMON-NEXT: st.param.b32 [func_retval0], %r1;954; COMMON-NEXT: ret;955 %r = bitcast <2 x i16> %a to i32956 ret i32 %r957}958 959define <2 x half> @test_bitcast_2xi16_to_2xhalf(i16 %a) #0 {960; COMMON-LABEL: test_bitcast_2xi16_to_2xhalf(961; COMMON: {962; COMMON-NEXT: .reg .b16 %rs<2>;963; COMMON-EMPTY:964; COMMON-NEXT: // %bb.0:965; COMMON-NEXT: ld.param.b16 %rs1, [test_bitcast_2xi16_to_2xhalf_param_0];966; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs1, 5};967; COMMON-NEXT: ret;968 %ins.0 = insertelement <2 x i16> undef, i16 %a, i32 0969 %ins.1 = insertelement <2 x i16> %ins.0, i16 5, i32 1970 %r = bitcast <2 x i16> %ins.1 to <2 x half>971 ret <2 x half> %r972}973 974 975define <2 x i16> @test_shufflevector(<2 x i16> %a) #0 {976; COMMON-LABEL: test_shufflevector(977; COMMON: {978; COMMON-NEXT: .reg .b16 %rs<3>;979; COMMON-NEXT: .reg .b32 %r<2>;980; COMMON-EMPTY:981; COMMON-NEXT: // %bb.0:982; COMMON-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];983; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;984; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};985; COMMON-NEXT: ret;986 %s = shufflevector <2 x i16> %a, <2 x i16> undef, <2 x i32> <i32 1, i32 0>987 ret <2 x i16> %s988}989 990define <2 x i16> @test_insertelement(<2 x i16> %a, i16 %x) #0 {991; I16x2-LABEL: test_insertelement(992; I16x2: {993; I16x2-NEXT: .reg .b16 %rs<3>;994; I16x2-NEXT: .reg .b32 %r<2>;995; I16x2-EMPTY:996; I16x2-NEXT: // %bb.0:997; I16x2-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];998; I16x2-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];999; I16x2-NEXT: mov.b32 {%rs2, _}, %r1;1000; I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};1001; I16x2-NEXT: ret;1002;1003; NO-I16x2-LABEL: test_insertelement(1004; NO-I16x2: {1005; NO-I16x2-NEXT: .reg .b16 %rs<3>;1006; NO-I16x2-NEXT: .reg .b32 %r<2>;1007; NO-I16x2-EMPTY:1008; NO-I16x2-NEXT: // %bb.0:1009; NO-I16x2-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];1010; NO-I16x2-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];1011; NO-I16x2-NEXT: { .reg .b16 tmp; mov.b32 {%rs2, tmp}, %r1; }1012; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};1013; NO-I16x2-NEXT: ret;1014 %i = insertelement <2 x i16> %a, i16 %x, i64 11015 ret <2 x i16> %i1016}1017 1018define <2 x i16> @test_fptosi_2xhalf_to_2xi16(<2 x half> %a) #0 {1019; COMMON-LABEL: test_fptosi_2xhalf_to_2xi16(1020; COMMON: {1021; COMMON-NEXT: .reg .b16 %rs<5>;1022; COMMON-NEXT: .reg .b32 %r<2>;1023; COMMON-EMPTY:1024; COMMON-NEXT: // %bb.0:1025; COMMON-NEXT: ld.param.b32 %r1, [test_fptosi_2xhalf_to_2xi16_param_0];1026; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;1027; COMMON-NEXT: cvt.rzi.s16.f16 %rs3, %rs2;1028; COMMON-NEXT: cvt.rzi.s16.f16 %rs4, %rs1;1029; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};1030; COMMON-NEXT: ret;1031 %r = fptosi <2 x half> %a to <2 x i16>1032 ret <2 x i16> %r1033}1034 1035define <2 x i16> @test_fptoui_2xhalf_to_2xi16(<2 x half> %a) #0 {1036; COMMON-LABEL: test_fptoui_2xhalf_to_2xi16(1037; COMMON: {1038; COMMON-NEXT: .reg .b16 %rs<5>;1039; COMMON-NEXT: .reg .b32 %r<2>;1040; COMMON-EMPTY:1041; COMMON-NEXT: // %bb.0:1042; COMMON-NEXT: ld.param.b32 %r1, [test_fptoui_2xhalf_to_2xi16_param_0];1043; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;1044; COMMON-NEXT: cvt.rzi.u16.f16 %rs3, %rs2;1045; COMMON-NEXT: cvt.rzi.u16.f16 %rs4, %rs1;1046; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};1047; COMMON-NEXT: ret;1048 %r = fptoui <2 x half> %a to <2 x i16>1049 ret <2 x i16> %r1050}1051 1052attributes #0 = { nounwind }1053