3008 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; ## Full FP32x2 support enabled by default.3; RUN: llc < %s -mcpu=sm_80 -O0 -disable-post-ra -frame-pointer=all \4; RUN: -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-NOF32X2 %s5; RUN: %if ptxas-sm_80 %{ \6; RUN: llc < %s -mcpu=sm_80 -O0 -disable-post-ra -frame-pointer=all \7; RUN: -verify-machineinstrs | %ptxas-verify -arch=sm_80 \8; RUN: %}9; RUN: llc < %s -mcpu=sm_100 -O0 -disable-post-ra -frame-pointer=all \10; RUN: -verify-machineinstrs | FileCheck --check-prefixes=CHECK,CHECK-F32X2 %s11; RUN: %if ptxas-sm_100 %{ \12; RUN: llc < %s -mcpu=sm_100 -O0 -disable-post-ra -frame-pointer=all \13; RUN: -verify-machineinstrs | %ptxas-verify -arch=sm_100 \14; RUN: %}15 16target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"17target triple = "nvptx64-nvidia-cuda"18 19define <2 x float> @test_ret_const() #0 {20; CHECK-LABEL: test_ret_const(21; CHECK: {22; CHECK-EMPTY:23; CHECK-EMPTY:24; CHECK-NEXT: // %bb.0:25; CHECK-NEXT: st.param.v2.b32 [func_retval0], {0f3F800000, 0f40000000};26; CHECK-NEXT: ret;27 ret <2 x float> <float 1.0, float 2.0>28}29 30define float @test_extract_0(<2 x float> %a) #0 {31; CHECK-NOF32X2-LABEL: test_extract_0(32; CHECK-NOF32X2: {33; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;34; CHECK-NOF32X2-EMPTY:35; CHECK-NOF32X2-NEXT: // %bb.0:36; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_0_param_0];37; CHECK-NOF32X2-NEXT: st.param.b32 [func_retval0], %r1;38; CHECK-NOF32X2-NEXT: ret;39;40; CHECK-F32X2-LABEL: test_extract_0(41; CHECK-F32X2: {42; CHECK-F32X2-NEXT: .reg .b32 %r<2>;43; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;44; CHECK-F32X2-EMPTY:45; CHECK-F32X2-NEXT: // %bb.0:46; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_extract_0_param_0];47; CHECK-F32X2-NEXT: mov.b64 {%r1, _}, %rd1;48; CHECK-F32X2-NEXT: st.param.b32 [func_retval0], %r1;49; CHECK-F32X2-NEXT: ret;50 %e = extractelement <2 x float> %a, i32 051 ret float %e52}53 54define float @test_extract_1(<2 x float> %a) #0 {55; CHECK-NOF32X2-LABEL: test_extract_1(56; CHECK-NOF32X2: {57; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;58; CHECK-NOF32X2-EMPTY:59; CHECK-NOF32X2-NEXT: // %bb.0:60; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_1_param_0];61; CHECK-NOF32X2-NEXT: st.param.b32 [func_retval0], %r2;62; CHECK-NOF32X2-NEXT: ret;63;64; CHECK-F32X2-LABEL: test_extract_1(65; CHECK-F32X2: {66; CHECK-F32X2-NEXT: .reg .b32 %r<2>;67; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;68; CHECK-F32X2-EMPTY:69; CHECK-F32X2-NEXT: // %bb.0:70; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_extract_1_param_0];71; CHECK-F32X2-NEXT: mov.b64 {_, %r1}, %rd1;72; CHECK-F32X2-NEXT: st.param.b32 [func_retval0], %r1;73; CHECK-F32X2-NEXT: ret;74 %e = extractelement <2 x float> %a, i32 175 ret float %e76}77 78define float @test_extract_i(<2 x float> %a, i64 %idx) #0 {79; CHECK-NOF32X2-LABEL: test_extract_i(80; CHECK-NOF32X2: {81; CHECK-NOF32X2-NEXT: .local .align 8 .b8 __local_depot3[8];82; CHECK-NOF32X2-NEXT: .reg .b64 %SP;83; CHECK-NOF32X2-NEXT: .reg .b64 %SPL;84; CHECK-NOF32X2-NEXT: .reg .b32 %r<4>;85; CHECK-NOF32X2-NEXT: .reg .b64 %rd<6>;86; CHECK-NOF32X2-EMPTY:87; CHECK-NOF32X2-NEXT: // %bb.0:88; CHECK-NOF32X2-NEXT: mov.b64 %SPL, __local_depot3;89; CHECK-NOF32X2-NEXT: cvta.local.u64 %SP, %SPL;90; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_i_param_0];91; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];92; CHECK-NOF32X2-NEXT: st.v2.b32 [%SP], {%r1, %r2};93; CHECK-NOF32X2-NEXT: and.b64 %rd2, %rd1, 1;94; CHECK-NOF32X2-NEXT: shl.b64 %rd3, %rd2, 2;95; CHECK-NOF32X2-NEXT: add.u64 %rd4, %SP, 0;96; CHECK-NOF32X2-NEXT: or.b64 %rd5, %rd4, %rd3;97; CHECK-NOF32X2-NEXT: ld.b32 %r3, [%rd5];98; CHECK-NOF32X2-NEXT: st.param.b32 [func_retval0], %r3;99; CHECK-NOF32X2-NEXT: ret;100;101; CHECK-F32X2-LABEL: test_extract_i(102; CHECK-F32X2: {103; CHECK-F32X2-NEXT: .reg .pred %p<2>;104; CHECK-F32X2-NEXT: .reg .b32 %r<4>;105; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;106; CHECK-F32X2-EMPTY:107; CHECK-F32X2-NEXT: // %bb.0:108; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_extract_i_param_1];109; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_extract_i_param_0];110; CHECK-F32X2-NEXT: setp.eq.b64 %p1, %rd2, 0;111; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;112; CHECK-F32X2-NEXT: selp.f32 %r3, %r1, %r2, %p1;113; CHECK-F32X2-NEXT: st.param.b32 [func_retval0], %r3;114; CHECK-F32X2-NEXT: ret;115 %e = extractelement <2 x float> %a, i64 %idx116 ret float %e117}118 119define <2 x float> @test_fadd(<2 x float> %a, <2 x float> %b) #0 {120; CHECK-NOF32X2-LABEL: test_fadd(121; CHECK-NOF32X2: {122; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;123; CHECK-NOF32X2-EMPTY:124; CHECK-NOF32X2-NEXT: // %bb.0:125; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fadd_param_1];126; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_param_0];127; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r2, %r4;128; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r1, %r3;129; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};130; CHECK-NOF32X2-NEXT: ret;131;132; CHECK-F32X2-LABEL: test_fadd(133; CHECK-F32X2: {134; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;135; CHECK-F32X2-EMPTY:136; CHECK-F32X2-NEXT: // %bb.0:137; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fadd_param_1];138; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fadd_param_0];139; CHECK-F32X2-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;140; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;141; CHECK-F32X2-NEXT: ret;142 %r = fadd <2 x float> %a, %b143 ret <2 x float> %r144}145 146define <2 x float> @test_fadd_imm_0(<2 x float> %a) #0 {147; CHECK-NOF32X2-LABEL: test_fadd_imm_0(148; CHECK-NOF32X2: {149; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;150; CHECK-NOF32X2-EMPTY:151; CHECK-NOF32X2-NEXT: // %bb.0:152; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_0_param_0];153; CHECK-NOF32X2-NEXT: add.rn.f32 %r3, %r2, 0f40000000;154; CHECK-NOF32X2-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;155; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};156; CHECK-NOF32X2-NEXT: ret;157;158; CHECK-F32X2-LABEL: test_fadd_imm_0(159; CHECK-F32X2: {160; CHECK-F32X2-NEXT: .reg .b32 %r<3>;161; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;162; CHECK-F32X2-EMPTY:163; CHECK-F32X2-NEXT: // %bb.0:164; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_param_0];165; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40000000;166; CHECK-F32X2-NEXT: mov.b32 %r2, 0f3F800000;167; CHECK-F32X2-NEXT: mov.b64 %rd2, {%r2, %r1};168; CHECK-F32X2-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;169; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;170; CHECK-F32X2-NEXT: ret;171 %r = fadd <2 x float> <float 1.0, float 2.0>, %a172 ret <2 x float> %r173}174 175define <2 x float> @test_fadd_imm_1(<2 x float> %a) #0 {176; CHECK-NOF32X2-LABEL: test_fadd_imm_1(177; CHECK-NOF32X2: {178; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;179; CHECK-NOF32X2-EMPTY:180; CHECK-NOF32X2-NEXT: // %bb.0:181; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_1_param_0];182; CHECK-NOF32X2-NEXT: add.rn.f32 %r3, %r2, 0f40000000;183; CHECK-NOF32X2-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;184; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};185; CHECK-NOF32X2-NEXT: ret;186;187; CHECK-F32X2-LABEL: test_fadd_imm_1(188; CHECK-F32X2: {189; CHECK-F32X2-NEXT: .reg .b32 %r<3>;190; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;191; CHECK-F32X2-EMPTY:192; CHECK-F32X2-NEXT: // %bb.0:193; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_param_0];194; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40000000;195; CHECK-F32X2-NEXT: mov.b32 %r2, 0f3F800000;196; CHECK-F32X2-NEXT: mov.b64 %rd2, {%r2, %r1};197; CHECK-F32X2-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;198; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;199; CHECK-F32X2-NEXT: ret;200 %r = fadd <2 x float> %a, <float 1.0, float 2.0>201 ret <2 x float> %r202}203 204define <4 x float> @test_fadd_v4(<4 x float> %a, <4 x float> %b) #0 {205; CHECK-NOF32X2-LABEL: test_fadd_v4(206; CHECK-NOF32X2: {207; CHECK-NOF32X2-NEXT: .reg .b32 %r<13>;208; CHECK-NOF32X2-EMPTY:209; CHECK-NOF32X2-NEXT: // %bb.0:210; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_fadd_v4_param_1];211; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_v4_param_0];212; CHECK-NOF32X2-NEXT: add.rn.f32 %r9, %r4, %r8;213; CHECK-NOF32X2-NEXT: add.rn.f32 %r10, %r3, %r7;214; CHECK-NOF32X2-NEXT: add.rn.f32 %r11, %r2, %r6;215; CHECK-NOF32X2-NEXT: add.rn.f32 %r12, %r1, %r5;216; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r12, %r11, %r10, %r9};217; CHECK-NOF32X2-NEXT: ret;218;219; CHECK-F32X2-LABEL: test_fadd_v4(220; CHECK-F32X2: {221; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;222; CHECK-F32X2-EMPTY:223; CHECK-F32X2-NEXT: // %bb.0:224; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_fadd_v4_param_1];225; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_v4_param_0];226; CHECK-F32X2-NEXT: add.rn.f32x2 %rd5, %rd2, %rd4;227; CHECK-F32X2-NEXT: add.rn.f32x2 %rd6, %rd1, %rd3;228; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd5};229; CHECK-F32X2-NEXT: ret;230 %r = fadd <4 x float> %a, %b231 ret <4 x float> %r232}233 234define <4 x float> @test_fadd_imm_0_v4(<4 x float> %a) #0 {235; CHECK-NOF32X2-LABEL: test_fadd_imm_0_v4(236; CHECK-NOF32X2: {237; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;238; CHECK-NOF32X2-EMPTY:239; CHECK-NOF32X2-NEXT: // %bb.0:240; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_0_v4_param_0];241; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r4, 0f40800000;242; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r3, 0f40400000;243; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r2, 0f40000000;244; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;245; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};246; CHECK-NOF32X2-NEXT: ret;247;248; CHECK-F32X2-LABEL: test_fadd_imm_0_v4(249; CHECK-F32X2: {250; CHECK-F32X2-NEXT: .reg .b32 %r<5>;251; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;252; CHECK-F32X2-EMPTY:253; CHECK-F32X2-NEXT: // %bb.0:254; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_param_0];255; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40800000;256; CHECK-F32X2-NEXT: mov.b32 %r2, 0f40400000;257; CHECK-F32X2-NEXT: mov.b64 %rd3, {%r2, %r1};258; CHECK-F32X2-NEXT: add.rn.f32x2 %rd4, %rd2, %rd3;259; CHECK-F32X2-NEXT: mov.b32 %r3, 0f40000000;260; CHECK-F32X2-NEXT: mov.b32 %r4, 0f3F800000;261; CHECK-F32X2-NEXT: mov.b64 %rd5, {%r4, %r3};262; CHECK-F32X2-NEXT: add.rn.f32x2 %rd6, %rd1, %rd5;263; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};264; CHECK-F32X2-NEXT: ret;265 %r = fadd <4 x float> <float 1.0, float 2.0, float 3.0, float 4.0>, %a266 ret <4 x float> %r267}268 269define <4 x float> @test_fadd_imm_1_v4(<4 x float> %a) #0 {270; CHECK-NOF32X2-LABEL: test_fadd_imm_1_v4(271; CHECK-NOF32X2: {272; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;273; CHECK-NOF32X2-EMPTY:274; CHECK-NOF32X2-NEXT: // %bb.0:275; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_1_v4_param_0];276; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r4, 0f40800000;277; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r3, 0f40400000;278; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r2, 0f40000000;279; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;280; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};281; CHECK-NOF32X2-NEXT: ret;282;283; CHECK-F32X2-LABEL: test_fadd_imm_1_v4(284; CHECK-F32X2: {285; CHECK-F32X2-NEXT: .reg .b32 %r<5>;286; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;287; CHECK-F32X2-EMPTY:288; CHECK-F32X2-NEXT: // %bb.0:289; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_1_v4_param_0];290; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40800000;291; CHECK-F32X2-NEXT: mov.b32 %r2, 0f40400000;292; CHECK-F32X2-NEXT: mov.b64 %rd3, {%r2, %r1};293; CHECK-F32X2-NEXT: add.rn.f32x2 %rd4, %rd2, %rd3;294; CHECK-F32X2-NEXT: mov.b32 %r3, 0f40000000;295; CHECK-F32X2-NEXT: mov.b32 %r4, 0f3F800000;296; CHECK-F32X2-NEXT: mov.b64 %rd5, {%r4, %r3};297; CHECK-F32X2-NEXT: add.rn.f32x2 %rd6, %rd1, %rd5;298; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};299; CHECK-F32X2-NEXT: ret;300 %r = fadd <4 x float> %a, <float 1.0, float 2.0, float 3.0, float 4.0>301 ret <4 x float> %r302}303 304define <2 x float> @test_fsub(<2 x float> %a, <2 x float> %b) #0 {305; CHECK-NOF32X2-LABEL: test_fsub(306; CHECK-NOF32X2: {307; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;308; CHECK-NOF32X2-EMPTY:309; CHECK-NOF32X2-NEXT: // %bb.0:310; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fsub_param_1];311; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fsub_param_0];312; CHECK-NOF32X2-NEXT: sub.rn.f32 %r5, %r2, %r4;313; CHECK-NOF32X2-NEXT: sub.rn.f32 %r6, %r1, %r3;314; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};315; CHECK-NOF32X2-NEXT: ret;316;317; CHECK-F32X2-LABEL: test_fsub(318; CHECK-F32X2: {319; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;320; CHECK-F32X2-EMPTY:321; CHECK-F32X2-NEXT: // %bb.0:322; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fsub_param_1];323; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fsub_param_0];324; CHECK-F32X2-NEXT: sub.rn.f32x2 %rd3, %rd1, %rd2;325; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;326; CHECK-F32X2-NEXT: ret;327 %r = fsub <2 x float> %a, %b328 ret <2 x float> %r329}330 331define <2 x float> @test_fneg(<2 x float> %a) #0 {332; CHECK-NOF32X2-LABEL: test_fneg(333; CHECK-NOF32X2: {334; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;335; CHECK-NOF32X2-EMPTY:336; CHECK-NOF32X2-NEXT: // %bb.0:337; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fneg_param_0];338; CHECK-NOF32X2-NEXT: neg.f32 %r3, %r2;339; CHECK-NOF32X2-NEXT: neg.f32 %r4, %r1;340; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};341; CHECK-NOF32X2-NEXT: ret;342;343; CHECK-F32X2-LABEL: test_fneg(344; CHECK-F32X2: {345; CHECK-F32X2-NEXT: .reg .b32 %r<5>;346; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;347; CHECK-F32X2-EMPTY:348; CHECK-F32X2-NEXT: // %bb.0:349; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fneg_param_0];350; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;351; CHECK-F32X2-NEXT: neg.f32 %r3, %r2;352; CHECK-F32X2-NEXT: neg.f32 %r4, %r1;353; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};354; CHECK-F32X2-NEXT: ret;355 %r = fneg <2 x float> %a356 ret <2 x float> %r357}358 359define <2 x float> @test_fmul(<2 x float> %a, <2 x float> %b) #0 {360; CHECK-NOF32X2-LABEL: test_fmul(361; CHECK-NOF32X2: {362; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;363; CHECK-NOF32X2-EMPTY:364; CHECK-NOF32X2-NEXT: // %bb.0:365; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmul_param_1];366; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fmul_param_0];367; CHECK-NOF32X2-NEXT: mul.rn.f32 %r5, %r2, %r4;368; CHECK-NOF32X2-NEXT: mul.rn.f32 %r6, %r1, %r3;369; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};370; CHECK-NOF32X2-NEXT: ret;371;372; CHECK-F32X2-LABEL: test_fmul(373; CHECK-F32X2: {374; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;375; CHECK-F32X2-EMPTY:376; CHECK-F32X2-NEXT: // %bb.0:377; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fmul_param_1];378; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fmul_param_0];379; CHECK-F32X2-NEXT: mul.rn.f32x2 %rd3, %rd1, %rd2;380; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;381; CHECK-F32X2-NEXT: ret;382 %r = fmul <2 x float> %a, %b383 ret <2 x float> %r384}385 386define <2 x float> @test_fdiv(<2 x float> %a, <2 x float> %b) #0 {387; CHECK-NOF32X2-LABEL: test_fdiv(388; CHECK-NOF32X2: {389; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;390; CHECK-NOF32X2-EMPTY:391; CHECK-NOF32X2-NEXT: // %bb.0:392; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fdiv_param_1];393; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fdiv_param_0];394; CHECK-NOF32X2-NEXT: div.rn.f32 %r5, %r2, %r4;395; CHECK-NOF32X2-NEXT: div.rn.f32 %r6, %r1, %r3;396; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};397; CHECK-NOF32X2-NEXT: ret;398;399; CHECK-F32X2-LABEL: test_fdiv(400; CHECK-F32X2: {401; CHECK-F32X2-NEXT: .reg .b32 %r<7>;402; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;403; CHECK-F32X2-EMPTY:404; CHECK-F32X2-NEXT: // %bb.0:405; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fdiv_param_1];406; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fdiv_param_0];407; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;408; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;409; CHECK-F32X2-NEXT: div.rn.f32 %r5, %r4, %r2;410; CHECK-F32X2-NEXT: div.rn.f32 %r6, %r3, %r1;411; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};412; CHECK-F32X2-NEXT: ret;413 %r = fdiv <2 x float> %a, %b414 ret <2 x float> %r415}416 417define <2 x float> @test_frem(<2 x float> %a, <2 x float> %b) #0 {418; CHECK-NOF32X2-LABEL: test_frem(419; CHECK-NOF32X2: {420; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;421; CHECK-NOF32X2-NEXT: .reg .b32 %r<15>;422; CHECK-NOF32X2-EMPTY:423; CHECK-NOF32X2-NEXT: // %bb.0:424; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_frem_param_1];425; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_frem_param_0];426; CHECK-NOF32X2-NEXT: div.rn.f32 %r5, %r2, %r4;427; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r6, %r5;428; CHECK-NOF32X2-NEXT: neg.f32 %r7, %r6;429; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r7, %r4, %r2;430; CHECK-NOF32X2-NEXT: testp.infinite.f32 %p1, %r4;431; CHECK-NOF32X2-NEXT: selp.f32 %r9, %r2, %r8, %p1;432; CHECK-NOF32X2-NEXT: div.rn.f32 %r10, %r1, %r3;433; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r11, %r10;434; CHECK-NOF32X2-NEXT: neg.f32 %r12, %r11;435; CHECK-NOF32X2-NEXT: fma.rn.f32 %r13, %r12, %r3, %r1;436; CHECK-NOF32X2-NEXT: testp.infinite.f32 %p2, %r3;437; CHECK-NOF32X2-NEXT: selp.f32 %r14, %r1, %r13, %p2;438; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r9};439; CHECK-NOF32X2-NEXT: ret;440;441; CHECK-F32X2-LABEL: test_frem(442; CHECK-F32X2: {443; CHECK-F32X2-NEXT: .reg .pred %p<3>;444; CHECK-F32X2-NEXT: .reg .b32 %r<15>;445; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;446; CHECK-F32X2-EMPTY:447; CHECK-F32X2-NEXT: // %bb.0:448; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_frem_param_1];449; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_frem_param_0];450; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;451; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;452; CHECK-F32X2-NEXT: div.rn.f32 %r5, %r4, %r2;453; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r6, %r5;454; CHECK-F32X2-NEXT: neg.f32 %r7, %r6;455; CHECK-F32X2-NEXT: fma.rn.f32 %r8, %r7, %r2, %r4;456; CHECK-F32X2-NEXT: testp.infinite.f32 %p1, %r2;457; CHECK-F32X2-NEXT: selp.f32 %r9, %r4, %r8, %p1;458; CHECK-F32X2-NEXT: div.rn.f32 %r10, %r3, %r1;459; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r11, %r10;460; CHECK-F32X2-NEXT: neg.f32 %r12, %r11;461; CHECK-F32X2-NEXT: fma.rn.f32 %r13, %r12, %r1, %r3;462; CHECK-F32X2-NEXT: testp.infinite.f32 %p2, %r1;463; CHECK-F32X2-NEXT: selp.f32 %r14, %r3, %r13, %p2;464; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r9};465; CHECK-F32X2-NEXT: ret;466 %r = frem <2 x float> %a, %b467 ret <2 x float> %r468}469 470define <2 x float> @test_fadd_ftz(<2 x float> %a, <2 x float> %b) #2 {471; CHECK-NOF32X2-LABEL: test_fadd_ftz(472; CHECK-NOF32X2: {473; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;474; CHECK-NOF32X2-EMPTY:475; CHECK-NOF32X2-NEXT: // %bb.0:476; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fadd_ftz_param_1];477; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_ftz_param_0];478; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r2, %r4;479; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r1, %r3;480; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};481; CHECK-NOF32X2-NEXT: ret;482;483; CHECK-F32X2-LABEL: test_fadd_ftz(484; CHECK-F32X2: {485; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;486; CHECK-F32X2-EMPTY:487; CHECK-F32X2-NEXT: // %bb.0:488; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fadd_ftz_param_1];489; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fadd_ftz_param_0];490; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;491; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;492; CHECK-F32X2-NEXT: ret;493 %r = fadd <2 x float> %a, %b494 ret <2 x float> %r495}496 497define <2 x float> @test_fadd_imm_0_ftz(<2 x float> %a) #2 {498; CHECK-NOF32X2-LABEL: test_fadd_imm_0_ftz(499; CHECK-NOF32X2: {500; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;501; CHECK-NOF32X2-EMPTY:502; CHECK-NOF32X2-NEXT: // %bb.0:503; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_0_ftz_param_0];504; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40000000;505; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r4, %r1, 0f3F800000;506; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};507; CHECK-NOF32X2-NEXT: ret;508;509; CHECK-F32X2-LABEL: test_fadd_imm_0_ftz(510; CHECK-F32X2: {511; CHECK-F32X2-NEXT: .reg .b32 %r<3>;512; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;513; CHECK-F32X2-EMPTY:514; CHECK-F32X2-NEXT: // %bb.0:515; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_ftz_param_0];516; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40000000;517; CHECK-F32X2-NEXT: mov.b32 %r2, 0f3F800000;518; CHECK-F32X2-NEXT: mov.b64 %rd2, {%r2, %r1};519; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;520; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;521; CHECK-F32X2-NEXT: ret;522 %r = fadd <2 x float> <float 1.0, float 2.0>, %a523 ret <2 x float> %r524}525 526define <2 x float> @test_fadd_imm_1_ftz(<2 x float> %a) #2 {527; CHECK-NOF32X2-LABEL: test_fadd_imm_1_ftz(528; CHECK-NOF32X2: {529; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;530; CHECK-NOF32X2-EMPTY:531; CHECK-NOF32X2-NEXT: // %bb.0:532; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_1_ftz_param_0];533; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40000000;534; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r4, %r1, 0f3F800000;535; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};536; CHECK-NOF32X2-NEXT: ret;537;538; CHECK-F32X2-LABEL: test_fadd_imm_1_ftz(539; CHECK-F32X2: {540; CHECK-F32X2-NEXT: .reg .b32 %r<3>;541; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;542; CHECK-F32X2-EMPTY:543; CHECK-F32X2-NEXT: // %bb.0:544; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_ftz_param_0];545; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40000000;546; CHECK-F32X2-NEXT: mov.b32 %r2, 0f3F800000;547; CHECK-F32X2-NEXT: mov.b64 %rd2, {%r2, %r1};548; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;549; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;550; CHECK-F32X2-NEXT: ret;551 %r = fadd <2 x float> %a, <float 1.0, float 2.0>552 ret <2 x float> %r553}554 555define <4 x float> @test_fadd_v4_ftz(<4 x float> %a, <4 x float> %b) #2 {556; CHECK-NOF32X2-LABEL: test_fadd_v4_ftz(557; CHECK-NOF32X2: {558; CHECK-NOF32X2-NEXT: .reg .b32 %r<13>;559; CHECK-NOF32X2-EMPTY:560; CHECK-NOF32X2-NEXT: // %bb.0:561; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_fadd_v4_ftz_param_1];562; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_v4_ftz_param_0];563; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r9, %r4, %r8;564; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r10, %r3, %r7;565; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r11, %r2, %r6;566; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r12, %r1, %r5;567; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r12, %r11, %r10, %r9};568; CHECK-NOF32X2-NEXT: ret;569;570; CHECK-F32X2-LABEL: test_fadd_v4_ftz(571; CHECK-F32X2: {572; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;573; CHECK-F32X2-EMPTY:574; CHECK-F32X2-NEXT: // %bb.0:575; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_fadd_v4_ftz_param_1];576; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_v4_ftz_param_0];577; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd5, %rd2, %rd4;578; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd6, %rd1, %rd3;579; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd5};580; CHECK-F32X2-NEXT: ret;581 %r = fadd <4 x float> %a, %b582 ret <4 x float> %r583}584 585define <4 x float> @test_fadd_imm_0_v4_ftz(<4 x float> %a) #2 {586; CHECK-NOF32X2-LABEL: test_fadd_imm_0_v4_ftz(587; CHECK-NOF32X2: {588; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;589; CHECK-NOF32X2-EMPTY:590; CHECK-NOF32X2-NEXT: // %bb.0:591; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_0_v4_ftz_param_0];592; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r4, 0f40800000;593; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r3, 0f40400000;594; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r7, %r2, 0f40000000;595; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r8, %r1, 0f3F800000;596; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};597; CHECK-NOF32X2-NEXT: ret;598;599; CHECK-F32X2-LABEL: test_fadd_imm_0_v4_ftz(600; CHECK-F32X2: {601; CHECK-F32X2-NEXT: .reg .b32 %r<5>;602; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;603; CHECK-F32X2-EMPTY:604; CHECK-F32X2-NEXT: // %bb.0:605; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_ftz_param_0];606; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40800000;607; CHECK-F32X2-NEXT: mov.b32 %r2, 0f40400000;608; CHECK-F32X2-NEXT: mov.b64 %rd3, {%r2, %r1};609; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd4, %rd2, %rd3;610; CHECK-F32X2-NEXT: mov.b32 %r3, 0f40000000;611; CHECK-F32X2-NEXT: mov.b32 %r4, 0f3F800000;612; CHECK-F32X2-NEXT: mov.b64 %rd5, {%r4, %r3};613; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd6, %rd1, %rd5;614; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};615; CHECK-F32X2-NEXT: ret;616 %r = fadd <4 x float> <float 1.0, float 2.0, float 3.0, float 4.0>, %a617 ret <4 x float> %r618}619 620define <4 x float> @test_fadd_imm_1_v4_ftz(<4 x float> %a) #2 {621; CHECK-NOF32X2-LABEL: test_fadd_imm_1_v4_ftz(622; CHECK-NOF32X2: {623; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;624; CHECK-NOF32X2-EMPTY:625; CHECK-NOF32X2-NEXT: // %bb.0:626; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_1_v4_ftz_param_0];627; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r4, 0f40800000;628; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r3, 0f40400000;629; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r7, %r2, 0f40000000;630; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r8, %r1, 0f3F800000;631; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};632; CHECK-NOF32X2-NEXT: ret;633;634; CHECK-F32X2-LABEL: test_fadd_imm_1_v4_ftz(635; CHECK-F32X2: {636; CHECK-F32X2-NEXT: .reg .b32 %r<5>;637; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;638; CHECK-F32X2-EMPTY:639; CHECK-F32X2-NEXT: // %bb.0:640; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_1_v4_ftz_param_0];641; CHECK-F32X2-NEXT: mov.b32 %r1, 0f40800000;642; CHECK-F32X2-NEXT: mov.b32 %r2, 0f40400000;643; CHECK-F32X2-NEXT: mov.b64 %rd3, {%r2, %r1};644; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd4, %rd2, %rd3;645; CHECK-F32X2-NEXT: mov.b32 %r3, 0f40000000;646; CHECK-F32X2-NEXT: mov.b32 %r4, 0f3F800000;647; CHECK-F32X2-NEXT: mov.b64 %rd5, {%r4, %r3};648; CHECK-F32X2-NEXT: add.rn.ftz.f32x2 %rd6, %rd1, %rd5;649; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};650; CHECK-F32X2-NEXT: ret;651 %r = fadd <4 x float> %a, <float 1.0, float 2.0, float 3.0, float 4.0>652 ret <4 x float> %r653}654 655define <2 x float> @test_fsub_ftz(<2 x float> %a, <2 x float> %b) #2 {656; CHECK-NOF32X2-LABEL: test_fsub_ftz(657; CHECK-NOF32X2: {658; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;659; CHECK-NOF32X2-EMPTY:660; CHECK-NOF32X2-NEXT: // %bb.0:661; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fsub_ftz_param_1];662; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fsub_ftz_param_0];663; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r5, %r2, %r4;664; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r6, %r1, %r3;665; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};666; CHECK-NOF32X2-NEXT: ret;667;668; CHECK-F32X2-LABEL: test_fsub_ftz(669; CHECK-F32X2: {670; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;671; CHECK-F32X2-EMPTY:672; CHECK-F32X2-NEXT: // %bb.0:673; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fsub_ftz_param_1];674; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fsub_ftz_param_0];675; CHECK-F32X2-NEXT: sub.rn.ftz.f32x2 %rd3, %rd1, %rd2;676; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;677; CHECK-F32X2-NEXT: ret;678 %r = fsub <2 x float> %a, %b679 ret <2 x float> %r680}681 682define <2 x float> @test_fneg_ftz(<2 x float> %a) #2 {683; CHECK-NOF32X2-LABEL: test_fneg_ftz(684; CHECK-NOF32X2: {685; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;686; CHECK-NOF32X2-EMPTY:687; CHECK-NOF32X2-NEXT: // %bb.0:688; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fneg_ftz_param_0];689; CHECK-NOF32X2-NEXT: neg.ftz.f32 %r3, %r2;690; CHECK-NOF32X2-NEXT: neg.ftz.f32 %r4, %r1;691; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};692; CHECK-NOF32X2-NEXT: ret;693;694; CHECK-F32X2-LABEL: test_fneg_ftz(695; CHECK-F32X2: {696; CHECK-F32X2-NEXT: .reg .b32 %r<5>;697; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;698; CHECK-F32X2-EMPTY:699; CHECK-F32X2-NEXT: // %bb.0:700; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fneg_ftz_param_0];701; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;702; CHECK-F32X2-NEXT: neg.ftz.f32 %r3, %r2;703; CHECK-F32X2-NEXT: neg.ftz.f32 %r4, %r1;704; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};705; CHECK-F32X2-NEXT: ret;706 %r = fneg <2 x float> %a707 ret <2 x float> %r708}709 710define <2 x float> @test_fmul_ftz(<2 x float> %a, <2 x float> %b) #2 {711; CHECK-NOF32X2-LABEL: test_fmul_ftz(712; CHECK-NOF32X2: {713; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;714; CHECK-NOF32X2-EMPTY:715; CHECK-NOF32X2-NEXT: // %bb.0:716; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmul_ftz_param_1];717; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fmul_ftz_param_0];718; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r5, %r2, %r4;719; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r6, %r1, %r3;720; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};721; CHECK-NOF32X2-NEXT: ret;722;723; CHECK-F32X2-LABEL: test_fmul_ftz(724; CHECK-F32X2: {725; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;726; CHECK-F32X2-EMPTY:727; CHECK-F32X2-NEXT: // %bb.0:728; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fmul_ftz_param_1];729; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fmul_ftz_param_0];730; CHECK-F32X2-NEXT: mul.rn.ftz.f32x2 %rd3, %rd1, %rd2;731; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;732; CHECK-F32X2-NEXT: ret;733 %r = fmul <2 x float> %a, %b734 ret <2 x float> %r735}736 737define <2 x float> @test_fma_ftz(<2 x float> %a, <2 x float> %b, <2 x float> %c) #2 {738; CHECK-NOF32X2-LABEL: test_fma_ftz(739; CHECK-NOF32X2: {740; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;741; CHECK-NOF32X2-EMPTY:742; CHECK-NOF32X2-NEXT: // %bb.0:743; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fma_ftz_param_2];744; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fma_ftz_param_1];745; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fma_ftz_param_0];746; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r7, %r2, %r4, %r6;747; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r8, %r1, %r3, %r5;748; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};749; CHECK-NOF32X2-NEXT: ret;750;751; CHECK-F32X2-LABEL: test_fma_ftz(752; CHECK-F32X2: {753; CHECK-F32X2-NEXT: .reg .b64 %rd<5>;754; CHECK-F32X2-EMPTY:755; CHECK-F32X2-NEXT: // %bb.0:756; CHECK-F32X2-NEXT: ld.param.b64 %rd3, [test_fma_ftz_param_2];757; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fma_ftz_param_1];758; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fma_ftz_param_0];759; CHECK-F32X2-NEXT: fma.rn.ftz.f32x2 %rd4, %rd1, %rd2, %rd3;760; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd4;761; CHECK-F32X2-NEXT: ret;762 %r = call <2 x float> @llvm.fma(<2 x float> %a, <2 x float> %b, <2 x float> %c)763 ret <2 x float> %r764}765 766define <2 x float> @test_fdiv_ftz(<2 x float> %a, <2 x float> %b) #2 {767; CHECK-NOF32X2-LABEL: test_fdiv_ftz(768; CHECK-NOF32X2: {769; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;770; CHECK-NOF32X2-EMPTY:771; CHECK-NOF32X2-NEXT: // %bb.0:772; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fdiv_ftz_param_1];773; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fdiv_ftz_param_0];774; CHECK-NOF32X2-NEXT: div.rn.ftz.f32 %r5, %r2, %r4;775; CHECK-NOF32X2-NEXT: div.rn.ftz.f32 %r6, %r1, %r3;776; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};777; CHECK-NOF32X2-NEXT: ret;778;779; CHECK-F32X2-LABEL: test_fdiv_ftz(780; CHECK-F32X2: {781; CHECK-F32X2-NEXT: .reg .b32 %r<7>;782; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;783; CHECK-F32X2-EMPTY:784; CHECK-F32X2-NEXT: // %bb.0:785; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fdiv_ftz_param_1];786; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fdiv_ftz_param_0];787; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;788; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;789; CHECK-F32X2-NEXT: div.rn.ftz.f32 %r5, %r4, %r2;790; CHECK-F32X2-NEXT: div.rn.ftz.f32 %r6, %r3, %r1;791; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};792; CHECK-F32X2-NEXT: ret;793 %r = fdiv <2 x float> %a, %b794 ret <2 x float> %r795}796 797define <2 x float> @test_frem_ftz(<2 x float> %a, <2 x float> %b) #2 {798; CHECK-NOF32X2-LABEL: test_frem_ftz(799; CHECK-NOF32X2: {800; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;801; CHECK-NOF32X2-NEXT: .reg .b32 %r<15>;802; CHECK-NOF32X2-EMPTY:803; CHECK-NOF32X2-NEXT: // %bb.0:804; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_frem_ftz_param_1];805; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_frem_ftz_param_0];806; CHECK-NOF32X2-NEXT: div.rn.ftz.f32 %r5, %r2, %r4;807; CHECK-NOF32X2-NEXT: cvt.rzi.ftz.f32.f32 %r6, %r5;808; CHECK-NOF32X2-NEXT: neg.ftz.f32 %r7, %r6;809; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r8, %r7, %r4, %r2;810; CHECK-NOF32X2-NEXT: testp.infinite.f32 %p1, %r4;811; CHECK-NOF32X2-NEXT: selp.f32 %r9, %r2, %r8, %p1;812; CHECK-NOF32X2-NEXT: div.rn.ftz.f32 %r10, %r1, %r3;813; CHECK-NOF32X2-NEXT: cvt.rzi.ftz.f32.f32 %r11, %r10;814; CHECK-NOF32X2-NEXT: neg.ftz.f32 %r12, %r11;815; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r13, %r12, %r3, %r1;816; CHECK-NOF32X2-NEXT: testp.infinite.f32 %p2, %r3;817; CHECK-NOF32X2-NEXT: selp.f32 %r14, %r1, %r13, %p2;818; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r9};819; CHECK-NOF32X2-NEXT: ret;820;821; CHECK-F32X2-LABEL: test_frem_ftz(822; CHECK-F32X2: {823; CHECK-F32X2-NEXT: .reg .pred %p<3>;824; CHECK-F32X2-NEXT: .reg .b32 %r<15>;825; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;826; CHECK-F32X2-EMPTY:827; CHECK-F32X2-NEXT: // %bb.0:828; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_frem_ftz_param_1];829; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_frem_ftz_param_0];830; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;831; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;832; CHECK-F32X2-NEXT: div.rn.ftz.f32 %r5, %r4, %r2;833; CHECK-F32X2-NEXT: cvt.rzi.ftz.f32.f32 %r6, %r5;834; CHECK-F32X2-NEXT: neg.ftz.f32 %r7, %r6;835; CHECK-F32X2-NEXT: fma.rn.ftz.f32 %r8, %r7, %r2, %r4;836; CHECK-F32X2-NEXT: testp.infinite.f32 %p1, %r2;837; CHECK-F32X2-NEXT: selp.f32 %r9, %r4, %r8, %p1;838; CHECK-F32X2-NEXT: div.rn.ftz.f32 %r10, %r3, %r1;839; CHECK-F32X2-NEXT: cvt.rzi.ftz.f32.f32 %r11, %r10;840; CHECK-F32X2-NEXT: neg.ftz.f32 %r12, %r11;841; CHECK-F32X2-NEXT: fma.rn.ftz.f32 %r13, %r12, %r1, %r3;842; CHECK-F32X2-NEXT: testp.infinite.f32 %p2, %r1;843; CHECK-F32X2-NEXT: selp.f32 %r14, %r3, %r13, %p2;844; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r9};845; CHECK-F32X2-NEXT: ret;846 %r = frem <2 x float> %a, %b847 ret <2 x float> %r848}849 850define void @test_ldst_v2f32(ptr %a, ptr %b) #0 {851; CHECK-NOF32X2-LABEL: test_ldst_v2f32(852; CHECK-NOF32X2: {853; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;854; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;855; CHECK-NOF32X2-EMPTY:856; CHECK-NOF32X2-NEXT: // %bb.0:857; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v2f32_param_1];858; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v2f32_param_0];859; CHECK-NOF32X2-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];860; CHECK-NOF32X2-NEXT: st.v2.b32 [%rd2], {%r1, %r2};861; CHECK-NOF32X2-NEXT: ret;862;863; CHECK-F32X2-LABEL: test_ldst_v2f32(864; CHECK-F32X2: {865; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;866; CHECK-F32X2-EMPTY:867; CHECK-F32X2-NEXT: // %bb.0:868; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v2f32_param_1];869; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v2f32_param_0];870; CHECK-F32X2-NEXT: ld.b64 %rd3, [%rd1];871; CHECK-F32X2-NEXT: st.b64 [%rd2], %rd3;872; CHECK-F32X2-NEXT: ret;873 %t1 = load <2 x float>, ptr %a874 store <2 x float> %t1, ptr %b, align 32875 ret void876}877 878define void @test_ldst_v3f32(ptr %a, ptr %b) #0 {879; CHECK-LABEL: test_ldst_v3f32(880; CHECK: {881; CHECK-NEXT: .reg .b32 %r<2>;882; CHECK-NEXT: .reg .b64 %rd<4>;883; CHECK-EMPTY:884; CHECK-NEXT: // %bb.0:885; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v3f32_param_1];886; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v3f32_param_0];887; CHECK-NEXT: ld.b64 %rd3, [%rd1];888; CHECK-NEXT: ld.b32 %r1, [%rd1+8];889; CHECK-NEXT: st.b32 [%rd2+8], %r1;890; CHECK-NEXT: st.b64 [%rd2], %rd3;891; CHECK-NEXT: ret;892 %t1 = load <3 x float>, ptr %a893 store <3 x float> %t1, ptr %b, align 32894 ret void895}896 897define void @test_ldst_v4f32(ptr %a, ptr %b) #0 {898; CHECK-NOF32X2-LABEL: test_ldst_v4f32(899; CHECK-NOF32X2: {900; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;901; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;902; CHECK-NOF32X2-EMPTY:903; CHECK-NOF32X2-NEXT: // %bb.0:904; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v4f32_param_1];905; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v4f32_param_0];906; CHECK-NOF32X2-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];907; CHECK-NOF32X2-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};908; CHECK-NOF32X2-NEXT: ret;909;910; CHECK-F32X2-LABEL: test_ldst_v4f32(911; CHECK-F32X2: {912; CHECK-F32X2-NEXT: .reg .b64 %rd<5>;913; CHECK-F32X2-EMPTY:914; CHECK-F32X2-NEXT: // %bb.0:915; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v4f32_param_1];916; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v4f32_param_0];917; CHECK-F32X2-NEXT: ld.v2.b64 {%rd3, %rd4}, [%rd1];918; CHECK-F32X2-NEXT: st.v2.b64 [%rd2], {%rd3, %rd4};919; CHECK-F32X2-NEXT: ret;920 %t1 = load <4 x float>, ptr %a921 store <4 x float> %t1, ptr %b, align 32922 ret void923}924 925define void @test_ldst_v8f32(ptr %a, ptr %b) #0 {926; CHECK-NOF32X2-LABEL: test_ldst_v8f32(927; CHECK-NOF32X2: {928; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;929; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;930; CHECK-NOF32X2-EMPTY:931; CHECK-NOF32X2-NEXT: // %bb.0:932; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v8f32_param_1];933; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v8f32_param_0];934; CHECK-NOF32X2-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];935; CHECK-NOF32X2-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];936; CHECK-NOF32X2-NEXT: st.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};937; CHECK-NOF32X2-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};938; CHECK-NOF32X2-NEXT: ret;939;940; CHECK-F32X2-LABEL: test_ldst_v8f32(941; CHECK-F32X2: {942; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;943; CHECK-F32X2-EMPTY:944; CHECK-F32X2-NEXT: // %bb.0:945; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_ldst_v8f32_param_1];946; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_ldst_v8f32_param_0];947; CHECK-F32X2-NEXT: ld.v2.b64 {%rd3, %rd4}, [%rd1];948; CHECK-F32X2-NEXT: ld.v2.b64 {%rd5, %rd6}, [%rd1+16];949; CHECK-F32X2-NEXT: st.v2.b64 [%rd2+16], {%rd5, %rd6};950; CHECK-F32X2-NEXT: st.v2.b64 [%rd2], {%rd3, %rd4};951; CHECK-F32X2-NEXT: ret;952 %t1 = load <8 x float>, ptr %a953 store <8 x float> %t1, ptr %b, align 32954 ret void955}956 957declare <2 x float> @test_callee(<2 x float> %a, <2 x float> %b) #0958 959define <2 x float> @test_call(<2 x float> %a, <2 x float> %b) #0 {960; CHECK-NOF32X2-LABEL: test_call(961; CHECK-NOF32X2: {962; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;963; CHECK-NOF32X2-EMPTY:964; CHECK-NOF32X2-NEXT: // %bb.0:965; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_call_param_1];966; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_call_param_0];967; CHECK-NOF32X2-NEXT: { // callseq 0, 0968; CHECK-NOF32X2-NEXT: .param .align 8 .b8 param0[8];969; CHECK-NOF32X2-NEXT: .param .align 8 .b8 param1[8];970; CHECK-NOF32X2-NEXT: .param .align 8 .b8 retval0[8];971; CHECK-NOF32X2-NEXT: st.param.v2.b32 [param1], {%r3, %r4};972; CHECK-NOF32X2-NEXT: st.param.v2.b32 [param0], {%r1, %r2};973; CHECK-NOF32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);974; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [retval0];975; CHECK-NOF32X2-NEXT: } // callseq 0976; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r5, %r6};977; CHECK-NOF32X2-NEXT: ret;978;979; CHECK-F32X2-LABEL: test_call(980; CHECK-F32X2: {981; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;982; CHECK-F32X2-EMPTY:983; CHECK-F32X2-NEXT: // %bb.0:984; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_call_param_1];985; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_call_param_0];986; CHECK-F32X2-NEXT: { // callseq 0, 0987; CHECK-F32X2-NEXT: .param .align 8 .b8 param0[8];988; CHECK-F32X2-NEXT: .param .align 8 .b8 param1[8];989; CHECK-F32X2-NEXT: .param .align 8 .b8 retval0[8];990; CHECK-F32X2-NEXT: st.param.b64 [param1], %rd2;991; CHECK-F32X2-NEXT: st.param.b64 [param0], %rd1;992; CHECK-F32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);993; CHECK-F32X2-NEXT: ld.param.b64 %rd3, [retval0];994; CHECK-F32X2-NEXT: } // callseq 0995; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;996; CHECK-F32X2-NEXT: ret;997 %r = call <2 x float> @test_callee(<2 x float> %a, <2 x float> %b)998 ret <2 x float> %r999}1000 1001define <2 x float> @test_call_flipped(<2 x float> %a, <2 x float> %b) #0 {1002; CHECK-NOF32X2-LABEL: test_call_flipped(1003; CHECK-NOF32X2: {1004; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;1005; CHECK-NOF32X2-EMPTY:1006; CHECK-NOF32X2-NEXT: // %bb.0:1007; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_call_flipped_param_1];1008; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_call_flipped_param_0];1009; CHECK-NOF32X2-NEXT: { // callseq 1, 01010; CHECK-NOF32X2-NEXT: .param .align 8 .b8 param0[8];1011; CHECK-NOF32X2-NEXT: .param .align 8 .b8 param1[8];1012; CHECK-NOF32X2-NEXT: .param .align 8 .b8 retval0[8];1013; CHECK-NOF32X2-NEXT: st.param.v2.b32 [param1], {%r1, %r2};1014; CHECK-NOF32X2-NEXT: st.param.v2.b32 [param0], {%r3, %r4};1015; CHECK-NOF32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);1016; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [retval0];1017; CHECK-NOF32X2-NEXT: } // callseq 11018; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r5, %r6};1019; CHECK-NOF32X2-NEXT: ret;1020;1021; CHECK-F32X2-LABEL: test_call_flipped(1022; CHECK-F32X2: {1023; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;1024; CHECK-F32X2-EMPTY:1025; CHECK-F32X2-NEXT: // %bb.0:1026; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_call_flipped_param_1];1027; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_call_flipped_param_0];1028; CHECK-F32X2-NEXT: { // callseq 1, 01029; CHECK-F32X2-NEXT: .param .align 8 .b8 param0[8];1030; CHECK-F32X2-NEXT: .param .align 8 .b8 param1[8];1031; CHECK-F32X2-NEXT: .param .align 8 .b8 retval0[8];1032; CHECK-F32X2-NEXT: st.param.b64 [param1], %rd1;1033; CHECK-F32X2-NEXT: st.param.b64 [param0], %rd2;1034; CHECK-F32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);1035; CHECK-F32X2-NEXT: ld.param.b64 %rd3, [retval0];1036; CHECK-F32X2-NEXT: } // callseq 11037; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;1038; CHECK-F32X2-NEXT: ret;1039 %r = call <2 x float> @test_callee(<2 x float> %b, <2 x float> %a)1040 ret <2 x float> %r1041}1042 1043define <2 x float> @test_tailcall_flipped(<2 x float> %a, <2 x float> %b) #0 {1044; CHECK-NOF32X2-LABEL: test_tailcall_flipped(1045; CHECK-NOF32X2: {1046; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;1047; CHECK-NOF32X2-EMPTY:1048; CHECK-NOF32X2-NEXT: // %bb.0:1049; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_tailcall_flipped_param_1];1050; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_tailcall_flipped_param_0];1051; CHECK-NOF32X2-NEXT: { // callseq 2, 01052; CHECK-NOF32X2-NEXT: .param .align 8 .b8 param0[8];1053; CHECK-NOF32X2-NEXT: .param .align 8 .b8 param1[8];1054; CHECK-NOF32X2-NEXT: .param .align 8 .b8 retval0[8];1055; CHECK-NOF32X2-NEXT: st.param.v2.b32 [param1], {%r1, %r2};1056; CHECK-NOF32X2-NEXT: st.param.v2.b32 [param0], {%r3, %r4};1057; CHECK-NOF32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);1058; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [retval0];1059; CHECK-NOF32X2-NEXT: } // callseq 21060; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r5, %r6};1061; CHECK-NOF32X2-NEXT: ret;1062;1063; CHECK-F32X2-LABEL: test_tailcall_flipped(1064; CHECK-F32X2: {1065; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;1066; CHECK-F32X2-EMPTY:1067; CHECK-F32X2-NEXT: // %bb.0:1068; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_tailcall_flipped_param_1];1069; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_tailcall_flipped_param_0];1070; CHECK-F32X2-NEXT: { // callseq 2, 01071; CHECK-F32X2-NEXT: .param .align 8 .b8 param0[8];1072; CHECK-F32X2-NEXT: .param .align 8 .b8 param1[8];1073; CHECK-F32X2-NEXT: .param .align 8 .b8 retval0[8];1074; CHECK-F32X2-NEXT: st.param.b64 [param1], %rd1;1075; CHECK-F32X2-NEXT: st.param.b64 [param0], %rd2;1076; CHECK-F32X2-NEXT: call.uni (retval0), test_callee, (param0, param1);1077; CHECK-F32X2-NEXT: ld.param.b64 %rd3, [retval0];1078; CHECK-F32X2-NEXT: } // callseq 21079; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;1080; CHECK-F32X2-NEXT: ret;1081 %r = tail call <2 x float> @test_callee(<2 x float> %b, <2 x float> %a)1082 ret <2 x float> %r1083}1084 1085define <2 x float> @test_select(<2 x float> %a, <2 x float> %b, i1 zeroext %c) #0 {1086; CHECK-NOF32X2-LABEL: test_select(1087; CHECK-NOF32X2: {1088; CHECK-NOF32X2-NEXT: .reg .pred %p<2>;1089; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1090; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;1091; CHECK-NOF32X2-EMPTY:1092; CHECK-NOF32X2-NEXT: // %bb.0:1093; CHECK-NOF32X2-NEXT: ld.param.b8 %rs1, [test_select_param_2];1094; CHECK-NOF32X2-NEXT: and.b16 %rs2, %rs1, 1;1095; CHECK-NOF32X2-NEXT: setp.ne.b16 %p1, %rs2, 0;1096; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_param_1];1097; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_param_0];1098; CHECK-NOF32X2-NEXT: selp.f32 %r5, %r2, %r4, %p1;1099; CHECK-NOF32X2-NEXT: selp.f32 %r6, %r1, %r3, %p1;1100; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};1101; CHECK-NOF32X2-NEXT: ret;1102;1103; CHECK-F32X2-LABEL: test_select(1104; CHECK-F32X2: {1105; CHECK-F32X2-NEXT: .reg .pred %p<2>;1106; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1107; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;1108; CHECK-F32X2-EMPTY:1109; CHECK-F32X2-NEXT: // %bb.0:1110; CHECK-F32X2-NEXT: ld.param.b8 %rs1, [test_select_param_2];1111; CHECK-F32X2-NEXT: and.b16 %rs2, %rs1, 1;1112; CHECK-F32X2-NEXT: setp.ne.b16 %p1, %rs2, 0;1113; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_select_param_1];1114; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_select_param_0];1115; CHECK-F32X2-NEXT: selp.b64 %rd3, %rd1, %rd2, %p1;1116; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;1117; CHECK-F32X2-NEXT: ret;1118 %r = select i1 %c, <2 x float> %a, <2 x float> %b1119 ret <2 x float> %r1120}1121 1122define <2 x float> @test_select_cc(<2 x float> %a, <2 x float> %b, <2 x float> %c, <2 x float> %d) #0 {1123; CHECK-NOF32X2-LABEL: test_select_cc(1124; CHECK-NOF32X2: {1125; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1126; CHECK-NOF32X2-NEXT: .reg .b32 %r<11>;1127; CHECK-NOF32X2-EMPTY:1128; CHECK-NOF32X2-NEXT: // %bb.0:1129; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r7, %r8}, [test_select_cc_param_3];1130; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_param_2];1131; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_param_1];1132; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_param_0];1133; CHECK-NOF32X2-NEXT: setp.neu.f32 %p1, %r5, %r7;1134; CHECK-NOF32X2-NEXT: setp.neu.f32 %p2, %r6, %r8;1135; CHECK-NOF32X2-NEXT: selp.f32 %r9, %r2, %r4, %p2;1136; CHECK-NOF32X2-NEXT: selp.f32 %r10, %r1, %r3, %p1;1137; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r10, %r9};1138; CHECK-NOF32X2-NEXT: ret;1139;1140; CHECK-F32X2-LABEL: test_select_cc(1141; CHECK-F32X2: {1142; CHECK-F32X2-NEXT: .reg .pred %p<3>;1143; CHECK-F32X2-NEXT: .reg .b32 %r<11>;1144; CHECK-F32X2-NEXT: .reg .b64 %rd<5>;1145; CHECK-F32X2-EMPTY:1146; CHECK-F32X2-NEXT: // %bb.0:1147; CHECK-F32X2-NEXT: ld.param.b64 %rd4, [test_select_cc_param_3];1148; CHECK-F32X2-NEXT: ld.param.b64 %rd3, [test_select_cc_param_2];1149; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_select_cc_param_1];1150; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_select_cc_param_0];1151; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd4;1152; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd3;1153; CHECK-F32X2-NEXT: setp.neu.f32 %p1, %r3, %r1;1154; CHECK-F32X2-NEXT: setp.neu.f32 %p2, %r4, %r2;1155; CHECK-F32X2-NEXT: mov.b64 {%r5, %r6}, %rd2;1156; CHECK-F32X2-NEXT: mov.b64 {%r7, %r8}, %rd1;1157; CHECK-F32X2-NEXT: selp.f32 %r9, %r8, %r6, %p2;1158; CHECK-F32X2-NEXT: selp.f32 %r10, %r7, %r5, %p1;1159; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r10, %r9};1160; CHECK-F32X2-NEXT: ret;1161 %cc = fcmp une <2 x float> %c, %d1162 %r = select <2 x i1> %cc, <2 x float> %a, <2 x float> %b1163 ret <2 x float> %r1164}1165 1166define <2 x double> @test_select_cc_f64_f32(<2 x double> %a, <2 x double> %b, <2 x float> %c, <2 x float> %d) #0 {1167; CHECK-NOF32X2-LABEL: test_select_cc_f64_f32(1168; CHECK-NOF32X2: {1169; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1170; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1171; CHECK-NOF32X2-NEXT: .reg .b64 %rd<7>;1172; CHECK-NOF32X2-EMPTY:1173; CHECK-NOF32X2-NEXT: // %bb.0:1174; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f64_f32_param_3];1175; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f64_f32_param_2];1176; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_select_cc_f64_f32_param_1];1177; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_select_cc_f64_f32_param_0];1178; CHECK-NOF32X2-NEXT: setp.neu.f32 %p1, %r1, %r3;1179; CHECK-NOF32X2-NEXT: setp.neu.f32 %p2, %r2, %r4;1180; CHECK-NOF32X2-NEXT: selp.f64 %rd5, %rd2, %rd4, %p2;1181; CHECK-NOF32X2-NEXT: selp.f64 %rd6, %rd1, %rd3, %p1;1182; CHECK-NOF32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd5};1183; CHECK-NOF32X2-NEXT: ret;1184;1185; CHECK-F32X2-LABEL: test_select_cc_f64_f32(1186; CHECK-F32X2: {1187; CHECK-F32X2-NEXT: .reg .pred %p<3>;1188; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1189; CHECK-F32X2-NEXT: .reg .b64 %rd<9>;1190; CHECK-F32X2-EMPTY:1191; CHECK-F32X2-NEXT: // %bb.0:1192; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_select_cc_f64_f32_param_1];1193; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_select_cc_f64_f32_param_0];1194; CHECK-F32X2-NEXT: ld.param.b64 %rd6, [test_select_cc_f64_f32_param_3];1195; CHECK-F32X2-NEXT: ld.param.b64 %rd5, [test_select_cc_f64_f32_param_2];1196; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd6;1197; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd5;1198; CHECK-F32X2-NEXT: setp.neu.f32 %p1, %r3, %r1;1199; CHECK-F32X2-NEXT: setp.neu.f32 %p2, %r4, %r2;1200; CHECK-F32X2-NEXT: selp.f64 %rd7, %rd2, %rd4, %p2;1201; CHECK-F32X2-NEXT: selp.f64 %rd8, %rd1, %rd3, %p1;1202; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd8, %rd7};1203; CHECK-F32X2-NEXT: ret;1204 %cc = fcmp une <2 x float> %c, %d1205 %r = select <2 x i1> %cc, <2 x double> %a, <2 x double> %b1206 ret <2 x double> %r1207}1208 1209define <2 x float> @test_select_cc_f32_f64(<2 x float> %a, <2 x float> %b, <2 x double> %c, <2 x double> %d) #0 {1210; CHECK-NOF32X2-LABEL: test_select_cc_f32_f64(1211; CHECK-NOF32X2: {1212; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1213; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;1214; CHECK-NOF32X2-NEXT: .reg .b64 %rd<5>;1215; CHECK-NOF32X2-EMPTY:1216; CHECK-NOF32X2-NEXT: // %bb.0:1217; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_select_cc_f32_f64_param_3];1218; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_select_cc_f32_f64_param_2];1219; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f32_f64_param_1];1220; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f32_f64_param_0];1221; CHECK-NOF32X2-NEXT: setp.neu.f64 %p1, %rd1, %rd3;1222; CHECK-NOF32X2-NEXT: setp.neu.f64 %p2, %rd2, %rd4;1223; CHECK-NOF32X2-NEXT: selp.f32 %r5, %r2, %r4, %p2;1224; CHECK-NOF32X2-NEXT: selp.f32 %r6, %r1, %r3, %p1;1225; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};1226; CHECK-NOF32X2-NEXT: ret;1227;1228; CHECK-F32X2-LABEL: test_select_cc_f32_f64(1229; CHECK-F32X2: {1230; CHECK-F32X2-NEXT: .reg .pred %p<3>;1231; CHECK-F32X2-NEXT: .reg .b32 %r<7>;1232; CHECK-F32X2-NEXT: .reg .b64 %rd<7>;1233; CHECK-F32X2-EMPTY:1234; CHECK-F32X2-NEXT: // %bb.0:1235; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd5, %rd6}, [test_select_cc_f32_f64_param_3];1236; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_select_cc_f32_f64_param_2];1237; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_select_cc_f32_f64_param_1];1238; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_select_cc_f32_f64_param_0];1239; CHECK-F32X2-NEXT: setp.neu.f64 %p1, %rd3, %rd5;1240; CHECK-F32X2-NEXT: setp.neu.f64 %p2, %rd4, %rd6;1241; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1242; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1243; CHECK-F32X2-NEXT: selp.f32 %r5, %r4, %r2, %p2;1244; CHECK-F32X2-NEXT: selp.f32 %r6, %r3, %r1, %p1;1245; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};1246; CHECK-F32X2-NEXT: ret;1247 %cc = fcmp une <2 x double> %c, %d1248 %r = select <2 x i1> %cc, <2 x float> %a, <2 x float> %b1249 ret <2 x float> %r1250}1251 1252define <2 x i1> @test_fcmp_une(<2 x float> %a, <2 x float> %b) #0 {1253; CHECK-NOF32X2-LABEL: test_fcmp_une(1254; CHECK-NOF32X2: {1255; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1256; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1257; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1258; CHECK-NOF32X2-EMPTY:1259; CHECK-NOF32X2-NEXT: // %bb.0:1260; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_une_param_1];1261; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_une_param_0];1262; CHECK-NOF32X2-NEXT: setp.neu.f32 %p1, %r2, %r4;1263; CHECK-NOF32X2-NEXT: setp.neu.f32 %p2, %r1, %r3;1264; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1265; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1266; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1267; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1268; CHECK-NOF32X2-NEXT: ret;1269;1270; CHECK-F32X2-LABEL: test_fcmp_une(1271; CHECK-F32X2: {1272; CHECK-F32X2-NEXT: .reg .pred %p<3>;1273; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1274; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1275; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1276; CHECK-F32X2-EMPTY:1277; CHECK-F32X2-NEXT: // %bb.0:1278; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_une_param_1];1279; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_une_param_0];1280; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1281; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1282; CHECK-F32X2-NEXT: setp.neu.f32 %p1, %r4, %r2;1283; CHECK-F32X2-NEXT: setp.neu.f32 %p2, %r3, %r1;1284; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1285; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1286; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1287; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1288; CHECK-F32X2-NEXT: ret;1289 %r = fcmp une <2 x float> %a, %b1290 ret <2 x i1> %r1291}1292 1293define <2 x i1> @test_fcmp_ueq(<2 x float> %a, <2 x float> %b) #0 {1294; CHECK-NOF32X2-LABEL: test_fcmp_ueq(1295; CHECK-NOF32X2: {1296; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1297; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1298; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1299; CHECK-NOF32X2-EMPTY:1300; CHECK-NOF32X2-NEXT: // %bb.0:1301; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ueq_param_1];1302; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ueq_param_0];1303; CHECK-NOF32X2-NEXT: setp.equ.f32 %p1, %r2, %r4;1304; CHECK-NOF32X2-NEXT: setp.equ.f32 %p2, %r1, %r3;1305; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1306; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1307; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1308; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1309; CHECK-NOF32X2-NEXT: ret;1310;1311; CHECK-F32X2-LABEL: test_fcmp_ueq(1312; CHECK-F32X2: {1313; CHECK-F32X2-NEXT: .reg .pred %p<3>;1314; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1315; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1316; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1317; CHECK-F32X2-EMPTY:1318; CHECK-F32X2-NEXT: // %bb.0:1319; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_ueq_param_1];1320; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_ueq_param_0];1321; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1322; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1323; CHECK-F32X2-NEXT: setp.equ.f32 %p1, %r4, %r2;1324; CHECK-F32X2-NEXT: setp.equ.f32 %p2, %r3, %r1;1325; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1326; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1327; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1328; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1329; CHECK-F32X2-NEXT: ret;1330 %r = fcmp ueq <2 x float> %a, %b1331 ret <2 x i1> %r1332}1333 1334define <2 x i1> @test_fcmp_ugt(<2 x float> %a, <2 x float> %b) #0 {1335; CHECK-NOF32X2-LABEL: test_fcmp_ugt(1336; CHECK-NOF32X2: {1337; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1338; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1339; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1340; CHECK-NOF32X2-EMPTY:1341; CHECK-NOF32X2-NEXT: // %bb.0:1342; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ugt_param_1];1343; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ugt_param_0];1344; CHECK-NOF32X2-NEXT: setp.gtu.f32 %p1, %r2, %r4;1345; CHECK-NOF32X2-NEXT: setp.gtu.f32 %p2, %r1, %r3;1346; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1347; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1348; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1349; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1350; CHECK-NOF32X2-NEXT: ret;1351;1352; CHECK-F32X2-LABEL: test_fcmp_ugt(1353; CHECK-F32X2: {1354; CHECK-F32X2-NEXT: .reg .pred %p<3>;1355; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1356; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1357; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1358; CHECK-F32X2-EMPTY:1359; CHECK-F32X2-NEXT: // %bb.0:1360; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_ugt_param_1];1361; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_ugt_param_0];1362; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1363; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1364; CHECK-F32X2-NEXT: setp.gtu.f32 %p1, %r4, %r2;1365; CHECK-F32X2-NEXT: setp.gtu.f32 %p2, %r3, %r1;1366; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1367; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1368; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1369; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1370; CHECK-F32X2-NEXT: ret;1371 %r = fcmp ugt <2 x float> %a, %b1372 ret <2 x i1> %r1373}1374 1375define <2 x i1> @test_fcmp_uge(<2 x float> %a, <2 x float> %b) #0 {1376; CHECK-NOF32X2-LABEL: test_fcmp_uge(1377; CHECK-NOF32X2: {1378; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1379; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1380; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1381; CHECK-NOF32X2-EMPTY:1382; CHECK-NOF32X2-NEXT: // %bb.0:1383; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_uge_param_1];1384; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_uge_param_0];1385; CHECK-NOF32X2-NEXT: setp.geu.f32 %p1, %r2, %r4;1386; CHECK-NOF32X2-NEXT: setp.geu.f32 %p2, %r1, %r3;1387; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1388; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1389; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1390; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1391; CHECK-NOF32X2-NEXT: ret;1392;1393; CHECK-F32X2-LABEL: test_fcmp_uge(1394; CHECK-F32X2: {1395; CHECK-F32X2-NEXT: .reg .pred %p<3>;1396; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1397; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1398; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1399; CHECK-F32X2-EMPTY:1400; CHECK-F32X2-NEXT: // %bb.0:1401; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_uge_param_1];1402; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_uge_param_0];1403; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1404; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1405; CHECK-F32X2-NEXT: setp.geu.f32 %p1, %r4, %r2;1406; CHECK-F32X2-NEXT: setp.geu.f32 %p2, %r3, %r1;1407; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1408; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1409; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1410; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1411; CHECK-F32X2-NEXT: ret;1412 %r = fcmp uge <2 x float> %a, %b1413 ret <2 x i1> %r1414}1415 1416define <2 x i1> @test_fcmp_ult(<2 x float> %a, <2 x float> %b) #0 {1417; CHECK-NOF32X2-LABEL: test_fcmp_ult(1418; CHECK-NOF32X2: {1419; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1420; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1421; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1422; CHECK-NOF32X2-EMPTY:1423; CHECK-NOF32X2-NEXT: // %bb.0:1424; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ult_param_1];1425; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ult_param_0];1426; CHECK-NOF32X2-NEXT: setp.ltu.f32 %p1, %r2, %r4;1427; CHECK-NOF32X2-NEXT: setp.ltu.f32 %p2, %r1, %r3;1428; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1429; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1430; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1431; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1432; CHECK-NOF32X2-NEXT: ret;1433;1434; CHECK-F32X2-LABEL: test_fcmp_ult(1435; CHECK-F32X2: {1436; CHECK-F32X2-NEXT: .reg .pred %p<3>;1437; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1438; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1439; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1440; CHECK-F32X2-EMPTY:1441; CHECK-F32X2-NEXT: // %bb.0:1442; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_ult_param_1];1443; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_ult_param_0];1444; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1445; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1446; CHECK-F32X2-NEXT: setp.ltu.f32 %p1, %r4, %r2;1447; CHECK-F32X2-NEXT: setp.ltu.f32 %p2, %r3, %r1;1448; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1449; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1450; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1451; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1452; CHECK-F32X2-NEXT: ret;1453 %r = fcmp ult <2 x float> %a, %b1454 ret <2 x i1> %r1455}1456 1457define <2 x i1> @test_fcmp_ule(<2 x float> %a, <2 x float> %b) #0 {1458; CHECK-NOF32X2-LABEL: test_fcmp_ule(1459; CHECK-NOF32X2: {1460; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1461; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1462; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1463; CHECK-NOF32X2-EMPTY:1464; CHECK-NOF32X2-NEXT: // %bb.0:1465; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ule_param_1];1466; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ule_param_0];1467; CHECK-NOF32X2-NEXT: setp.leu.f32 %p1, %r2, %r4;1468; CHECK-NOF32X2-NEXT: setp.leu.f32 %p2, %r1, %r3;1469; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1470; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1471; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1472; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1473; CHECK-NOF32X2-NEXT: ret;1474;1475; CHECK-F32X2-LABEL: test_fcmp_ule(1476; CHECK-F32X2: {1477; CHECK-F32X2-NEXT: .reg .pred %p<3>;1478; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1479; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1480; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1481; CHECK-F32X2-EMPTY:1482; CHECK-F32X2-NEXT: // %bb.0:1483; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_ule_param_1];1484; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_ule_param_0];1485; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1486; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1487; CHECK-F32X2-NEXT: setp.leu.f32 %p1, %r4, %r2;1488; CHECK-F32X2-NEXT: setp.leu.f32 %p2, %r3, %r1;1489; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1490; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1491; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1492; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1493; CHECK-F32X2-NEXT: ret;1494 %r = fcmp ule <2 x float> %a, %b1495 ret <2 x i1> %r1496}1497 1498define <2 x i1> @test_fcmp_uno(<2 x float> %a, <2 x float> %b) #0 {1499; CHECK-NOF32X2-LABEL: test_fcmp_uno(1500; CHECK-NOF32X2: {1501; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1502; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1503; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1504; CHECK-NOF32X2-EMPTY:1505; CHECK-NOF32X2-NEXT: // %bb.0:1506; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_uno_param_1];1507; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_uno_param_0];1508; CHECK-NOF32X2-NEXT: setp.nan.f32 %p1, %r2, %r4;1509; CHECK-NOF32X2-NEXT: setp.nan.f32 %p2, %r1, %r3;1510; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1511; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1512; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1513; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1514; CHECK-NOF32X2-NEXT: ret;1515;1516; CHECK-F32X2-LABEL: test_fcmp_uno(1517; CHECK-F32X2: {1518; CHECK-F32X2-NEXT: .reg .pred %p<3>;1519; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1520; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1521; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1522; CHECK-F32X2-EMPTY:1523; CHECK-F32X2-NEXT: // %bb.0:1524; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_uno_param_1];1525; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_uno_param_0];1526; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1527; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1528; CHECK-F32X2-NEXT: setp.nan.f32 %p1, %r4, %r2;1529; CHECK-F32X2-NEXT: setp.nan.f32 %p2, %r3, %r1;1530; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1531; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1532; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1533; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1534; CHECK-F32X2-NEXT: ret;1535 %r = fcmp uno <2 x float> %a, %b1536 ret <2 x i1> %r1537}1538 1539define <2 x i1> @test_fcmp_one(<2 x float> %a, <2 x float> %b) #0 {1540; CHECK-NOF32X2-LABEL: test_fcmp_one(1541; CHECK-NOF32X2: {1542; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1543; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1544; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1545; CHECK-NOF32X2-EMPTY:1546; CHECK-NOF32X2-NEXT: // %bb.0:1547; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_one_param_1];1548; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_one_param_0];1549; CHECK-NOF32X2-NEXT: setp.ne.f32 %p1, %r2, %r4;1550; CHECK-NOF32X2-NEXT: setp.ne.f32 %p2, %r1, %r3;1551; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1552; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1553; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1554; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1555; CHECK-NOF32X2-NEXT: ret;1556;1557; CHECK-F32X2-LABEL: test_fcmp_one(1558; CHECK-F32X2: {1559; CHECK-F32X2-NEXT: .reg .pred %p<3>;1560; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1561; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1562; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1563; CHECK-F32X2-EMPTY:1564; CHECK-F32X2-NEXT: // %bb.0:1565; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_one_param_1];1566; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_one_param_0];1567; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1568; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1569; CHECK-F32X2-NEXT: setp.ne.f32 %p1, %r4, %r2;1570; CHECK-F32X2-NEXT: setp.ne.f32 %p2, %r3, %r1;1571; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1572; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1573; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1574; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1575; CHECK-F32X2-NEXT: ret;1576 %r = fcmp one <2 x float> %a, %b1577 ret <2 x i1> %r1578}1579 1580define <2 x i1> @test_fcmp_oeq(<2 x float> %a, <2 x float> %b) #0 {1581; CHECK-NOF32X2-LABEL: test_fcmp_oeq(1582; CHECK-NOF32X2: {1583; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1584; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1585; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1586; CHECK-NOF32X2-EMPTY:1587; CHECK-NOF32X2-NEXT: // %bb.0:1588; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_oeq_param_1];1589; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_oeq_param_0];1590; CHECK-NOF32X2-NEXT: setp.eq.f32 %p1, %r2, %r4;1591; CHECK-NOF32X2-NEXT: setp.eq.f32 %p2, %r1, %r3;1592; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1593; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1594; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1595; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1596; CHECK-NOF32X2-NEXT: ret;1597;1598; CHECK-F32X2-LABEL: test_fcmp_oeq(1599; CHECK-F32X2: {1600; CHECK-F32X2-NEXT: .reg .pred %p<3>;1601; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1602; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1603; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1604; CHECK-F32X2-EMPTY:1605; CHECK-F32X2-NEXT: // %bb.0:1606; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_oeq_param_1];1607; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_oeq_param_0];1608; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1609; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1610; CHECK-F32X2-NEXT: setp.eq.f32 %p1, %r4, %r2;1611; CHECK-F32X2-NEXT: setp.eq.f32 %p2, %r3, %r1;1612; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1613; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1614; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1615; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1616; CHECK-F32X2-NEXT: ret;1617 %r = fcmp oeq <2 x float> %a, %b1618 ret <2 x i1> %r1619}1620 1621define <2 x i1> @test_fcmp_ogt(<2 x float> %a, <2 x float> %b) #0 {1622; CHECK-NOF32X2-LABEL: test_fcmp_ogt(1623; CHECK-NOF32X2: {1624; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1625; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1626; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1627; CHECK-NOF32X2-EMPTY:1628; CHECK-NOF32X2-NEXT: // %bb.0:1629; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ogt_param_1];1630; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ogt_param_0];1631; CHECK-NOF32X2-NEXT: setp.gt.f32 %p1, %r2, %r4;1632; CHECK-NOF32X2-NEXT: setp.gt.f32 %p2, %r1, %r3;1633; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1634; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1635; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1636; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1637; CHECK-NOF32X2-NEXT: ret;1638;1639; CHECK-F32X2-LABEL: test_fcmp_ogt(1640; CHECK-F32X2: {1641; CHECK-F32X2-NEXT: .reg .pred %p<3>;1642; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1643; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1644; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1645; CHECK-F32X2-EMPTY:1646; CHECK-F32X2-NEXT: // %bb.0:1647; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_ogt_param_1];1648; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_ogt_param_0];1649; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1650; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1651; CHECK-F32X2-NEXT: setp.gt.f32 %p1, %r4, %r2;1652; CHECK-F32X2-NEXT: setp.gt.f32 %p2, %r3, %r1;1653; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1654; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1655; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1656; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1657; CHECK-F32X2-NEXT: ret;1658 %r = fcmp ogt <2 x float> %a, %b1659 ret <2 x i1> %r1660}1661 1662define <2 x i1> @test_fcmp_oge(<2 x float> %a, <2 x float> %b) #0 {1663; CHECK-NOF32X2-LABEL: test_fcmp_oge(1664; CHECK-NOF32X2: {1665; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1666; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1667; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1668; CHECK-NOF32X2-EMPTY:1669; CHECK-NOF32X2-NEXT: // %bb.0:1670; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_oge_param_1];1671; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_oge_param_0];1672; CHECK-NOF32X2-NEXT: setp.ge.f32 %p1, %r2, %r4;1673; CHECK-NOF32X2-NEXT: setp.ge.f32 %p2, %r1, %r3;1674; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1675; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1676; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1677; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1678; CHECK-NOF32X2-NEXT: ret;1679;1680; CHECK-F32X2-LABEL: test_fcmp_oge(1681; CHECK-F32X2: {1682; CHECK-F32X2-NEXT: .reg .pred %p<3>;1683; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1684; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1685; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1686; CHECK-F32X2-EMPTY:1687; CHECK-F32X2-NEXT: // %bb.0:1688; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_oge_param_1];1689; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_oge_param_0];1690; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1691; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1692; CHECK-F32X2-NEXT: setp.ge.f32 %p1, %r4, %r2;1693; CHECK-F32X2-NEXT: setp.ge.f32 %p2, %r3, %r1;1694; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1695; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1696; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1697; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1698; CHECK-F32X2-NEXT: ret;1699 %r = fcmp oge <2 x float> %a, %b1700 ret <2 x i1> %r1701}1702 1703define <2 x i1> @test_fcmp_olt(<2 x float> %a, <2 x float> %b) #0 {1704; CHECK-NOF32X2-LABEL: test_fcmp_olt(1705; CHECK-NOF32X2: {1706; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1707; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1708; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1709; CHECK-NOF32X2-EMPTY:1710; CHECK-NOF32X2-NEXT: // %bb.0:1711; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_olt_param_1];1712; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_olt_param_0];1713; CHECK-NOF32X2-NEXT: setp.lt.f32 %p1, %r2, %r4;1714; CHECK-NOF32X2-NEXT: setp.lt.f32 %p2, %r1, %r3;1715; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1716; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1717; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1718; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1719; CHECK-NOF32X2-NEXT: ret;1720;1721; CHECK-F32X2-LABEL: test_fcmp_olt(1722; CHECK-F32X2: {1723; CHECK-F32X2-NEXT: .reg .pred %p<3>;1724; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1725; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1726; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1727; CHECK-F32X2-EMPTY:1728; CHECK-F32X2-NEXT: // %bb.0:1729; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_olt_param_1];1730; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_olt_param_0];1731; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1732; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1733; CHECK-F32X2-NEXT: setp.lt.f32 %p1, %r4, %r2;1734; CHECK-F32X2-NEXT: setp.lt.f32 %p2, %r3, %r1;1735; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1736; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1737; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1738; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1739; CHECK-F32X2-NEXT: ret;1740 %r = fcmp olt <2 x float> %a, %b1741 ret <2 x i1> %r1742}1743 1744define <2 x i1> @test_fcmp_ole(<2 x float> %a, <2 x float> %b) #0 {1745; CHECK-NOF32X2-LABEL: test_fcmp_ole(1746; CHECK-NOF32X2: {1747; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1748; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1749; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1750; CHECK-NOF32X2-EMPTY:1751; CHECK-NOF32X2-NEXT: // %bb.0:1752; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ole_param_1];1753; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ole_param_0];1754; CHECK-NOF32X2-NEXT: setp.le.f32 %p1, %r2, %r4;1755; CHECK-NOF32X2-NEXT: setp.le.f32 %p2, %r1, %r3;1756; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1757; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1758; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1759; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1760; CHECK-NOF32X2-NEXT: ret;1761;1762; CHECK-F32X2-LABEL: test_fcmp_ole(1763; CHECK-F32X2: {1764; CHECK-F32X2-NEXT: .reg .pred %p<3>;1765; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1766; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1767; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1768; CHECK-F32X2-EMPTY:1769; CHECK-F32X2-NEXT: // %bb.0:1770; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_ole_param_1];1771; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_ole_param_0];1772; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1773; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1774; CHECK-F32X2-NEXT: setp.le.f32 %p1, %r4, %r2;1775; CHECK-F32X2-NEXT: setp.le.f32 %p2, %r3, %r1;1776; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1777; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1778; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1779; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1780; CHECK-F32X2-NEXT: ret;1781 %r = fcmp ole <2 x float> %a, %b1782 ret <2 x i1> %r1783}1784 1785define <2 x i1> @test_fcmp_ord(<2 x float> %a, <2 x float> %b) #0 {1786; CHECK-NOF32X2-LABEL: test_fcmp_ord(1787; CHECK-NOF32X2: {1788; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;1789; CHECK-NOF32X2-NEXT: .reg .b16 %rs<3>;1790; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1791; CHECK-NOF32X2-EMPTY:1792; CHECK-NOF32X2-NEXT: // %bb.0:1793; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ord_param_1];1794; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ord_param_0];1795; CHECK-NOF32X2-NEXT: setp.num.f32 %p1, %r2, %r4;1796; CHECK-NOF32X2-NEXT: setp.num.f32 %p2, %r1, %r3;1797; CHECK-NOF32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1798; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0], %rs1;1799; CHECK-NOF32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1800; CHECK-NOF32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1801; CHECK-NOF32X2-NEXT: ret;1802;1803; CHECK-F32X2-LABEL: test_fcmp_ord(1804; CHECK-F32X2: {1805; CHECK-F32X2-NEXT: .reg .pred %p<3>;1806; CHECK-F32X2-NEXT: .reg .b16 %rs<3>;1807; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1808; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;1809; CHECK-F32X2-EMPTY:1810; CHECK-F32X2-NEXT: // %bb.0:1811; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fcmp_ord_param_1];1812; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fcmp_ord_param_0];1813; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;1814; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;1815; CHECK-F32X2-NEXT: setp.num.f32 %p1, %r4, %r2;1816; CHECK-F32X2-NEXT: setp.num.f32 %p2, %r3, %r1;1817; CHECK-F32X2-NEXT: selp.b16 %rs1, -1, 0, %p2;1818; CHECK-F32X2-NEXT: st.param.b8 [func_retval0], %rs1;1819; CHECK-F32X2-NEXT: selp.b16 %rs2, -1, 0, %p1;1820; CHECK-F32X2-NEXT: st.param.b8 [func_retval0+1], %rs2;1821; CHECK-F32X2-NEXT: ret;1822 %r = fcmp ord <2 x float> %a, %b1823 ret <2 x i1> %r1824}1825 1826define <2 x i32> @test_fptosi_i32(<2 x float> %a) #0 {1827; CHECK-NOF32X2-LABEL: test_fptosi_i32(1828; CHECK-NOF32X2: {1829; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1830; CHECK-NOF32X2-EMPTY:1831; CHECK-NOF32X2-NEXT: // %bb.0:1832; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_i32_param_0];1833; CHECK-NOF32X2-NEXT: cvt.rzi.s32.f32 %r3, %r2;1834; CHECK-NOF32X2-NEXT: cvt.rzi.s32.f32 %r4, %r1;1835; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1836; CHECK-NOF32X2-NEXT: ret;1837;1838; CHECK-F32X2-LABEL: test_fptosi_i32(1839; CHECK-F32X2: {1840; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1841; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;1842; CHECK-F32X2-EMPTY:1843; CHECK-F32X2-NEXT: // %bb.0:1844; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fptosi_i32_param_0];1845; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1846; CHECK-F32X2-NEXT: cvt.rzi.s32.f32 %r3, %r2;1847; CHECK-F32X2-NEXT: cvt.rzi.s32.f32 %r4, %r1;1848; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1849; CHECK-F32X2-NEXT: ret;1850 %r = fptosi <2 x float> %a to <2 x i32>1851 ret <2 x i32> %r1852}1853 1854define <2 x i64> @test_fptosi_i64(<2 x float> %a) #0 {1855; CHECK-NOF32X2-LABEL: test_fptosi_i64(1856; CHECK-NOF32X2: {1857; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;1858; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;1859; CHECK-NOF32X2-EMPTY:1860; CHECK-NOF32X2-NEXT: // %bb.0:1861; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_i64_param_0];1862; CHECK-NOF32X2-NEXT: cvt.rzi.s64.f32 %rd1, %r2;1863; CHECK-NOF32X2-NEXT: cvt.rzi.s64.f32 %rd2, %r1;1864; CHECK-NOF32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};1865; CHECK-NOF32X2-NEXT: ret;1866;1867; CHECK-F32X2-LABEL: test_fptosi_i64(1868; CHECK-F32X2: {1869; CHECK-F32X2-NEXT: .reg .b32 %r<3>;1870; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;1871; CHECK-F32X2-EMPTY:1872; CHECK-F32X2-NEXT: // %bb.0:1873; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fptosi_i64_param_0];1874; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1875; CHECK-F32X2-NEXT: cvt.rzi.s64.f32 %rd2, %r2;1876; CHECK-F32X2-NEXT: cvt.rzi.s64.f32 %rd3, %r1;1877; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};1878; CHECK-F32X2-NEXT: ret;1879 %r = fptosi <2 x float> %a to <2 x i64>1880 ret <2 x i64> %r1881}1882 1883define <2 x i32> @test_fptoui_2xi32(<2 x float> %a) #0 {1884; CHECK-NOF32X2-LABEL: test_fptoui_2xi32(1885; CHECK-NOF32X2: {1886; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1887; CHECK-NOF32X2-EMPTY:1888; CHECK-NOF32X2-NEXT: // %bb.0:1889; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_2xi32_param_0];1890; CHECK-NOF32X2-NEXT: cvt.rzi.u32.f32 %r3, %r2;1891; CHECK-NOF32X2-NEXT: cvt.rzi.u32.f32 %r4, %r1;1892; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1893; CHECK-NOF32X2-NEXT: ret;1894;1895; CHECK-F32X2-LABEL: test_fptoui_2xi32(1896; CHECK-F32X2: {1897; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1898; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;1899; CHECK-F32X2-EMPTY:1900; CHECK-F32X2-NEXT: // %bb.0:1901; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fptoui_2xi32_param_0];1902; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1903; CHECK-F32X2-NEXT: cvt.rzi.u32.f32 %r3, %r2;1904; CHECK-F32X2-NEXT: cvt.rzi.u32.f32 %r4, %r1;1905; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1906; CHECK-F32X2-NEXT: ret;1907 %r = fptoui <2 x float> %a to <2 x i32>1908 ret <2 x i32> %r1909}1910 1911define <2 x i64> @test_fptoui_2xi64(<2 x float> %a) #0 {1912; CHECK-NOF32X2-LABEL: test_fptoui_2xi64(1913; CHECK-NOF32X2: {1914; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;1915; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;1916; CHECK-NOF32X2-EMPTY:1917; CHECK-NOF32X2-NEXT: // %bb.0:1918; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_2xi64_param_0];1919; CHECK-NOF32X2-NEXT: cvt.rzi.u64.f32 %rd1, %r2;1920; CHECK-NOF32X2-NEXT: cvt.rzi.u64.f32 %rd2, %r1;1921; CHECK-NOF32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};1922; CHECK-NOF32X2-NEXT: ret;1923;1924; CHECK-F32X2-LABEL: test_fptoui_2xi64(1925; CHECK-F32X2: {1926; CHECK-F32X2-NEXT: .reg .b32 %r<3>;1927; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;1928; CHECK-F32X2-EMPTY:1929; CHECK-F32X2-NEXT: // %bb.0:1930; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fptoui_2xi64_param_0];1931; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1932; CHECK-F32X2-NEXT: cvt.rzi.u64.f32 %rd2, %r2;1933; CHECK-F32X2-NEXT: cvt.rzi.u64.f32 %rd3, %r1;1934; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};1935; CHECK-F32X2-NEXT: ret;1936 %r = fptoui <2 x float> %a to <2 x i64>1937 ret <2 x i64> %r1938}1939 1940define <2 x float> @test_uitofp_2xi32(<2 x i32> %a) #0 {1941; CHECK-NOF32X2-LABEL: test_uitofp_2xi32(1942; CHECK-NOF32X2: {1943; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1944; CHECK-NOF32X2-EMPTY:1945; CHECK-NOF32X2-NEXT: // %bb.0:1946; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_param_0];1947; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r3, %r2;1948; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r4, %r1;1949; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1950; CHECK-NOF32X2-NEXT: ret;1951;1952; CHECK-F32X2-LABEL: test_uitofp_2xi32(1953; CHECK-F32X2: {1954; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1955; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;1956; CHECK-F32X2-EMPTY:1957; CHECK-F32X2-NEXT: // %bb.0:1958; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_uitofp_2xi32_param_0];1959; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;1960; CHECK-F32X2-NEXT: cvt.rn.f32.u32 %r3, %r2;1961; CHECK-F32X2-NEXT: cvt.rn.f32.u32 %r4, %r1;1962; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1963; CHECK-F32X2-NEXT: ret;1964 %r = uitofp <2 x i32> %a to <2 x float>1965 ret <2 x float> %r1966}1967 1968define <2 x float> @test_uitofp_2xi64(<2 x i64> %a) #0 {1969; CHECK-LABEL: test_uitofp_2xi64(1970; CHECK: {1971; CHECK-NEXT: .reg .b32 %r<3>;1972; CHECK-NEXT: .reg .b64 %rd<3>;1973; CHECK-EMPTY:1974; CHECK-NEXT: // %bb.0:1975; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_uitofp_2xi64_param_0];1976; CHECK-NEXT: cvt.rn.f32.u64 %r1, %rd2;1977; CHECK-NEXT: cvt.rn.f32.u64 %r2, %rd1;1978; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};1979; CHECK-NEXT: ret;1980 %r = uitofp <2 x i64> %a to <2 x float>1981 ret <2 x float> %r1982}1983 1984define <2 x float> @test_sitofp_2xi32(<2 x i32> %a) #0 {1985; CHECK-NOF32X2-LABEL: test_sitofp_2xi32(1986; CHECK-NOF32X2: {1987; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;1988; CHECK-NOF32X2-EMPTY:1989; CHECK-NOF32X2-NEXT: // %bb.0:1990; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sitofp_2xi32_param_0];1991; CHECK-NOF32X2-NEXT: cvt.rn.f32.s32 %r3, %r2;1992; CHECK-NOF32X2-NEXT: cvt.rn.f32.s32 %r4, %r1;1993; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};1994; CHECK-NOF32X2-NEXT: ret;1995;1996; CHECK-F32X2-LABEL: test_sitofp_2xi32(1997; CHECK-F32X2: {1998; CHECK-F32X2-NEXT: .reg .b32 %r<5>;1999; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2000; CHECK-F32X2-EMPTY:2001; CHECK-F32X2-NEXT: // %bb.0:2002; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_sitofp_2xi32_param_0];2003; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2004; CHECK-F32X2-NEXT: cvt.rn.f32.s32 %r3, %r2;2005; CHECK-F32X2-NEXT: cvt.rn.f32.s32 %r4, %r1;2006; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2007; CHECK-F32X2-NEXT: ret;2008 %r = sitofp <2 x i32> %a to <2 x float>2009 ret <2 x float> %r2010}2011 2012define <2 x float> @test_sitofp_2xi64(<2 x i64> %a) #0 {2013; CHECK-LABEL: test_sitofp_2xi64(2014; CHECK: {2015; CHECK-NEXT: .reg .b32 %r<3>;2016; CHECK-NEXT: .reg .b64 %rd<3>;2017; CHECK-EMPTY:2018; CHECK-NEXT: // %bb.0:2019; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_sitofp_2xi64_param_0];2020; CHECK-NEXT: cvt.rn.f32.s64 %r1, %rd2;2021; CHECK-NEXT: cvt.rn.f32.s64 %r2, %rd1;2022; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};2023; CHECK-NEXT: ret;2024 %r = sitofp <2 x i64> %a to <2 x float>2025 ret <2 x float> %r2026}2027 2028define <2 x float> @test_uitofp_2xi32_fadd(<2 x i32> %a, <2 x float> %b) #0 {2029; CHECK-NOF32X2-LABEL: test_uitofp_2xi32_fadd(2030; CHECK-NOF32X2: {2031; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;2032; CHECK-NOF32X2-EMPTY:2033; CHECK-NOF32X2-NEXT: // %bb.0:2034; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_uitofp_2xi32_fadd_param_1];2035; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_fadd_param_0];2036; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r5, %r1;2037; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r6, %r2;2038; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r4, %r6;2039; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r3, %r5;2040; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};2041; CHECK-NOF32X2-NEXT: ret;2042;2043; CHECK-F32X2-LABEL: test_uitofp_2xi32_fadd(2044; CHECK-F32X2: {2045; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2046; CHECK-F32X2-NEXT: .reg .b64 %rd<5>;2047; CHECK-F32X2-EMPTY:2048; CHECK-F32X2-NEXT: // %bb.0:2049; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_uitofp_2xi32_fadd_param_1];2050; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_uitofp_2xi32_fadd_param_0];2051; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2052; CHECK-F32X2-NEXT: cvt.rn.f32.u32 %r3, %r2;2053; CHECK-F32X2-NEXT: cvt.rn.f32.u32 %r4, %r1;2054; CHECK-F32X2-NEXT: mov.b64 %rd3, {%r4, %r3};2055; CHECK-F32X2-NEXT: add.rn.f32x2 %rd4, %rd2, %rd3;2056; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd4;2057; CHECK-F32X2-NEXT: ret;2058 %c = uitofp <2 x i32> %a to <2 x float>2059 %r = fadd <2 x float> %b, %c2060 ret <2 x float> %r2061}2062 2063define <2 x float> @test_fptrunc_2xdouble(<2 x double> %a) #0 {2064; CHECK-NOF32X2-LABEL: test_fptrunc_2xdouble(2065; CHECK-NOF32X2: {2066; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;2067; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;2068; CHECK-NOF32X2-EMPTY:2069; CHECK-NOF32X2-NEXT: // %bb.0:2070; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fptrunc_2xdouble_param_0];2071; CHECK-NOF32X2-NEXT: cvt.rn.f32.f64 %r1, %rd2;2072; CHECK-NOF32X2-NEXT: cvt.rn.f32.f64 %r2, %rd1;2073; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};2074; CHECK-NOF32X2-NEXT: ret;2075;2076; CHECK-F32X2-LABEL: test_fptrunc_2xdouble(2077; CHECK-F32X2: {2078; CHECK-F32X2-NEXT: .reg .b32 %r<3>;2079; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;2080; CHECK-F32X2-EMPTY:2081; CHECK-F32X2-NEXT: // %bb.0:2082; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fptrunc_2xdouble_param_0];2083; CHECK-F32X2-NEXT: cvt.rn.f32.f64 %r1, %rd2;2084; CHECK-F32X2-NEXT: cvt.rn.f32.f64 %r2, %rd1;2085; CHECK-F32X2-NEXT: mov.b64 %rd3, {%r2, %r1};2086; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd3;2087; CHECK-F32X2-NEXT: ret;2088 %r = fptrunc <2 x double> %a to <2 x float>2089 ret <2 x float> %r2090}2091 2092define <2 x double> @test_fpext_2xdouble(<2 x float> %a) #0 {2093; CHECK-NOF32X2-LABEL: test_fpext_2xdouble(2094; CHECK-NOF32X2: {2095; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;2096; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;2097; CHECK-NOF32X2-EMPTY:2098; CHECK-NOF32X2-NEXT: // %bb.0:2099; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fpext_2xdouble_param_0];2100; CHECK-NOF32X2-NEXT: cvt.f64.f32 %rd1, %r2;2101; CHECK-NOF32X2-NEXT: cvt.f64.f32 %rd2, %r1;2102; CHECK-NOF32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};2103; CHECK-NOF32X2-NEXT: ret;2104;2105; CHECK-F32X2-LABEL: test_fpext_2xdouble(2106; CHECK-F32X2: {2107; CHECK-F32X2-NEXT: .reg .b32 %r<3>;2108; CHECK-F32X2-NEXT: .reg .b64 %rd<4>;2109; CHECK-F32X2-EMPTY:2110; CHECK-F32X2-NEXT: // %bb.0:2111; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fpext_2xdouble_param_0];2112; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2113; CHECK-F32X2-NEXT: cvt.f64.f32 %rd2, %r2;2114; CHECK-F32X2-NEXT: cvt.f64.f32 %rd3, %r1;2115; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};2116; CHECK-F32X2-NEXT: ret;2117 %r = fpext <2 x float> %a to <2 x double>2118 ret <2 x double> %r2119}2120 2121define <2 x i32> @test_bitcast_2xfloat_to_2xi32(<2 x float> %a) #0 {2122; CHECK-NOF32X2-LABEL: test_bitcast_2xfloat_to_2xi32(2123; CHECK-NOF32X2: {2124; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;2125; CHECK-NOF32X2-EMPTY:2126; CHECK-NOF32X2-NEXT: // %bb.0:2127; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_bitcast_2xfloat_to_2xi32_param_0];2128; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};2129; CHECK-NOF32X2-NEXT: ret;2130;2131; CHECK-F32X2-LABEL: test_bitcast_2xfloat_to_2xi32(2132; CHECK-F32X2: {2133; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2134; CHECK-F32X2-EMPTY:2135; CHECK-F32X2-NEXT: // %bb.0:2136; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_bitcast_2xfloat_to_2xi32_param_0];2137; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd1;2138; CHECK-F32X2-NEXT: ret;2139 %r = bitcast <2 x float> %a to <2 x i32>2140 ret <2 x i32> %r2141}2142 2143define <2 x float> @test_bitcast_2xi32_to_2xfloat(<2 x i32> %a) #0 {2144; CHECK-NOF32X2-LABEL: test_bitcast_2xi32_to_2xfloat(2145; CHECK-NOF32X2: {2146; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;2147; CHECK-NOF32X2-EMPTY:2148; CHECK-NOF32X2-NEXT: // %bb.0:2149; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_bitcast_2xi32_to_2xfloat_param_0];2150; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};2151; CHECK-NOF32X2-NEXT: ret;2152;2153; CHECK-F32X2-LABEL: test_bitcast_2xi32_to_2xfloat(2154; CHECK-F32X2: {2155; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2156; CHECK-F32X2-EMPTY:2157; CHECK-F32X2-NEXT: // %bb.0:2158; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_bitcast_2xi32_to_2xfloat_param_0];2159; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd1;2160; CHECK-F32X2-NEXT: ret;2161 %r = bitcast <2 x i32> %a to <2 x float>2162 ret <2 x float> %r2163}2164 2165define <2 x float> @test_bitcast_double_to_2xfloat(double %a) #0 {2166; CHECK-LABEL: test_bitcast_double_to_2xfloat(2167; CHECK: {2168; CHECK-NEXT: .reg .b64 %rd<2>;2169; CHECK-EMPTY:2170; CHECK-NEXT: // %bb.0:2171; CHECK-NEXT: ld.param.b64 %rd1, [test_bitcast_double_to_2xfloat_param_0];2172; CHECK-NEXT: st.param.b64 [func_retval0], %rd1;2173; CHECK-NEXT: ret;2174 %r = bitcast double %a to <2 x float>2175 ret <2 x float> %r2176}2177 2178define double @test_bitcast_2xfloat_to_double(<2 x float> %a) #0 {2179; CHECK-NOF32X2-LABEL: test_bitcast_2xfloat_to_double(2180; CHECK-NOF32X2: {2181; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;2182; CHECK-NOF32X2-EMPTY:2183; CHECK-NOF32X2-NEXT: // %bb.0:2184; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_bitcast_2xfloat_to_double_param_0];2185; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r2};2186; CHECK-NOF32X2-NEXT: ret;2187;2188; CHECK-F32X2-LABEL: test_bitcast_2xfloat_to_double(2189; CHECK-F32X2: {2190; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2191; CHECK-F32X2-EMPTY:2192; CHECK-F32X2-NEXT: // %bb.0:2193; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_bitcast_2xfloat_to_double_param_0];2194; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd1;2195; CHECK-F32X2-NEXT: ret;2196 %r = bitcast <2 x float> %a to double2197 ret double %r2198}2199 2200define <2 x float> @test_sqrt(<2 x float> %a) #0 {2201; CHECK-NOF32X2-LABEL: test_sqrt(2202; CHECK-NOF32X2: {2203; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2204; CHECK-NOF32X2-EMPTY:2205; CHECK-NOF32X2-NEXT: // %bb.0:2206; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sqrt_param_0];2207; CHECK-NOF32X2-NEXT: sqrt.rn.f32 %r3, %r2;2208; CHECK-NOF32X2-NEXT: sqrt.rn.f32 %r4, %r1;2209; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2210; CHECK-NOF32X2-NEXT: ret;2211;2212; CHECK-F32X2-LABEL: test_sqrt(2213; CHECK-F32X2: {2214; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2215; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2216; CHECK-F32X2-EMPTY:2217; CHECK-F32X2-NEXT: // %bb.0:2218; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_sqrt_param_0];2219; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2220; CHECK-F32X2-NEXT: sqrt.rn.f32 %r3, %r2;2221; CHECK-F32X2-NEXT: sqrt.rn.f32 %r4, %r1;2222; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2223; CHECK-F32X2-NEXT: ret;2224 %r = call <2 x float> @llvm.sqrt(<2 x float> %a)2225 ret <2 x float> %r2226}2227 2228;;; Can't do this yet: requires libcall.2229; XCHECK-LABEL: test_powi(2230;define <2 x float> @test_powi(<2 x float> %a, <2 x i32> %b) #0 {2231; %r = call <2 x float> @llvm.powi.i32(<2 x float> %a, <2 x i32> %b)2232; ret <2 x float> %r2233;}2234 2235define <2 x float> @test_sin(<2 x float> %a) #0 {2236; CHECK-NOF32X2-LABEL: test_sin(2237; CHECK-NOF32X2: {2238; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2239; CHECK-NOF32X2-EMPTY:2240; CHECK-NOF32X2-NEXT: // %bb.0:2241; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sin_param_0];2242; CHECK-NOF32X2-NEXT: sin.approx.f32 %r3, %r2;2243; CHECK-NOF32X2-NEXT: sin.approx.f32 %r4, %r1;2244; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2245; CHECK-NOF32X2-NEXT: ret;2246;2247; CHECK-F32X2-LABEL: test_sin(2248; CHECK-F32X2: {2249; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2250; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2251; CHECK-F32X2-EMPTY:2252; CHECK-F32X2-NEXT: // %bb.0:2253; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_sin_param_0];2254; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2255; CHECK-F32X2-NEXT: sin.approx.f32 %r3, %r2;2256; CHECK-F32X2-NEXT: sin.approx.f32 %r4, %r1;2257; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2258; CHECK-F32X2-NEXT: ret;2259 %r = call afn <2 x float> @llvm.sin(<2 x float> %a)2260 ret <2 x float> %r2261}2262 2263define <2 x float> @test_cos(<2 x float> %a) #0 {2264; CHECK-NOF32X2-LABEL: test_cos(2265; CHECK-NOF32X2: {2266; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2267; CHECK-NOF32X2-EMPTY:2268; CHECK-NOF32X2-NEXT: // %bb.0:2269; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_cos_param_0];2270; CHECK-NOF32X2-NEXT: cos.approx.f32 %r3, %r2;2271; CHECK-NOF32X2-NEXT: cos.approx.f32 %r4, %r1;2272; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2273; CHECK-NOF32X2-NEXT: ret;2274;2275; CHECK-F32X2-LABEL: test_cos(2276; CHECK-F32X2: {2277; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2278; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2279; CHECK-F32X2-EMPTY:2280; CHECK-F32X2-NEXT: // %bb.0:2281; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_cos_param_0];2282; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2283; CHECK-F32X2-NEXT: cos.approx.f32 %r3, %r2;2284; CHECK-F32X2-NEXT: cos.approx.f32 %r4, %r1;2285; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2286; CHECK-F32X2-NEXT: ret;2287 %r = call afn <2 x float> @llvm.cos(<2 x float> %a)2288 ret <2 x float> %r2289}2290 2291;;; Can't do this yet: requires libcall.2292; XCHECK-LABEL: test_pow(2293;define <2 x float> @test_pow(<2 x float> %a, <2 x float> %b) #0 {2294; %r = call <2 x float> @llvm.pow(<2 x float> %a, <2 x float> %b)2295; ret <2 x float> %r2296;}2297 2298;;; Can't do this yet: requires libcall.2299; XCHECK-LABEL: test_exp(2300;define <2 x float> @test_exp(<2 x float> %a) #0 {2301; %r = call <2 x float> @llvm.exp(<2 x float> %a)2302; ret <2 x float> %r2303;}2304 2305;;; Can't do this yet: requires libcall.2306; XCHECK-LABEL: test_exp2(2307;define <2 x float> @test_exp2(<2 x float> %a) #0 {2308; %r = call <2 x float> @llvm.exp2(<2 x float> %a)2309; ret <2 x float> %r2310;}2311 2312;;; Can't do this yet: requires libcall.2313; XCHECK-LABEL: test_log(2314;define <2 x float> @test_log(<2 x float> %a) #0 {2315; %r = call <2 x float> @llvm.log(<2 x float> %a)2316; ret <2 x float> %r2317;}2318 2319;;; Can't do this yet: requires libcall.2320; XCHECK-LABEL: test_log10(2321;define <2 x float> @test_log10(<2 x float> %a) #0 {2322; %r = call <2 x float> @llvm.log10(<2 x float> %a)2323; ret <2 x float> %r2324;}2325 2326;;; Can't do this yet: requires libcall.2327; XCHECK-LABEL: test_log2(2328;define <2 x float> @test_log2(<2 x float> %a) #0 {2329; %r = call <2 x float> @llvm.log2(<2 x float> %a)2330; ret <2 x float> %r2331;}2332 2333 2334define <2 x float> @test_fma(<2 x float> %a, <2 x float> %b, <2 x float> %c) #0 {2335; CHECK-NOF32X2-LABEL: test_fma(2336; CHECK-NOF32X2: {2337; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;2338; CHECK-NOF32X2-EMPTY:2339; CHECK-NOF32X2-NEXT: // %bb.0:2340; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fma_param_2];2341; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fma_param_1];2342; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fma_param_0];2343; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r2, %r4, %r6;2344; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r1, %r3, %r5;2345; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};2346; CHECK-NOF32X2-NEXT: ret;2347;2348; CHECK-F32X2-LABEL: test_fma(2349; CHECK-F32X2: {2350; CHECK-F32X2-NEXT: .reg .b64 %rd<5>;2351; CHECK-F32X2-EMPTY:2352; CHECK-F32X2-NEXT: // %bb.0:2353; CHECK-F32X2-NEXT: ld.param.b64 %rd3, [test_fma_param_2];2354; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fma_param_1];2355; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fma_param_0];2356; CHECK-F32X2-NEXT: fma.rn.f32x2 %rd4, %rd1, %rd2, %rd3;2357; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd4;2358; CHECK-F32X2-NEXT: ret;2359 %r = call <2 x float> @llvm.fma(<2 x float> %a, <2 x float> %b, <2 x float> %c)2360 ret <2 x float> %r2361}2362 2363define <2 x float> @test_fabs(<2 x float> %a) #0 {2364; CHECK-NOF32X2-LABEL: test_fabs(2365; CHECK-NOF32X2: {2366; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2367; CHECK-NOF32X2-EMPTY:2368; CHECK-NOF32X2-NEXT: // %bb.0:2369; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fabs_param_0];2370; CHECK-NOF32X2-NEXT: abs.f32 %r3, %r2;2371; CHECK-NOF32X2-NEXT: abs.f32 %r4, %r1;2372; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2373; CHECK-NOF32X2-NEXT: ret;2374;2375; CHECK-F32X2-LABEL: test_fabs(2376; CHECK-F32X2: {2377; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2378; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2379; CHECK-F32X2-EMPTY:2380; CHECK-F32X2-NEXT: // %bb.0:2381; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fabs_param_0];2382; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2383; CHECK-F32X2-NEXT: abs.f32 %r3, %r2;2384; CHECK-F32X2-NEXT: abs.f32 %r4, %r1;2385; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2386; CHECK-F32X2-NEXT: ret;2387 %r = call <2 x float> @llvm.fabs(<2 x float> %a)2388 ret <2 x float> %r2389}2390 2391define <2 x float> @test_minnum(<2 x float> %a, <2 x float> %b) #0 {2392; CHECK-NOF32X2-LABEL: test_minnum(2393; CHECK-NOF32X2: {2394; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;2395; CHECK-NOF32X2-EMPTY:2396; CHECK-NOF32X2-NEXT: // %bb.0:2397; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_minnum_param_1];2398; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_minnum_param_0];2399; CHECK-NOF32X2-NEXT: min.f32 %r5, %r2, %r4;2400; CHECK-NOF32X2-NEXT: min.f32 %r6, %r1, %r3;2401; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};2402; CHECK-NOF32X2-NEXT: ret;2403;2404; CHECK-F32X2-LABEL: test_minnum(2405; CHECK-F32X2: {2406; CHECK-F32X2-NEXT: .reg .b32 %r<7>;2407; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;2408; CHECK-F32X2-EMPTY:2409; CHECK-F32X2-NEXT: // %bb.0:2410; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_minnum_param_1];2411; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_minnum_param_0];2412; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;2413; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;2414; CHECK-F32X2-NEXT: min.f32 %r5, %r4, %r2;2415; CHECK-F32X2-NEXT: min.f32 %r6, %r3, %r1;2416; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};2417; CHECK-F32X2-NEXT: ret;2418 %r = call <2 x float> @llvm.minnum(<2 x float> %a, <2 x float> %b)2419 ret <2 x float> %r2420}2421 2422define <2 x float> @test_maxnum(<2 x float> %a, <2 x float> %b) #0 {2423; CHECK-NOF32X2-LABEL: test_maxnum(2424; CHECK-NOF32X2: {2425; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;2426; CHECK-NOF32X2-EMPTY:2427; CHECK-NOF32X2-NEXT: // %bb.0:2428; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_maxnum_param_1];2429; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_maxnum_param_0];2430; CHECK-NOF32X2-NEXT: max.f32 %r5, %r2, %r4;2431; CHECK-NOF32X2-NEXT: max.f32 %r6, %r1, %r3;2432; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};2433; CHECK-NOF32X2-NEXT: ret;2434;2435; CHECK-F32X2-LABEL: test_maxnum(2436; CHECK-F32X2: {2437; CHECK-F32X2-NEXT: .reg .b32 %r<7>;2438; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;2439; CHECK-F32X2-EMPTY:2440; CHECK-F32X2-NEXT: // %bb.0:2441; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_maxnum_param_1];2442; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_maxnum_param_0];2443; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;2444; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;2445; CHECK-F32X2-NEXT: max.f32 %r5, %r4, %r2;2446; CHECK-F32X2-NEXT: max.f32 %r6, %r3, %r1;2447; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};2448; CHECK-F32X2-NEXT: ret;2449 %r = call <2 x float> @llvm.maxnum(<2 x float> %a, <2 x float> %b)2450 ret <2 x float> %r2451}2452 2453define <2 x float> @test_copysign(<2 x float> %a, <2 x float> %b) #0 {2454; CHECK-NOF32X2-LABEL: test_copysign(2455; CHECK-NOF32X2: {2456; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;2457; CHECK-NOF32X2-EMPTY:2458; CHECK-NOF32X2-NEXT: // %bb.0:2459; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_copysign_param_1];2460; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_param_0];2461; CHECK-NOF32X2-NEXT: copysign.f32 %r5, %r4, %r2;2462; CHECK-NOF32X2-NEXT: copysign.f32 %r6, %r3, %r1;2463; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};2464; CHECK-NOF32X2-NEXT: ret;2465;2466; CHECK-F32X2-LABEL: test_copysign(2467; CHECK-F32X2: {2468; CHECK-F32X2-NEXT: .reg .b32 %r<7>;2469; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;2470; CHECK-F32X2-EMPTY:2471; CHECK-F32X2-NEXT: // %bb.0:2472; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_copysign_param_1];2473; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_copysign_param_0];2474; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2475; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd2;2476; CHECK-F32X2-NEXT: copysign.f32 %r5, %r4, %r2;2477; CHECK-F32X2-NEXT: copysign.f32 %r6, %r3, %r1;2478; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};2479; CHECK-F32X2-NEXT: ret;2480 %r = call <2 x float> @llvm.copysign(<2 x float> %a, <2 x float> %b)2481 ret <2 x float> %r2482}2483 2484define <2 x float> @test_copysign_f64(<2 x float> %a, <2 x double> %b) #0 {2485; CHECK-NOF32X2-LABEL: test_copysign_f64(2486; CHECK-NOF32X2: {2487; CHECK-NOF32X2-NEXT: .reg .pred %p<3>;2488; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;2489; CHECK-NOF32X2-NEXT: .reg .b64 %rd<7>;2490; CHECK-NOF32X2-EMPTY:2491; CHECK-NOF32X2-NEXT: // %bb.0:2492; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_copysign_f64_param_1];2493; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_f64_param_0];2494; CHECK-NOF32X2-NEXT: abs.f32 %r3, %r2;2495; CHECK-NOF32X2-NEXT: neg.f32 %r4, %r3;2496; CHECK-NOF32X2-NEXT: shr.u64 %rd3, %rd2, 63;2497; CHECK-NOF32X2-NEXT: and.b64 %rd4, %rd3, 1;2498; CHECK-NOF32X2-NEXT: setp.ne.b64 %p1, %rd4, 0;2499; CHECK-NOF32X2-NEXT: selp.f32 %r5, %r4, %r3, %p1;2500; CHECK-NOF32X2-NEXT: abs.f32 %r6, %r1;2501; CHECK-NOF32X2-NEXT: neg.f32 %r7, %r6;2502; CHECK-NOF32X2-NEXT: shr.u64 %rd5, %rd1, 63;2503; CHECK-NOF32X2-NEXT: and.b64 %rd6, %rd5, 1;2504; CHECK-NOF32X2-NEXT: setp.ne.b64 %p2, %rd6, 0;2505; CHECK-NOF32X2-NEXT: selp.f32 %r8, %r7, %r6, %p2;2506; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r5};2507; CHECK-NOF32X2-NEXT: ret;2508;2509; CHECK-F32X2-LABEL: test_copysign_f64(2510; CHECK-F32X2: {2511; CHECK-F32X2-NEXT: .reg .pred %p<3>;2512; CHECK-F32X2-NEXT: .reg .b32 %r<9>;2513; CHECK-F32X2-NEXT: .reg .b64 %rd<8>;2514; CHECK-F32X2-EMPTY:2515; CHECK-F32X2-NEXT: // %bb.0:2516; CHECK-F32X2-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_copysign_f64_param_1];2517; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_copysign_f64_param_0];2518; CHECK-F32X2-NEXT: shr.u64 %rd4, %rd3, 63;2519; CHECK-F32X2-NEXT: and.b64 %rd5, %rd4, 1;2520; CHECK-F32X2-NEXT: setp.ne.b64 %p1, %rd5, 0;2521; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2522; CHECK-F32X2-NEXT: abs.f32 %r3, %r2;2523; CHECK-F32X2-NEXT: neg.f32 %r4, %r3;2524; CHECK-F32X2-NEXT: selp.f32 %r5, %r4, %r3, %p1;2525; CHECK-F32X2-NEXT: shr.u64 %rd6, %rd2, 63;2526; CHECK-F32X2-NEXT: and.b64 %rd7, %rd6, 1;2527; CHECK-F32X2-NEXT: setp.ne.b64 %p2, %rd7, 0;2528; CHECK-F32X2-NEXT: abs.f32 %r6, %r1;2529; CHECK-F32X2-NEXT: neg.f32 %r7, %r6;2530; CHECK-F32X2-NEXT: selp.f32 %r8, %r7, %r6, %p2;2531; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r5};2532; CHECK-F32X2-NEXT: ret;2533 %tb = fptrunc <2 x double> %b to <2 x float>2534 %r = call <2 x float> @llvm.copysign(<2 x float> %a, <2 x float> %tb)2535 ret <2 x float> %r2536}2537 2538define <2 x double> @test_copysign_extended(<2 x float> %a, <2 x float> %b) #0 {2539; CHECK-NOF32X2-LABEL: test_copysign_extended(2540; CHECK-NOF32X2: {2541; CHECK-NOF32X2-NEXT: .reg .b32 %r<7>;2542; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;2543; CHECK-NOF32X2-EMPTY:2544; CHECK-NOF32X2-NEXT: // %bb.0:2545; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_copysign_extended_param_1];2546; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_extended_param_0];2547; CHECK-NOF32X2-NEXT: copysign.f32 %r5, %r3, %r1;2548; CHECK-NOF32X2-NEXT: copysign.f32 %r6, %r4, %r2;2549; CHECK-NOF32X2-NEXT: cvt.f64.f32 %rd1, %r6;2550; CHECK-NOF32X2-NEXT: cvt.f64.f32 %rd2, %r5;2551; CHECK-NOF32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};2552; CHECK-NOF32X2-NEXT: ret;2553;2554; CHECK-F32X2-LABEL: test_copysign_extended(2555; CHECK-F32X2: {2556; CHECK-F32X2-NEXT: .reg .b32 %r<7>;2557; CHECK-F32X2-NEXT: .reg .b64 %rd<5>;2558; CHECK-F32X2-EMPTY:2559; CHECK-F32X2-NEXT: // %bb.0:2560; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_copysign_extended_param_1];2561; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_copysign_extended_param_0];2562; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2563; CHECK-F32X2-NEXT: mov.b64 {%r3, %r4}, %rd2;2564; CHECK-F32X2-NEXT: copysign.f32 %r5, %r3, %r1;2565; CHECK-F32X2-NEXT: copysign.f32 %r6, %r4, %r2;2566; CHECK-F32X2-NEXT: cvt.f64.f32 %rd3, %r6;2567; CHECK-F32X2-NEXT: cvt.f64.f32 %rd4, %r5;2568; CHECK-F32X2-NEXT: st.param.v2.b64 [func_retval0], {%rd4, %rd3};2569; CHECK-F32X2-NEXT: ret;2570 %r = call <2 x float> @llvm.copysign(<2 x float> %a, <2 x float> %b)2571 %xr = fpext <2 x float> %r to <2 x double>2572 ret <2 x double> %xr2573}2574 2575define <2 x float> @test_floor(<2 x float> %a) #0 {2576; CHECK-NOF32X2-LABEL: test_floor(2577; CHECK-NOF32X2: {2578; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2579; CHECK-NOF32X2-EMPTY:2580; CHECK-NOF32X2-NEXT: // %bb.0:2581; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_floor_param_0];2582; CHECK-NOF32X2-NEXT: cvt.rmi.f32.f32 %r3, %r2;2583; CHECK-NOF32X2-NEXT: cvt.rmi.f32.f32 %r4, %r1;2584; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2585; CHECK-NOF32X2-NEXT: ret;2586;2587; CHECK-F32X2-LABEL: test_floor(2588; CHECK-F32X2: {2589; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2590; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2591; CHECK-F32X2-EMPTY:2592; CHECK-F32X2-NEXT: // %bb.0:2593; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_floor_param_0];2594; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2595; CHECK-F32X2-NEXT: cvt.rmi.f32.f32 %r3, %r2;2596; CHECK-F32X2-NEXT: cvt.rmi.f32.f32 %r4, %r1;2597; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2598; CHECK-F32X2-NEXT: ret;2599 %r = call <2 x float> @llvm.floor(<2 x float> %a)2600 ret <2 x float> %r2601}2602 2603define <2 x float> @test_ceil(<2 x float> %a) #0 {2604; CHECK-NOF32X2-LABEL: test_ceil(2605; CHECK-NOF32X2: {2606; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2607; CHECK-NOF32X2-EMPTY:2608; CHECK-NOF32X2-NEXT: // %bb.0:2609; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_ceil_param_0];2610; CHECK-NOF32X2-NEXT: cvt.rpi.f32.f32 %r3, %r2;2611; CHECK-NOF32X2-NEXT: cvt.rpi.f32.f32 %r4, %r1;2612; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2613; CHECK-NOF32X2-NEXT: ret;2614;2615; CHECK-F32X2-LABEL: test_ceil(2616; CHECK-F32X2: {2617; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2618; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2619; CHECK-F32X2-EMPTY:2620; CHECK-F32X2-NEXT: // %bb.0:2621; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_ceil_param_0];2622; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2623; CHECK-F32X2-NEXT: cvt.rpi.f32.f32 %r3, %r2;2624; CHECK-F32X2-NEXT: cvt.rpi.f32.f32 %r4, %r1;2625; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2626; CHECK-F32X2-NEXT: ret;2627 %r = call <2 x float> @llvm.ceil(<2 x float> %a)2628 ret <2 x float> %r2629}2630 2631define <2 x float> @test_trunc(<2 x float> %a) #0 {2632; CHECK-NOF32X2-LABEL: test_trunc(2633; CHECK-NOF32X2: {2634; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2635; CHECK-NOF32X2-EMPTY:2636; CHECK-NOF32X2-NEXT: // %bb.0:2637; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_param_0];2638; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r3, %r2;2639; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r4, %r1;2640; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2641; CHECK-NOF32X2-NEXT: ret;2642;2643; CHECK-F32X2-LABEL: test_trunc(2644; CHECK-F32X2: {2645; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2646; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2647; CHECK-F32X2-EMPTY:2648; CHECK-F32X2-NEXT: // %bb.0:2649; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_trunc_param_0];2650; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2651; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r3, %r2;2652; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r4, %r1;2653; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2654; CHECK-F32X2-NEXT: ret;2655 %r = call <2 x float> @llvm.trunc(<2 x float> %a)2656 ret <2 x float> %r2657}2658 2659define <2 x float> @test_rint(<2 x float> %a) #0 {2660; CHECK-NOF32X2-LABEL: test_rint(2661; CHECK-NOF32X2: {2662; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2663; CHECK-NOF32X2-EMPTY:2664; CHECK-NOF32X2-NEXT: // %bb.0:2665; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_rint_param_0];2666; CHECK-NOF32X2-NEXT: cvt.rni.f32.f32 %r3, %r2;2667; CHECK-NOF32X2-NEXT: cvt.rni.f32.f32 %r4, %r1;2668; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2669; CHECK-NOF32X2-NEXT: ret;2670;2671; CHECK-F32X2-LABEL: test_rint(2672; CHECK-F32X2: {2673; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2674; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2675; CHECK-F32X2-EMPTY:2676; CHECK-F32X2-NEXT: // %bb.0:2677; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_rint_param_0];2678; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2679; CHECK-F32X2-NEXT: cvt.rni.f32.f32 %r3, %r2;2680; CHECK-F32X2-NEXT: cvt.rni.f32.f32 %r4, %r1;2681; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2682; CHECK-F32X2-NEXT: ret;2683 %r = call <2 x float> @llvm.rint(<2 x float> %a)2684 ret <2 x float> %r2685}2686 2687define <2 x float> @test_nearbyint(<2 x float> %a) #0 {2688; CHECK-NOF32X2-LABEL: test_nearbyint(2689; CHECK-NOF32X2: {2690; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2691; CHECK-NOF32X2-EMPTY:2692; CHECK-NOF32X2-NEXT: // %bb.0:2693; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_nearbyint_param_0];2694; CHECK-NOF32X2-NEXT: cvt.rni.f32.f32 %r3, %r2;2695; CHECK-NOF32X2-NEXT: cvt.rni.f32.f32 %r4, %r1;2696; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2697; CHECK-NOF32X2-NEXT: ret;2698;2699; CHECK-F32X2-LABEL: test_nearbyint(2700; CHECK-F32X2: {2701; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2702; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2703; CHECK-F32X2-EMPTY:2704; CHECK-F32X2-NEXT: // %bb.0:2705; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_nearbyint_param_0];2706; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2707; CHECK-F32X2-NEXT: cvt.rni.f32.f32 %r3, %r2;2708; CHECK-F32X2-NEXT: cvt.rni.f32.f32 %r4, %r1;2709; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2710; CHECK-F32X2-NEXT: ret;2711 %r = call <2 x float> @llvm.nearbyint(<2 x float> %a)2712 ret <2 x float> %r2713}2714 2715define <2 x float> @test_roundeven(<2 x float> %a) #0 {2716; CHECK-NOF32X2-LABEL: test_roundeven(2717; CHECK-NOF32X2: {2718; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2719; CHECK-NOF32X2-EMPTY:2720; CHECK-NOF32X2-NEXT: // %bb.0:2721; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_roundeven_param_0];2722; CHECK-NOF32X2-NEXT: cvt.rni.f32.f32 %r3, %r2;2723; CHECK-NOF32X2-NEXT: cvt.rni.f32.f32 %r4, %r1;2724; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2725; CHECK-NOF32X2-NEXT: ret;2726;2727; CHECK-F32X2-LABEL: test_roundeven(2728; CHECK-F32X2: {2729; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2730; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2731; CHECK-F32X2-EMPTY:2732; CHECK-F32X2-NEXT: // %bb.0:2733; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_roundeven_param_0];2734; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2735; CHECK-F32X2-NEXT: cvt.rni.f32.f32 %r3, %r2;2736; CHECK-F32X2-NEXT: cvt.rni.f32.f32 %r4, %r1;2737; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2738; CHECK-F32X2-NEXT: ret;2739 %r = call <2 x float> @llvm.roundeven(<2 x float> %a)2740 ret <2 x float> %r2741}2742 2743; check the use of sign mask and 0.5 to implement round2744define <2 x float> @test_round(<2 x float> %a) #0 {2745; CHECK-NOF32X2-LABEL: test_round(2746; CHECK-NOF32X2: {2747; CHECK-NOF32X2-NEXT: .reg .pred %p<5>;2748; CHECK-NOF32X2-NEXT: .reg .b32 %r<19>;2749; CHECK-NOF32X2-EMPTY:2750; CHECK-NOF32X2-NEXT: // %bb.0:2751; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_round_param_0];2752; CHECK-NOF32X2-NEXT: and.b32 %r3, %r2, -2147483648;2753; CHECK-NOF32X2-NEXT: or.b32 %r4, %r3, 1056964608;2754; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r2, %r4;2755; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r6, %r5;2756; CHECK-NOF32X2-NEXT: abs.f32 %r7, %r2;2757; CHECK-NOF32X2-NEXT: setp.gt.f32 %p1, %r7, 0f4B000000;2758; CHECK-NOF32X2-NEXT: selp.f32 %r8, %r2, %r6, %p1;2759; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r9, %r2;2760; CHECK-NOF32X2-NEXT: setp.lt.f32 %p2, %r7, 0f3F000000;2761; CHECK-NOF32X2-NEXT: selp.f32 %r10, %r9, %r8, %p2;2762; CHECK-NOF32X2-NEXT: and.b32 %r11, %r1, -2147483648;2763; CHECK-NOF32X2-NEXT: or.b32 %r12, %r11, 1056964608;2764; CHECK-NOF32X2-NEXT: add.rn.f32 %r13, %r1, %r12;2765; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r14, %r13;2766; CHECK-NOF32X2-NEXT: abs.f32 %r15, %r1;2767; CHECK-NOF32X2-NEXT: setp.gt.f32 %p3, %r15, 0f4B000000;2768; CHECK-NOF32X2-NEXT: selp.f32 %r16, %r1, %r14, %p3;2769; CHECK-NOF32X2-NEXT: cvt.rzi.f32.f32 %r17, %r1;2770; CHECK-NOF32X2-NEXT: setp.lt.f32 %p4, %r15, 0f3F000000;2771; CHECK-NOF32X2-NEXT: selp.f32 %r18, %r17, %r16, %p4;2772; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r18, %r10};2773; CHECK-NOF32X2-NEXT: ret;2774;2775; CHECK-F32X2-LABEL: test_round(2776; CHECK-F32X2: {2777; CHECK-F32X2-NEXT: .reg .pred %p<5>;2778; CHECK-F32X2-NEXT: .reg .b32 %r<19>;2779; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2780; CHECK-F32X2-EMPTY:2781; CHECK-F32X2-NEXT: // %bb.0:2782; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_round_param_0];2783; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2784; CHECK-F32X2-NEXT: and.b32 %r3, %r2, -2147483648;2785; CHECK-F32X2-NEXT: or.b32 %r4, %r3, 1056964608;2786; CHECK-F32X2-NEXT: add.rn.f32 %r5, %r2, %r4;2787; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r6, %r5;2788; CHECK-F32X2-NEXT: abs.f32 %r7, %r2;2789; CHECK-F32X2-NEXT: setp.gt.f32 %p1, %r7, 0f4B000000;2790; CHECK-F32X2-NEXT: selp.f32 %r8, %r2, %r6, %p1;2791; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r9, %r2;2792; CHECK-F32X2-NEXT: setp.lt.f32 %p2, %r7, 0f3F000000;2793; CHECK-F32X2-NEXT: selp.f32 %r10, %r9, %r8, %p2;2794; CHECK-F32X2-NEXT: and.b32 %r11, %r1, -2147483648;2795; CHECK-F32X2-NEXT: or.b32 %r12, %r11, 1056964608;2796; CHECK-F32X2-NEXT: add.rn.f32 %r13, %r1, %r12;2797; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r14, %r13;2798; CHECK-F32X2-NEXT: abs.f32 %r15, %r1;2799; CHECK-F32X2-NEXT: setp.gt.f32 %p3, %r15, 0f4B000000;2800; CHECK-F32X2-NEXT: selp.f32 %r16, %r1, %r14, %p3;2801; CHECK-F32X2-NEXT: cvt.rzi.f32.f32 %r17, %r1;2802; CHECK-F32X2-NEXT: setp.lt.f32 %p4, %r15, 0f3F000000;2803; CHECK-F32X2-NEXT: selp.f32 %r18, %r17, %r16, %p4;2804; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r18, %r10};2805; CHECK-F32X2-NEXT: ret;2806 %r = call <2 x float> @llvm.round(<2 x float> %a)2807 ret <2 x float> %r2808}2809 2810define <2 x float> @test_fmuladd(<2 x float> %a, <2 x float> %b, <2 x float> %c) #0 {2811; CHECK-NOF32X2-LABEL: test_fmuladd(2812; CHECK-NOF32X2: {2813; CHECK-NOF32X2-NEXT: .reg .b32 %r<9>;2814; CHECK-NOF32X2-EMPTY:2815; CHECK-NOF32X2-NEXT: // %bb.0:2816; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fmuladd_param_2];2817; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmuladd_param_1];2818; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fmuladd_param_0];2819; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r2, %r4, %r6;2820; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r1, %r3, %r5;2821; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};2822; CHECK-NOF32X2-NEXT: ret;2823;2824; CHECK-F32X2-LABEL: test_fmuladd(2825; CHECK-F32X2: {2826; CHECK-F32X2-NEXT: .reg .b64 %rd<5>;2827; CHECK-F32X2-EMPTY:2828; CHECK-F32X2-NEXT: // %bb.0:2829; CHECK-F32X2-NEXT: ld.param.b64 %rd3, [test_fmuladd_param_2];2830; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_fmuladd_param_1];2831; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_fmuladd_param_0];2832; CHECK-F32X2-NEXT: fma.rn.f32x2 %rd4, %rd1, %rd2, %rd3;2833; CHECK-F32X2-NEXT: st.param.b64 [func_retval0], %rd4;2834; CHECK-F32X2-NEXT: ret;2835 %r = call <2 x float> @llvm.fmuladd(<2 x float> %a, <2 x float> %b, <2 x float> %c)2836 ret <2 x float> %r2837}2838 2839define <2 x float> @test_shufflevector(<2 x float> %a) #0 {2840; CHECK-NOF32X2-LABEL: test_shufflevector(2841; CHECK-NOF32X2: {2842; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;2843; CHECK-NOF32X2-EMPTY:2844; CHECK-NOF32X2-NEXT: // %bb.0:2845; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_shufflevector_param_0];2846; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};2847; CHECK-NOF32X2-NEXT: ret;2848;2849; CHECK-F32X2-LABEL: test_shufflevector(2850; CHECK-F32X2: {2851; CHECK-F32X2-NEXT: .reg .b32 %r<3>;2852; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2853; CHECK-F32X2-EMPTY:2854; CHECK-F32X2-NEXT: // %bb.0:2855; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_shufflevector_param_0];2856; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2857; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};2858; CHECK-F32X2-NEXT: ret;2859 %s = shufflevector <2 x float> %a, <2 x float> poison, <2 x i32> <i32 1, i32 0>2860 ret <2 x float> %s2861}2862 2863define <2 x float> @test_insertelement(<2 x float> %a, float %x) #0 {2864; CHECK-NOF32X2-LABEL: test_insertelement(2865; CHECK-NOF32X2: {2866; CHECK-NOF32X2-NEXT: .reg .b32 %r<4>;2867; CHECK-NOF32X2-EMPTY:2868; CHECK-NOF32X2-NEXT: // %bb.0:2869; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_insertelement_param_0];2870; CHECK-NOF32X2-NEXT: ld.param.b32 %r3, [test_insertelement_param_1];2871; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r1, %r3};2872; CHECK-NOF32X2-NEXT: ret;2873;2874; CHECK-F32X2-LABEL: test_insertelement(2875; CHECK-F32X2: {2876; CHECK-F32X2-NEXT: .reg .b32 %r<3>;2877; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2878; CHECK-F32X2-EMPTY:2879; CHECK-F32X2-NEXT: // %bb.0:2880; CHECK-F32X2-NEXT: ld.param.b32 %r1, [test_insertelement_param_1];2881; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_insertelement_param_0];2882; CHECK-F32X2-NEXT: mov.b64 {%r2, _}, %rd1;2883; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};2884; CHECK-F32X2-NEXT: ret;2885 %i = insertelement <2 x float> %a, float %x, i64 12886 ret <2 x float> %i2887}2888 2889define <2 x float> @test_sitofp_2xi32_to_2xfloat(<2 x i32> %a) #0 {2890; CHECK-NOF32X2-LABEL: test_sitofp_2xi32_to_2xfloat(2891; CHECK-NOF32X2: {2892; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2893; CHECK-NOF32X2-EMPTY:2894; CHECK-NOF32X2-NEXT: // %bb.0:2895; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sitofp_2xi32_to_2xfloat_param_0];2896; CHECK-NOF32X2-NEXT: cvt.rn.f32.s32 %r3, %r2;2897; CHECK-NOF32X2-NEXT: cvt.rn.f32.s32 %r4, %r1;2898; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2899; CHECK-NOF32X2-NEXT: ret;2900;2901; CHECK-F32X2-LABEL: test_sitofp_2xi32_to_2xfloat(2902; CHECK-F32X2: {2903; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2904; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2905; CHECK-F32X2-EMPTY:2906; CHECK-F32X2-NEXT: // %bb.0:2907; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_sitofp_2xi32_to_2xfloat_param_0];2908; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2909; CHECK-F32X2-NEXT: cvt.rn.f32.s32 %r3, %r2;2910; CHECK-F32X2-NEXT: cvt.rn.f32.s32 %r4, %r1;2911; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2912; CHECK-F32X2-NEXT: ret;2913 %r = sitofp <2 x i32> %a to <2 x float>2914 ret <2 x float> %r2915}2916 2917define <2 x float> @test_uitofp_2xi32_to_2xfloat(<2 x i32> %a) #0 {2918; CHECK-NOF32X2-LABEL: test_uitofp_2xi32_to_2xfloat(2919; CHECK-NOF32X2: {2920; CHECK-NOF32X2-NEXT: .reg .b32 %r<5>;2921; CHECK-NOF32X2-EMPTY:2922; CHECK-NOF32X2-NEXT: // %bb.0:2923; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_to_2xfloat_param_0];2924; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r3, %r2;2925; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r4, %r1;2926; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2927; CHECK-NOF32X2-NEXT: ret;2928;2929; CHECK-F32X2-LABEL: test_uitofp_2xi32_to_2xfloat(2930; CHECK-F32X2: {2931; CHECK-F32X2-NEXT: .reg .b32 %r<5>;2932; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;2933; CHECK-F32X2-EMPTY:2934; CHECK-F32X2-NEXT: // %bb.0:2935; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_uitofp_2xi32_to_2xfloat_param_0];2936; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2937; CHECK-F32X2-NEXT: cvt.rn.f32.u32 %r3, %r2;2938; CHECK-F32X2-NEXT: cvt.rn.f32.u32 %r4, %r1;2939; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};2940; CHECK-F32X2-NEXT: ret;2941 %r = uitofp <2 x i32> %a to <2 x float>2942 ret <2 x float> %r2943}2944 2945define void @test_trunc_to_v2bf16(<2 x float> %a, ptr %p) {2946; CHECK-NOF32X2-LABEL: test_trunc_to_v2bf16(2947; CHECK-NOF32X2: {2948; CHECK-NOF32X2-NEXT: .reg .b32 %r<4>;2949; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;2950; CHECK-NOF32X2-EMPTY:2951; CHECK-NOF32X2-NEXT: // %bb.0:2952; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_to_v2bf16_param_0];2953; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_trunc_to_v2bf16_param_1];2954; CHECK-NOF32X2-NEXT: cvt.rn.bf16x2.f32 %r3, %r2, %r1;2955; CHECK-NOF32X2-NEXT: st.b32 [%rd1], %r3;2956; CHECK-NOF32X2-NEXT: ret;2957;2958; CHECK-F32X2-LABEL: test_trunc_to_v2bf16(2959; CHECK-F32X2: {2960; CHECK-F32X2-NEXT: .reg .b32 %r<4>;2961; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;2962; CHECK-F32X2-EMPTY:2963; CHECK-F32X2-NEXT: // %bb.0:2964; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_trunc_to_v2bf16_param_1];2965; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_trunc_to_v2bf16_param_0];2966; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2967; CHECK-F32X2-NEXT: cvt.rn.bf16x2.f32 %r3, %r2, %r1;2968; CHECK-F32X2-NEXT: st.b32 [%rd2], %r3;2969; CHECK-F32X2-NEXT: ret;2970 %trunc = fptrunc <2 x float> %a to <2 x bfloat>2971 store <2 x bfloat> %trunc, ptr %p2972 ret void2973}2974 2975define void @test_trunc_to_v2f16(<2 x float> %a, ptr %p) {2976; CHECK-NOF32X2-LABEL: test_trunc_to_v2f16(2977; CHECK-NOF32X2: {2978; CHECK-NOF32X2-NEXT: .reg .b32 %r<4>;2979; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;2980; CHECK-NOF32X2-EMPTY:2981; CHECK-NOF32X2-NEXT: // %bb.0:2982; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_to_v2f16_param_0];2983; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_trunc_to_v2f16_param_1];2984; CHECK-NOF32X2-NEXT: cvt.rn.f16x2.f32 %r3, %r2, %r1;2985; CHECK-NOF32X2-NEXT: st.b32 [%rd1], %r3;2986; CHECK-NOF32X2-NEXT: ret;2987;2988; CHECK-F32X2-LABEL: test_trunc_to_v2f16(2989; CHECK-F32X2: {2990; CHECK-F32X2-NEXT: .reg .b32 %r<4>;2991; CHECK-F32X2-NEXT: .reg .b64 %rd<3>;2992; CHECK-F32X2-EMPTY:2993; CHECK-F32X2-NEXT: // %bb.0:2994; CHECK-F32X2-NEXT: ld.param.b64 %rd2, [test_trunc_to_v2f16_param_1];2995; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_trunc_to_v2f16_param_0];2996; CHECK-F32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;2997; CHECK-F32X2-NEXT: cvt.rn.f16x2.f32 %r3, %r2, %r1;2998; CHECK-F32X2-NEXT: st.b32 [%rd2], %r3;2999; CHECK-F32X2-NEXT: ret;3000 %trunc = fptrunc <2 x float> %a to <2 x half>3001 store <2 x half> %trunc, ptr %p3002 ret void3003}3004 3005 3006attributes #0 = { nounwind }3007attributes #2 = { "denormal-fp-math"="preserve-sign" }3008