2318 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; ## Full FP16 support enabled by default.3; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \4; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \5; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes CHECK,CHECK-F16 %s6; RUN: %if ptxas-sm_53 %{ \7; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \8; RUN: -O0 -disable-post-ra -frame-pointer=all -verify-machineinstrs \9; RUN: | %ptxas-verify -arch=sm_53 \10; RUN: %}11; ## FP16 support explicitly disabled.12; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \13; RUN: -O0 -disable-post-ra -frame-pointer=all --nvptx-no-f16-math \14; RUN: -verify-machineinstrs \15; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes CHECK,CHECK-NOF16 %s16; RUN: %if ptxas-sm_53 %{ \17; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -mcpu=sm_53 \18; RUN: -O0 -disable-post-ra -frame-pointer=all --nvptx-no-f16-math \19; RUN: -verify-machineinstrs \20; RUN: | %ptxas-verify -arch=sm_53 \21; RUN: %}22; ## FP16 is not supported by hardware.23; RUN: llc < %s -O0 -mtriple=nvptx64-nvidia-cuda -mcpu=sm_52 \24; RUN: -disable-post-ra -frame-pointer=all -verify-machineinstrs \25; RUN: | FileCheck -allow-deprecated-dag-overlap -check-prefixes CHECK,CHECK-NOF16 %s26; RUN: %if ptxas-sm_52 %{ \27; RUN: llc < %s -O0 -mtriple=nvptx64-nvidia-cuda -mcpu=sm_52 \28; RUN: -disable-post-ra -frame-pointer=all -verify-machineinstrs \29; RUN: | %ptxas-verify -arch=sm_52 \30; RUN: %}31 32target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"33 34define <2 x half> @test_ret_const() #0 {35; CHECK-LABEL: test_ret_const(36; CHECK: {37; CHECK-EMPTY:38; CHECK-EMPTY:39; CHECK-NEXT: // %bb.0:40; CHECK-NEXT: st.param.b32 [func_retval0], 1073757184;41; CHECK-NEXT: ret;42 ret <2 x half> <half 1.0, half 2.0>43}44 45define half @test_extract_0(<2 x half> %a) #0 {46; CHECK-LABEL: test_extract_0(47; CHECK: {48; CHECK-NEXT: .reg .b16 %rs<2>;49; CHECK-NEXT: .reg .b32 %r<2>;50; CHECK-EMPTY:51; CHECK-NEXT: // %bb.0:52; CHECK-NEXT: ld.param.b32 %r1, [test_extract_0_param_0];53; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r1; }54; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;55; CHECK-NEXT: ret;56 %e = extractelement <2 x half> %a, i32 057 ret half %e58}59 60define half @test_extract_1(<2 x half> %a) #0 {61; CHECK-LABEL: test_extract_1(62; CHECK: {63; CHECK-NEXT: .reg .b16 %rs<2>;64; CHECK-NEXT: .reg .b32 %r<2>;65; CHECK-EMPTY:66; CHECK-NEXT: // %bb.0:67; CHECK-NEXT: ld.param.b32 %r1, [test_extract_1_param_0];68; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }69; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;70; CHECK-NEXT: ret;71 %e = extractelement <2 x half> %a, i32 172 ret half %e73}74 75define half @test_extract_i(<2 x half> %a, i64 %idx) #0 {76; CHECK-LABEL: test_extract_i(77; CHECK: {78; CHECK-NEXT: .reg .pred %p<2>;79; CHECK-NEXT: .reg .b16 %rs<4>;80; CHECK-NEXT: .reg .b32 %r<2>;81; CHECK-NEXT: .reg .b64 %rd<2>;82; CHECK-EMPTY:83; CHECK-NEXT: // %bb.0:84; CHECK-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];85; CHECK-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];86; CHECK-NEXT: setp.eq.b64 %p1, %rd1, 0;87; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;88; CHECK-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;89; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;90; CHECK-NEXT: ret;91 %e = extractelement <2 x half> %a, i64 %idx92 ret half %e93}94 95define <2 x half> @test_fadd(<2 x half> %a, <2 x half> %b) #0 {96; CHECK-F16-LABEL: test_fadd(97; CHECK-F16: {98; CHECK-F16-NEXT: .reg .b32 %r<4>;99; CHECK-F16-EMPTY:100; CHECK-F16-NEXT: // %bb.0:101; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fadd_param_1];102; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fadd_param_0];103; CHECK-F16-NEXT: add.rn.f16x2 %r3, %r1, %r2;104; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;105; CHECK-F16-NEXT: ret;106;107; CHECK-NOF16-LABEL: test_fadd(108; CHECK-NOF16: {109; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;110; CHECK-NOF16-NEXT: .reg .b32 %r<10>;111; CHECK-NOF16-EMPTY:112; CHECK-NOF16-NEXT: // %bb.0:113; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fadd_param_1];114; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_param_0];115; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;116; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;117; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;118; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;119; CHECK-NOF16-NEXT: add.rn.f32 %r5, %r4, %r3;120; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r5;121; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;122; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;123; CHECK-NOF16-NEXT: add.rn.f32 %r8, %r7, %r6;124; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r8;125; CHECK-NOF16-NEXT: mov.b32 %r9, {%rs6, %rs5};126; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r9;127; CHECK-NOF16-NEXT: ret;128 %r = fadd <2 x half> %a, %b129 ret <2 x half> %r130}131 132; Check that we can lower fadd with immediate arguments.133define <2 x half> @test_fadd_imm_0(<2 x half> %a) #0 {134; CHECK-F16-LABEL: test_fadd_imm_0(135; CHECK-F16: {136; CHECK-F16-NEXT: .reg .b32 %r<4>;137; CHECK-F16-EMPTY:138; CHECK-F16-NEXT: // %bb.0:139; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fadd_imm_0_param_0];140; CHECK-F16-NEXT: mov.b32 %r2, 1073757184;141; CHECK-F16-NEXT: add.rn.f16x2 %r3, %r1, %r2;142; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;143; CHECK-F16-NEXT: ret;144;145; CHECK-NOF16-LABEL: test_fadd_imm_0(146; CHECK-NOF16: {147; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;148; CHECK-NOF16-NEXT: .reg .b32 %r<7>;149; CHECK-NOF16-EMPTY:150; CHECK-NOF16-NEXT: // %bb.0:151; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_imm_0_param_0];152; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;153; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;154; CHECK-NOF16-NEXT: add.rn.f32 %r3, %r2, 0f40000000;155; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;156; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs1;157; CHECK-NOF16-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;158; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs4, %r5;159; CHECK-NOF16-NEXT: mov.b32 %r6, {%rs4, %rs3};160; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r6;161; CHECK-NOF16-NEXT: ret;162 %r = fadd <2 x half> <half 1.0, half 2.0>, %a163 ret <2 x half> %r164}165 166define <2 x half> @test_fadd_imm_1(<2 x half> %a) #0 {167; CHECK-F16-LABEL: test_fadd_imm_1(168; CHECK-F16: {169; CHECK-F16-NEXT: .reg .b32 %r<4>;170; CHECK-F16-EMPTY:171; CHECK-F16-NEXT: // %bb.0:172; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fadd_imm_1_param_0];173; CHECK-F16-NEXT: mov.b32 %r2, 1073757184;174; CHECK-F16-NEXT: add.rn.f16x2 %r3, %r1, %r2;175; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;176; CHECK-F16-NEXT: ret;177;178; CHECK-NOF16-LABEL: test_fadd_imm_1(179; CHECK-NOF16: {180; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;181; CHECK-NOF16-NEXT: .reg .b32 %r<7>;182; CHECK-NOF16-EMPTY:183; CHECK-NOF16-NEXT: // %bb.0:184; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_imm_1_param_0];185; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;186; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;187; CHECK-NOF16-NEXT: add.rn.f32 %r3, %r2, 0f40000000;188; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;189; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs1;190; CHECK-NOF16-NEXT: add.rn.f32 %r5, %r4, 0f3F800000;191; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs4, %r5;192; CHECK-NOF16-NEXT: mov.b32 %r6, {%rs4, %rs3};193; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r6;194; CHECK-NOF16-NEXT: ret;195 %r = fadd <2 x half> %a, <half 1.0, half 2.0>196 ret <2 x half> %r197}198 199define <2 x half> @test_fsub(<2 x half> %a, <2 x half> %b) #0 {200; CHECK-F16-LABEL: test_fsub(201; CHECK-F16: {202; CHECK-F16-NEXT: .reg .b32 %r<4>;203; CHECK-F16-EMPTY:204; CHECK-F16-NEXT: // %bb.0:205; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fsub_param_1];206; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fsub_param_0];207; CHECK-F16-NEXT: sub.rn.f16x2 %r3, %r1, %r2;208; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;209; CHECK-F16-NEXT: ret;210;211; CHECK-NOF16-LABEL: test_fsub(212; CHECK-NOF16: {213; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;214; CHECK-NOF16-NEXT: .reg .b32 %r<10>;215; CHECK-NOF16-EMPTY:216; CHECK-NOF16-NEXT: // %bb.0:217; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fsub_param_1];218; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fsub_param_0];219; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;220; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;221; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;222; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;223; CHECK-NOF16-NEXT: sub.rn.f32 %r5, %r4, %r3;224; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r5;225; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;226; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;227; CHECK-NOF16-NEXT: sub.rn.f32 %r8, %r7, %r6;228; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r8;229; CHECK-NOF16-NEXT: mov.b32 %r9, {%rs6, %rs5};230; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r9;231; CHECK-NOF16-NEXT: ret;232 %r = fsub <2 x half> %a, %b233 ret <2 x half> %r234}235 236define <2 x half> @test_fneg(<2 x half> %a) #0 {237; CHECK-F16-LABEL: test_fneg(238; CHECK-F16: {239; CHECK-F16-NEXT: .reg .b32 %r<4>;240; CHECK-F16-EMPTY:241; CHECK-F16-NEXT: // %bb.0:242; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fneg_param_0];243; CHECK-F16-NEXT: mov.b32 %r2, 0;244; CHECK-F16-NEXT: sub.rn.f16x2 %r3, %r2, %r1;245; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;246; CHECK-F16-NEXT: ret;247;248; CHECK-NOF16-LABEL: test_fneg(249; CHECK-NOF16: {250; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;251; CHECK-NOF16-NEXT: .reg .b32 %r<8>;252; CHECK-NOF16-EMPTY:253; CHECK-NOF16-NEXT: // %bb.0:254; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fneg_param_0];255; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;256; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;257; CHECK-NOF16-NEXT: mov.b32 %r3, 0f00000000;258; CHECK-NOF16-NEXT: sub.rn.f32 %r4, %r3, %r2;259; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r4;260; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;261; CHECK-NOF16-NEXT: sub.rn.f32 %r6, %r3, %r5;262; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs4, %r6;263; CHECK-NOF16-NEXT: mov.b32 %r7, {%rs4, %rs3};264; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;265; CHECK-NOF16-NEXT: ret;266 %r = fsub <2 x half> <half 0.0, half 0.0>, %a267 ret <2 x half> %r268}269 270define <2 x half> @test_fmul(<2 x half> %a, <2 x half> %b) #0 {271; CHECK-F16-LABEL: test_fmul(272; CHECK-F16: {273; CHECK-F16-NEXT: .reg .b32 %r<4>;274; CHECK-F16-EMPTY:275; CHECK-F16-NEXT: // %bb.0:276; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fmul_param_1];277; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fmul_param_0];278; CHECK-F16-NEXT: mul.rn.f16x2 %r3, %r1, %r2;279; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;280; CHECK-F16-NEXT: ret;281;282; CHECK-NOF16-LABEL: test_fmul(283; CHECK-NOF16: {284; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;285; CHECK-NOF16-NEXT: .reg .b32 %r<10>;286; CHECK-NOF16-EMPTY:287; CHECK-NOF16-NEXT: // %bb.0:288; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fmul_param_1];289; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fmul_param_0];290; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;291; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;292; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;293; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;294; CHECK-NOF16-NEXT: mul.rn.f32 %r5, %r4, %r3;295; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r5;296; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;297; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;298; CHECK-NOF16-NEXT: mul.rn.f32 %r8, %r7, %r6;299; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r8;300; CHECK-NOF16-NEXT: mov.b32 %r9, {%rs6, %rs5};301; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r9;302; CHECK-NOF16-NEXT: ret;303 %r = fmul <2 x half> %a, %b304 ret <2 x half> %r305}306 307define <2 x half> @test_fdiv(<2 x half> %a, <2 x half> %b) #0 {308; CHECK-LABEL: test_fdiv(309; CHECK: {310; CHECK-NEXT: .reg .b16 %rs<7>;311; CHECK-NEXT: .reg .b32 %r<10>;312; CHECK-EMPTY:313; CHECK-NEXT: // %bb.0:314; CHECK-NEXT: ld.param.b32 %r2, [test_fdiv_param_1];315; CHECK-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];316; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;317; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;318; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;319; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;320; CHECK-NEXT: div.rn.f32 %r5, %r4, %r3;321; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r5;322; CHECK-NEXT: cvt.f32.f16 %r6, %rs1;323; CHECK-NEXT: cvt.f32.f16 %r7, %rs3;324; CHECK-NEXT: div.rn.f32 %r8, %r7, %r6;325; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r8;326; CHECK-NEXT: mov.b32 %r9, {%rs6, %rs5};327; CHECK-NEXT: st.param.b32 [func_retval0], %r9;328; CHECK-NEXT: ret;329 %r = fdiv <2 x half> %a, %b330 ret <2 x half> %r331}332 333; -- Load two 16x2 inputs and split them into f16 elements334; -- Split into elements335; -- promote to f32.336; -- frem(a[0],b[0]).337; -- frem(a[1],b[1]).338; -- convert back to f16.339; -- merge into f16x2 and return it.340define <2 x half> @test_frem(<2 x half> %a, <2 x half> %b) #0 {341; CHECK-LABEL: test_frem(342; CHECK: {343; CHECK-NEXT: .reg .pred %p<3>;344; CHECK-NEXT: .reg .b16 %rs<7>;345; CHECK-NEXT: .reg .b32 %r<18>;346; CHECK-EMPTY:347; CHECK-NEXT: // %bb.0:348; CHECK-NEXT: ld.param.b32 %r2, [test_frem_param_1];349; CHECK-NEXT: ld.param.b32 %r1, [test_frem_param_0];350; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;351; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;352; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;353; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;354; CHECK-NEXT: div.rn.f32 %r5, %r4, %r3;355; CHECK-NEXT: cvt.rzi.f32.f32 %r6, %r5;356; CHECK-NEXT: neg.f32 %r7, %r6;357; CHECK-NEXT: fma.rn.f32 %r8, %r7, %r3, %r4;358; CHECK-NEXT: testp.infinite.f32 %p1, %r3;359; CHECK-NEXT: selp.f32 %r9, %r4, %r8, %p1;360; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r9;361; CHECK-NEXT: cvt.f32.f16 %r10, %rs1;362; CHECK-NEXT: cvt.f32.f16 %r11, %rs3;363; CHECK-NEXT: div.rn.f32 %r12, %r11, %r10;364; CHECK-NEXT: cvt.rzi.f32.f32 %r13, %r12;365; CHECK-NEXT: neg.f32 %r14, %r13;366; CHECK-NEXT: fma.rn.f32 %r15, %r14, %r10, %r11;367; CHECK-NEXT: testp.infinite.f32 %p2, %r10;368; CHECK-NEXT: selp.f32 %r16, %r11, %r15, %p2;369; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r16;370; CHECK-NEXT: mov.b32 %r17, {%rs6, %rs5};371; CHECK-NEXT: st.param.b32 [func_retval0], %r17;372; CHECK-NEXT: ret;373 %r = frem <2 x half> %a, %b374 ret <2 x half> %r375}376 377define void @test_ldst_v2f16(ptr %a, ptr %b) {378; CHECK-LABEL: test_ldst_v2f16(379; CHECK: {380; CHECK-NEXT: .reg .b32 %r<2>;381; CHECK-NEXT: .reg .b64 %rd<3>;382; CHECK-EMPTY:383; CHECK-NEXT: // %bb.0:384; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v2f16_param_1];385; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v2f16_param_0];386; CHECK-NEXT: ld.b32 %r1, [%rd1];387; CHECK-NEXT: st.b32 [%rd2], %r1;388; CHECK-NEXT: ret;389 %t1 = load <2 x half>, ptr %a390 store <2 x half> %t1, ptr %b, align 16391 ret void392}393 394; -- v3 is inconvenient to capture as it's lowered as ld.b64 + fair395; number of bitshifting instructions that may change at llvm's whim.396; So we only verify that we only issue correct number of writes using397; correct offset, but not the values we write.398define void @test_ldst_v3f16(ptr %a, ptr %b) {399; CHECK-LABEL: test_ldst_v3f16(400; CHECK: {401; CHECK-NEXT: .reg .b16 %rs<2>;402; CHECK-NEXT: .reg .b32 %r<2>;403; CHECK-NEXT: .reg .b64 %rd<4>;404; CHECK-EMPTY:405; CHECK-NEXT: // %bb.0:406; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v3f16_param_1];407; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v3f16_param_0];408; CHECK-NEXT: ld.b64 %rd3, [%rd1];409; CHECK-NEXT: { .reg .b32 tmp; mov.b64 {tmp, %r1}, %rd3; }410; CHECK-NEXT: st.b32 [%rd2], %rd3;411; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r1; }412; CHECK-NEXT: st.b16 [%rd2+4], %rs1;413; CHECK-NEXT: ret;414 %t1 = load <3 x half>, ptr %a415 store <3 x half> %t1, ptr %b, align 16416 ret void417}418 419define void @test_ldst_v4f16(ptr %a, ptr %b) {420; CHECK-LABEL: test_ldst_v4f16(421; CHECK: {422; CHECK-NEXT: .reg .b32 %r<3>;423; CHECK-NEXT: .reg .b64 %rd<3>;424; CHECK-EMPTY:425; CHECK-NEXT: // %bb.0:426; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v4f16_param_1];427; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v4f16_param_0];428; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];429; CHECK-NEXT: st.v2.b32 [%rd2], {%r1, %r2};430; CHECK-NEXT: ret;431 %t1 = load <4 x half>, ptr %a432 store <4 x half> %t1, ptr %b, align 16433 ret void434}435 436define void @test_ldst_v8f16(ptr %a, ptr %b) {437; CHECK-LABEL: test_ldst_v8f16(438; CHECK: {439; CHECK-NEXT: .reg .b32 %r<5>;440; CHECK-NEXT: .reg .b64 %rd<3>;441; CHECK-EMPTY:442; CHECK-NEXT: // %bb.0:443; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v8f16_param_1];444; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v8f16_param_0];445; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];446; CHECK-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};447; CHECK-NEXT: ret;448 %t1 = load <8 x half>, ptr %a449 store <8 x half> %t1, ptr %b, align 16450 ret void451}452 453declare <2 x half> @test_callee(<2 x half> %a, <2 x half> %b) #0454 455define <2 x half> @test_call(<2 x half> %a, <2 x half> %b) #0 {456; CHECK-LABEL: test_call(457; CHECK: {458; CHECK-NEXT: .reg .b32 %r<4>;459; CHECK-EMPTY:460; CHECK-NEXT: // %bb.0:461; CHECK-NEXT: ld.param.b32 %r2, [test_call_param_1];462; CHECK-NEXT: ld.param.b32 %r1, [test_call_param_0];463; CHECK-NEXT: { // callseq 0, 0464; CHECK-NEXT: .param .align 4 .b8 param0[4];465; CHECK-NEXT: .param .align 4 .b8 param1[4];466; CHECK-NEXT: .param .align 4 .b8 retval0[4];467; CHECK-NEXT: st.param.b32 [param1], %r2;468; CHECK-NEXT: st.param.b32 [param0], %r1;469; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);470; CHECK-NEXT: ld.param.b32 %r3, [retval0];471; CHECK-NEXT: } // callseq 0472; CHECK-NEXT: st.param.b32 [func_retval0], %r3;473; CHECK-NEXT: ret;474 %r = call <2 x half> @test_callee(<2 x half> %a, <2 x half> %b)475 ret <2 x half> %r476}477 478define <2 x half> @test_call_flipped(<2 x half> %a, <2 x half> %b) #0 {479; CHECK-LABEL: test_call_flipped(480; CHECK: {481; CHECK-NEXT: .reg .b32 %r<4>;482; CHECK-EMPTY:483; CHECK-NEXT: // %bb.0:484; CHECK-NEXT: ld.param.b32 %r2, [test_call_flipped_param_1];485; CHECK-NEXT: ld.param.b32 %r1, [test_call_flipped_param_0];486; CHECK-NEXT: { // callseq 1, 0487; CHECK-NEXT: .param .align 4 .b8 param0[4];488; CHECK-NEXT: .param .align 4 .b8 param1[4];489; CHECK-NEXT: .param .align 4 .b8 retval0[4];490; CHECK-NEXT: st.param.b32 [param1], %r1;491; CHECK-NEXT: st.param.b32 [param0], %r2;492; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);493; CHECK-NEXT: ld.param.b32 %r3, [retval0];494; CHECK-NEXT: } // callseq 1495; CHECK-NEXT: st.param.b32 [func_retval0], %r3;496; CHECK-NEXT: ret;497 %r = call <2 x half> @test_callee(<2 x half> %b, <2 x half> %a)498 ret <2 x half> %r499}500 501define <2 x half> @test_tailcall_flipped(<2 x half> %a, <2 x half> %b) #0 {502; CHECK-LABEL: test_tailcall_flipped(503; CHECK: {504; CHECK-NEXT: .reg .b32 %r<4>;505; CHECK-EMPTY:506; CHECK-NEXT: // %bb.0:507; CHECK-NEXT: ld.param.b32 %r2, [test_tailcall_flipped_param_1];508; CHECK-NEXT: ld.param.b32 %r1, [test_tailcall_flipped_param_0];509; CHECK-NEXT: { // callseq 2, 0510; CHECK-NEXT: .param .align 4 .b8 param0[4];511; CHECK-NEXT: .param .align 4 .b8 param1[4];512; CHECK-NEXT: .param .align 4 .b8 retval0[4];513; CHECK-NEXT: st.param.b32 [param1], %r1;514; CHECK-NEXT: st.param.b32 [param0], %r2;515; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);516; CHECK-NEXT: ld.param.b32 %r3, [retval0];517; CHECK-NEXT: } // callseq 2518; CHECK-NEXT: st.param.b32 [func_retval0], %r3;519; CHECK-NEXT: ret;520 %r = tail call <2 x half> @test_callee(<2 x half> %b, <2 x half> %a)521 ret <2 x half> %r522}523 524define <2 x half> @test_select(<2 x half> %a, <2 x half> %b, i1 zeroext %c) #0 {525; CHECK-LABEL: test_select(526; CHECK: {527; CHECK-NEXT: .reg .pred %p<2>;528; CHECK-NEXT: .reg .b16 %rs<3>;529; CHECK-NEXT: .reg .b32 %r<4>;530; CHECK-EMPTY:531; CHECK-NEXT: // %bb.0:532; CHECK-NEXT: ld.param.b8 %rs1, [test_select_param_2];533; CHECK-NEXT: and.b16 %rs2, %rs1, 1;534; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;535; CHECK-NEXT: ld.param.b32 %r2, [test_select_param_1];536; CHECK-NEXT: ld.param.b32 %r1, [test_select_param_0];537; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p1;538; CHECK-NEXT: st.param.b32 [func_retval0], %r3;539; CHECK-NEXT: ret;540 %r = select i1 %c, <2 x half> %a, <2 x half> %b541 ret <2 x half> %r542}543 544define <2 x half> @test_select_cc(<2 x half> %a, <2 x half> %b, <2 x half> %c, <2 x half> %d) #0 {545; CHECK-F16-LABEL: test_select_cc(546; CHECK-F16: {547; CHECK-F16-NEXT: .reg .pred %p<3>;548; CHECK-F16-NEXT: .reg .b16 %rs<7>;549; CHECK-F16-NEXT: .reg .b32 %r<5>;550; CHECK-F16-EMPTY:551; CHECK-F16-NEXT: // %bb.0:552; CHECK-F16-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];553; CHECK-F16-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];554; CHECK-F16-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];555; CHECK-F16-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];556; CHECK-F16-NEXT: setp.neu.f16x2 %p1|%p2, %r3, %r4;557; CHECK-F16-NEXT: mov.b32 {%rs1, %rs2}, %r2;558; CHECK-F16-NEXT: mov.b32 {%rs3, %rs4}, %r1;559; CHECK-F16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p2;560; CHECK-F16-NEXT: selp.b16 %rs6, %rs3, %rs1, %p1;561; CHECK-F16-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};562; CHECK-F16-NEXT: ret;563;564; CHECK-NOF16-LABEL: test_select_cc(565; CHECK-NOF16: {566; CHECK-NOF16-NEXT: .reg .pred %p<3>;567; CHECK-NOF16-NEXT: .reg .b16 %rs<11>;568; CHECK-NOF16-NEXT: .reg .b32 %r<9>;569; CHECK-NOF16-EMPTY:570; CHECK-NOF16-NEXT: // %bb.0:571; CHECK-NOF16-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];572; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];573; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];574; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];575; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r4;576; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;577; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r3;578; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;579; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r6, %r5;580; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs2;581; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs4;582; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r8, %r7;583; CHECK-NOF16-NEXT: mov.b32 {%rs5, %rs6}, %r2;584; CHECK-NOF16-NEXT: mov.b32 {%rs7, %rs8}, %r1;585; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs6, %p2;586; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs7, %rs5, %p1;587; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};588; CHECK-NOF16-NEXT: ret;589 %cc = fcmp une <2 x half> %c, %d590 %r = select <2 x i1> %cc, <2 x half> %a, <2 x half> %b591 ret <2 x half> %r592}593 594define <2 x float> @test_select_cc_f32_f16(<2 x float> %a, <2 x float> %b,595; CHECK-F16-LABEL: test_select_cc_f32_f16(596; CHECK-F16: {597; CHECK-F16-NEXT: .reg .pred %p<3>;598; CHECK-F16-NEXT: .reg .b32 %r<9>;599; CHECK-F16-EMPTY:600; CHECK-F16-NEXT: // %bb.0:601; CHECK-F16-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f32_f16_param_1];602; CHECK-F16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f32_f16_param_0];603; CHECK-F16-NEXT: ld.param.b32 %r6, [test_select_cc_f32_f16_param_3];604; CHECK-F16-NEXT: ld.param.b32 %r5, [test_select_cc_f32_f16_param_2];605; CHECK-F16-NEXT: setp.neu.f16x2 %p1|%p2, %r5, %r6;606; CHECK-F16-NEXT: selp.f32 %r7, %r2, %r4, %p2;607; CHECK-F16-NEXT: selp.f32 %r8, %r1, %r3, %p1;608; CHECK-F16-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};609; CHECK-F16-NEXT: ret;610;611; CHECK-NOF16-LABEL: test_select_cc_f32_f16(612; CHECK-NOF16: {613; CHECK-NOF16-NEXT: .reg .pred %p<3>;614; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;615; CHECK-NOF16-NEXT: .reg .b32 %r<13>;616; CHECK-NOF16-EMPTY:617; CHECK-NOF16-NEXT: // %bb.0:618; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f32_f16_param_1];619; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f32_f16_param_0];620; CHECK-NOF16-NEXT: ld.param.b32 %r6, [test_select_cc_f32_f16_param_3];621; CHECK-NOF16-NEXT: ld.param.b32 %r5, [test_select_cc_f32_f16_param_2];622; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r6;623; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;624; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r5;625; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;626; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r8, %r7;627; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs2;628; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs4;629; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r10, %r9;630; CHECK-NOF16-NEXT: selp.f32 %r11, %r2, %r4, %p2;631; CHECK-NOF16-NEXT: selp.f32 %r12, %r1, %r3, %p1;632; CHECK-NOF16-NEXT: st.param.v2.b32 [func_retval0], {%r12, %r11};633; CHECK-NOF16-NEXT: ret;634 <2 x half> %c, <2 x half> %d) #0 {635 %cc = fcmp une <2 x half> %c, %d636 %r = select <2 x i1> %cc, <2 x float> %a, <2 x float> %b637 ret <2 x float> %r638}639 640define <2 x half> @test_select_cc_f16_f32(<2 x half> %a, <2 x half> %b,641; CHECK-LABEL: test_select_cc_f16_f32(642; CHECK: {643; CHECK-NEXT: .reg .pred %p<3>;644; CHECK-NEXT: .reg .b16 %rs<7>;645; CHECK-NEXT: .reg .b32 %r<7>;646; CHECK-EMPTY:647; CHECK-NEXT: // %bb.0:648; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_f16_f32_param_3];649; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f16_f32_param_2];650; CHECK-NEXT: ld.param.b32 %r2, [test_select_cc_f16_f32_param_1];651; CHECK-NEXT: ld.param.b32 %r1, [test_select_cc_f16_f32_param_0];652; CHECK-NEXT: setp.neu.f32 %p1, %r3, %r5;653; CHECK-NEXT: setp.neu.f32 %p2, %r4, %r6;654; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;655; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;656; CHECK-NEXT: selp.b16 %rs5, %rs4, %rs2, %p2;657; CHECK-NEXT: selp.b16 %rs6, %rs3, %rs1, %p1;658; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};659; CHECK-NEXT: ret;660 <2 x float> %c, <2 x float> %d) #0 {661 %cc = fcmp une <2 x float> %c, %d662 %r = select <2 x i1> %cc, <2 x half> %a, <2 x half> %b663 ret <2 x half> %r664}665 666define <2 x i1> @test_fcmp_une(<2 x half> %a, <2 x half> %b) #0 {667; CHECK-F16-LABEL: test_fcmp_une(668; CHECK-F16: {669; CHECK-F16-NEXT: .reg .pred %p<3>;670; CHECK-F16-NEXT: .reg .b16 %rs<3>;671; CHECK-F16-NEXT: .reg .b32 %r<3>;672; CHECK-F16-EMPTY:673; CHECK-F16-NEXT: // %bb.0:674; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_une_param_1];675; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_une_param_0];676; CHECK-F16-NEXT: setp.neu.f16x2 %p1|%p2, %r1, %r2;677; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;678; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;679; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;680; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;681; CHECK-F16-NEXT: ret;682;683; CHECK-NOF16-LABEL: test_fcmp_une(684; CHECK-NOF16: {685; CHECK-NOF16-NEXT: .reg .pred %p<3>;686; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;687; CHECK-NOF16-NEXT: .reg .b32 %r<7>;688; CHECK-NOF16-EMPTY:689; CHECK-NOF16-NEXT: // %bb.0:690; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_une_param_1];691; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_une_param_0];692; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;693; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;694; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;695; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;696; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r4, %r3;697; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;698; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;699; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r6, %r5;700; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;701; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;702; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;703; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;704; CHECK-NOF16-NEXT: ret;705 %r = fcmp une <2 x half> %a, %b706 ret <2 x i1> %r707}708 709define <2 x i1> @test_fcmp_ueq(<2 x half> %a, <2 x half> %b) #0 {710; CHECK-F16-LABEL: test_fcmp_ueq(711; CHECK-F16: {712; CHECK-F16-NEXT: .reg .pred %p<3>;713; CHECK-F16-NEXT: .reg .b16 %rs<3>;714; CHECK-F16-NEXT: .reg .b32 %r<3>;715; CHECK-F16-EMPTY:716; CHECK-F16-NEXT: // %bb.0:717; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_ueq_param_1];718; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_ueq_param_0];719; CHECK-F16-NEXT: setp.equ.f16x2 %p1|%p2, %r1, %r2;720; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;721; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;722; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;723; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;724; CHECK-F16-NEXT: ret;725;726; CHECK-NOF16-LABEL: test_fcmp_ueq(727; CHECK-NOF16: {728; CHECK-NOF16-NEXT: .reg .pred %p<3>;729; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;730; CHECK-NOF16-NEXT: .reg .b32 %r<7>;731; CHECK-NOF16-EMPTY:732; CHECK-NOF16-NEXT: // %bb.0:733; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ueq_param_1];734; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ueq_param_0];735; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;736; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;737; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;738; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;739; CHECK-NOF16-NEXT: setp.equ.f32 %p1, %r4, %r3;740; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;741; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;742; CHECK-NOF16-NEXT: setp.equ.f32 %p2, %r6, %r5;743; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;744; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;745; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;746; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;747; CHECK-NOF16-NEXT: ret;748 %r = fcmp ueq <2 x half> %a, %b749 ret <2 x i1> %r750}751 752define <2 x i1> @test_fcmp_ugt(<2 x half> %a, <2 x half> %b) #0 {753; CHECK-F16-LABEL: test_fcmp_ugt(754; CHECK-F16: {755; CHECK-F16-NEXT: .reg .pred %p<3>;756; CHECK-F16-NEXT: .reg .b16 %rs<3>;757; CHECK-F16-NEXT: .reg .b32 %r<3>;758; CHECK-F16-EMPTY:759; CHECK-F16-NEXT: // %bb.0:760; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_ugt_param_1];761; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_ugt_param_0];762; CHECK-F16-NEXT: setp.gtu.f16x2 %p1|%p2, %r1, %r2;763; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;764; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;765; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;766; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;767; CHECK-F16-NEXT: ret;768;769; CHECK-NOF16-LABEL: test_fcmp_ugt(770; CHECK-NOF16: {771; CHECK-NOF16-NEXT: .reg .pred %p<3>;772; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;773; CHECK-NOF16-NEXT: .reg .b32 %r<7>;774; CHECK-NOF16-EMPTY:775; CHECK-NOF16-NEXT: // %bb.0:776; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ugt_param_1];777; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ugt_param_0];778; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;779; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;780; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;781; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;782; CHECK-NOF16-NEXT: setp.gtu.f32 %p1, %r4, %r3;783; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;784; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;785; CHECK-NOF16-NEXT: setp.gtu.f32 %p2, %r6, %r5;786; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;787; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;788; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;789; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;790; CHECK-NOF16-NEXT: ret;791 %r = fcmp ugt <2 x half> %a, %b792 ret <2 x i1> %r793}794 795define <2 x i1> @test_fcmp_uge(<2 x half> %a, <2 x half> %b) #0 {796; CHECK-F16-LABEL: test_fcmp_uge(797; CHECK-F16: {798; CHECK-F16-NEXT: .reg .pred %p<3>;799; CHECK-F16-NEXT: .reg .b16 %rs<3>;800; CHECK-F16-NEXT: .reg .b32 %r<3>;801; CHECK-F16-EMPTY:802; CHECK-F16-NEXT: // %bb.0:803; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_uge_param_1];804; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_uge_param_0];805; CHECK-F16-NEXT: setp.geu.f16x2 %p1|%p2, %r1, %r2;806; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;807; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;808; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;809; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;810; CHECK-F16-NEXT: ret;811;812; CHECK-NOF16-LABEL: test_fcmp_uge(813; CHECK-NOF16: {814; CHECK-NOF16-NEXT: .reg .pred %p<3>;815; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;816; CHECK-NOF16-NEXT: .reg .b32 %r<7>;817; CHECK-NOF16-EMPTY:818; CHECK-NOF16-NEXT: // %bb.0:819; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_uge_param_1];820; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_uge_param_0];821; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;822; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;823; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;824; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;825; CHECK-NOF16-NEXT: setp.geu.f32 %p1, %r4, %r3;826; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;827; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;828; CHECK-NOF16-NEXT: setp.geu.f32 %p2, %r6, %r5;829; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;830; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;831; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;832; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;833; CHECK-NOF16-NEXT: ret;834 %r = fcmp uge <2 x half> %a, %b835 ret <2 x i1> %r836}837 838define <2 x i1> @test_fcmp_ult(<2 x half> %a, <2 x half> %b) #0 {839; CHECK-F16-LABEL: test_fcmp_ult(840; CHECK-F16: {841; CHECK-F16-NEXT: .reg .pred %p<3>;842; CHECK-F16-NEXT: .reg .b16 %rs<3>;843; CHECK-F16-NEXT: .reg .b32 %r<3>;844; CHECK-F16-EMPTY:845; CHECK-F16-NEXT: // %bb.0:846; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_ult_param_1];847; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_ult_param_0];848; CHECK-F16-NEXT: setp.ltu.f16x2 %p1|%p2, %r1, %r2;849; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;850; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;851; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;852; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;853; CHECK-F16-NEXT: ret;854;855; CHECK-NOF16-LABEL: test_fcmp_ult(856; CHECK-NOF16: {857; CHECK-NOF16-NEXT: .reg .pred %p<3>;858; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;859; CHECK-NOF16-NEXT: .reg .b32 %r<7>;860; CHECK-NOF16-EMPTY:861; CHECK-NOF16-NEXT: // %bb.0:862; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ult_param_1];863; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ult_param_0];864; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;865; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;866; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;867; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;868; CHECK-NOF16-NEXT: setp.ltu.f32 %p1, %r4, %r3;869; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;870; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;871; CHECK-NOF16-NEXT: setp.ltu.f32 %p2, %r6, %r5;872; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;873; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;874; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;875; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;876; CHECK-NOF16-NEXT: ret;877 %r = fcmp ult <2 x half> %a, %b878 ret <2 x i1> %r879}880 881define <2 x i1> @test_fcmp_ule(<2 x half> %a, <2 x half> %b) #0 {882; CHECK-F16-LABEL: test_fcmp_ule(883; CHECK-F16: {884; CHECK-F16-NEXT: .reg .pred %p<3>;885; CHECK-F16-NEXT: .reg .b16 %rs<3>;886; CHECK-F16-NEXT: .reg .b32 %r<3>;887; CHECK-F16-EMPTY:888; CHECK-F16-NEXT: // %bb.0:889; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_ule_param_1];890; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_ule_param_0];891; CHECK-F16-NEXT: setp.leu.f16x2 %p1|%p2, %r1, %r2;892; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;893; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;894; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;895; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;896; CHECK-F16-NEXT: ret;897;898; CHECK-NOF16-LABEL: test_fcmp_ule(899; CHECK-NOF16: {900; CHECK-NOF16-NEXT: .reg .pred %p<3>;901; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;902; CHECK-NOF16-NEXT: .reg .b32 %r<7>;903; CHECK-NOF16-EMPTY:904; CHECK-NOF16-NEXT: // %bb.0:905; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ule_param_1];906; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ule_param_0];907; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;908; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;909; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;910; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;911; CHECK-NOF16-NEXT: setp.leu.f32 %p1, %r4, %r3;912; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;913; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;914; CHECK-NOF16-NEXT: setp.leu.f32 %p2, %r6, %r5;915; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;916; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;917; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;918; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;919; CHECK-NOF16-NEXT: ret;920 %r = fcmp ule <2 x half> %a, %b921 ret <2 x i1> %r922}923 924 925define <2 x i1> @test_fcmp_uno(<2 x half> %a, <2 x half> %b) #0 {926; CHECK-F16-LABEL: test_fcmp_uno(927; CHECK-F16: {928; CHECK-F16-NEXT: .reg .pred %p<3>;929; CHECK-F16-NEXT: .reg .b16 %rs<3>;930; CHECK-F16-NEXT: .reg .b32 %r<3>;931; CHECK-F16-EMPTY:932; CHECK-F16-NEXT: // %bb.0:933; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_uno_param_1];934; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_uno_param_0];935; CHECK-F16-NEXT: setp.nan.f16x2 %p1|%p2, %r1, %r2;936; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;937; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;938; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;939; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;940; CHECK-F16-NEXT: ret;941;942; CHECK-NOF16-LABEL: test_fcmp_uno(943; CHECK-NOF16: {944; CHECK-NOF16-NEXT: .reg .pred %p<3>;945; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;946; CHECK-NOF16-NEXT: .reg .b32 %r<7>;947; CHECK-NOF16-EMPTY:948; CHECK-NOF16-NEXT: // %bb.0:949; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_uno_param_1];950; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_uno_param_0];951; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;952; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;953; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;954; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;955; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r4, %r3;956; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;957; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;958; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r6, %r5;959; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;960; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;961; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;962; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;963; CHECK-NOF16-NEXT: ret;964 %r = fcmp uno <2 x half> %a, %b965 ret <2 x i1> %r966}967 968define <2 x i1> @test_fcmp_one(<2 x half> %a, <2 x half> %b) #0 {969; CHECK-F16-LABEL: test_fcmp_one(970; CHECK-F16: {971; CHECK-F16-NEXT: .reg .pred %p<3>;972; CHECK-F16-NEXT: .reg .b16 %rs<3>;973; CHECK-F16-NEXT: .reg .b32 %r<3>;974; CHECK-F16-EMPTY:975; CHECK-F16-NEXT: // %bb.0:976; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_one_param_1];977; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_one_param_0];978; CHECK-F16-NEXT: setp.ne.f16x2 %p1|%p2, %r1, %r2;979; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;980; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;981; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;982; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;983; CHECK-F16-NEXT: ret;984;985; CHECK-NOF16-LABEL: test_fcmp_one(986; CHECK-NOF16: {987; CHECK-NOF16-NEXT: .reg .pred %p<3>;988; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;989; CHECK-NOF16-NEXT: .reg .b32 %r<7>;990; CHECK-NOF16-EMPTY:991; CHECK-NOF16-NEXT: // %bb.0:992; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_one_param_1];993; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_one_param_0];994; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;995; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;996; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;997; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;998; CHECK-NOF16-NEXT: setp.ne.f32 %p1, %r4, %r3;999; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1000; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;1001; CHECK-NOF16-NEXT: setp.ne.f32 %p2, %r6, %r5;1002; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;1003; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;1004; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;1005; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;1006; CHECK-NOF16-NEXT: ret;1007 %r = fcmp one <2 x half> %a, %b1008 ret <2 x i1> %r1009}1010 1011define <2 x i1> @test_fcmp_oeq(<2 x half> %a, <2 x half> %b) #0 {1012; CHECK-F16-LABEL: test_fcmp_oeq(1013; CHECK-F16: {1014; CHECK-F16-NEXT: .reg .pred %p<3>;1015; CHECK-F16-NEXT: .reg .b16 %rs<3>;1016; CHECK-F16-NEXT: .reg .b32 %r<3>;1017; CHECK-F16-EMPTY:1018; CHECK-F16-NEXT: // %bb.0:1019; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_oeq_param_1];1020; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_oeq_param_0];1021; CHECK-F16-NEXT: setp.eq.f16x2 %p1|%p2, %r1, %r2;1022; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;1023; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;1024; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;1025; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;1026; CHECK-F16-NEXT: ret;1027;1028; CHECK-NOF16-LABEL: test_fcmp_oeq(1029; CHECK-NOF16: {1030; CHECK-NOF16-NEXT: .reg .pred %p<3>;1031; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1032; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1033; CHECK-NOF16-EMPTY:1034; CHECK-NOF16-NEXT: // %bb.0:1035; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_oeq_param_1];1036; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_oeq_param_0];1037; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;1038; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;1039; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;1040; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;1041; CHECK-NOF16-NEXT: setp.eq.f32 %p1, %r4, %r3;1042; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1043; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;1044; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %r6, %r5;1045; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;1046; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;1047; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;1048; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;1049; CHECK-NOF16-NEXT: ret;1050 %r = fcmp oeq <2 x half> %a, %b1051 ret <2 x i1> %r1052}1053 1054define <2 x i1> @test_fcmp_ogt(<2 x half> %a, <2 x half> %b) #0 {1055; CHECK-F16-LABEL: test_fcmp_ogt(1056; CHECK-F16: {1057; CHECK-F16-NEXT: .reg .pred %p<3>;1058; CHECK-F16-NEXT: .reg .b16 %rs<3>;1059; CHECK-F16-NEXT: .reg .b32 %r<3>;1060; CHECK-F16-EMPTY:1061; CHECK-F16-NEXT: // %bb.0:1062; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_ogt_param_1];1063; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_ogt_param_0];1064; CHECK-F16-NEXT: setp.gt.f16x2 %p1|%p2, %r1, %r2;1065; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;1066; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;1067; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;1068; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;1069; CHECK-F16-NEXT: ret;1070;1071; CHECK-NOF16-LABEL: test_fcmp_ogt(1072; CHECK-NOF16: {1073; CHECK-NOF16-NEXT: .reg .pred %p<3>;1074; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1075; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1076; CHECK-NOF16-EMPTY:1077; CHECK-NOF16-NEXT: // %bb.0:1078; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ogt_param_1];1079; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ogt_param_0];1080; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;1081; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;1082; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;1083; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;1084; CHECK-NOF16-NEXT: setp.gt.f32 %p1, %r4, %r3;1085; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1086; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;1087; CHECK-NOF16-NEXT: setp.gt.f32 %p2, %r6, %r5;1088; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;1089; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;1090; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;1091; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;1092; CHECK-NOF16-NEXT: ret;1093 %r = fcmp ogt <2 x half> %a, %b1094 ret <2 x i1> %r1095}1096 1097define <2 x i1> @test_fcmp_oge(<2 x half> %a, <2 x half> %b) #0 {1098; CHECK-F16-LABEL: test_fcmp_oge(1099; CHECK-F16: {1100; CHECK-F16-NEXT: .reg .pred %p<3>;1101; CHECK-F16-NEXT: .reg .b16 %rs<3>;1102; CHECK-F16-NEXT: .reg .b32 %r<3>;1103; CHECK-F16-EMPTY:1104; CHECK-F16-NEXT: // %bb.0:1105; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_oge_param_1];1106; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_oge_param_0];1107; CHECK-F16-NEXT: setp.ge.f16x2 %p1|%p2, %r1, %r2;1108; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;1109; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;1110; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;1111; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;1112; CHECK-F16-NEXT: ret;1113;1114; CHECK-NOF16-LABEL: test_fcmp_oge(1115; CHECK-NOF16: {1116; CHECK-NOF16-NEXT: .reg .pred %p<3>;1117; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1118; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1119; CHECK-NOF16-EMPTY:1120; CHECK-NOF16-NEXT: // %bb.0:1121; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_oge_param_1];1122; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_oge_param_0];1123; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;1124; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;1125; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;1126; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;1127; CHECK-NOF16-NEXT: setp.ge.f32 %p1, %r4, %r3;1128; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1129; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;1130; CHECK-NOF16-NEXT: setp.ge.f32 %p2, %r6, %r5;1131; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;1132; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;1133; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;1134; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;1135; CHECK-NOF16-NEXT: ret;1136 %r = fcmp oge <2 x half> %a, %b1137 ret <2 x i1> %r1138}1139 1140define <2 x i1> @test_fcmp_olt(<2 x half> %a, <2 x half> %b) #0 {1141; CHECK-F16-LABEL: test_fcmp_olt(1142; CHECK-F16: {1143; CHECK-F16-NEXT: .reg .pred %p<3>;1144; CHECK-F16-NEXT: .reg .b16 %rs<3>;1145; CHECK-F16-NEXT: .reg .b32 %r<3>;1146; CHECK-F16-EMPTY:1147; CHECK-F16-NEXT: // %bb.0:1148; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_olt_param_1];1149; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_olt_param_0];1150; CHECK-F16-NEXT: setp.lt.f16x2 %p1|%p2, %r1, %r2;1151; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;1152; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;1153; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;1154; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;1155; CHECK-F16-NEXT: ret;1156;1157; CHECK-NOF16-LABEL: test_fcmp_olt(1158; CHECK-NOF16: {1159; CHECK-NOF16-NEXT: .reg .pred %p<3>;1160; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1161; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1162; CHECK-NOF16-EMPTY:1163; CHECK-NOF16-NEXT: // %bb.0:1164; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_olt_param_1];1165; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_olt_param_0];1166; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;1167; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;1168; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;1169; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;1170; CHECK-NOF16-NEXT: setp.lt.f32 %p1, %r4, %r3;1171; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1172; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;1173; CHECK-NOF16-NEXT: setp.lt.f32 %p2, %r6, %r5;1174; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;1175; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;1176; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;1177; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;1178; CHECK-NOF16-NEXT: ret;1179 %r = fcmp olt <2 x half> %a, %b1180 ret <2 x i1> %r1181}1182 1183define <2 x i1> @test_fcmp_ole(<2 x half> %a, <2 x half> %b) #0 {1184; CHECK-F16-LABEL: test_fcmp_ole(1185; CHECK-F16: {1186; CHECK-F16-NEXT: .reg .pred %p<3>;1187; CHECK-F16-NEXT: .reg .b16 %rs<3>;1188; CHECK-F16-NEXT: .reg .b32 %r<3>;1189; CHECK-F16-EMPTY:1190; CHECK-F16-NEXT: // %bb.0:1191; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_ole_param_1];1192; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_ole_param_0];1193; CHECK-F16-NEXT: setp.le.f16x2 %p1|%p2, %r1, %r2;1194; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;1195; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;1196; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;1197; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;1198; CHECK-F16-NEXT: ret;1199;1200; CHECK-NOF16-LABEL: test_fcmp_ole(1201; CHECK-NOF16: {1202; CHECK-NOF16-NEXT: .reg .pred %p<3>;1203; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1204; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1205; CHECK-NOF16-EMPTY:1206; CHECK-NOF16-NEXT: // %bb.0:1207; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ole_param_1];1208; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ole_param_0];1209; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;1210; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;1211; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;1212; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;1213; CHECK-NOF16-NEXT: setp.le.f32 %p1, %r4, %r3;1214; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1215; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;1216; CHECK-NOF16-NEXT: setp.le.f32 %p2, %r6, %r5;1217; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;1218; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;1219; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;1220; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;1221; CHECK-NOF16-NEXT: ret;1222 %r = fcmp ole <2 x half> %a, %b1223 ret <2 x i1> %r1224}1225 1226define <2 x i1> @test_fcmp_ord(<2 x half> %a, <2 x half> %b) #0 {1227; CHECK-F16-LABEL: test_fcmp_ord(1228; CHECK-F16: {1229; CHECK-F16-NEXT: .reg .pred %p<3>;1230; CHECK-F16-NEXT: .reg .b16 %rs<3>;1231; CHECK-F16-NEXT: .reg .b32 %r<3>;1232; CHECK-F16-EMPTY:1233; CHECK-F16-NEXT: // %bb.0:1234; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fcmp_ord_param_1];1235; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fcmp_ord_param_0];1236; CHECK-F16-NEXT: setp.num.f16x2 %p1|%p2, %r1, %r2;1237; CHECK-F16-NEXT: selp.b16 %rs1, -1, 0, %p1;1238; CHECK-F16-NEXT: st.param.b8 [func_retval0], %rs1;1239; CHECK-F16-NEXT: selp.b16 %rs2, -1, 0, %p2;1240; CHECK-F16-NEXT: st.param.b8 [func_retval0+1], %rs2;1241; CHECK-F16-NEXT: ret;1242;1243; CHECK-NOF16-LABEL: test_fcmp_ord(1244; CHECK-NOF16: {1245; CHECK-NOF16-NEXT: .reg .pred %p<3>;1246; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1247; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1248; CHECK-NOF16-EMPTY:1249; CHECK-NOF16-NEXT: // %bb.0:1250; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ord_param_1];1251; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ord_param_0];1252; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;1253; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;1254; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;1255; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;1256; CHECK-NOF16-NEXT: setp.num.f32 %p1, %r4, %r3;1257; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;1258; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;1259; CHECK-NOF16-NEXT: setp.num.f32 %p2, %r6, %r5;1260; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;1261; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;1262; CHECK-NOF16-NEXT: selp.b16 %rs6, -1, 0, %p1;1263; CHECK-NOF16-NEXT: st.param.b8 [func_retval0+1], %rs6;1264; CHECK-NOF16-NEXT: ret;1265 %r = fcmp ord <2 x half> %a, %b1266 ret <2 x i1> %r1267}1268 1269define <2 x i32> @test_fptosi_i32(<2 x half> %a) #0 {1270; CHECK-LABEL: test_fptosi_i32(1271; CHECK: {1272; CHECK-NEXT: .reg .b16 %rs<3>;1273; CHECK-NEXT: .reg .b32 %r<4>;1274; CHECK-EMPTY:1275; CHECK-NEXT: // %bb.0:1276; CHECK-NEXT: ld.param.b32 %r1, [test_fptosi_i32_param_0];1277; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1278; CHECK-NEXT: cvt.rzi.s32.f16 %r2, %rs2;1279; CHECK-NEXT: cvt.rzi.s32.f16 %r3, %rs1;1280; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};1281; CHECK-NEXT: ret;1282 %r = fptosi <2 x half> %a to <2 x i32>1283 ret <2 x i32> %r1284}1285 1286define <2 x i64> @test_fptosi_i64(<2 x half> %a) #0 {1287; CHECK-LABEL: test_fptosi_i64(1288; CHECK: {1289; CHECK-NEXT: .reg .b16 %rs<3>;1290; CHECK-NEXT: .reg .b32 %r<2>;1291; CHECK-NEXT: .reg .b64 %rd<3>;1292; CHECK-EMPTY:1293; CHECK-NEXT: // %bb.0:1294; CHECK-NEXT: ld.param.b32 %r1, [test_fptosi_i64_param_0];1295; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1296; CHECK-NEXT: cvt.rzi.s64.f16 %rd1, %rs2;1297; CHECK-NEXT: cvt.rzi.s64.f16 %rd2, %rs1;1298; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};1299; CHECK-NEXT: ret;1300 %r = fptosi <2 x half> %a to <2 x i64>1301 ret <2 x i64> %r1302}1303 1304define <2 x i32> @test_fptoui_2xi32(<2 x half> %a) #0 {1305; CHECK-LABEL: test_fptoui_2xi32(1306; CHECK: {1307; CHECK-NEXT: .reg .b16 %rs<3>;1308; CHECK-NEXT: .reg .b32 %r<4>;1309; CHECK-EMPTY:1310; CHECK-NEXT: // %bb.0:1311; CHECK-NEXT: ld.param.b32 %r1, [test_fptoui_2xi32_param_0];1312; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1313; CHECK-NEXT: cvt.rzi.u32.f16 %r2, %rs2;1314; CHECK-NEXT: cvt.rzi.u32.f16 %r3, %rs1;1315; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};1316; CHECK-NEXT: ret;1317 %r = fptoui <2 x half> %a to <2 x i32>1318 ret <2 x i32> %r1319}1320 1321define <2 x i64> @test_fptoui_2xi64(<2 x half> %a) #0 {1322; CHECK-LABEL: test_fptoui_2xi64(1323; CHECK: {1324; CHECK-NEXT: .reg .b16 %rs<3>;1325; CHECK-NEXT: .reg .b32 %r<2>;1326; CHECK-NEXT: .reg .b64 %rd<3>;1327; CHECK-EMPTY:1328; CHECK-NEXT: // %bb.0:1329; CHECK-NEXT: ld.param.b32 %r1, [test_fptoui_2xi64_param_0];1330; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1331; CHECK-NEXT: cvt.rzi.u64.f16 %rd1, %rs2;1332; CHECK-NEXT: cvt.rzi.u64.f16 %rd2, %rs1;1333; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};1334; CHECK-NEXT: ret;1335 %r = fptoui <2 x half> %a to <2 x i64>1336 ret <2 x i64> %r1337}1338 1339define <2 x half> @test_uitofp_2xi32(<2 x i32> %a) #0 {1340; CHECK-LABEL: test_uitofp_2xi32(1341; CHECK: {1342; CHECK-NEXT: .reg .b16 %rs<3>;1343; CHECK-NEXT: .reg .b32 %r<3>;1344; CHECK-EMPTY:1345; CHECK-NEXT: // %bb.0:1346; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_param_0];1347; CHECK-NEXT: cvt.rn.f16.u32 %rs1, %r2;1348; CHECK-NEXT: cvt.rn.f16.u32 %rs2, %r1;1349; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};1350; CHECK-NEXT: ret;1351 %r = uitofp <2 x i32> %a to <2 x half>1352 ret <2 x half> %r1353}1354 1355define <2 x half> @test_uitofp_2xi64(<2 x i64> %a) #0 {1356; CHECK-LABEL: test_uitofp_2xi64(1357; CHECK: {1358; CHECK-NEXT: .reg .b16 %rs<3>;1359; CHECK-NEXT: .reg .b64 %rd<3>;1360; CHECK-EMPTY:1361; CHECK-NEXT: // %bb.0:1362; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_uitofp_2xi64_param_0];1363; CHECK-NEXT: cvt.rn.f16.u64 %rs1, %rd2;1364; CHECK-NEXT: cvt.rn.f16.u64 %rs2, %rd1;1365; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};1366; CHECK-NEXT: ret;1367 %r = uitofp <2 x i64> %a to <2 x half>1368 ret <2 x half> %r1369}1370 1371define <2 x half> @test_sitofp_2xi32(<2 x i32> %a) #0 {1372; CHECK-LABEL: test_sitofp_2xi32(1373; CHECK: {1374; CHECK-NEXT: .reg .b16 %rs<3>;1375; CHECK-NEXT: .reg .b32 %r<3>;1376; CHECK-EMPTY:1377; CHECK-NEXT: // %bb.0:1378; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sitofp_2xi32_param_0];1379; CHECK-NEXT: cvt.rn.f16.s32 %rs1, %r2;1380; CHECK-NEXT: cvt.rn.f16.s32 %rs2, %r1;1381; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};1382; CHECK-NEXT: ret;1383 %r = sitofp <2 x i32> %a to <2 x half>1384 ret <2 x half> %r1385}1386 1387define <2 x half> @test_sitofp_2xi64(<2 x i64> %a) #0 {1388; CHECK-LABEL: test_sitofp_2xi64(1389; CHECK: {1390; CHECK-NEXT: .reg .b16 %rs<3>;1391; CHECK-NEXT: .reg .b64 %rd<3>;1392; CHECK-EMPTY:1393; CHECK-NEXT: // %bb.0:1394; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_sitofp_2xi64_param_0];1395; CHECK-NEXT: cvt.rn.f16.s64 %rs1, %rd2;1396; CHECK-NEXT: cvt.rn.f16.s64 %rs2, %rd1;1397; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};1398; CHECK-NEXT: ret;1399 %r = sitofp <2 x i64> %a to <2 x half>1400 ret <2 x half> %r1401}1402 1403 1404define <2 x half> @test_uitofp_2xi32_fadd(<2 x i32> %a, <2 x half> %b) #0 {1405; CHECK-F16-LABEL: test_uitofp_2xi32_fadd(1406; CHECK-F16: {1407; CHECK-F16-NEXT: .reg .b16 %rs<3>;1408; CHECK-F16-NEXT: .reg .b32 %r<6>;1409; CHECK-F16-EMPTY:1410; CHECK-F16-NEXT: // %bb.0:1411; CHECK-F16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_fadd_param_0];1412; CHECK-F16-NEXT: ld.param.b32 %r3, [test_uitofp_2xi32_fadd_param_1];1413; CHECK-F16-NEXT: cvt.rn.f16.u32 %rs1, %r2;1414; CHECK-F16-NEXT: cvt.rn.f16.u32 %rs2, %r1;1415; CHECK-F16-NEXT: mov.b32 %r4, {%rs2, %rs1};1416; CHECK-F16-NEXT: add.rn.f16x2 %r5, %r3, %r4;1417; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r5;1418; CHECK-F16-NEXT: ret;1419;1420; CHECK-NOF16-LABEL: test_uitofp_2xi32_fadd(1421; CHECK-NOF16: {1422; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1423; CHECK-NOF16-NEXT: .reg .b32 %r<11>;1424; CHECK-NOF16-EMPTY:1425; CHECK-NOF16-NEXT: // %bb.0:1426; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_fadd_param_0];1427; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_uitofp_2xi32_fadd_param_1];1428; CHECK-NOF16-NEXT: cvt.rn.f16.u32 %rs1, %r1;1429; CHECK-NOF16-NEXT: cvt.rn.f16.u32 %rs2, %r2;1430; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;1431; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r3;1432; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;1433; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r5, %r4;1434; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r6;1435; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;1436; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;1437; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r8, %r7;1438; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r9;1439; CHECK-NOF16-NEXT: mov.b32 %r10, {%rs6, %rs5};1440; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r10;1441; CHECK-NOF16-NEXT: ret;1442 %c = uitofp <2 x i32> %a to <2 x half>1443 %r = fadd <2 x half> %b, %c1444 ret <2 x half> %r1445}1446 1447define <2 x half> @test_sitofp_2xi32_fadd(<2 x i32> %a, <2 x half> %b) #0 {1448; CHECK-F16-LABEL: test_sitofp_2xi32_fadd(1449; CHECK-F16: {1450; CHECK-F16-NEXT: .reg .b16 %rs<3>;1451; CHECK-F16-NEXT: .reg .b32 %r<6>;1452; CHECK-F16-EMPTY:1453; CHECK-F16-NEXT: // %bb.0:1454; CHECK-F16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sitofp_2xi32_fadd_param_0];1455; CHECK-F16-NEXT: ld.param.b32 %r3, [test_sitofp_2xi32_fadd_param_1];1456; CHECK-F16-NEXT: cvt.rn.f16.s32 %rs1, %r2;1457; CHECK-F16-NEXT: cvt.rn.f16.s32 %rs2, %r1;1458; CHECK-F16-NEXT: mov.b32 %r4, {%rs2, %rs1};1459; CHECK-F16-NEXT: add.rn.f16x2 %r5, %r3, %r4;1460; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r5;1461; CHECK-F16-NEXT: ret;1462;1463; CHECK-NOF16-LABEL: test_sitofp_2xi32_fadd(1464; CHECK-NOF16: {1465; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;1466; CHECK-NOF16-NEXT: .reg .b32 %r<11>;1467; CHECK-NOF16-EMPTY:1468; CHECK-NOF16-NEXT: // %bb.0:1469; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sitofp_2xi32_fadd_param_0];1470; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_sitofp_2xi32_fadd_param_1];1471; CHECK-NOF16-NEXT: cvt.rn.f16.s32 %rs1, %r1;1472; CHECK-NOF16-NEXT: cvt.rn.f16.s32 %rs2, %r2;1473; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;1474; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r3;1475; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;1476; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r5, %r4;1477; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r6;1478; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;1479; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;1480; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r8, %r7;1481; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r9;1482; CHECK-NOF16-NEXT: mov.b32 %r10, {%rs6, %rs5};1483; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r10;1484; CHECK-NOF16-NEXT: ret;1485 %c = sitofp <2 x i32> %a to <2 x half>1486 %r = fadd <2 x half> %b, %c1487 ret <2 x half> %r1488}1489 1490define <2 x half> @test_fptrunc_2xfloat(<2 x float> %a) #0 {1491; CHECK-LABEL: test_fptrunc_2xfloat(1492; CHECK: {1493; CHECK-NEXT: .reg .b16 %rs<3>;1494; CHECK-NEXT: .reg .b32 %r<4>;1495; CHECK-EMPTY:1496; CHECK-NEXT: // %bb.0:1497; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptrunc_2xfloat_param_0];1498; CHECK-NEXT: cvt.rn.f16.f32 %rs1, %r2;1499; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r1;1500; CHECK-NEXT: mov.b32 %r3, {%rs2, %rs1};1501; CHECK-NEXT: st.param.b32 [func_retval0], %r3;1502; CHECK-NEXT: ret;1503 %r = fptrunc <2 x float> %a to <2 x half>1504 ret <2 x half> %r1505}1506 1507define <2 x half> @test_fptrunc_2xdouble(<2 x double> %a) #0 {1508; CHECK-LABEL: test_fptrunc_2xdouble(1509; CHECK: {1510; CHECK-NEXT: .reg .b16 %rs<3>;1511; CHECK-NEXT: .reg .b32 %r<2>;1512; CHECK-NEXT: .reg .b64 %rd<3>;1513; CHECK-EMPTY:1514; CHECK-NEXT: // %bb.0:1515; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fptrunc_2xdouble_param_0];1516; CHECK-NEXT: cvt.rn.f16.f64 %rs1, %rd2;1517; CHECK-NEXT: cvt.rn.f16.f64 %rs2, %rd1;1518; CHECK-NEXT: mov.b32 %r1, {%rs2, %rs1};1519; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1520; CHECK-NEXT: ret;1521 %r = fptrunc <2 x double> %a to <2 x half>1522 ret <2 x half> %r1523}1524 1525define <2 x float> @test_fpext_2xfloat(<2 x half> %a) #0 {1526; CHECK-LABEL: test_fpext_2xfloat(1527; CHECK: {1528; CHECK-NEXT: .reg .b16 %rs<3>;1529; CHECK-NEXT: .reg .b32 %r<4>;1530; CHECK-EMPTY:1531; CHECK-NEXT: // %bb.0:1532; CHECK-NEXT: ld.param.b32 %r1, [test_fpext_2xfloat_param_0];1533; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1534; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;1535; CHECK-NEXT: cvt.f32.f16 %r3, %rs1;1536; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};1537; CHECK-NEXT: ret;1538 %r = fpext <2 x half> %a to <2 x float>1539 ret <2 x float> %r1540}1541 1542define <2 x double> @test_fpext_2xdouble(<2 x half> %a) #0 {1543; CHECK-LABEL: test_fpext_2xdouble(1544; CHECK: {1545; CHECK-NEXT: .reg .b16 %rs<3>;1546; CHECK-NEXT: .reg .b32 %r<2>;1547; CHECK-NEXT: .reg .b64 %rd<3>;1548; CHECK-EMPTY:1549; CHECK-NEXT: // %bb.0:1550; CHECK-NEXT: ld.param.b32 %r1, [test_fpext_2xdouble_param_0];1551; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1552; CHECK-NEXT: cvt.f64.f16 %rd1, %rs2;1553; CHECK-NEXT: cvt.f64.f16 %rd2, %rs1;1554; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};1555; CHECK-NEXT: ret;1556 %r = fpext <2 x half> %a to <2 x double>1557 ret <2 x double> %r1558}1559 1560 1561define <2 x i16> @test_bitcast_2xhalf_to_2xi16(<2 x half> %a) #0 {1562; CHECK-LABEL: test_bitcast_2xhalf_to_2xi16(1563; CHECK: {1564; CHECK-NEXT: .reg .b32 %r<2>;1565; CHECK-EMPTY:1566; CHECK-NEXT: // %bb.0:1567; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_2xhalf_to_2xi16_param_0];1568; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1569; CHECK-NEXT: ret;1570 %r = bitcast <2 x half> %a to <2 x i16>1571 ret <2 x i16> %r1572}1573 1574define <2 x half> @test_bitcast_2xi16_to_2xhalf(<2 x i16> %a) #0 {1575; CHECK-LABEL: test_bitcast_2xi16_to_2xhalf(1576; CHECK: {1577; CHECK-NEXT: .reg .b32 %r<2>;1578; CHECK-EMPTY:1579; CHECK-NEXT: // %bb.0:1580; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_2xi16_to_2xhalf_param_0];1581; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1582; CHECK-NEXT: ret;1583 %r = bitcast <2 x i16> %a to <2 x half>1584 ret <2 x half> %r1585}1586 1587define <2 x half> @test_bitcast_float_to_2xhalf(float %a) #0 {1588; CHECK-LABEL: test_bitcast_float_to_2xhalf(1589; CHECK: {1590; CHECK-NEXT: .reg .b32 %r<2>;1591; CHECK-EMPTY:1592; CHECK-NEXT: // %bb.0:1593; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_float_to_2xhalf_param_0];1594; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1595; CHECK-NEXT: ret;1596 %r = bitcast float %a to <2 x half>1597 ret <2 x half> %r1598}1599 1600define float @test_bitcast_2xhalf_to_float(<2 x half> %a) #0 {1601; CHECK-LABEL: test_bitcast_2xhalf_to_float(1602; CHECK: {1603; CHECK-NEXT: .reg .b32 %r<2>;1604; CHECK-EMPTY:1605; CHECK-NEXT: // %bb.0:1606; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_2xhalf_to_float_param_0];1607; CHECK-NEXT: st.param.b32 [func_retval0], %r1;1608; CHECK-NEXT: ret;1609 %r = bitcast <2 x half> %a to float1610 ret float %r1611}1612 1613declare <2 x half> @llvm.sqrt.f16(<2 x half> %a) #01614declare <2 x half> @llvm.powi.f16.i32(<2 x half> %a, <2 x i32> %b) #01615declare <2 x half> @llvm.sin.f16(<2 x half> %a) #01616declare <2 x half> @llvm.cos.f16(<2 x half> %a) #01617declare <2 x half> @llvm.pow.f16(<2 x half> %a, <2 x half> %b) #01618declare <2 x half> @llvm.exp.f16(<2 x half> %a) #01619declare <2 x half> @llvm.exp2.f16(<2 x half> %a) #01620declare <2 x half> @llvm.log.f16(<2 x half> %a) #01621declare <2 x half> @llvm.log10.f16(<2 x half> %a) #01622declare <2 x half> @llvm.log2.f16(<2 x half> %a) #01623declare <2 x half> @llvm.fma.f16(<2 x half> %a, <2 x half> %b, <2 x half> %c) #01624declare <2 x half> @llvm.fabs.f16(<2 x half> %a) #01625declare <2 x half> @llvm.minnum.f16(<2 x half> %a, <2 x half> %b) #01626declare <2 x half> @llvm.maxnum.f16(<2 x half> %a, <2 x half> %b) #01627declare <2 x half> @llvm.copysign.f16(<2 x half> %a, <2 x half> %b) #01628declare <2 x half> @llvm.floor.f16(<2 x half> %a) #01629declare <2 x half> @llvm.ceil.f16(<2 x half> %a) #01630declare <2 x half> @llvm.trunc.f16(<2 x half> %a) #01631declare <2 x half> @llvm.rint.f16(<2 x half> %a) #01632declare <2 x half> @llvm.nearbyint.f16(<2 x half> %a) #01633declare <2 x half> @llvm.round.f16(<2 x half> %a) #01634declare <2 x half> @llvm.roundeven.f16(<2 x half> %a) #01635declare <2 x half> @llvm.fmuladd.f16(<2 x half> %a, <2 x half> %b, <2 x half> %c) #01636 1637define <2 x half> @test_sqrt(<2 x half> %a) #0 {1638; CHECK-LABEL: test_sqrt(1639; CHECK: {1640; CHECK-NEXT: .reg .b16 %rs<5>;1641; CHECK-NEXT: .reg .b32 %r<7>;1642; CHECK-EMPTY:1643; CHECK-NEXT: // %bb.0:1644; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_param_0];1645; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1646; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;1647; CHECK-NEXT: sqrt.rn.f32 %r3, %r2;1648; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;1649; CHECK-NEXT: cvt.f32.f16 %r4, %rs1;1650; CHECK-NEXT: sqrt.rn.f32 %r5, %r4;1651; CHECK-NEXT: cvt.rn.f16.f32 %rs4, %r5;1652; CHECK-NEXT: mov.b32 %r6, {%rs4, %rs3};1653; CHECK-NEXT: st.param.b32 [func_retval0], %r6;1654; CHECK-NEXT: ret;1655 %r = call <2 x half> @llvm.sqrt.f16(<2 x half> %a)1656 ret <2 x half> %r1657}1658 1659;;; Can't do this yet: requires libcall.1660; XCHECK-LABEL: test_powi(1661;define <2 x half> @test_powi(<2 x half> %a, <2 x i32> %b) #0 {1662; %r = call <2 x half> @llvm.powi.f16.i32(<2 x half> %a, <2 x i32> %b)1663; ret <2 x half> %r1664;}1665 1666define <2 x half> @test_sin(<2 x half> %a) #0 {1667; CHECK-LABEL: test_sin(1668; CHECK: {1669; CHECK-NEXT: .reg .b16 %rs<5>;1670; CHECK-NEXT: .reg .b32 %r<7>;1671; CHECK-EMPTY:1672; CHECK-NEXT: // %bb.0:1673; CHECK-NEXT: ld.param.b32 %r1, [test_sin_param_0];1674; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1675; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;1676; CHECK-NEXT: sin.approx.f32 %r3, %r2;1677; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;1678; CHECK-NEXT: cvt.f32.f16 %r4, %rs1;1679; CHECK-NEXT: sin.approx.f32 %r5, %r4;1680; CHECK-NEXT: cvt.rn.f16.f32 %rs4, %r5;1681; CHECK-NEXT: mov.b32 %r6, {%rs4, %rs3};1682; CHECK-NEXT: st.param.b32 [func_retval0], %r6;1683; CHECK-NEXT: ret;1684 %r = call afn <2 x half> @llvm.sin.f16(<2 x half> %a)1685 ret <2 x half> %r1686}1687 1688define <2 x half> @test_cos(<2 x half> %a) #0 {1689; CHECK-LABEL: test_cos(1690; CHECK: {1691; CHECK-NEXT: .reg .b16 %rs<5>;1692; CHECK-NEXT: .reg .b32 %r<7>;1693; CHECK-EMPTY:1694; CHECK-NEXT: // %bb.0:1695; CHECK-NEXT: ld.param.b32 %r1, [test_cos_param_0];1696; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;1697; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;1698; CHECK-NEXT: cos.approx.f32 %r3, %r2;1699; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;1700; CHECK-NEXT: cvt.f32.f16 %r4, %rs1;1701; CHECK-NEXT: cos.approx.f32 %r5, %r4;1702; CHECK-NEXT: cvt.rn.f16.f32 %rs4, %r5;1703; CHECK-NEXT: mov.b32 %r6, {%rs4, %rs3};1704; CHECK-NEXT: st.param.b32 [func_retval0], %r6;1705; CHECK-NEXT: ret;1706 %r = call afn <2 x half> @llvm.cos.f16(<2 x half> %a)1707 ret <2 x half> %r1708}1709 1710;;; Can't do this yet: requires libcall.1711; XCHECK-LABEL: test_pow(1712;define <2 x half> @test_pow(<2 x half> %a, <2 x half> %b) #0 {1713; %r = call <2 x half> @llvm.pow.f16(<2 x half> %a, <2 x half> %b)1714; ret <2 x half> %r1715;}1716 1717;;; Can't do this yet: requires libcall.1718; XCHECK-LABEL: test_exp(1719;define <2 x half> @test_exp(<2 x half> %a) #0 {1720; %r = call <2 x half> @llvm.exp.f16(<2 x half> %a)1721; ret <2 x half> %r1722;}1723 1724;;; Can't do this yet: requires libcall.1725; XCHECK-LABEL: test_exp2(1726;define <2 x half> @test_exp2(<2 x half> %a) #0 {1727; %r = call <2 x half> @llvm.exp2.f16(<2 x half> %a)1728; ret <2 x half> %r1729;}1730 1731;;; Can't do this yet: requires libcall.1732; XCHECK-LABEL: test_log(1733;define <2 x half> @test_log(<2 x half> %a) #0 {1734; %r = call <2 x half> @llvm.log.f16(<2 x half> %a)1735; ret <2 x half> %r1736;}1737 1738;;; Can't do this yet: requires libcall.1739; XCHECK-LABEL: test_log10(1740;define <2 x half> @test_log10(<2 x half> %a) #0 {1741; %r = call <2 x half> @llvm.log10.f16(<2 x half> %a)1742; ret <2 x half> %r1743;}1744 1745;;; Can't do this yet: requires libcall.1746; XCHECK-LABEL: test_log2(1747;define <2 x half> @test_log2(<2 x half> %a) #0 {1748; %r = call <2 x half> @llvm.log2.f16(<2 x half> %a)1749; ret <2 x half> %r1750;}1751 1752 1753define <2 x half> @test_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 {1754; CHECK-F16-LABEL: test_fma(1755; CHECK-F16: {1756; CHECK-F16-NEXT: .reg .b32 %r<5>;1757; CHECK-F16-EMPTY:1758; CHECK-F16-NEXT: // %bb.0:1759; CHECK-F16-NEXT: ld.param.b32 %r3, [test_fma_param_2];1760; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fma_param_1];1761; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fma_param_0];1762; CHECK-F16-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;1763; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r4;1764; CHECK-F16-NEXT: ret;1765;1766; CHECK-NOF16-LABEL: test_fma(1767; CHECK-NOF16: {1768; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;1769; CHECK-NOF16-NEXT: .reg .b32 %r<13>;1770; CHECK-NOF16-EMPTY:1771; CHECK-NOF16-NEXT: // %bb.0:1772; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_fma_param_2];1773; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fma_param_1];1774; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fma_param_0];1775; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r3;1776; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;1777; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r2;1778; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;1779; CHECK-NOF16-NEXT: mov.b32 {%rs5, %rs6}, %r1;1780; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs6;1781; CHECK-NOF16-NEXT: fma.rn.f32 %r7, %r6, %r5, %r4;1782; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs7, %r7;1783; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs1;1784; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs3;1785; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs5;1786; CHECK-NOF16-NEXT: fma.rn.f32 %r11, %r10, %r9, %r8;1787; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs8, %r11;1788; CHECK-NOF16-NEXT: mov.b32 %r12, {%rs8, %rs7};1789; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r12;1790; CHECK-NOF16-NEXT: ret;1791 %r = call <2 x half> @llvm.fma.f16(<2 x half> %a, <2 x half> %b, <2 x half> %c)1792 ret <2 x half> %r1793}1794 1795define <2 x half> @test_fabs(<2 x half> %a) #0 {1796; CHECK-F16-LABEL: test_fabs(1797; CHECK-F16: {1798; CHECK-F16-NEXT: .reg .b32 %r<3>;1799; CHECK-F16-EMPTY:1800; CHECK-F16-NEXT: // %bb.0:1801; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fabs_param_0];1802; CHECK-F16-NEXT: and.b32 %r2, %r1, 2147450879;1803; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r2;1804; CHECK-F16-NEXT: ret;1805;1806; CHECK-NOF16-LABEL: test_fabs(1807; CHECK-NOF16: {1808; CHECK-NOF16-NEXT: .reg .b16 %rs<5>;1809; CHECK-NOF16-NEXT: .reg .b32 %r<7>;1810; CHECK-NOF16-EMPTY:1811; CHECK-NOF16-NEXT: // %bb.0:1812; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fabs_param_0];1813; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;1814; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;1815; CHECK-NOF16-NEXT: abs.f32 %r3, %r2;1816; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;1817; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs1;1818; CHECK-NOF16-NEXT: abs.f32 %r5, %r4;1819; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs4, %r5;1820; CHECK-NOF16-NEXT: mov.b32 %r6, {%rs4, %rs3};1821; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r6;1822; CHECK-NOF16-NEXT: ret;1823 %r = call <2 x half> @llvm.fabs.f16(<2 x half> %a)1824 ret <2 x half> %r1825}1826 1827define <2 x half> @test_minnum(<2 x half> %a, <2 x half> %b) #0 {1828; CHECK-LABEL: test_minnum(1829; CHECK: {1830; CHECK-NEXT: .reg .b16 %rs<7>;1831; CHECK-NEXT: .reg .b32 %r<10>;1832; CHECK-EMPTY:1833; CHECK-NEXT: // %bb.0:1834; CHECK-NEXT: ld.param.b32 %r2, [test_minnum_param_1];1835; CHECK-NEXT: ld.param.b32 %r1, [test_minnum_param_0];1836; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;1837; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;1838; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;1839; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;1840; CHECK-NEXT: min.f32 %r5, %r4, %r3;1841; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r5;1842; CHECK-NEXT: cvt.f32.f16 %r6, %rs1;1843; CHECK-NEXT: cvt.f32.f16 %r7, %rs3;1844; CHECK-NEXT: min.f32 %r8, %r7, %r6;1845; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r8;1846; CHECK-NEXT: mov.b32 %r9, {%rs6, %rs5};1847; CHECK-NEXT: st.param.b32 [func_retval0], %r9;1848; CHECK-NEXT: ret;1849 %r = call <2 x half> @llvm.minnum.f16(<2 x half> %a, <2 x half> %b)1850 ret <2 x half> %r1851}1852 1853define <2 x half> @test_maxnum(<2 x half> %a, <2 x half> %b) #0 {1854; CHECK-LABEL: test_maxnum(1855; CHECK: {1856; CHECK-NEXT: .reg .b16 %rs<7>;1857; CHECK-NEXT: .reg .b32 %r<10>;1858; CHECK-EMPTY:1859; CHECK-NEXT: // %bb.0:1860; CHECK-NEXT: ld.param.b32 %r2, [test_maxnum_param_1];1861; CHECK-NEXT: ld.param.b32 %r1, [test_maxnum_param_0];1862; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;1863; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;1864; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;1865; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;1866; CHECK-NEXT: max.f32 %r5, %r4, %r3;1867; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r5;1868; CHECK-NEXT: cvt.f32.f16 %r6, %rs1;1869; CHECK-NEXT: cvt.f32.f16 %r7, %rs3;1870; CHECK-NEXT: max.f32 %r8, %r7, %r6;1871; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r8;1872; CHECK-NEXT: mov.b32 %r9, {%rs6, %rs5};1873; CHECK-NEXT: st.param.b32 [func_retval0], %r9;1874; CHECK-NEXT: ret;1875 %r = call <2 x half> @llvm.maxnum.f16(<2 x half> %a, <2 x half> %b)1876 ret <2 x half> %r1877}1878 1879define <2 x half> @test_copysign(<2 x half> %a, <2 x half> %b) #0 {1880; CHECK-F16-LABEL: test_copysign(1881; CHECK-F16: {1882; CHECK-F16-NEXT: .reg .b32 %r<6>;1883; CHECK-F16-EMPTY:1884; CHECK-F16-NEXT: // %bb.0:1885; CHECK-F16-NEXT: ld.param.b32 %r2, [test_copysign_param_1];1886; CHECK-F16-NEXT: ld.param.b32 %r1, [test_copysign_param_0];1887; CHECK-F16-NEXT: and.b32 %r3, %r2, -2147450880;1888; CHECK-F16-NEXT: and.b32 %r4, %r1, 2147450879;1889; CHECK-F16-NEXT: or.b32 %r5, %r4, %r3;1890; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r5;1891; CHECK-F16-NEXT: ret;1892;1893; CHECK-NOF16-LABEL: test_copysign(1894; CHECK-NOF16: {1895; CHECK-NOF16-NEXT: .reg .b16 %rs<11>;1896; CHECK-NOF16-NEXT: .reg .b32 %r<3>;1897; CHECK-NOF16-EMPTY:1898; CHECK-NOF16-NEXT: // %bb.0:1899; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_copysign_param_1];1900; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_param_0];1901; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;1902; CHECK-NOF16-NEXT: and.b16 %rs3, %rs2, -32768;1903; CHECK-NOF16-NEXT: mov.b32 {%rs4, %rs5}, %r1;1904; CHECK-NOF16-NEXT: and.b16 %rs6, %rs5, 32767;1905; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs3;1906; CHECK-NOF16-NEXT: and.b16 %rs8, %rs1, -32768;1907; CHECK-NOF16-NEXT: and.b16 %rs9, %rs4, 32767;1908; CHECK-NOF16-NEXT: or.b16 %rs10, %rs9, %rs8;1909; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs7};1910; CHECK-NOF16-NEXT: ret;1911 %r = call <2 x half> @llvm.copysign.f16(<2 x half> %a, <2 x half> %b)1912 ret <2 x half> %r1913}1914 1915define <2 x half> @test_copysign_f32(<2 x half> %a, <2 x float> %b) #0 {1916; CHECK-F16-LABEL: test_copysign_f32(1917; CHECK-F16: {1918; CHECK-F16-NEXT: .reg .b16 %rs<3>;1919; CHECK-F16-NEXT: .reg .b32 %r<8>;1920; CHECK-F16-EMPTY:1921; CHECK-F16-NEXT: // %bb.0:1922; CHECK-F16-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_copysign_f32_param_1];1923; CHECK-F16-NEXT: ld.param.b32 %r1, [test_copysign_f32_param_0];1924; CHECK-F16-NEXT: cvt.rn.f16.f32 %rs1, %r3;1925; CHECK-F16-NEXT: cvt.rn.f16.f32 %rs2, %r2;1926; CHECK-F16-NEXT: mov.b32 %r4, {%rs2, %rs1};1927; CHECK-F16-NEXT: and.b32 %r5, %r4, -2147450880;1928; CHECK-F16-NEXT: and.b32 %r6, %r1, 2147450879;1929; CHECK-F16-NEXT: or.b32 %r7, %r6, %r5;1930; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r7;1931; CHECK-F16-NEXT: ret;1932;1933; CHECK-NOF16-LABEL: test_copysign_f32(1934; CHECK-NOF16: {1935; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;1936; CHECK-NOF16-NEXT: .reg .b32 %r<6>;1937; CHECK-NOF16-EMPTY:1938; CHECK-NOF16-NEXT: // %bb.0:1939; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_copysign_f32_param_1];1940; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_f32_param_0];1941; CHECK-NOF16-NEXT: and.b32 %r4, %r3, -2147483648;1942; CHECK-NOF16-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r4; }1943; CHECK-NOF16-NEXT: mov.b32 {%rs2, %rs3}, %r1;1944; CHECK-NOF16-NEXT: and.b16 %rs4, %rs3, 32767;1945; CHECK-NOF16-NEXT: or.b16 %rs5, %rs4, %rs1;1946; CHECK-NOF16-NEXT: and.b32 %r5, %r2, -2147483648;1947; CHECK-NOF16-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs6}, %r5; }1948; CHECK-NOF16-NEXT: and.b16 %rs7, %rs2, 32767;1949; CHECK-NOF16-NEXT: or.b16 %rs8, %rs7, %rs6;1950; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs8, %rs5};1951; CHECK-NOF16-NEXT: ret;1952 %tb = fptrunc <2 x float> %b to <2 x half>1953 %r = call <2 x half> @llvm.copysign.f16(<2 x half> %a, <2 x half> %tb)1954 ret <2 x half> %r1955}1956 1957define <2 x half> @test_copysign_f64(<2 x half> %a, <2 x double> %b) #0 {1958; CHECK-F16-LABEL: test_copysign_f64(1959; CHECK-F16: {1960; CHECK-F16-NEXT: .reg .b16 %rs<3>;1961; CHECK-F16-NEXT: .reg .b32 %r<6>;1962; CHECK-F16-NEXT: .reg .b64 %rd<3>;1963; CHECK-F16-EMPTY:1964; CHECK-F16-NEXT: // %bb.0:1965; CHECK-F16-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_copysign_f64_param_1];1966; CHECK-F16-NEXT: ld.param.b32 %r1, [test_copysign_f64_param_0];1967; CHECK-F16-NEXT: cvt.rn.f16.f64 %rs1, %rd2;1968; CHECK-F16-NEXT: cvt.rn.f16.f64 %rs2, %rd1;1969; CHECK-F16-NEXT: mov.b32 %r2, {%rs2, %rs1};1970; CHECK-F16-NEXT: and.b32 %r3, %r2, -2147450880;1971; CHECK-F16-NEXT: and.b32 %r4, %r1, 2147450879;1972; CHECK-F16-NEXT: or.b32 %r5, %r4, %r3;1973; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r5;1974; CHECK-F16-NEXT: ret;1975;1976; CHECK-NOF16-LABEL: test_copysign_f64(1977; CHECK-NOF16: {1978; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;1979; CHECK-NOF16-NEXT: .reg .b32 %r<2>;1980; CHECK-NOF16-NEXT: .reg .b64 %rd<7>;1981; CHECK-NOF16-EMPTY:1982; CHECK-NOF16-NEXT: // %bb.0:1983; CHECK-NOF16-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_copysign_f64_param_1];1984; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_f64_param_0];1985; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;1986; CHECK-NOF16-NEXT: and.b16 %rs3, %rs2, 32767;1987; CHECK-NOF16-NEXT: and.b64 %rd3, %rd2, -9223372036854775808;1988; CHECK-NOF16-NEXT: shr.u64 %rd4, %rd3, 48;1989; CHECK-NOF16-NEXT: cvt.u16.u64 %rs4, %rd4;1990; CHECK-NOF16-NEXT: or.b16 %rs5, %rs3, %rs4;1991; CHECK-NOF16-NEXT: and.b16 %rs6, %rs1, 32767;1992; CHECK-NOF16-NEXT: and.b64 %rd5, %rd1, -9223372036854775808;1993; CHECK-NOF16-NEXT: shr.u64 %rd6, %rd5, 48;1994; CHECK-NOF16-NEXT: cvt.u16.u64 %rs7, %rd6;1995; CHECK-NOF16-NEXT: or.b16 %rs8, %rs6, %rs7;1996; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs8, %rs5};1997; CHECK-NOF16-NEXT: ret;1998 %tb = fptrunc <2 x double> %b to <2 x half>1999 %r = call <2 x half> @llvm.copysign.f16(<2 x half> %a, <2 x half> %tb)2000 ret <2 x half> %r2001}2002 2003define <2 x float> @test_copysign_extended(<2 x half> %a, <2 x half> %b) #0 {2004; CHECK-F16-LABEL: test_copysign_extended(2005; CHECK-F16: {2006; CHECK-F16-NEXT: .reg .b16 %rs<3>;2007; CHECK-F16-NEXT: .reg .b32 %r<8>;2008; CHECK-F16-EMPTY:2009; CHECK-F16-NEXT: // %bb.0:2010; CHECK-F16-NEXT: ld.param.b32 %r2, [test_copysign_extended_param_1];2011; CHECK-F16-NEXT: ld.param.b32 %r1, [test_copysign_extended_param_0];2012; CHECK-F16-NEXT: and.b32 %r3, %r2, -2147450880;2013; CHECK-F16-NEXT: and.b32 %r4, %r1, 2147450879;2014; CHECK-F16-NEXT: or.b32 %r5, %r4, %r3;2015; CHECK-F16-NEXT: mov.b32 {%rs1, %rs2}, %r5;2016; CHECK-F16-NEXT: cvt.f32.f16 %r6, %rs2;2017; CHECK-F16-NEXT: cvt.f32.f16 %r7, %rs1;2018; CHECK-F16-NEXT: st.param.v2.b32 [func_retval0], {%r7, %r6};2019; CHECK-F16-NEXT: ret;2020;2021; CHECK-NOF16-LABEL: test_copysign_extended(2022; CHECK-NOF16: {2023; CHECK-NOF16-NEXT: .reg .b16 %rs<11>;2024; CHECK-NOF16-NEXT: .reg .b32 %r<5>;2025; CHECK-NOF16-EMPTY:2026; CHECK-NOF16-NEXT: // %bb.0:2027; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_copysign_extended_param_1];2028; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_extended_param_0];2029; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;2030; CHECK-NOF16-NEXT: and.b16 %rs3, %rs1, -32768;2031; CHECK-NOF16-NEXT: mov.b32 {%rs4, %rs5}, %r1;2032; CHECK-NOF16-NEXT: and.b16 %rs6, %rs4, 32767;2033; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs3;2034; CHECK-NOF16-NEXT: and.b16 %rs8, %rs2, -32768;2035; CHECK-NOF16-NEXT: and.b16 %rs9, %rs5, 32767;2036; CHECK-NOF16-NEXT: or.b16 %rs10, %rs9, %rs8;2037; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs10;2038; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs7;2039; CHECK-NOF16-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2040; CHECK-NOF16-NEXT: ret;2041 %r = call <2 x half> @llvm.copysign.f16(<2 x half> %a, <2 x half> %b)2042 %xr = fpext <2 x half> %r to <2 x float>2043 ret <2 x float> %xr2044}2045 2046define <2 x half> @test_floor(<2 x half> %a) #0 {2047; CHECK-LABEL: test_floor(2048; CHECK: {2049; CHECK-NEXT: .reg .b16 %rs<5>;2050; CHECK-NEXT: .reg .b32 %r<2>;2051; CHECK-EMPTY:2052; CHECK-NEXT: // %bb.0:2053; CHECK-NEXT: ld.param.b32 %r1, [test_floor_param_0];2054; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2055; CHECK-NEXT: cvt.rmi.f16.f16 %rs3, %rs2;2056; CHECK-NEXT: cvt.rmi.f16.f16 %rs4, %rs1;2057; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2058; CHECK-NEXT: ret;2059 %r = call <2 x half> @llvm.floor.f16(<2 x half> %a)2060 ret <2 x half> %r2061}2062 2063define <2 x half> @test_ceil(<2 x half> %a) #0 {2064; CHECK-LABEL: test_ceil(2065; CHECK: {2066; CHECK-NEXT: .reg .b16 %rs<5>;2067; CHECK-NEXT: .reg .b32 %r<2>;2068; CHECK-EMPTY:2069; CHECK-NEXT: // %bb.0:2070; CHECK-NEXT: ld.param.b32 %r1, [test_ceil_param_0];2071; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2072; CHECK-NEXT: cvt.rpi.f16.f16 %rs3, %rs2;2073; CHECK-NEXT: cvt.rpi.f16.f16 %rs4, %rs1;2074; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2075; CHECK-NEXT: ret;2076 %r = call <2 x half> @llvm.ceil.f16(<2 x half> %a)2077 ret <2 x half> %r2078}2079 2080define <2 x half> @test_trunc(<2 x half> %a) #0 {2081; CHECK-LABEL: test_trunc(2082; CHECK: {2083; CHECK-NEXT: .reg .b16 %rs<5>;2084; CHECK-NEXT: .reg .b32 %r<2>;2085; CHECK-EMPTY:2086; CHECK-NEXT: // %bb.0:2087; CHECK-NEXT: ld.param.b32 %r1, [test_trunc_param_0];2088; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2089; CHECK-NEXT: cvt.rzi.f16.f16 %rs3, %rs2;2090; CHECK-NEXT: cvt.rzi.f16.f16 %rs4, %rs1;2091; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2092; CHECK-NEXT: ret;2093 %r = call <2 x half> @llvm.trunc.f16(<2 x half> %a)2094 ret <2 x half> %r2095}2096 2097define <2 x half> @test_rint(<2 x half> %a) #0 {2098; CHECK-LABEL: test_rint(2099; CHECK: {2100; CHECK-NEXT: .reg .b16 %rs<5>;2101; CHECK-NEXT: .reg .b32 %r<2>;2102; CHECK-EMPTY:2103; CHECK-NEXT: // %bb.0:2104; CHECK-NEXT: ld.param.b32 %r1, [test_rint_param_0];2105; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2106; CHECK-NEXT: cvt.rni.f16.f16 %rs3, %rs2;2107; CHECK-NEXT: cvt.rni.f16.f16 %rs4, %rs1;2108; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2109; CHECK-NEXT: ret;2110 %r = call <2 x half> @llvm.rint.f16(<2 x half> %a)2111 ret <2 x half> %r2112}2113 2114define <2 x half> @test_nearbyint(<2 x half> %a) #0 {2115; CHECK-LABEL: test_nearbyint(2116; CHECK: {2117; CHECK-NEXT: .reg .b16 %rs<5>;2118; CHECK-NEXT: .reg .b32 %r<2>;2119; CHECK-EMPTY:2120; CHECK-NEXT: // %bb.0:2121; CHECK-NEXT: ld.param.b32 %r1, [test_nearbyint_param_0];2122; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2123; CHECK-NEXT: cvt.rni.f16.f16 %rs3, %rs2;2124; CHECK-NEXT: cvt.rni.f16.f16 %rs4, %rs1;2125; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2126; CHECK-NEXT: ret;2127 %r = call <2 x half> @llvm.nearbyint.f16(<2 x half> %a)2128 ret <2 x half> %r2129}2130 2131define <2 x half> @test_roundeven(<2 x half> %a) #0 {2132; CHECK-LABEL: test_roundeven(2133; CHECK: {2134; CHECK-NEXT: .reg .b16 %rs<5>;2135; CHECK-NEXT: .reg .b32 %r<2>;2136; CHECK-EMPTY:2137; CHECK-NEXT: // %bb.0:2138; CHECK-NEXT: ld.param.b32 %r1, [test_roundeven_param_0];2139; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2140; CHECK-NEXT: cvt.rni.f16.f16 %rs3, %rs2;2141; CHECK-NEXT: cvt.rni.f16.f16 %rs4, %rs1;2142; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2143; CHECK-NEXT: ret;2144 %r = call <2 x half> @llvm.roundeven.f16(<2 x half> %a)2145 ret <2 x half> %r2146}2147 2148; check the use of sign mask and 0.5 to implement round2149define <2 x half> @test_round(<2 x half> %a) #0 {2150; CHECK-LABEL: test_round(2151; CHECK: {2152; CHECK-NEXT: .reg .pred %p<5>;2153; CHECK-NEXT: .reg .b16 %rs<5>;2154; CHECK-NEXT: .reg .b32 %r<21>;2155; CHECK-EMPTY:2156; CHECK-NEXT: // %bb.0:2157; CHECK-NEXT: ld.param.b32 %r1, [test_round_param_0];2158; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2159; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;2160; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;2161; CHECK-NEXT: or.b32 %r4, %r3, 1056964608;2162; CHECK-NEXT: add.rn.f32 %r5, %r2, %r4;2163; CHECK-NEXT: cvt.rzi.f32.f32 %r6, %r5;2164; CHECK-NEXT: abs.f32 %r7, %r2;2165; CHECK-NEXT: setp.gt.f32 %p1, %r7, 0f4B000000;2166; CHECK-NEXT: selp.f32 %r8, %r2, %r6, %p1;2167; CHECK-NEXT: cvt.rzi.f32.f32 %r9, %r2;2168; CHECK-NEXT: setp.lt.f32 %p2, %r7, 0f3F000000;2169; CHECK-NEXT: selp.f32 %r10, %r9, %r8, %p2;2170; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r10;2171; CHECK-NEXT: cvt.f32.f16 %r11, %rs1;2172; CHECK-NEXT: and.b32 %r12, %r11, -2147483648;2173; CHECK-NEXT: or.b32 %r13, %r12, 1056964608;2174; CHECK-NEXT: add.rn.f32 %r14, %r11, %r13;2175; CHECK-NEXT: cvt.rzi.f32.f32 %r15, %r14;2176; CHECK-NEXT: abs.f32 %r16, %r11;2177; CHECK-NEXT: setp.gt.f32 %p3, %r16, 0f4B000000;2178; CHECK-NEXT: selp.f32 %r17, %r11, %r15, %p3;2179; CHECK-NEXT: cvt.rzi.f32.f32 %r18, %r11;2180; CHECK-NEXT: setp.lt.f32 %p4, %r16, 0f3F000000;2181; CHECK-NEXT: selp.f32 %r19, %r18, %r17, %p4;2182; CHECK-NEXT: cvt.rn.f16.f32 %rs4, %r19;2183; CHECK-NEXT: mov.b32 %r20, {%rs4, %rs3};2184; CHECK-NEXT: st.param.b32 [func_retval0], %r20;2185; CHECK-NEXT: ret;2186 %r = call <2 x half> @llvm.round.f16(<2 x half> %a)2187 ret <2 x half> %r2188}2189 2190define <2 x half> @test_fmuladd(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 {2191; CHECK-F16-LABEL: test_fmuladd(2192; CHECK-F16: {2193; CHECK-F16-NEXT: .reg .b32 %r<5>;2194; CHECK-F16-EMPTY:2195; CHECK-F16-NEXT: // %bb.0:2196; CHECK-F16-NEXT: ld.param.b32 %r3, [test_fmuladd_param_2];2197; CHECK-F16-NEXT: ld.param.b32 %r2, [test_fmuladd_param_1];2198; CHECK-F16-NEXT: ld.param.b32 %r1, [test_fmuladd_param_0];2199; CHECK-F16-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;2200; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r4;2201; CHECK-F16-NEXT: ret;2202;2203; CHECK-NOF16-LABEL: test_fmuladd(2204; CHECK-NOF16: {2205; CHECK-NOF16-NEXT: .reg .b16 %rs<9>;2206; CHECK-NOF16-NEXT: .reg .b32 %r<13>;2207; CHECK-NOF16-EMPTY:2208; CHECK-NOF16-NEXT: // %bb.0:2209; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_fmuladd_param_2];2210; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fmuladd_param_1];2211; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fmuladd_param_0];2212; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r3;2213; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;2214; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r2;2215; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;2216; CHECK-NOF16-NEXT: mov.b32 {%rs5, %rs6}, %r1;2217; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs6;2218; CHECK-NOF16-NEXT: fma.rn.f32 %r7, %r6, %r5, %r4;2219; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs7, %r7;2220; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs1;2221; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs3;2222; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs5;2223; CHECK-NOF16-NEXT: fma.rn.f32 %r11, %r10, %r9, %r8;2224; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs8, %r11;2225; CHECK-NOF16-NEXT: mov.b32 %r12, {%rs8, %rs7};2226; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r12;2227; CHECK-NOF16-NEXT: ret;2228 %r = call <2 x half> @llvm.fmuladd.f16(<2 x half> %a, <2 x half> %b, <2 x half> %c)2229 ret <2 x half> %r2230}2231 2232define <2 x half> @test_shufflevector(<2 x half> %a) #0 {2233; CHECK-LABEL: test_shufflevector(2234; CHECK: {2235; CHECK-NEXT: .reg .b16 %rs<3>;2236; CHECK-NEXT: .reg .b32 %r<2>;2237; CHECK-EMPTY:2238; CHECK-NEXT: // %bb.0:2239; CHECK-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];2240; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2241; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};2242; CHECK-NEXT: ret;2243 %s = shufflevector <2 x half> %a, <2 x half> undef, <2 x i32> <i32 1, i32 0>2244 ret <2 x half> %s2245}2246 2247define <2 x half> @test_insertelement(<2 x half> %a, half %x) #0 {2248; CHECK-LABEL: test_insertelement(2249; CHECK: {2250; CHECK-NEXT: .reg .b16 %rs<3>;2251; CHECK-NEXT: .reg .b32 %r<2>;2252; CHECK-EMPTY:2253; CHECK-NEXT: // %bb.0:2254; CHECK-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];2255; CHECK-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];2256; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs2, tmp}, %r1; }2257; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};2258; CHECK-NEXT: ret;2259 %i = insertelement <2 x half> %a, half %x, i64 12260 ret <2 x half> %i2261}2262 2263define <2 x half> @test_sitofp_2xi16_to_2xhalf(<2 x i16> %a) #0 {2264; CHECK-LABEL: test_sitofp_2xi16_to_2xhalf(2265; CHECK: {2266; CHECK-NEXT: .reg .b16 %rs<5>;2267; CHECK-NEXT: .reg .b32 %r<2>;2268; CHECK-EMPTY:2269; CHECK-NEXT: // %bb.0:2270; CHECK-NEXT: ld.param.b32 %r1, [test_sitofp_2xi16_to_2xhalf_param_0];2271; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2272; CHECK-NEXT: cvt.rn.f16.s16 %rs3, %rs2;2273; CHECK-NEXT: cvt.rn.f16.s16 %rs4, %rs1;2274; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2275; CHECK-NEXT: ret;2276 %r = sitofp <2 x i16> %a to <2 x half>2277 ret <2 x half> %r2278}2279 2280define <2 x half> @test_uitofp_2xi16_to_2xhalf(<2 x i16> %a) #0 {2281; CHECK-LABEL: test_uitofp_2xi16_to_2xhalf(2282; CHECK: {2283; CHECK-NEXT: .reg .b16 %rs<5>;2284; CHECK-NEXT: .reg .b32 %r<2>;2285; CHECK-EMPTY:2286; CHECK-NEXT: // %bb.0:2287; CHECK-NEXT: ld.param.b32 %r1, [test_uitofp_2xi16_to_2xhalf_param_0];2288; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;2289; CHECK-NEXT: cvt.rn.f16.u16 %rs3, %rs2;2290; CHECK-NEXT: cvt.rn.f16.u16 %rs4, %rs1;2291; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};2292; CHECK-NEXT: ret;2293 %r = uitofp <2 x i16> %a to <2 x half>2294 ret <2 x half> %r2295}2296 2297define void @test_store_2xhalf(ptr %p1, ptr %p2, <2 x half> %v) {2298; CHECK-LABEL: test_store_2xhalf(2299; CHECK: {2300; CHECK-NEXT: .reg .b32 %r<2>;2301; CHECK-NEXT: .reg .b64 %rd<3>;2302; CHECK-EMPTY:2303; CHECK-NEXT: // %bb.0:2304; CHECK-NEXT: ld.param.b32 %r1, [test_store_2xhalf_param_2];2305; CHECK-NEXT: ld.param.b64 %rd2, [test_store_2xhalf_param_1];2306; CHECK-NEXT: ld.param.b64 %rd1, [test_store_2xhalf_param_0];2307; CHECK-NEXT: st.b32 [%rd1], %r1;2308; CHECK-NEXT: st.b32 [%rd2], 1006648320;2309; CHECK-NEXT: ret;2310 store <2 x half> %v, ptr %p12311 store <2 x half> <half 1.0, half 1.0>, ptr %p22312 ret void2313}2314 2315 2316 2317attributes #0 = { nounwind }2318