744 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | FileCheck --check-prefixes=CHECK,SM80 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | FileCheck --check-prefixes=CHECK,SM90 %s4; RUN: %if ptxas-sm_80 && ptxas-isa-7.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | %ptxas-verify -arch=sm_80 %}5; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}6 7target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"8 9define <2 x bfloat> @test_ret_const() #0 {10; CHECK-LABEL: test_ret_const(11; CHECK: {12; CHECK-EMPTY:13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: st.param.b32 [func_retval0], 1073758080;16; CHECK-NEXT: ret;17 ret <2 x bfloat> <bfloat 1.0, bfloat 2.0>18}19 20; Check that we can lower fadd with immediate arguments.21define <2 x bfloat> @test_fadd_imm_0(<2 x bfloat> %a) #0 {22; SM80-LABEL: test_fadd_imm_0(23; SM80: {24; SM80-NEXT: .reg .b32 %r<5>;25; SM80-EMPTY:26; SM80-NEXT: // %bb.0:27; SM80-NEXT: ld.param.b32 %r1, [test_fadd_imm_0_param_0];28; SM80-NEXT: mov.b32 %r2, 1065369472;29; SM80-NEXT: mov.b32 %r3, 1073758080;30; SM80-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;31; SM80-NEXT: st.param.b32 [func_retval0], %r4;32; SM80-NEXT: ret;33;34; SM90-LABEL: test_fadd_imm_0(35; SM90: {36; SM90-NEXT: .reg .b32 %r<4>;37; SM90-EMPTY:38; SM90-NEXT: // %bb.0:39; SM90-NEXT: ld.param.b32 %r1, [test_fadd_imm_0_param_0];40; SM90-NEXT: mov.b32 %r2, 1073758080;41; SM90-NEXT: add.rn.bf16x2 %r3, %r1, %r2;42; SM90-NEXT: st.param.b32 [func_retval0], %r3;43; SM90-NEXT: ret;44 %r = fadd <2 x bfloat> <bfloat 1.0, bfloat 2.0>, %a45 ret <2 x bfloat> %r46}47 48define bfloat @test_fadd_imm_1(bfloat %a) #0 {49; SM80-LABEL: test_fadd_imm_1(50; SM80: {51; SM80-NEXT: .reg .b16 %rs<4>;52; SM80-EMPTY:53; SM80-NEXT: // %bb.0:54; SM80-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];55; SM80-NEXT: mov.b16 %rs2, 0x3F80;56; SM80-NEXT: fma.rn.bf16 %rs3, %rs1, %rs2, %rs2;57; SM80-NEXT: st.param.b16 [func_retval0], %rs3;58; SM80-NEXT: ret;59;60; SM90-LABEL: test_fadd_imm_1(61; SM90: {62; SM90-NEXT: .reg .b16 %rs<4>;63; SM90-EMPTY:64; SM90-NEXT: // %bb.0:65; SM90-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];66; SM90-NEXT: mov.b16 %rs2, 0x3F80;67; SM90-NEXT: add.rn.bf16 %rs3, %rs1, %rs2;68; SM90-NEXT: st.param.b16 [func_retval0], %rs3;69; SM90-NEXT: ret;70 %r = fadd bfloat %a, 1.071 ret bfloat %r72}73 74define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {75; SM80-LABEL: test_fsubx2(76; SM80: {77; SM80-NEXT: .reg .b32 %r<5>;78; SM80-EMPTY:79; SM80-NEXT: // %bb.0:80; SM80-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];81; SM80-NEXT: mov.b32 %r2, -1082081408;82; SM80-NEXT: ld.param.b32 %r3, [test_fsubx2_param_1];83; SM80-NEXT: fma.rn.bf16x2 %r4, %r3, %r2, %r1;84; SM80-NEXT: st.param.b32 [func_retval0], %r4;85; SM80-NEXT: ret;86;87; SM90-LABEL: test_fsubx2(88; SM90: {89; SM90-NEXT: .reg .b32 %r<4>;90; SM90-EMPTY:91; SM90-NEXT: // %bb.0:92; SM90-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];93; SM90-NEXT: ld.param.b32 %r2, [test_fsubx2_param_1];94; SM90-NEXT: sub.rn.bf16x2 %r3, %r1, %r2;95; SM90-NEXT: st.param.b32 [func_retval0], %r3;96; SM90-NEXT: ret;97 %r = fsub <2 x bfloat> %a, %b98 ret <2 x bfloat> %r99}100 101define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {102; SM80-LABEL: test_fmulx2(103; SM80: {104; SM80-NEXT: .reg .b32 %r<5>;105; SM80-EMPTY:106; SM80-NEXT: // %bb.0:107; SM80-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];108; SM80-NEXT: mov.b32 %r2, -2147450880;109; SM80-NEXT: ld.param.b32 %r3, [test_fmulx2_param_1];110; SM80-NEXT: fma.rn.bf16x2 %r4, %r1, %r3, %r2;111; SM80-NEXT: st.param.b32 [func_retval0], %r4;112; SM80-NEXT: ret;113;114; SM90-LABEL: test_fmulx2(115; SM90: {116; SM90-NEXT: .reg .b32 %r<4>;117; SM90-EMPTY:118; SM90-NEXT: // %bb.0:119; SM90-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];120; SM90-NEXT: ld.param.b32 %r2, [test_fmulx2_param_1];121; SM90-NEXT: mul.rn.bf16x2 %r3, %r1, %r2;122; SM90-NEXT: st.param.b32 [func_retval0], %r3;123; SM90-NEXT: ret;124 %r = fmul <2 x bfloat> %a, %b125 ret <2 x bfloat> %r126}127 128define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 {129; CHECK-LABEL: test_fdiv(130; CHECK: {131; CHECK-NEXT: .reg .b16 %rs<5>;132; CHECK-NEXT: .reg .b32 %r<8>;133; CHECK-EMPTY:134; CHECK-NEXT: // %bb.0:135; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];136; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];137; CHECK-NEXT: cvt.f32.bf16 %r1, %rs3;138; CHECK-NEXT: cvt.f32.bf16 %r2, %rs1;139; CHECK-NEXT: div.rn.f32 %r3, %r2, %r1;140; CHECK-NEXT: cvt.f32.bf16 %r4, %rs4;141; CHECK-NEXT: cvt.f32.bf16 %r5, %rs2;142; CHECK-NEXT: div.rn.f32 %r6, %r5, %r4;143; CHECK-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;144; CHECK-NEXT: st.param.b32 [func_retval0], %r7;145; CHECK-NEXT: ret;146 %r = fdiv <2 x bfloat> %a, %b147 ret <2 x bfloat> %r148}149 150define <2 x bfloat> @test_fneg(<2 x bfloat> %a) #0 {151; CHECK-LABEL: test_fneg(152; CHECK: {153; CHECK-NEXT: .reg .b32 %r<3>;154; CHECK-EMPTY:155; CHECK-NEXT: // %bb.0:156; CHECK-NEXT: ld.param.b32 %r1, [test_fneg_param_0];157; CHECK-NEXT: neg.bf16x2 %r2, %r1;158; CHECK-NEXT: st.param.b32 [func_retval0], %r2;159; CHECK-NEXT: ret;160 %r = fneg <2 x bfloat> %a161 ret <2 x bfloat> %r162}163 164define void @test_ldst_v2bf16(ptr %a, ptr %b) {165; CHECK-LABEL: test_ldst_v2bf16(166; CHECK: {167; CHECK-NEXT: .reg .b32 %r<2>;168; CHECK-NEXT: .reg .b64 %rd<3>;169; CHECK-EMPTY:170; CHECK-NEXT: // %bb.0:171; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v2bf16_param_0];172; CHECK-NEXT: ld.b32 %r1, [%rd1];173; CHECK-NEXT: ld.param.b64 %rd2, [test_ldst_v2bf16_param_1];174; CHECK-NEXT: st.b32 [%rd2], %r1;175; CHECK-NEXT: ret;176 %t1 = load <2 x bfloat>, ptr %a177 store <2 x bfloat> %t1, ptr %b, align 16178 ret void179}180 181define void @test_ldst_v3bf16(ptr %a, ptr %b) {182; CHECK-LABEL: test_ldst_v3bf16(183; CHECK: {184; CHECK-NEXT: .reg .b16 %rs<2>;185; CHECK-NEXT: .reg .b32 %r<2>;186; CHECK-NEXT: .reg .b64 %rd<4>;187; CHECK-EMPTY:188; CHECK-NEXT: // %bb.0:189; CHECK-NEXT: ld.param.b64 %rd1, [test_ldst_v3bf16_param_0];190; CHECK-NEXT: ld.b64 %rd2, [%rd1];191; CHECK-NEXT: mov.b64 {_, %r1}, %rd2;192; CHECK-NEXT: ld.param.b64 %rd3, [test_ldst_v3bf16_param_1];193; CHECK-NEXT: st.b32 [%rd3], %rd2;194; CHECK-NEXT: mov.b32 {%rs1, _}, %r1;195; CHECK-NEXT: st.b16 [%rd3+4], %rs1;196; CHECK-NEXT: ret;197 %t1 = load <3 x bfloat>, ptr %a198 store <3 x bfloat> %t1, ptr %b, align 16199 ret void200}201 202declare <2 x bfloat> @test_callee(<2 x bfloat> %a, <2 x bfloat> %b) #0203 204define <2 x bfloat> @test_call(<2 x bfloat> %a, <2 x bfloat> %b) #0 {205; CHECK-LABEL: test_call(206; CHECK: {207; CHECK-NEXT: .reg .b32 %r<4>;208; CHECK-EMPTY:209; CHECK-NEXT: // %bb.0:210; CHECK-NEXT: ld.param.b32 %r1, [test_call_param_0];211; CHECK-NEXT: { // callseq 0, 0212; CHECK-NEXT: .param .align 4 .b8 param0[4];213; CHECK-NEXT: .param .align 4 .b8 param1[4];214; CHECK-NEXT: .param .align 4 .b8 retval0[4];215; CHECK-NEXT: ld.param.b32 %r2, [test_call_param_1];216; CHECK-NEXT: st.param.b32 [param1], %r2;217; CHECK-NEXT: st.param.b32 [param0], %r1;218; CHECK-NEXT: call.uni (retval0), test_callee, (param0, param1);219; CHECK-NEXT: ld.param.b32 %r3, [retval0];220; CHECK-NEXT: } // callseq 0221; CHECK-NEXT: st.param.b32 [func_retval0], %r3;222; CHECK-NEXT: ret;223 %r = call <2 x bfloat> @test_callee(<2 x bfloat> %a, <2 x bfloat> %b)224 ret <2 x bfloat> %r225}226 227define <2 x bfloat> @test_select(<2 x bfloat> %a, <2 x bfloat> %b, i1 zeroext %c) #0 {228; CHECK-LABEL: test_select(229; CHECK: {230; CHECK-NEXT: .reg .pred %p<2>;231; CHECK-NEXT: .reg .b16 %rs<3>;232; CHECK-NEXT: .reg .b32 %r<2>;233; CHECK-NEXT: .reg .b64 %rd<4>;234; CHECK-EMPTY:235; CHECK-NEXT: // %bb.0:236; CHECK-NEXT: ld.param.b8 %rs1, [test_select_param_2];237; CHECK-NEXT: and.b16 %rs2, %rs1, 1;238; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;239; CHECK-NEXT: mov.b64 %rd1, test_select_param_1;240; CHECK-NEXT: mov.b64 %rd2, test_select_param_0;241; CHECK-NEXT: selp.b64 %rd3, %rd2, %rd1, %p1;242; CHECK-NEXT: ld.param.b32 %r1, [%rd3];243; CHECK-NEXT: st.param.b32 [func_retval0], %r1;244; CHECK-NEXT: ret;245 %r = select i1 %c, <2 x bfloat> %a, <2 x bfloat> %b246 ret <2 x bfloat> %r247}248 249define <2 x bfloat> @test_select_cc(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c, <2 x bfloat> %d) #0 {250; SM80-LABEL: test_select_cc(251; SM80: {252; SM80-NEXT: .reg .pred %p<3>;253; SM80-NEXT: .reg .b16 %rs<11>;254; SM80-NEXT: .reg .b32 %r<5>;255; SM80-EMPTY:256; SM80-NEXT: // %bb.0:257; SM80-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_param_0];258; SM80-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_param_3];259; SM80-NEXT: cvt.f32.bf16 %r1, %rs3;260; SM80-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_select_cc_param_2];261; SM80-NEXT: cvt.f32.bf16 %r2, %rs5;262; SM80-NEXT: setp.neu.f32 %p1, %r2, %r1;263; SM80-NEXT: cvt.f32.bf16 %r3, %rs4;264; SM80-NEXT: cvt.f32.bf16 %r4, %rs6;265; SM80-NEXT: setp.neu.f32 %p2, %r4, %r3;266; SM80-NEXT: ld.param.v2.b16 {%rs7, %rs8}, [test_select_cc_param_1];267; SM80-NEXT: selp.b16 %rs9, %rs2, %rs8, %p2;268; SM80-NEXT: selp.b16 %rs10, %rs1, %rs7, %p1;269; SM80-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};270; SM80-NEXT: ret;271;272; SM90-LABEL: test_select_cc(273; SM90: {274; SM90-NEXT: .reg .pred %p<3>;275; SM90-NEXT: .reg .b16 %rs<7>;276; SM90-NEXT: .reg .b32 %r<3>;277; SM90-EMPTY:278; SM90-NEXT: // %bb.0:279; SM90-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_param_0];280; SM90-NEXT: ld.param.b32 %r1, [test_select_cc_param_2];281; SM90-NEXT: ld.param.b32 %r2, [test_select_cc_param_3];282; SM90-NEXT: setp.neu.bf16x2 %p1|%p2, %r1, %r2;283; SM90-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_param_1];284; SM90-NEXT: selp.b16 %rs5, %rs2, %rs4, %p2;285; SM90-NEXT: selp.b16 %rs6, %rs1, %rs3, %p1;286; SM90-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};287; SM90-NEXT: ret;288 %cc = fcmp une <2 x bfloat> %c, %d289 %r = select <2 x i1> %cc, <2 x bfloat> %a, <2 x bfloat> %b290 ret <2 x bfloat> %r291}292 293define <2 x float> @test_select_cc_f32_bf16(<2 x float> %a, <2 x float> %b,294; SM80-LABEL: test_select_cc_f32_bf16(295; SM80: {296; SM80-NEXT: .reg .pred %p<3>;297; SM80-NEXT: .reg .b16 %rs<5>;298; SM80-NEXT: .reg .b32 %r<11>;299; SM80-EMPTY:300; SM80-NEXT: // %bb.0:301; SM80-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f32_bf16_param_0];302; SM80-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_f32_bf16_param_3];303; SM80-NEXT: cvt.f32.bf16 %r3, %rs1;304; SM80-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_f32_bf16_param_2];305; SM80-NEXT: cvt.f32.bf16 %r4, %rs3;306; SM80-NEXT: setp.neu.f32 %p1, %r4, %r3;307; SM80-NEXT: cvt.f32.bf16 %r5, %rs2;308; SM80-NEXT: cvt.f32.bf16 %r6, %rs4;309; SM80-NEXT: setp.neu.f32 %p2, %r6, %r5;310; SM80-NEXT: ld.param.v2.b32 {%r7, %r8}, [test_select_cc_f32_bf16_param_1];311; SM80-NEXT: selp.f32 %r9, %r2, %r8, %p2;312; SM80-NEXT: selp.f32 %r10, %r1, %r7, %p1;313; SM80-NEXT: st.param.v2.b32 [func_retval0], {%r10, %r9};314; SM80-NEXT: ret;315;316; SM90-LABEL: test_select_cc_f32_bf16(317; SM90: {318; SM90-NEXT: .reg .pred %p<3>;319; SM90-NEXT: .reg .b32 %r<9>;320; SM90-EMPTY:321; SM90-NEXT: // %bb.0:322; SM90-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f32_bf16_param_0];323; SM90-NEXT: ld.param.b32 %r3, [test_select_cc_f32_bf16_param_2];324; SM90-NEXT: ld.param.b32 %r4, [test_select_cc_f32_bf16_param_3];325; SM90-NEXT: setp.neu.bf16x2 %p1|%p2, %r3, %r4;326; SM90-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_f32_bf16_param_1];327; SM90-NEXT: selp.f32 %r7, %r2, %r6, %p2;328; SM90-NEXT: selp.f32 %r8, %r1, %r5, %p1;329; SM90-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};330; SM90-NEXT: ret;331 <2 x bfloat> %c, <2 x bfloat> %d) #0 {332 %cc = fcmp une <2 x bfloat> %c, %d333 %r = select <2 x i1> %cc, <2 x float> %a, <2 x float> %b334 ret <2 x float> %r335}336 337define <2 x bfloat> @test_select_cc_bf16_f32(<2 x bfloat> %a, <2 x bfloat> %b,338; CHECK-LABEL: test_select_cc_bf16_f32(339; CHECK: {340; CHECK-NEXT: .reg .pred %p<3>;341; CHECK-NEXT: .reg .b16 %rs<7>;342; CHECK-NEXT: .reg .b32 %r<5>;343; CHECK-EMPTY:344; CHECK-NEXT: // %bb.0:345; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_bf16_f32_param_0];346; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_bf16_f32_param_2];347; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_bf16_f32_param_3];348; CHECK-NEXT: setp.neu.f32 %p1, %r1, %r3;349; CHECK-NEXT: setp.neu.f32 %p2, %r2, %r4;350; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_bf16_f32_param_1];351; CHECK-NEXT: selp.b16 %rs5, %rs2, %rs4, %p2;352; CHECK-NEXT: selp.b16 %rs6, %rs1, %rs3, %p1;353; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};354; CHECK-NEXT: ret;355 <2 x float> %c, <2 x float> %d) #0 {356 %cc = fcmp une <2 x float> %c, %d357 %r = select <2 x i1> %cc, <2 x bfloat> %a, <2 x bfloat> %b358 ret <2 x bfloat> %r359}360 361define <2 x bfloat> @test_fptrunc_2xfloat(<2 x float> %a) #0 {362; CHECK-LABEL: test_fptrunc_2xfloat(363; CHECK: {364; CHECK-NEXT: .reg .b32 %r<4>;365; CHECK-EMPTY:366; CHECK-NEXT: // %bb.0:367; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptrunc_2xfloat_param_0];368; CHECK-NEXT: cvt.rn.bf16x2.f32 %r3, %r2, %r1;369; CHECK-NEXT: st.param.b32 [func_retval0], %r3;370; CHECK-NEXT: ret;371 %r = fptrunc <2 x float> %a to <2 x bfloat>372 ret <2 x bfloat> %r373}374 375define <2 x float> @test_fpext_2xfloat(<2 x bfloat> %a) #0 {376; CHECK-LABEL: test_fpext_2xfloat(377; CHECK: {378; CHECK-NEXT: .reg .b16 %rs<3>;379; CHECK-NEXT: .reg .b32 %r<3>;380; CHECK-EMPTY:381; CHECK-NEXT: // %bb.0:382; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fpext_2xfloat_param_0];383; CHECK-NEXT: cvt.f32.bf16 %r1, %rs2;384; CHECK-NEXT: cvt.f32.bf16 %r2, %rs1;385; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};386; CHECK-NEXT: ret;387 %r = fpext <2 x bfloat> %a to <2 x float>388 ret <2 x float> %r389}390 391define <2 x i16> @test_bitcast_2xbf16_to_2xi16(<2 x bfloat> %a) #0 {392; CHECK-LABEL: test_bitcast_2xbf16_to_2xi16(393; CHECK: {394; CHECK-NEXT: .reg .b32 %r<2>;395; CHECK-EMPTY:396; CHECK-NEXT: // %bb.0:397; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_2xbf16_to_2xi16_param_0];398; CHECK-NEXT: st.param.b32 [func_retval0], %r1;399; CHECK-NEXT: ret;400 %r = bitcast <2 x bfloat> %a to <2 x i16>401 ret <2 x i16> %r402}403 404define <2 x bfloat> @test_bitcast_2xi16_to_2xbf16(<2 x i16> %a) #0 {405; CHECK-LABEL: test_bitcast_2xi16_to_2xbf16(406; CHECK: {407; CHECK-NEXT: .reg .b32 %r<2>;408; CHECK-EMPTY:409; CHECK-NEXT: // %bb.0:410; CHECK-NEXT: ld.param.b32 %r1, [test_bitcast_2xi16_to_2xbf16_param_0];411; CHECK-NEXT: st.param.b32 [func_retval0], %r1;412; CHECK-NEXT: ret;413 %r = bitcast <2 x i16> %a to <2 x bfloat>414 ret <2 x bfloat> %r415}416 417declare <2 x bfloat> @llvm.sqrt.f16(<2 x bfloat> %a) #0418declare <2 x bfloat> @llvm.powi.f16(<2 x bfloat> %a, <2 x i32> %b) #0419declare <2 x bfloat> @llvm.sin.f16(<2 x bfloat> %a) #0420declare <2 x bfloat> @llvm.cos.f16(<2 x bfloat> %a) #0421declare <2 x bfloat> @llvm.pow.f16(<2 x bfloat> %a, <2 x bfloat> %b) #0422declare <2 x bfloat> @llvm.exp.f16(<2 x bfloat> %a) #0423declare <2 x bfloat> @llvm.exp2.f16(<2 x bfloat> %a) #0424declare <2 x bfloat> @llvm.log.f16(<2 x bfloat> %a) #0425declare <2 x bfloat> @llvm.log10.f16(<2 x bfloat> %a) #0426declare <2 x bfloat> @llvm.log2.f16(<2 x bfloat> %a) #0427declare <2 x bfloat> @llvm.fma.f16(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) #0428declare <2 x bfloat> @llvm.fabs.f16(<2 x bfloat> %a) #0429declare <2 x bfloat> @llvm.minnum.f16(<2 x bfloat> %a, <2 x bfloat> %b) #0430declare <2 x bfloat> @llvm.maxnum.f16(<2 x bfloat> %a, <2 x bfloat> %b) #0431declare <2 x bfloat> @llvm.copysign.f16(<2 x bfloat> %a, <2 x bfloat> %b) #0432declare <2 x bfloat> @llvm.floor.f16(<2 x bfloat> %a) #0433declare <2 x bfloat> @llvm.ceil.f16(<2 x bfloat> %a) #0434declare <2 x bfloat> @llvm.trunc.f16(<2 x bfloat> %a) #0435declare <2 x bfloat> @llvm.rint.f16(<2 x bfloat> %a) #0436declare <2 x bfloat> @llvm.nearbyint.f16(<2 x bfloat> %a) #0437declare <2 x bfloat> @llvm.round.f16(<2 x bfloat> %a) #0438declare <2 x bfloat> @llvm.fmuladd.f16(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) #0439 440define <2 x bfloat> @test_sqrt(<2 x bfloat> %a) #0 {441; CHECK-LABEL: test_sqrt(442; CHECK: {443; CHECK-NEXT: .reg .b16 %rs<3>;444; CHECK-NEXT: .reg .b32 %r<6>;445; CHECK-EMPTY:446; CHECK-NEXT: // %bb.0:447; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sqrt_param_0];448; CHECK-NEXT: cvt.f32.bf16 %r1, %rs1;449; CHECK-NEXT: sqrt.rn.f32 %r2, %r1;450; CHECK-NEXT: cvt.f32.bf16 %r3, %rs2;451; CHECK-NEXT: sqrt.rn.f32 %r4, %r3;452; CHECK-NEXT: cvt.rn.bf16x2.f32 %r5, %r4, %r2;453; CHECK-NEXT: st.param.b32 [func_retval0], %r5;454; CHECK-NEXT: ret;455 %r = call <2 x bfloat> @llvm.sqrt.f16(<2 x bfloat> %a)456 ret <2 x bfloat> %r457}458 459define <2 x bfloat> @test_fmuladd(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) #0 {460; CHECK-LABEL: test_fmuladd(461; CHECK: {462; CHECK-NEXT: .reg .b32 %r<5>;463; CHECK-EMPTY:464; CHECK-NEXT: // %bb.0:465; CHECK-NEXT: ld.param.b32 %r1, [test_fmuladd_param_0];466; CHECK-NEXT: ld.param.b32 %r2, [test_fmuladd_param_1];467; CHECK-NEXT: ld.param.b32 %r3, [test_fmuladd_param_2];468; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;469; CHECK-NEXT: st.param.b32 [func_retval0], %r4;470; CHECK-NEXT: ret;471 %r = call <2 x bfloat> @llvm.fmuladd.f16(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)472 ret <2 x bfloat> %r473}474 475define <2 x bfloat> @test_fabs(<2 x bfloat> %a) #0 {476; CHECK-LABEL: test_fabs(477; CHECK: {478; CHECK-NEXT: .reg .b32 %r<3>;479; CHECK-EMPTY:480; CHECK-NEXT: // %bb.0:481; CHECK-NEXT: ld.param.b32 %r1, [test_fabs_param_0];482; CHECK-NEXT: abs.bf16x2 %r2, %r1;483; CHECK-NEXT: st.param.b32 [func_retval0], %r2;484; CHECK-NEXT: ret;485 %r = call <2 x bfloat> @llvm.fabs.f16(<2 x bfloat> %a)486 ret <2 x bfloat> %r487}488 489define <2 x bfloat> @test_fabs_add(<2 x bfloat> %a, <2 x bfloat> %b) #0 {490; SM80-LABEL: test_fabs_add(491; SM80: {492; SM80-NEXT: .reg .b32 %r<7>;493; SM80-EMPTY:494; SM80-NEXT: // %bb.0:495; SM80-NEXT: ld.param.b32 %r1, [test_fabs_add_param_0];496; SM80-NEXT: mov.b32 %r2, 1065369472;497; SM80-NEXT: fma.rn.bf16x2 %r3, %r1, %r2, %r1;498; SM80-NEXT: ld.param.b32 %r4, [test_fabs_add_param_1];499; SM80-NEXT: abs.bf16x2 %r5, %r3;500; SM80-NEXT: fma.rn.bf16x2 %r6, %r5, %r2, %r4;501; SM80-NEXT: st.param.b32 [func_retval0], %r6;502; SM80-NEXT: ret;503;504; SM90-LABEL: test_fabs_add(505; SM90: {506; SM90-NEXT: .reg .b32 %r<6>;507; SM90-EMPTY:508; SM90-NEXT: // %bb.0:509; SM90-NEXT: ld.param.b32 %r1, [test_fabs_add_param_0];510; SM90-NEXT: add.rn.bf16x2 %r2, %r1, %r1;511; SM90-NEXT: ld.param.b32 %r3, [test_fabs_add_param_1];512; SM90-NEXT: abs.bf16x2 %r4, %r2;513; SM90-NEXT: add.rn.bf16x2 %r5, %r4, %r3;514; SM90-NEXT: st.param.b32 [func_retval0], %r5;515; SM90-NEXT: ret;516 %s = fadd <2 x bfloat> %a, %a517 %r = call <2 x bfloat> @llvm.fabs.f16(<2 x bfloat> %s)518 %d = fadd <2 x bfloat> %r, %b519 ret <2 x bfloat> %d520}521 522define <2 x bfloat> @test_minnum(<2 x bfloat> %a, <2 x bfloat> %b) #0 {523; CHECK-LABEL: test_minnum(524; CHECK: {525; CHECK-NEXT: .reg .b32 %r<4>;526; CHECK-EMPTY:527; CHECK-NEXT: // %bb.0:528; CHECK-NEXT: ld.param.b32 %r1, [test_minnum_param_0];529; CHECK-NEXT: ld.param.b32 %r2, [test_minnum_param_1];530; CHECK-NEXT: min.bf16x2 %r3, %r1, %r2;531; CHECK-NEXT: st.param.b32 [func_retval0], %r3;532; CHECK-NEXT: ret;533 %r = call <2 x bfloat> @llvm.minnum.f16(<2 x bfloat> %a, <2 x bfloat> %b)534 ret <2 x bfloat> %r535}536 537define <2 x bfloat> @test_maxnum(<2 x bfloat> %a, <2 x bfloat> %b) #0 {538; CHECK-LABEL: test_maxnum(539; CHECK: {540; CHECK-NEXT: .reg .b32 %r<4>;541; CHECK-EMPTY:542; CHECK-NEXT: // %bb.0:543; CHECK-NEXT: ld.param.b32 %r1, [test_maxnum_param_0];544; CHECK-NEXT: ld.param.b32 %r2, [test_maxnum_param_1];545; CHECK-NEXT: max.bf16x2 %r3, %r1, %r2;546; CHECK-NEXT: st.param.b32 [func_retval0], %r3;547; CHECK-NEXT: ret;548 %r = call <2 x bfloat> @llvm.maxnum.f16(<2 x bfloat> %a, <2 x bfloat> %b)549 ret <2 x bfloat> %r550}551 552define <2 x bfloat> @test_floor(<2 x bfloat> %a) #0 {553; SM80-LABEL: test_floor(554; SM80: {555; SM80-NEXT: .reg .b16 %rs<3>;556; SM80-NEXT: .reg .b32 %r<6>;557; SM80-EMPTY:558; SM80-NEXT: // %bb.0:559; SM80-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_floor_param_0];560; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;561; SM80-NEXT: cvt.rmi.f32.f32 %r2, %r1;562; SM80-NEXT: cvt.f32.bf16 %r3, %rs2;563; SM80-NEXT: cvt.rmi.f32.f32 %r4, %r3;564; SM80-NEXT: cvt.rn.bf16x2.f32 %r5, %r4, %r2;565; SM80-NEXT: st.param.b32 [func_retval0], %r5;566; SM80-NEXT: ret;567;568; SM90-LABEL: test_floor(569; SM90: {570; SM90-NEXT: .reg .b16 %rs<5>;571; SM90-EMPTY:572; SM90-NEXT: // %bb.0:573; SM90-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_floor_param_0];574; SM90-NEXT: cvt.rmi.bf16.bf16 %rs3, %rs2;575; SM90-NEXT: cvt.rmi.bf16.bf16 %rs4, %rs1;576; SM90-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};577; SM90-NEXT: ret;578 %r = call <2 x bfloat> @llvm.floor.f16(<2 x bfloat> %a)579 ret <2 x bfloat> %r580}581 582define <2 x bfloat> @test_ceil(<2 x bfloat> %a) #0 {583; SM80-LABEL: test_ceil(584; SM80: {585; SM80-NEXT: .reg .b16 %rs<3>;586; SM80-NEXT: .reg .b32 %r<6>;587; SM80-EMPTY:588; SM80-NEXT: // %bb.0:589; SM80-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_ceil_param_0];590; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;591; SM80-NEXT: cvt.rpi.f32.f32 %r2, %r1;592; SM80-NEXT: cvt.f32.bf16 %r3, %rs2;593; SM80-NEXT: cvt.rpi.f32.f32 %r4, %r3;594; SM80-NEXT: cvt.rn.bf16x2.f32 %r5, %r4, %r2;595; SM80-NEXT: st.param.b32 [func_retval0], %r5;596; SM80-NEXT: ret;597;598; SM90-LABEL: test_ceil(599; SM90: {600; SM90-NEXT: .reg .b16 %rs<5>;601; SM90-EMPTY:602; SM90-NEXT: // %bb.0:603; SM90-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_ceil_param_0];604; SM90-NEXT: cvt.rpi.bf16.bf16 %rs3, %rs2;605; SM90-NEXT: cvt.rpi.bf16.bf16 %rs4, %rs1;606; SM90-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};607; SM90-NEXT: ret;608 %r = call <2 x bfloat> @llvm.ceil.f16(<2 x bfloat> %a)609 ret <2 x bfloat> %r610}611 612define <2 x bfloat> @test_trunc(<2 x bfloat> %a) #0 {613; SM80-LABEL: test_trunc(614; SM80: {615; SM80-NEXT: .reg .b16 %rs<3>;616; SM80-NEXT: .reg .b32 %r<6>;617; SM80-EMPTY:618; SM80-NEXT: // %bb.0:619; SM80-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_trunc_param_0];620; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;621; SM80-NEXT: cvt.rzi.f32.f32 %r2, %r1;622; SM80-NEXT: cvt.f32.bf16 %r3, %rs2;623; SM80-NEXT: cvt.rzi.f32.f32 %r4, %r3;624; SM80-NEXT: cvt.rn.bf16x2.f32 %r5, %r4, %r2;625; SM80-NEXT: st.param.b32 [func_retval0], %r5;626; SM80-NEXT: ret;627;628; SM90-LABEL: test_trunc(629; SM90: {630; SM90-NEXT: .reg .b16 %rs<5>;631; SM90-EMPTY:632; SM90-NEXT: // %bb.0:633; SM90-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_trunc_param_0];634; SM90-NEXT: cvt.rzi.bf16.bf16 %rs3, %rs2;635; SM90-NEXT: cvt.rzi.bf16.bf16 %rs4, %rs1;636; SM90-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};637; SM90-NEXT: ret;638 %r = call <2 x bfloat> @llvm.trunc.f16(<2 x bfloat> %a)639 ret <2 x bfloat> %r640}641 642define <2 x bfloat> @test_rint(<2 x bfloat> %a) #0 {643; SM80-LABEL: test_rint(644; SM80: {645; SM80-NEXT: .reg .b16 %rs<3>;646; SM80-NEXT: .reg .b32 %r<6>;647; SM80-EMPTY:648; SM80-NEXT: // %bb.0:649; SM80-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_rint_param_0];650; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;651; SM80-NEXT: cvt.rni.f32.f32 %r2, %r1;652; SM80-NEXT: cvt.f32.bf16 %r3, %rs2;653; SM80-NEXT: cvt.rni.f32.f32 %r4, %r3;654; SM80-NEXT: cvt.rn.bf16x2.f32 %r5, %r4, %r2;655; SM80-NEXT: st.param.b32 [func_retval0], %r5;656; SM80-NEXT: ret;657;658; SM90-LABEL: test_rint(659; SM90: {660; SM90-NEXT: .reg .b16 %rs<5>;661; SM90-EMPTY:662; SM90-NEXT: // %bb.0:663; SM90-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_rint_param_0];664; SM90-NEXT: cvt.rni.bf16.bf16 %rs3, %rs2;665; SM90-NEXT: cvt.rni.bf16.bf16 %rs4, %rs1;666; SM90-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};667; SM90-NEXT: ret;668 %r = call <2 x bfloat> @llvm.rint.f16(<2 x bfloat> %a)669 ret <2 x bfloat> %r670}671 672define <2 x bfloat> @test_round(<2 x bfloat> %a) #0 {673; CHECK-LABEL: test_round(674; CHECK: {675; CHECK-NEXT: .reg .pred %p<5>;676; CHECK-NEXT: .reg .b16 %rs<3>;677; CHECK-NEXT: .reg .b32 %r<20>;678; CHECK-EMPTY:679; CHECK-NEXT: // %bb.0:680; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_round_param_0];681; CHECK-NEXT: cvt.f32.bf16 %r1, %rs1;682; CHECK-NEXT: and.b32 %r2, %r1, -2147483648;683; CHECK-NEXT: or.b32 %r3, %r2, 1056964608;684; CHECK-NEXT: add.rn.f32 %r4, %r1, %r3;685; CHECK-NEXT: cvt.rzi.f32.f32 %r5, %r4;686; CHECK-NEXT: abs.f32 %r6, %r1;687; CHECK-NEXT: setp.gt.f32 %p1, %r6, 0f4B000000;688; CHECK-NEXT: selp.f32 %r7, %r1, %r5, %p1;689; CHECK-NEXT: cvt.rzi.f32.f32 %r8, %r1;690; CHECK-NEXT: setp.lt.f32 %p2, %r6, 0f3F000000;691; CHECK-NEXT: selp.f32 %r9, %r8, %r7, %p2;692; CHECK-NEXT: cvt.f32.bf16 %r10, %rs2;693; CHECK-NEXT: and.b32 %r11, %r10, -2147483648;694; CHECK-NEXT: or.b32 %r12, %r11, 1056964608;695; CHECK-NEXT: add.rn.f32 %r13, %r10, %r12;696; CHECK-NEXT: cvt.rzi.f32.f32 %r14, %r13;697; CHECK-NEXT: abs.f32 %r15, %r10;698; CHECK-NEXT: setp.gt.f32 %p3, %r15, 0f4B000000;699; CHECK-NEXT: selp.f32 %r16, %r10, %r14, %p3;700; CHECK-NEXT: cvt.rzi.f32.f32 %r17, %r10;701; CHECK-NEXT: setp.lt.f32 %p4, %r15, 0f3F000000;702; CHECK-NEXT: selp.f32 %r18, %r17, %r16, %p4;703; CHECK-NEXT: cvt.rn.bf16x2.f32 %r19, %r18, %r9;704; CHECK-NEXT: st.param.b32 [func_retval0], %r19;705; CHECK-NEXT: ret;706 %r = call <2 x bfloat> @llvm.round.f16(<2 x bfloat> %a)707 ret <2 x bfloat> %r708}709 710define <2 x bfloat> @test_copysign(<2 x bfloat> %a, <2 x bfloat> %b) #0 {711; CHECK-LABEL: test_copysign(712; CHECK: {713; CHECK-NEXT: .reg .b32 %r<6>;714; CHECK-EMPTY:715; CHECK-NEXT: // %bb.0:716; CHECK-NEXT: ld.param.b32 %r1, [test_copysign_param_1];717; CHECK-NEXT: and.b32 %r2, %r1, -2147450880;718; CHECK-NEXT: ld.param.b32 %r3, [test_copysign_param_0];719; CHECK-NEXT: and.b32 %r4, %r3, 2147450879;720; CHECK-NEXT: or.b32 %r5, %r4, %r2;721; CHECK-NEXT: st.param.b32 [func_retval0], %r5;722; CHECK-NEXT: ret;723 %r = call <2 x bfloat> @llvm.copysign.f16(<2 x bfloat> %a, <2 x bfloat> %b)724 ret <2 x bfloat> %r725}726 727define void @test_store_bf16x2(ptr %p1, ptr %p2, <2 x bfloat> %v) {728; CHECK-LABEL: test_store_bf16x2(729; CHECK: {730; CHECK-NEXT: .reg .b32 %r<2>;731; CHECK-NEXT: .reg .b64 %rd<3>;732; CHECK-EMPTY:733; CHECK-NEXT: // %bb.0:734; CHECK-NEXT: ld.param.b64 %rd1, [test_store_bf16x2_param_0];735; CHECK-NEXT: ld.param.b32 %r1, [test_store_bf16x2_param_2];736; CHECK-NEXT: st.b32 [%rd1], %r1;737; CHECK-NEXT: ld.param.b64 %rd2, [test_store_bf16x2_param_1];738; CHECK-NEXT: st.b32 [%rd2], 1065369472;739; CHECK-NEXT: ret;740 store <2 x bfloat> %v, ptr %p1741 store <2 x bfloat> <bfloat 1.0, bfloat 1.0>, ptr %p2742 ret void743}744