1898 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_70 -mattr=+ptx71 | FileCheck --check-prefixes=CHECK,SM70 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | FileCheck --check-prefixes=CHECK,SM80 %s4; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 -denormal-fp-math-f32=preserve-sign | FileCheck --check-prefixes=CHECK,SM80-FTZ %s5; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 -denormal-fp-math-f32=preserve-sign | FileCheck --check-prefixes=CHECK,SM90-FTZ %s6; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | FileCheck --check-prefixes=CHECK,SM90 %s7; RUN: %if ptxas-sm_80 && ptxas-isa-7.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | %ptxas-verify -arch=sm_80 %}8; RUN: %if ptxas-sm_80 && ptxas-isa-7.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 -denormal-fp-math-f32=preserve-sign | %ptxas-verify -arch=sm_80 %}9; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}10 11target triple = "nvptx64-nvidia-cuda"12 13; LDST: .b8 bfloat_array[8] = {1, 2, 3, 4, 5, 6, 7, 8};14@"bfloat_array" = addrspace(1) constant [4 x bfloat]15 [bfloat 0xR0201, bfloat 0xR0403, bfloat 0xR0605, bfloat 0xR0807]16 17define bfloat @test_fadd(bfloat %0, bfloat %1) {18; SM70-LABEL: test_fadd(19; SM70: {20; SM70-NEXT: .reg .pred %p<2>;21; SM70-NEXT: .reg .b32 %r<12>;22; SM70-EMPTY:23; SM70-NEXT: // %bb.0:24; SM70-NEXT: ld.param.b16 %r1, [test_fadd_param_1];25; SM70-NEXT: shl.b32 %r2, %r1, 16;26; SM70-NEXT: ld.param.b16 %r3, [test_fadd_param_0];27; SM70-NEXT: shl.b32 %r4, %r3, 16;28; SM70-NEXT: add.rn.f32 %r5, %r4, %r2;29; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;30; SM70-NEXT: add.s32 %r7, %r6, %r5;31; SM70-NEXT: add.s32 %r8, %r7, 32767;32; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;33; SM70-NEXT: or.b32 %r9, %r5, 4194304;34; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;35; SM70-NEXT: shr.u32 %r11, %r10, 16;36; SM70-NEXT: st.param.b16 [func_retval0], %r11;37; SM70-NEXT: ret;38;39; SM80-LABEL: test_fadd(40; SM80: {41; SM80-NEXT: .reg .b16 %rs<5>;42; SM80-EMPTY:43; SM80-NEXT: // %bb.0:44; SM80-NEXT: ld.param.b16 %rs1, [test_fadd_param_0];45; SM80-NEXT: ld.param.b16 %rs2, [test_fadd_param_1];46; SM80-NEXT: mov.b16 %rs3, 0x3F80;47; SM80-NEXT: fma.rn.bf16 %rs4, %rs1, %rs3, %rs2;48; SM80-NEXT: st.param.b16 [func_retval0], %rs4;49; SM80-NEXT: ret;50;51; SM80-FTZ-LABEL: test_fadd(52; SM80-FTZ: {53; SM80-FTZ-NEXT: .reg .b16 %rs<4>;54; SM80-FTZ-NEXT: .reg .b32 %r<4>;55; SM80-FTZ-EMPTY:56; SM80-FTZ-NEXT: // %bb.0:57; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_param_0];58; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_fadd_param_1];59; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs2;60; SM80-FTZ-NEXT: cvt.f32.bf16 %r2, %rs1;61; SM80-FTZ-NEXT: add.rn.ftz.f32 %r3, %r2, %r1;62; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r3;63; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;64; SM80-FTZ-NEXT: ret;65;66; SM90-FTZ-LABEL: test_fadd(67; SM90-FTZ: {68; SM90-FTZ-NEXT: .reg .b16 %rs<4>;69; SM90-FTZ-EMPTY:70; SM90-FTZ-NEXT: // %bb.0:71; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_param_0];72; SM90-FTZ-NEXT: ld.param.b16 %rs2, [test_fadd_param_1];73; SM90-FTZ-NEXT: add.rn.bf16 %rs3, %rs1, %rs2;74; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;75; SM90-FTZ-NEXT: ret;76;77; SM90-LABEL: test_fadd(78; SM90: {79; SM90-NEXT: .reg .b16 %rs<4>;80; SM90-EMPTY:81; SM90-NEXT: // %bb.0:82; SM90-NEXT: ld.param.b16 %rs1, [test_fadd_param_0];83; SM90-NEXT: ld.param.b16 %rs2, [test_fadd_param_1];84; SM90-NEXT: add.rn.bf16 %rs3, %rs1, %rs2;85; SM90-NEXT: st.param.b16 [func_retval0], %rs3;86; SM90-NEXT: ret;87 %3 = fadd bfloat %0, %188 ret bfloat %389}90 91define bfloat @test_fsub(bfloat %0, bfloat %1) {92; SM70-LABEL: test_fsub(93; SM70: {94; SM70-NEXT: .reg .pred %p<2>;95; SM70-NEXT: .reg .b32 %r<12>;96; SM70-EMPTY:97; SM70-NEXT: // %bb.0:98; SM70-NEXT: ld.param.b16 %r1, [test_fsub_param_1];99; SM70-NEXT: shl.b32 %r2, %r1, 16;100; SM70-NEXT: ld.param.b16 %r3, [test_fsub_param_0];101; SM70-NEXT: shl.b32 %r4, %r3, 16;102; SM70-NEXT: sub.rn.f32 %r5, %r4, %r2;103; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;104; SM70-NEXT: add.s32 %r7, %r6, %r5;105; SM70-NEXT: add.s32 %r8, %r7, 32767;106; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;107; SM70-NEXT: or.b32 %r9, %r5, 4194304;108; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;109; SM70-NEXT: shr.u32 %r11, %r10, 16;110; SM70-NEXT: st.param.b16 [func_retval0], %r11;111; SM70-NEXT: ret;112;113; SM80-LABEL: test_fsub(114; SM80: {115; SM80-NEXT: .reg .b16 %rs<5>;116; SM80-EMPTY:117; SM80-NEXT: // %bb.0:118; SM80-NEXT: ld.param.b16 %rs1, [test_fsub_param_0];119; SM80-NEXT: mov.b16 %rs2, 0xBF80;120; SM80-NEXT: ld.param.b16 %rs3, [test_fsub_param_1];121; SM80-NEXT: fma.rn.bf16 %rs4, %rs3, %rs2, %rs1;122; SM80-NEXT: st.param.b16 [func_retval0], %rs4;123; SM80-NEXT: ret;124;125; SM80-FTZ-LABEL: test_fsub(126; SM80-FTZ: {127; SM80-FTZ-NEXT: .reg .b16 %rs<4>;128; SM80-FTZ-NEXT: .reg .b32 %r<4>;129; SM80-FTZ-EMPTY:130; SM80-FTZ-NEXT: // %bb.0:131; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fsub_param_0];132; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_fsub_param_1];133; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs2;134; SM80-FTZ-NEXT: cvt.f32.bf16 %r2, %rs1;135; SM80-FTZ-NEXT: sub.rn.ftz.f32 %r3, %r2, %r1;136; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r3;137; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;138; SM80-FTZ-NEXT: ret;139;140; SM90-FTZ-LABEL: test_fsub(141; SM90-FTZ: {142; SM90-FTZ-NEXT: .reg .b16 %rs<4>;143; SM90-FTZ-EMPTY:144; SM90-FTZ-NEXT: // %bb.0:145; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_fsub_param_0];146; SM90-FTZ-NEXT: ld.param.b16 %rs2, [test_fsub_param_1];147; SM90-FTZ-NEXT: sub.rn.bf16 %rs3, %rs1, %rs2;148; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;149; SM90-FTZ-NEXT: ret;150;151; SM90-LABEL: test_fsub(152; SM90: {153; SM90-NEXT: .reg .b16 %rs<4>;154; SM90-EMPTY:155; SM90-NEXT: // %bb.0:156; SM90-NEXT: ld.param.b16 %rs1, [test_fsub_param_0];157; SM90-NEXT: ld.param.b16 %rs2, [test_fsub_param_1];158; SM90-NEXT: sub.rn.bf16 %rs3, %rs1, %rs2;159; SM90-NEXT: st.param.b16 [func_retval0], %rs3;160; SM90-NEXT: ret;161 %3 = fsub bfloat %0, %1162 ret bfloat %3163}164 165define <2 x bfloat> @test_faddx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {166; SM70-LABEL: test_faddx2(167; SM70: {168; SM70-NEXT: .reg .pred %p<3>;169; SM70-NEXT: .reg .b16 %rs<5>;170; SM70-NEXT: .reg .b32 %r<22>;171; SM70-EMPTY:172; SM70-NEXT: // %bb.0:173; SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_faddx2_param_0];174; SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_faddx2_param_1];175; SM70-NEXT: cvt.u32.u16 %r1, %rs4;176; SM70-NEXT: shl.b32 %r2, %r1, 16;177; SM70-NEXT: cvt.u32.u16 %r3, %rs2;178; SM70-NEXT: shl.b32 %r4, %r3, 16;179; SM70-NEXT: add.rn.f32 %r5, %r4, %r2;180; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;181; SM70-NEXT: add.s32 %r7, %r6, %r5;182; SM70-NEXT: add.s32 %r8, %r7, 32767;183; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;184; SM70-NEXT: or.b32 %r9, %r5, 4194304;185; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;186; SM70-NEXT: cvt.u32.u16 %r11, %rs3;187; SM70-NEXT: shl.b32 %r12, %r11, 16;188; SM70-NEXT: cvt.u32.u16 %r13, %rs1;189; SM70-NEXT: shl.b32 %r14, %r13, 16;190; SM70-NEXT: add.rn.f32 %r15, %r14, %r12;191; SM70-NEXT: bfe.u32 %r16, %r15, 16, 1;192; SM70-NEXT: add.s32 %r17, %r16, %r15;193; SM70-NEXT: add.s32 %r18, %r17, 32767;194; SM70-NEXT: setp.nan.f32 %p2, %r15, %r15;195; SM70-NEXT: or.b32 %r19, %r15, 4194304;196; SM70-NEXT: selp.b32 %r20, %r19, %r18, %p2;197; SM70-NEXT: prmt.b32 %r21, %r20, %r10, 0x7632U;198; SM70-NEXT: st.param.b32 [func_retval0], %r21;199; SM70-NEXT: ret;200;201; SM80-LABEL: test_faddx2(202; SM80: {203; SM80-NEXT: .reg .b32 %r<5>;204; SM80-EMPTY:205; SM80-NEXT: // %bb.0:206; SM80-NEXT: ld.param.b32 %r1, [test_faddx2_param_0];207; SM80-NEXT: ld.param.b32 %r2, [test_faddx2_param_1];208; SM80-NEXT: mov.b32 %r3, 1065369472;209; SM80-NEXT: fma.rn.bf16x2 %r4, %r1, %r3, %r2;210; SM80-NEXT: st.param.b32 [func_retval0], %r4;211; SM80-NEXT: ret;212;213; SM80-FTZ-LABEL: test_faddx2(214; SM80-FTZ: {215; SM80-FTZ-NEXT: .reg .b16 %rs<5>;216; SM80-FTZ-NEXT: .reg .b32 %r<8>;217; SM80-FTZ-EMPTY:218; SM80-FTZ-NEXT: // %bb.0:219; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_faddx2_param_0];220; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_faddx2_param_1];221; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs3;222; SM80-FTZ-NEXT: cvt.f32.bf16 %r2, %rs1;223; SM80-FTZ-NEXT: add.rn.ftz.f32 %r3, %r2, %r1;224; SM80-FTZ-NEXT: cvt.f32.bf16 %r4, %rs4;225; SM80-FTZ-NEXT: cvt.f32.bf16 %r5, %rs2;226; SM80-FTZ-NEXT: add.rn.ftz.f32 %r6, %r5, %r4;227; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;228; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r7;229; SM80-FTZ-NEXT: ret;230;231; SM90-FTZ-LABEL: test_faddx2(232; SM90-FTZ: {233; SM90-FTZ-NEXT: .reg .b32 %r<4>;234; SM90-FTZ-EMPTY:235; SM90-FTZ-NEXT: // %bb.0:236; SM90-FTZ-NEXT: ld.param.b32 %r1, [test_faddx2_param_0];237; SM90-FTZ-NEXT: ld.param.b32 %r2, [test_faddx2_param_1];238; SM90-FTZ-NEXT: add.rn.bf16x2 %r3, %r1, %r2;239; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r3;240; SM90-FTZ-NEXT: ret;241;242; SM90-LABEL: test_faddx2(243; SM90: {244; SM90-NEXT: .reg .b32 %r<4>;245; SM90-EMPTY:246; SM90-NEXT: // %bb.0:247; SM90-NEXT: ld.param.b32 %r1, [test_faddx2_param_0];248; SM90-NEXT: ld.param.b32 %r2, [test_faddx2_param_1];249; SM90-NEXT: add.rn.bf16x2 %r3, %r1, %r2;250; SM90-NEXT: st.param.b32 [func_retval0], %r3;251; SM90-NEXT: ret;252 %r = fadd <2 x bfloat> %a, %b253 ret <2 x bfloat> %r254}255 256define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {257; SM70-LABEL: test_fsubx2(258; SM70: {259; SM70-NEXT: .reg .pred %p<3>;260; SM70-NEXT: .reg .b16 %rs<5>;261; SM70-NEXT: .reg .b32 %r<22>;262; SM70-EMPTY:263; SM70-NEXT: // %bb.0:264; SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fsubx2_param_0];265; SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fsubx2_param_1];266; SM70-NEXT: cvt.u32.u16 %r1, %rs4;267; SM70-NEXT: shl.b32 %r2, %r1, 16;268; SM70-NEXT: cvt.u32.u16 %r3, %rs2;269; SM70-NEXT: shl.b32 %r4, %r3, 16;270; SM70-NEXT: sub.rn.f32 %r5, %r4, %r2;271; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;272; SM70-NEXT: add.s32 %r7, %r6, %r5;273; SM70-NEXT: add.s32 %r8, %r7, 32767;274; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;275; SM70-NEXT: or.b32 %r9, %r5, 4194304;276; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;277; SM70-NEXT: cvt.u32.u16 %r11, %rs3;278; SM70-NEXT: shl.b32 %r12, %r11, 16;279; SM70-NEXT: cvt.u32.u16 %r13, %rs1;280; SM70-NEXT: shl.b32 %r14, %r13, 16;281; SM70-NEXT: sub.rn.f32 %r15, %r14, %r12;282; SM70-NEXT: bfe.u32 %r16, %r15, 16, 1;283; SM70-NEXT: add.s32 %r17, %r16, %r15;284; SM70-NEXT: add.s32 %r18, %r17, 32767;285; SM70-NEXT: setp.nan.f32 %p2, %r15, %r15;286; SM70-NEXT: or.b32 %r19, %r15, 4194304;287; SM70-NEXT: selp.b32 %r20, %r19, %r18, %p2;288; SM70-NEXT: prmt.b32 %r21, %r20, %r10, 0x7632U;289; SM70-NEXT: st.param.b32 [func_retval0], %r21;290; SM70-NEXT: ret;291;292; SM80-LABEL: test_fsubx2(293; SM80: {294; SM80-NEXT: .reg .b32 %r<5>;295; SM80-EMPTY:296; SM80-NEXT: // %bb.0:297; SM80-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];298; SM80-NEXT: mov.b32 %r2, -1082081408;299; SM80-NEXT: ld.param.b32 %r3, [test_fsubx2_param_1];300; SM80-NEXT: fma.rn.bf16x2 %r4, %r3, %r2, %r1;301; SM80-NEXT: st.param.b32 [func_retval0], %r4;302; SM80-NEXT: ret;303;304; SM80-FTZ-LABEL: test_fsubx2(305; SM80-FTZ: {306; SM80-FTZ-NEXT: .reg .b16 %rs<5>;307; SM80-FTZ-NEXT: .reg .b32 %r<8>;308; SM80-FTZ-EMPTY:309; SM80-FTZ-NEXT: // %bb.0:310; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fsubx2_param_0];311; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fsubx2_param_1];312; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs3;313; SM80-FTZ-NEXT: cvt.f32.bf16 %r2, %rs1;314; SM80-FTZ-NEXT: sub.rn.ftz.f32 %r3, %r2, %r1;315; SM80-FTZ-NEXT: cvt.f32.bf16 %r4, %rs4;316; SM80-FTZ-NEXT: cvt.f32.bf16 %r5, %rs2;317; SM80-FTZ-NEXT: sub.rn.ftz.f32 %r6, %r5, %r4;318; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;319; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r7;320; SM80-FTZ-NEXT: ret;321;322; SM90-FTZ-LABEL: test_fsubx2(323; SM90-FTZ: {324; SM90-FTZ-NEXT: .reg .b32 %r<4>;325; SM90-FTZ-EMPTY:326; SM90-FTZ-NEXT: // %bb.0:327; SM90-FTZ-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];328; SM90-FTZ-NEXT: ld.param.b32 %r2, [test_fsubx2_param_1];329; SM90-FTZ-NEXT: sub.rn.bf16x2 %r3, %r1, %r2;330; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r3;331; SM90-FTZ-NEXT: ret;332;333; SM90-LABEL: test_fsubx2(334; SM90: {335; SM90-NEXT: .reg .b32 %r<4>;336; SM90-EMPTY:337; SM90-NEXT: // %bb.0:338; SM90-NEXT: ld.param.b32 %r1, [test_fsubx2_param_0];339; SM90-NEXT: ld.param.b32 %r2, [test_fsubx2_param_1];340; SM90-NEXT: sub.rn.bf16x2 %r3, %r1, %r2;341; SM90-NEXT: st.param.b32 [func_retval0], %r3;342; SM90-NEXT: ret;343 %r = fsub <2 x bfloat> %a, %b344 ret <2 x bfloat> %r345}346 347define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {348; SM70-LABEL: test_fmulx2(349; SM70: {350; SM70-NEXT: .reg .pred %p<3>;351; SM70-NEXT: .reg .b16 %rs<5>;352; SM70-NEXT: .reg .b32 %r<22>;353; SM70-EMPTY:354; SM70-NEXT: // %bb.0:355; SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fmulx2_param_0];356; SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fmulx2_param_1];357; SM70-NEXT: cvt.u32.u16 %r1, %rs4;358; SM70-NEXT: shl.b32 %r2, %r1, 16;359; SM70-NEXT: cvt.u32.u16 %r3, %rs2;360; SM70-NEXT: shl.b32 %r4, %r3, 16;361; SM70-NEXT: mul.rn.f32 %r5, %r4, %r2;362; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;363; SM70-NEXT: add.s32 %r7, %r6, %r5;364; SM70-NEXT: add.s32 %r8, %r7, 32767;365; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;366; SM70-NEXT: or.b32 %r9, %r5, 4194304;367; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;368; SM70-NEXT: cvt.u32.u16 %r11, %rs3;369; SM70-NEXT: shl.b32 %r12, %r11, 16;370; SM70-NEXT: cvt.u32.u16 %r13, %rs1;371; SM70-NEXT: shl.b32 %r14, %r13, 16;372; SM70-NEXT: mul.rn.f32 %r15, %r14, %r12;373; SM70-NEXT: bfe.u32 %r16, %r15, 16, 1;374; SM70-NEXT: add.s32 %r17, %r16, %r15;375; SM70-NEXT: add.s32 %r18, %r17, 32767;376; SM70-NEXT: setp.nan.f32 %p2, %r15, %r15;377; SM70-NEXT: or.b32 %r19, %r15, 4194304;378; SM70-NEXT: selp.b32 %r20, %r19, %r18, %p2;379; SM70-NEXT: prmt.b32 %r21, %r20, %r10, 0x7632U;380; SM70-NEXT: st.param.b32 [func_retval0], %r21;381; SM70-NEXT: ret;382;383; SM80-LABEL: test_fmulx2(384; SM80: {385; SM80-NEXT: .reg .b32 %r<5>;386; SM80-EMPTY:387; SM80-NEXT: // %bb.0:388; SM80-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];389; SM80-NEXT: mov.b32 %r2, -2147450880;390; SM80-NEXT: ld.param.b32 %r3, [test_fmulx2_param_1];391; SM80-NEXT: fma.rn.bf16x2 %r4, %r1, %r3, %r2;392; SM80-NEXT: st.param.b32 [func_retval0], %r4;393; SM80-NEXT: ret;394;395; SM80-FTZ-LABEL: test_fmulx2(396; SM80-FTZ: {397; SM80-FTZ-NEXT: .reg .b16 %rs<5>;398; SM80-FTZ-NEXT: .reg .b32 %r<8>;399; SM80-FTZ-EMPTY:400; SM80-FTZ-NEXT: // %bb.0:401; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fmulx2_param_0];402; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fmulx2_param_1];403; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs3;404; SM80-FTZ-NEXT: cvt.f32.bf16 %r2, %rs1;405; SM80-FTZ-NEXT: mul.rn.ftz.f32 %r3, %r2, %r1;406; SM80-FTZ-NEXT: cvt.f32.bf16 %r4, %rs4;407; SM80-FTZ-NEXT: cvt.f32.bf16 %r5, %rs2;408; SM80-FTZ-NEXT: mul.rn.ftz.f32 %r6, %r5, %r4;409; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;410; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r7;411; SM80-FTZ-NEXT: ret;412;413; SM90-FTZ-LABEL: test_fmulx2(414; SM90-FTZ: {415; SM90-FTZ-NEXT: .reg .b32 %r<4>;416; SM90-FTZ-EMPTY:417; SM90-FTZ-NEXT: // %bb.0:418; SM90-FTZ-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];419; SM90-FTZ-NEXT: ld.param.b32 %r2, [test_fmulx2_param_1];420; SM90-FTZ-NEXT: mul.rn.bf16x2 %r3, %r1, %r2;421; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r3;422; SM90-FTZ-NEXT: ret;423;424; SM90-LABEL: test_fmulx2(425; SM90: {426; SM90-NEXT: .reg .b32 %r<4>;427; SM90-EMPTY:428; SM90-NEXT: // %bb.0:429; SM90-NEXT: ld.param.b32 %r1, [test_fmulx2_param_0];430; SM90-NEXT: ld.param.b32 %r2, [test_fmulx2_param_1];431; SM90-NEXT: mul.rn.bf16x2 %r3, %r1, %r2;432; SM90-NEXT: st.param.b32 [func_retval0], %r3;433; SM90-NEXT: ret;434 %r = fmul <2 x bfloat> %a, %b435 ret <2 x bfloat> %r436}437 438define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 {439; SM70-LABEL: test_fdiv(440; SM70: {441; SM70-NEXT: .reg .pred %p<3>;442; SM70-NEXT: .reg .b16 %rs<5>;443; SM70-NEXT: .reg .b32 %r<22>;444; SM70-EMPTY:445; SM70-NEXT: // %bb.0:446; SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];447; SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];448; SM70-NEXT: cvt.u32.u16 %r1, %rs4;449; SM70-NEXT: shl.b32 %r2, %r1, 16;450; SM70-NEXT: cvt.u32.u16 %r3, %rs2;451; SM70-NEXT: shl.b32 %r4, %r3, 16;452; SM70-NEXT: div.rn.f32 %r5, %r4, %r2;453; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;454; SM70-NEXT: add.s32 %r7, %r6, %r5;455; SM70-NEXT: add.s32 %r8, %r7, 32767;456; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;457; SM70-NEXT: or.b32 %r9, %r5, 4194304;458; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;459; SM70-NEXT: cvt.u32.u16 %r11, %rs3;460; SM70-NEXT: shl.b32 %r12, %r11, 16;461; SM70-NEXT: cvt.u32.u16 %r13, %rs1;462; SM70-NEXT: shl.b32 %r14, %r13, 16;463; SM70-NEXT: div.rn.f32 %r15, %r14, %r12;464; SM70-NEXT: bfe.u32 %r16, %r15, 16, 1;465; SM70-NEXT: add.s32 %r17, %r16, %r15;466; SM70-NEXT: add.s32 %r18, %r17, 32767;467; SM70-NEXT: setp.nan.f32 %p2, %r15, %r15;468; SM70-NEXT: or.b32 %r19, %r15, 4194304;469; SM70-NEXT: selp.b32 %r20, %r19, %r18, %p2;470; SM70-NEXT: prmt.b32 %r21, %r20, %r10, 0x7632U;471; SM70-NEXT: st.param.b32 [func_retval0], %r21;472; SM70-NEXT: ret;473;474; SM80-LABEL: test_fdiv(475; SM80: {476; SM80-NEXT: .reg .b16 %rs<5>;477; SM80-NEXT: .reg .b32 %r<8>;478; SM80-EMPTY:479; SM80-NEXT: // %bb.0:480; SM80-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];481; SM80-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];482; SM80-NEXT: cvt.f32.bf16 %r1, %rs3;483; SM80-NEXT: cvt.f32.bf16 %r2, %rs1;484; SM80-NEXT: div.rn.f32 %r3, %r2, %r1;485; SM80-NEXT: cvt.f32.bf16 %r4, %rs4;486; SM80-NEXT: cvt.f32.bf16 %r5, %rs2;487; SM80-NEXT: div.rn.f32 %r6, %r5, %r4;488; SM80-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;489; SM80-NEXT: st.param.b32 [func_retval0], %r7;490; SM80-NEXT: ret;491;492; SM80-FTZ-LABEL: test_fdiv(493; SM80-FTZ: {494; SM80-FTZ-NEXT: .reg .b16 %rs<5>;495; SM80-FTZ-NEXT: .reg .b32 %r<8>;496; SM80-FTZ-EMPTY:497; SM80-FTZ-NEXT: // %bb.0:498; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];499; SM80-FTZ-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];500; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs3;501; SM80-FTZ-NEXT: cvt.f32.bf16 %r2, %rs1;502; SM80-FTZ-NEXT: div.rn.ftz.f32 %r3, %r2, %r1;503; SM80-FTZ-NEXT: cvt.f32.bf16 %r4, %rs4;504; SM80-FTZ-NEXT: cvt.f32.bf16 %r5, %rs2;505; SM80-FTZ-NEXT: div.rn.ftz.f32 %r6, %r5, %r4;506; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;507; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r7;508; SM80-FTZ-NEXT: ret;509;510; SM90-FTZ-LABEL: test_fdiv(511; SM90-FTZ: {512; SM90-FTZ-NEXT: .reg .b16 %rs<5>;513; SM90-FTZ-NEXT: .reg .b32 %r<8>;514; SM90-FTZ-EMPTY:515; SM90-FTZ-NEXT: // %bb.0:516; SM90-FTZ-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];517; SM90-FTZ-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];518; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r1, %rs3;519; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r2, %rs1;520; SM90-FTZ-NEXT: div.rn.ftz.f32 %r3, %r2, %r1;521; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r4, %rs4;522; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r5, %rs2;523; SM90-FTZ-NEXT: div.rn.ftz.f32 %r6, %r5, %r4;524; SM90-FTZ-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;525; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r7;526; SM90-FTZ-NEXT: ret;527;528; SM90-LABEL: test_fdiv(529; SM90: {530; SM90-NEXT: .reg .b16 %rs<5>;531; SM90-NEXT: .reg .b32 %r<8>;532; SM90-EMPTY:533; SM90-NEXT: // %bb.0:534; SM90-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];535; SM90-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];536; SM90-NEXT: cvt.f32.bf16 %r1, %rs3;537; SM90-NEXT: cvt.f32.bf16 %r2, %rs1;538; SM90-NEXT: div.rn.f32 %r3, %r2, %r1;539; SM90-NEXT: cvt.f32.bf16 %r4, %rs4;540; SM90-NEXT: cvt.f32.bf16 %r5, %rs2;541; SM90-NEXT: div.rn.f32 %r6, %r5, %r4;542; SM90-NEXT: cvt.rn.bf16x2.f32 %r7, %r6, %r3;543; SM90-NEXT: st.param.b32 [func_retval0], %r7;544; SM90-NEXT: ret;545 %r = fdiv <2 x bfloat> %a, %b546 ret <2 x bfloat> %r547}548 549define bfloat @test_extract_0(<2 x bfloat> %a) #0 {550; CHECK-LABEL: test_extract_0(551; CHECK: {552; CHECK-NEXT: .reg .b16 %rs<2>;553; CHECK-EMPTY:554; CHECK-NEXT: // %bb.0:555; CHECK-NEXT: ld.param.b16 %rs1, [test_extract_0_param_0];556; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;557; CHECK-NEXT: ret;558 %e = extractelement <2 x bfloat> %a, i32 0559 ret bfloat %e560}561 562define bfloat @test_extract_1(<2 x bfloat> %a) #0 {563; CHECK-LABEL: test_extract_1(564; CHECK: {565; CHECK-NEXT: .reg .b16 %rs<2>;566; CHECK-EMPTY:567; CHECK-NEXT: // %bb.0:568; CHECK-NEXT: ld.param.b16 %rs1, [test_extract_1_param_0+2];569; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;570; CHECK-NEXT: ret;571 %e = extractelement <2 x bfloat> %a, i32 1572 ret bfloat %e573}574 575define float @test_fpext_float(bfloat %a) #0 {576; SM70-LABEL: test_fpext_float(577; SM70: {578; SM70-NEXT: .reg .b32 %r<3>;579; SM70-EMPTY:580; SM70-NEXT: // %bb.0:581; SM70-NEXT: ld.param.b16 %r1, [test_fpext_float_param_0];582; SM70-NEXT: shl.b32 %r2, %r1, 16;583; SM70-NEXT: st.param.b32 [func_retval0], %r2;584; SM70-NEXT: ret;585;586; SM80-LABEL: test_fpext_float(587; SM80: {588; SM80-NEXT: .reg .b16 %rs<2>;589; SM80-NEXT: .reg .b32 %r<2>;590; SM80-EMPTY:591; SM80-NEXT: // %bb.0:592; SM80-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];593; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;594; SM80-NEXT: st.param.b32 [func_retval0], %r1;595; SM80-NEXT: ret;596;597; SM80-FTZ-LABEL: test_fpext_float(598; SM80-FTZ: {599; SM80-FTZ-NEXT: .reg .b16 %rs<2>;600; SM80-FTZ-NEXT: .reg .b32 %r<2>;601; SM80-FTZ-EMPTY:602; SM80-FTZ-NEXT: // %bb.0:603; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];604; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs1;605; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r1;606; SM80-FTZ-NEXT: ret;607;608; SM90-FTZ-LABEL: test_fpext_float(609; SM90-FTZ: {610; SM90-FTZ-NEXT: .reg .b16 %rs<2>;611; SM90-FTZ-NEXT: .reg .b32 %r<2>;612; SM90-FTZ-EMPTY:613; SM90-FTZ-NEXT: // %bb.0:614; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];615; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r1, %rs1;616; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r1;617; SM90-FTZ-NEXT: ret;618;619; SM90-LABEL: test_fpext_float(620; SM90: {621; SM90-NEXT: .reg .b16 %rs<2>;622; SM90-NEXT: .reg .b32 %r<2>;623; SM90-EMPTY:624; SM90-NEXT: // %bb.0:625; SM90-NEXT: ld.param.b16 %rs1, [test_fpext_float_param_0];626; SM90-NEXT: cvt.f32.bf16 %r1, %rs1;627; SM90-NEXT: st.param.b32 [func_retval0], %r1;628; SM90-NEXT: ret;629 %r = fpext bfloat %a to float630 ret float %r631}632 633define bfloat @test_fptrunc_float(float %a) #0 {634; SM70-LABEL: test_fptrunc_float(635; SM70: {636; SM70-NEXT: .reg .pred %p<2>;637; SM70-NEXT: .reg .b32 %r<8>;638; SM70-EMPTY:639; SM70-NEXT: // %bb.0:640; SM70-NEXT: ld.param.b32 %r1, [test_fptrunc_float_param_0];641; SM70-NEXT: bfe.u32 %r2, %r1, 16, 1;642; SM70-NEXT: add.s32 %r3, %r2, %r1;643; SM70-NEXT: add.s32 %r4, %r3, 32767;644; SM70-NEXT: setp.nan.f32 %p1, %r1, %r1;645; SM70-NEXT: or.b32 %r5, %r1, 4194304;646; SM70-NEXT: selp.b32 %r6, %r5, %r4, %p1;647; SM70-NEXT: shr.u32 %r7, %r6, 16;648; SM70-NEXT: st.param.b16 [func_retval0], %r7;649; SM70-NEXT: ret;650;651; SM80-LABEL: test_fptrunc_float(652; SM80: {653; SM80-NEXT: .reg .b16 %rs<2>;654; SM80-NEXT: .reg .b32 %r<2>;655; SM80-EMPTY:656; SM80-NEXT: // %bb.0:657; SM80-NEXT: ld.param.b32 %r1, [test_fptrunc_float_param_0];658; SM80-NEXT: cvt.rn.bf16.f32 %rs1, %r1;659; SM80-NEXT: st.param.b16 [func_retval0], %rs1;660; SM80-NEXT: ret;661;662; SM80-FTZ-LABEL: test_fptrunc_float(663; SM80-FTZ: {664; SM80-FTZ-NEXT: .reg .b16 %rs<2>;665; SM80-FTZ-NEXT: .reg .b32 %r<2>;666; SM80-FTZ-EMPTY:667; SM80-FTZ-NEXT: // %bb.0:668; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_fptrunc_float_param_0];669; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs1, %r1;670; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs1;671; SM80-FTZ-NEXT: ret;672;673; SM90-FTZ-LABEL: test_fptrunc_float(674; SM90-FTZ: {675; SM90-FTZ-NEXT: .reg .b16 %rs<2>;676; SM90-FTZ-NEXT: .reg .b32 %r<2>;677; SM90-FTZ-EMPTY:678; SM90-FTZ-NEXT: // %bb.0:679; SM90-FTZ-NEXT: ld.param.b32 %r1, [test_fptrunc_float_param_0];680; SM90-FTZ-NEXT: cvt.rn.bf16.f32 %rs1, %r1;681; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs1;682; SM90-FTZ-NEXT: ret;683;684; SM90-LABEL: test_fptrunc_float(685; SM90: {686; SM90-NEXT: .reg .b16 %rs<2>;687; SM90-NEXT: .reg .b32 %r<2>;688; SM90-EMPTY:689; SM90-NEXT: // %bb.0:690; SM90-NEXT: ld.param.b32 %r1, [test_fptrunc_float_param_0];691; SM90-NEXT: cvt.rn.bf16.f32 %rs1, %r1;692; SM90-NEXT: st.param.b16 [func_retval0], %rs1;693; SM90-NEXT: ret;694 %r = fptrunc float %a to bfloat695 ret bfloat %r696}697 698define bfloat @test_fadd_imm_1(bfloat %a) #0 {699; SM70-LABEL: test_fadd_imm_1(700; SM70: {701; SM70-NEXT: .reg .pred %p<2>;702; SM70-NEXT: .reg .b32 %r<10>;703; SM70-EMPTY:704; SM70-NEXT: // %bb.0:705; SM70-NEXT: ld.param.b16 %r1, [test_fadd_imm_1_param_0];706; SM70-NEXT: shl.b32 %r2, %r1, 16;707; SM70-NEXT: add.rn.f32 %r3, %r2, 0f3F800000;708; SM70-NEXT: bfe.u32 %r4, %r3, 16, 1;709; SM70-NEXT: add.s32 %r5, %r4, %r3;710; SM70-NEXT: add.s32 %r6, %r5, 32767;711; SM70-NEXT: setp.nan.f32 %p1, %r3, %r3;712; SM70-NEXT: or.b32 %r7, %r3, 4194304;713; SM70-NEXT: selp.b32 %r8, %r7, %r6, %p1;714; SM70-NEXT: shr.u32 %r9, %r8, 16;715; SM70-NEXT: st.param.b16 [func_retval0], %r9;716; SM70-NEXT: ret;717;718; SM80-LABEL: test_fadd_imm_1(719; SM80: {720; SM80-NEXT: .reg .b16 %rs<4>;721; SM80-EMPTY:722; SM80-NEXT: // %bb.0:723; SM80-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];724; SM80-NEXT: mov.b16 %rs2, 0x3F80;725; SM80-NEXT: fma.rn.bf16 %rs3, %rs1, %rs2, %rs2;726; SM80-NEXT: st.param.b16 [func_retval0], %rs3;727; SM80-NEXT: ret;728;729; SM80-FTZ-LABEL: test_fadd_imm_1(730; SM80-FTZ: {731; SM80-FTZ-NEXT: .reg .b16 %rs<3>;732; SM80-FTZ-NEXT: .reg .b32 %r<3>;733; SM80-FTZ-EMPTY:734; SM80-FTZ-NEXT: // %bb.0:735; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];736; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs1;737; SM80-FTZ-NEXT: add.rn.ftz.f32 %r2, %r1, 0f3F800000;738; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %r2;739; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;740; SM80-FTZ-NEXT: ret;741;742; SM90-FTZ-LABEL: test_fadd_imm_1(743; SM90-FTZ: {744; SM90-FTZ-NEXT: .reg .b16 %rs<4>;745; SM90-FTZ-EMPTY:746; SM90-FTZ-NEXT: // %bb.0:747; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];748; SM90-FTZ-NEXT: mov.b16 %rs2, 0x3F80;749; SM90-FTZ-NEXT: add.rn.bf16 %rs3, %rs1, %rs2;750; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;751; SM90-FTZ-NEXT: ret;752;753; SM90-LABEL: test_fadd_imm_1(754; SM90: {755; SM90-NEXT: .reg .b16 %rs<4>;756; SM90-EMPTY:757; SM90-NEXT: // %bb.0:758; SM90-NEXT: ld.param.b16 %rs1, [test_fadd_imm_1_param_0];759; SM90-NEXT: mov.b16 %rs2, 0x3F80;760; SM90-NEXT: add.rn.bf16 %rs3, %rs1, %rs2;761; SM90-NEXT: st.param.b16 [func_retval0], %rs3;762; SM90-NEXT: ret;763 %r = fadd bfloat %a, 1.0764 ret bfloat %r765}766 767define bfloat @test_select_cc_bf16_f64(double %a, double %b, bfloat %c, bfloat %d) #0 {768; CHECK-LABEL: test_select_cc_bf16_f64(769; CHECK: {770; CHECK-NEXT: .reg .pred %p<2>;771; CHECK-NEXT: .reg .b16 %rs<2>;772; CHECK-NEXT: .reg .b64 %rd<6>;773; CHECK-EMPTY:774; CHECK-NEXT: // %bb.0:775; CHECK-NEXT: ld.param.b64 %rd1, [test_select_cc_bf16_f64_param_0];776; CHECK-NEXT: ld.param.b64 %rd2, [test_select_cc_bf16_f64_param_1];777; CHECK-NEXT: setp.lt.f64 %p1, %rd1, %rd2;778; CHECK-NEXT: mov.b64 %rd3, test_select_cc_bf16_f64_param_3;779; CHECK-NEXT: mov.b64 %rd4, test_select_cc_bf16_f64_param_2;780; CHECK-NEXT: selp.b64 %rd5, %rd4, %rd3, %p1;781; CHECK-NEXT: ld.param.b16 %rs1, [%rd5];782; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;783; CHECK-NEXT: ret;784 %cc = fcmp olt double %a, %b785 %r = select i1 %cc, bfloat %c, bfloat %d786 ret bfloat %r787}788 789define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 {790; SM70-LABEL: test_extload_bf16x8(791; SM70: {792; SM70-NEXT: .reg .b16 %rs<9>;793; SM70-NEXT: .reg .b32 %r<21>;794; SM70-NEXT: .reg .b64 %rd<2>;795; SM70-EMPTY:796; SM70-NEXT: // %bb.0:797; SM70-NEXT: ld.param.b64 %rd1, [test_extload_bf16x8_param_0];798; SM70-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];799; SM70-NEXT: mov.b32 {%rs1, %rs2}, %r3;800; SM70-NEXT: mov.b32 {%rs3, %rs4}, %r4;801; SM70-NEXT: mov.b32 {%rs5, %rs6}, %r1;802; SM70-NEXT: mov.b32 {%rs7, %rs8}, %r2;803; SM70-NEXT: cvt.u32.u16 %r5, %rs8;804; SM70-NEXT: shl.b32 %r6, %r5, 16;805; SM70-NEXT: cvt.u32.u16 %r7, %rs7;806; SM70-NEXT: shl.b32 %r8, %r7, 16;807; SM70-NEXT: cvt.u32.u16 %r9, %rs6;808; SM70-NEXT: shl.b32 %r10, %r9, 16;809; SM70-NEXT: cvt.u32.u16 %r11, %rs5;810; SM70-NEXT: shl.b32 %r12, %r11, 16;811; SM70-NEXT: cvt.u32.u16 %r13, %rs4;812; SM70-NEXT: shl.b32 %r14, %r13, 16;813; SM70-NEXT: cvt.u32.u16 %r15, %rs3;814; SM70-NEXT: shl.b32 %r16, %r15, 16;815; SM70-NEXT: cvt.u32.u16 %r17, %rs2;816; SM70-NEXT: shl.b32 %r18, %r17, 16;817; SM70-NEXT: cvt.u32.u16 %r19, %rs1;818; SM70-NEXT: shl.b32 %r20, %r19, 16;819; SM70-NEXT: st.param.v4.b32 [func_retval0+16], {%r20, %r18, %r16, %r14};820; SM70-NEXT: st.param.v4.b32 [func_retval0], {%r12, %r10, %r8, %r6};821; SM70-NEXT: ret;822;823; SM80-LABEL: test_extload_bf16x8(824; SM80: {825; SM80-NEXT: .reg .b16 %rs<9>;826; SM80-NEXT: .reg .b32 %r<13>;827; SM80-NEXT: .reg .b64 %rd<2>;828; SM80-EMPTY:829; SM80-NEXT: // %bb.0:830; SM80-NEXT: ld.param.b64 %rd1, [test_extload_bf16x8_param_0];831; SM80-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];832; SM80-NEXT: mov.b32 {%rs1, %rs2}, %r3;833; SM80-NEXT: mov.b32 {%rs3, %rs4}, %r4;834; SM80-NEXT: mov.b32 {%rs5, %rs6}, %r1;835; SM80-NEXT: mov.b32 {%rs7, %rs8}, %r2;836; SM80-NEXT: cvt.f32.bf16 %r5, %rs8;837; SM80-NEXT: cvt.f32.bf16 %r6, %rs7;838; SM80-NEXT: cvt.f32.bf16 %r7, %rs6;839; SM80-NEXT: cvt.f32.bf16 %r8, %rs5;840; SM80-NEXT: cvt.f32.bf16 %r9, %rs4;841; SM80-NEXT: cvt.f32.bf16 %r10, %rs3;842; SM80-NEXT: cvt.f32.bf16 %r11, %rs2;843; SM80-NEXT: cvt.f32.bf16 %r12, %rs1;844; SM80-NEXT: st.param.v4.b32 [func_retval0+16], {%r12, %r11, %r10, %r9};845; SM80-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};846; SM80-NEXT: ret;847;848; SM80-FTZ-LABEL: test_extload_bf16x8(849; SM80-FTZ: {850; SM80-FTZ-NEXT: .reg .b16 %rs<9>;851; SM80-FTZ-NEXT: .reg .b32 %r<13>;852; SM80-FTZ-NEXT: .reg .b64 %rd<2>;853; SM80-FTZ-EMPTY:854; SM80-FTZ-NEXT: // %bb.0:855; SM80-FTZ-NEXT: ld.param.b64 %rd1, [test_extload_bf16x8_param_0];856; SM80-FTZ-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];857; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r3;858; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r4;859; SM80-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r1;860; SM80-FTZ-NEXT: mov.b32 {%rs7, %rs8}, %r2;861; SM80-FTZ-NEXT: cvt.f32.bf16 %r5, %rs8;862; SM80-FTZ-NEXT: cvt.f32.bf16 %r6, %rs7;863; SM80-FTZ-NEXT: cvt.f32.bf16 %r7, %rs6;864; SM80-FTZ-NEXT: cvt.f32.bf16 %r8, %rs5;865; SM80-FTZ-NEXT: cvt.f32.bf16 %r9, %rs4;866; SM80-FTZ-NEXT: cvt.f32.bf16 %r10, %rs3;867; SM80-FTZ-NEXT: cvt.f32.bf16 %r11, %rs2;868; SM80-FTZ-NEXT: cvt.f32.bf16 %r12, %rs1;869; SM80-FTZ-NEXT: st.param.v4.b32 [func_retval0+16], {%r12, %r11, %r10, %r9};870; SM80-FTZ-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};871; SM80-FTZ-NEXT: ret;872;873; SM90-FTZ-LABEL: test_extload_bf16x8(874; SM90-FTZ: {875; SM90-FTZ-NEXT: .reg .b16 %rs<9>;876; SM90-FTZ-NEXT: .reg .b32 %r<13>;877; SM90-FTZ-NEXT: .reg .b64 %rd<2>;878; SM90-FTZ-EMPTY:879; SM90-FTZ-NEXT: // %bb.0:880; SM90-FTZ-NEXT: ld.param.b64 %rd1, [test_extload_bf16x8_param_0];881; SM90-FTZ-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];882; SM90-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r3;883; SM90-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r4;884; SM90-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r1;885; SM90-FTZ-NEXT: mov.b32 {%rs7, %rs8}, %r2;886; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r5, %rs8;887; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r6, %rs7;888; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r7, %rs6;889; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r8, %rs5;890; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r9, %rs4;891; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r10, %rs3;892; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r11, %rs2;893; SM90-FTZ-NEXT: cvt.ftz.f32.bf16 %r12, %rs1;894; SM90-FTZ-NEXT: st.param.v4.b32 [func_retval0+16], {%r12, %r11, %r10, %r9};895; SM90-FTZ-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};896; SM90-FTZ-NEXT: ret;897;898; SM90-LABEL: test_extload_bf16x8(899; SM90: {900; SM90-NEXT: .reg .b16 %rs<9>;901; SM90-NEXT: .reg .b32 %r<13>;902; SM90-NEXT: .reg .b64 %rd<2>;903; SM90-EMPTY:904; SM90-NEXT: // %bb.0:905; SM90-NEXT: ld.param.b64 %rd1, [test_extload_bf16x8_param_0];906; SM90-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];907; SM90-NEXT: mov.b32 {%rs1, %rs2}, %r3;908; SM90-NEXT: mov.b32 {%rs3, %rs4}, %r4;909; SM90-NEXT: mov.b32 {%rs5, %rs6}, %r1;910; SM90-NEXT: mov.b32 {%rs7, %rs8}, %r2;911; SM90-NEXT: cvt.f32.bf16 %r5, %rs8;912; SM90-NEXT: cvt.f32.bf16 %r6, %rs7;913; SM90-NEXT: cvt.f32.bf16 %r7, %rs6;914; SM90-NEXT: cvt.f32.bf16 %r8, %rs5;915; SM90-NEXT: cvt.f32.bf16 %r9, %rs4;916; SM90-NEXT: cvt.f32.bf16 %r10, %rs3;917; SM90-NEXT: cvt.f32.bf16 %r11, %rs2;918; SM90-NEXT: cvt.f32.bf16 %r12, %rs1;919; SM90-NEXT: st.param.v4.b32 [func_retval0+16], {%r12, %r11, %r10, %r9};920; SM90-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};921; SM90-NEXT: ret;922 %load = load <8 x bfloat>, ptr addrspace(3) %arg, align 16923 %res = fpext <8 x bfloat> %load to <8 x float>924 ret <8 x float> %res925}926 927define i16 @test_fptosi_i16(bfloat %a) {928; SM70-LABEL: test_fptosi_i16(929; SM70: {930; SM70-NEXT: .reg .b16 %rs<2>;931; SM70-NEXT: .reg .b32 %r<4>;932; SM70-EMPTY:933; SM70-NEXT: // %bb.0:934; SM70-NEXT: ld.param.b16 %r1, [test_fptosi_i16_param_0];935; SM70-NEXT: shl.b32 %r2, %r1, 16;936; SM70-NEXT: cvt.rzi.s16.f32 %rs1, %r2;937; SM70-NEXT: cvt.u32.u16 %r3, %rs1;938; SM70-NEXT: st.param.b32 [func_retval0], %r3;939; SM70-NEXT: ret;940;941; SM80-LABEL: test_fptosi_i16(942; SM80: {943; SM80-NEXT: .reg .b16 %rs<3>;944; SM80-NEXT: .reg .b32 %r<3>;945; SM80-EMPTY:946; SM80-NEXT: // %bb.0:947; SM80-NEXT: ld.param.b16 %rs1, [test_fptosi_i16_param_0];948; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;949; SM80-NEXT: cvt.rzi.s16.f32 %rs2, %r1;950; SM80-NEXT: cvt.u32.u16 %r2, %rs2;951; SM80-NEXT: st.param.b32 [func_retval0], %r2;952; SM80-NEXT: ret;953;954; SM80-FTZ-LABEL: test_fptosi_i16(955; SM80-FTZ: {956; SM80-FTZ-NEXT: .reg .b16 %rs<3>;957; SM80-FTZ-NEXT: .reg .b32 %r<3>;958; SM80-FTZ-EMPTY:959; SM80-FTZ-NEXT: // %bb.0:960; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fptosi_i16_param_0];961; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs1;962; SM80-FTZ-NEXT: cvt.rzi.ftz.s16.f32 %rs2, %r1;963; SM80-FTZ-NEXT: cvt.u32.u16 %r2, %rs2;964; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r2;965; SM80-FTZ-NEXT: ret;966;967; SM90-FTZ-LABEL: test_fptosi_i16(968; SM90-FTZ: {969; SM90-FTZ-NEXT: .reg .b16 %rs<3>;970; SM90-FTZ-NEXT: .reg .b32 %r<2>;971; SM90-FTZ-EMPTY:972; SM90-FTZ-NEXT: // %bb.0:973; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_fptosi_i16_param_0];974; SM90-FTZ-NEXT: cvt.rzi.s16.bf16 %rs2, %rs1;975; SM90-FTZ-NEXT: cvt.u32.u16 %r1, %rs2;976; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r1;977; SM90-FTZ-NEXT: ret;978;979; SM90-LABEL: test_fptosi_i16(980; SM90: {981; SM90-NEXT: .reg .b16 %rs<3>;982; SM90-NEXT: .reg .b32 %r<2>;983; SM90-EMPTY:984; SM90-NEXT: // %bb.0:985; SM90-NEXT: ld.param.b16 %rs1, [test_fptosi_i16_param_0];986; SM90-NEXT: cvt.rzi.s16.bf16 %rs2, %rs1;987; SM90-NEXT: cvt.u32.u16 %r1, %rs2;988; SM90-NEXT: st.param.b32 [func_retval0], %r1;989; SM90-NEXT: ret;990 %r = fptosi bfloat %a to i16991 ret i16 %r992}993 994define i16 @test_fptoui_i16(bfloat %a) {995; SM70-LABEL: test_fptoui_i16(996; SM70: {997; SM70-NEXT: .reg .b16 %rs<2>;998; SM70-NEXT: .reg .b32 %r<4>;999; SM70-EMPTY:1000; SM70-NEXT: // %bb.0:1001; SM70-NEXT: ld.param.b16 %r1, [test_fptoui_i16_param_0];1002; SM70-NEXT: shl.b32 %r2, %r1, 16;1003; SM70-NEXT: cvt.rzi.u16.f32 %rs1, %r2;1004; SM70-NEXT: cvt.u32.u16 %r3, %rs1;1005; SM70-NEXT: st.param.b32 [func_retval0], %r3;1006; SM70-NEXT: ret;1007;1008; SM80-LABEL: test_fptoui_i16(1009; SM80: {1010; SM80-NEXT: .reg .b16 %rs<3>;1011; SM80-NEXT: .reg .b32 %r<3>;1012; SM80-EMPTY:1013; SM80-NEXT: // %bb.0:1014; SM80-NEXT: ld.param.b16 %rs1, [test_fptoui_i16_param_0];1015; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;1016; SM80-NEXT: cvt.rzi.u16.f32 %rs2, %r1;1017; SM80-NEXT: cvt.u32.u16 %r2, %rs2;1018; SM80-NEXT: st.param.b32 [func_retval0], %r2;1019; SM80-NEXT: ret;1020;1021; SM80-FTZ-LABEL: test_fptoui_i16(1022; SM80-FTZ: {1023; SM80-FTZ-NEXT: .reg .b16 %rs<3>;1024; SM80-FTZ-NEXT: .reg .b32 %r<3>;1025; SM80-FTZ-EMPTY:1026; SM80-FTZ-NEXT: // %bb.0:1027; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_fptoui_i16_param_0];1028; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs1;1029; SM80-FTZ-NEXT: cvt.rzi.ftz.u16.f32 %rs2, %r1;1030; SM80-FTZ-NEXT: cvt.u32.u16 %r2, %rs2;1031; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r2;1032; SM80-FTZ-NEXT: ret;1033;1034; SM90-FTZ-LABEL: test_fptoui_i16(1035; SM90-FTZ: {1036; SM90-FTZ-NEXT: .reg .b16 %rs<3>;1037; SM90-FTZ-NEXT: .reg .b32 %r<2>;1038; SM90-FTZ-EMPTY:1039; SM90-FTZ-NEXT: // %bb.0:1040; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_fptoui_i16_param_0];1041; SM90-FTZ-NEXT: cvt.rzi.u16.bf16 %rs2, %rs1;1042; SM90-FTZ-NEXT: cvt.u32.u16 %r1, %rs2;1043; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r1;1044; SM90-FTZ-NEXT: ret;1045;1046; SM90-LABEL: test_fptoui_i16(1047; SM90: {1048; SM90-NEXT: .reg .b16 %rs<3>;1049; SM90-NEXT: .reg .b32 %r<2>;1050; SM90-EMPTY:1051; SM90-NEXT: // %bb.0:1052; SM90-NEXT: ld.param.b16 %rs1, [test_fptoui_i16_param_0];1053; SM90-NEXT: cvt.rzi.u16.bf16 %rs2, %rs1;1054; SM90-NEXT: cvt.u32.u16 %r1, %rs2;1055; SM90-NEXT: st.param.b32 [func_retval0], %r1;1056; SM90-NEXT: ret;1057 %r = fptoui bfloat %a to i161058 ret i16 %r1059}1060 1061define bfloat @test_sitofp_i16(i16 %a) {1062; SM70-LABEL: test_sitofp_i16(1063; SM70: {1064; SM70-NEXT: .reg .pred %p<2>;1065; SM70-NEXT: .reg .b16 %rs<2>;1066; SM70-NEXT: .reg .b32 %r<8>;1067; SM70-EMPTY:1068; SM70-NEXT: // %bb.0:1069; SM70-NEXT: ld.param.b16 %rs1, [test_sitofp_i16_param_0];1070; SM70-NEXT: cvt.rn.f32.s16 %r1, %rs1;1071; SM70-NEXT: bfe.u32 %r2, %r1, 16, 1;1072; SM70-NEXT: add.s32 %r3, %r2, %r1;1073; SM70-NEXT: add.s32 %r4, %r3, 32767;1074; SM70-NEXT: setp.nan.f32 %p1, %r1, %r1;1075; SM70-NEXT: or.b32 %r5, %r1, 4194304;1076; SM70-NEXT: selp.b32 %r6, %r5, %r4, %p1;1077; SM70-NEXT: shr.u32 %r7, %r6, 16;1078; SM70-NEXT: st.param.b16 [func_retval0], %r7;1079; SM70-NEXT: ret;1080;1081; SM80-LABEL: test_sitofp_i16(1082; SM80: {1083; SM80-NEXT: .reg .b16 %rs<3>;1084; SM80-NEXT: .reg .b32 %r<2>;1085; SM80-EMPTY:1086; SM80-NEXT: // %bb.0:1087; SM80-NEXT: ld.param.b16 %rs1, [test_sitofp_i16_param_0];1088; SM80-NEXT: cvt.rn.f32.s16 %r1, %rs1;1089; SM80-NEXT: cvt.rn.bf16.f32 %rs2, %r1;1090; SM80-NEXT: st.param.b16 [func_retval0], %rs2;1091; SM80-NEXT: ret;1092;1093; SM80-FTZ-LABEL: test_sitofp_i16(1094; SM80-FTZ: {1095; SM80-FTZ-NEXT: .reg .b16 %rs<3>;1096; SM80-FTZ-NEXT: .reg .b32 %r<2>;1097; SM80-FTZ-EMPTY:1098; SM80-FTZ-NEXT: // %bb.0:1099; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_sitofp_i16_param_0];1100; SM80-FTZ-NEXT: cvt.rn.f32.s16 %r1, %rs1;1101; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %r1;1102; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1103; SM80-FTZ-NEXT: ret;1104;1105; SM90-FTZ-LABEL: test_sitofp_i16(1106; SM90-FTZ: {1107; SM90-FTZ-NEXT: .reg .b16 %rs<3>;1108; SM90-FTZ-EMPTY:1109; SM90-FTZ-NEXT: // %bb.0:1110; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_sitofp_i16_param_0];1111; SM90-FTZ-NEXT: cvt.rn.bf16.s16 %rs2, %rs1;1112; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1113; SM90-FTZ-NEXT: ret;1114;1115; SM90-LABEL: test_sitofp_i16(1116; SM90: {1117; SM90-NEXT: .reg .b16 %rs<3>;1118; SM90-EMPTY:1119; SM90-NEXT: // %bb.0:1120; SM90-NEXT: ld.param.b16 %rs1, [test_sitofp_i16_param_0];1121; SM90-NEXT: cvt.rn.bf16.s16 %rs2, %rs1;1122; SM90-NEXT: st.param.b16 [func_retval0], %rs2;1123; SM90-NEXT: ret;1124 %r = sitofp i16 %a to bfloat1125 ret bfloat %r1126}1127 1128define bfloat @test_uitofp_i8(i8 %a) {1129; SM70-LABEL: test_uitofp_i8(1130; SM70: {1131; SM70-NEXT: .reg .pred %p<2>;1132; SM70-NEXT: .reg .b16 %rs<2>;1133; SM70-NEXT: .reg .b32 %r<8>;1134; SM70-EMPTY:1135; SM70-NEXT: // %bb.0:1136; SM70-NEXT: ld.param.b8 %rs1, [test_uitofp_i8_param_0];1137; SM70-NEXT: cvt.rn.f32.u16 %r1, %rs1;1138; SM70-NEXT: bfe.u32 %r2, %r1, 16, 1;1139; SM70-NEXT: add.s32 %r3, %r2, %r1;1140; SM70-NEXT: add.s32 %r4, %r3, 32767;1141; SM70-NEXT: setp.nan.f32 %p1, %r1, %r1;1142; SM70-NEXT: or.b32 %r5, %r1, 4194304;1143; SM70-NEXT: selp.b32 %r6, %r5, %r4, %p1;1144; SM70-NEXT: shr.u32 %r7, %r6, 16;1145; SM70-NEXT: st.param.b16 [func_retval0], %r7;1146; SM70-NEXT: ret;1147;1148; SM80-LABEL: test_uitofp_i8(1149; SM80: {1150; SM80-NEXT: .reg .b16 %rs<3>;1151; SM80-NEXT: .reg .b32 %r<2>;1152; SM80-EMPTY:1153; SM80-NEXT: // %bb.0:1154; SM80-NEXT: ld.param.b8 %rs1, [test_uitofp_i8_param_0];1155; SM80-NEXT: cvt.rn.f32.u16 %r1, %rs1;1156; SM80-NEXT: cvt.rn.bf16.f32 %rs2, %r1;1157; SM80-NEXT: st.param.b16 [func_retval0], %rs2;1158; SM80-NEXT: ret;1159;1160; SM80-FTZ-LABEL: test_uitofp_i8(1161; SM80-FTZ: {1162; SM80-FTZ-NEXT: .reg .b16 %rs<3>;1163; SM80-FTZ-NEXT: .reg .b32 %r<2>;1164; SM80-FTZ-EMPTY:1165; SM80-FTZ-NEXT: // %bb.0:1166; SM80-FTZ-NEXT: ld.param.b8 %rs1, [test_uitofp_i8_param_0];1167; SM80-FTZ-NEXT: cvt.rn.f32.u16 %r1, %rs1;1168; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %r1;1169; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1170; SM80-FTZ-NEXT: ret;1171;1172; SM90-FTZ-LABEL: test_uitofp_i8(1173; SM90-FTZ: {1174; SM90-FTZ-NEXT: .reg .b16 %rs<3>;1175; SM90-FTZ-EMPTY:1176; SM90-FTZ-NEXT: // %bb.0:1177; SM90-FTZ-NEXT: ld.param.b8 %rs1, [test_uitofp_i8_param_0];1178; SM90-FTZ-NEXT: cvt.rn.bf16.u16 %rs2, %rs1;1179; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1180; SM90-FTZ-NEXT: ret;1181;1182; SM90-LABEL: test_uitofp_i8(1183; SM90: {1184; SM90-NEXT: .reg .b16 %rs<3>;1185; SM90-EMPTY:1186; SM90-NEXT: // %bb.0:1187; SM90-NEXT: ld.param.b8 %rs1, [test_uitofp_i8_param_0];1188; SM90-NEXT: cvt.rn.bf16.u16 %rs2, %rs1;1189; SM90-NEXT: st.param.b16 [func_retval0], %rs2;1190; SM90-NEXT: ret;1191 %r = uitofp i8 %a to bfloat1192 ret bfloat %r1193}1194 1195define bfloat @test_uitofp_i1(i1 %a) {1196; SM70-LABEL: test_uitofp_i1(1197; SM70: {1198; SM70-NEXT: .reg .pred %p<3>;1199; SM70-NEXT: .reg .b16 %rs<3>;1200; SM70-NEXT: .reg .b32 %r<9>;1201; SM70-EMPTY:1202; SM70-NEXT: // %bb.0:1203; SM70-NEXT: ld.param.b8 %rs1, [test_uitofp_i1_param_0];1204; SM70-NEXT: and.b16 %rs2, %rs1, 1;1205; SM70-NEXT: setp.ne.b16 %p1, %rs2, 0;1206; SM70-NEXT: selp.b32 %r1, 1, 0, %p1;1207; SM70-NEXT: cvt.rn.f32.u32 %r2, %r1;1208; SM70-NEXT: bfe.u32 %r3, %r2, 16, 1;1209; SM70-NEXT: add.s32 %r4, %r3, %r2;1210; SM70-NEXT: add.s32 %r5, %r4, 32767;1211; SM70-NEXT: setp.nan.f32 %p2, %r2, %r2;1212; SM70-NEXT: or.b32 %r6, %r2, 4194304;1213; SM70-NEXT: selp.b32 %r7, %r6, %r5, %p2;1214; SM70-NEXT: shr.u32 %r8, %r7, 16;1215; SM70-NEXT: st.param.b16 [func_retval0], %r8;1216; SM70-NEXT: ret;1217;1218; SM80-LABEL: test_uitofp_i1(1219; SM80: {1220; SM80-NEXT: .reg .pred %p<2>;1221; SM80-NEXT: .reg .b16 %rs<4>;1222; SM80-NEXT: .reg .b32 %r<3>;1223; SM80-EMPTY:1224; SM80-NEXT: // %bb.0:1225; SM80-NEXT: ld.param.b8 %rs1, [test_uitofp_i1_param_0];1226; SM80-NEXT: and.b16 %rs2, %rs1, 1;1227; SM80-NEXT: setp.ne.b16 %p1, %rs2, 0;1228; SM80-NEXT: selp.b32 %r1, 1, 0, %p1;1229; SM80-NEXT: cvt.rn.f32.u32 %r2, %r1;1230; SM80-NEXT: cvt.rn.bf16.f32 %rs3, %r2;1231; SM80-NEXT: st.param.b16 [func_retval0], %rs3;1232; SM80-NEXT: ret;1233;1234; SM80-FTZ-LABEL: test_uitofp_i1(1235; SM80-FTZ: {1236; SM80-FTZ-NEXT: .reg .pred %p<2>;1237; SM80-FTZ-NEXT: .reg .b16 %rs<4>;1238; SM80-FTZ-NEXT: .reg .b32 %r<3>;1239; SM80-FTZ-EMPTY:1240; SM80-FTZ-NEXT: // %bb.0:1241; SM80-FTZ-NEXT: ld.param.b8 %rs1, [test_uitofp_i1_param_0];1242; SM80-FTZ-NEXT: and.b16 %rs2, %rs1, 1;1243; SM80-FTZ-NEXT: setp.ne.b16 %p1, %rs2, 0;1244; SM80-FTZ-NEXT: selp.b32 %r1, 1, 0, %p1;1245; SM80-FTZ-NEXT: cvt.rn.f32.u32 %r2, %r1;1246; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r2;1247; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;1248; SM80-FTZ-NEXT: ret;1249;1250; SM90-FTZ-LABEL: test_uitofp_i1(1251; SM90-FTZ: {1252; SM90-FTZ-NEXT: .reg .pred %p<2>;1253; SM90-FTZ-NEXT: .reg .b16 %rs<4>;1254; SM90-FTZ-NEXT: .reg .b32 %r<2>;1255; SM90-FTZ-EMPTY:1256; SM90-FTZ-NEXT: // %bb.0:1257; SM90-FTZ-NEXT: ld.param.b8 %rs1, [test_uitofp_i1_param_0];1258; SM90-FTZ-NEXT: and.b16 %rs2, %rs1, 1;1259; SM90-FTZ-NEXT: setp.ne.b16 %p1, %rs2, 0;1260; SM90-FTZ-NEXT: selp.b32 %r1, 1, 0, %p1;1261; SM90-FTZ-NEXT: cvt.rn.bf16.u32 %rs3, %r1;1262; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;1263; SM90-FTZ-NEXT: ret;1264;1265; SM90-LABEL: test_uitofp_i1(1266; SM90: {1267; SM90-NEXT: .reg .pred %p<2>;1268; SM90-NEXT: .reg .b16 %rs<4>;1269; SM90-NEXT: .reg .b32 %r<2>;1270; SM90-EMPTY:1271; SM90-NEXT: // %bb.0:1272; SM90-NEXT: ld.param.b8 %rs1, [test_uitofp_i1_param_0];1273; SM90-NEXT: and.b16 %rs2, %rs1, 1;1274; SM90-NEXT: setp.ne.b16 %p1, %rs2, 0;1275; SM90-NEXT: selp.b32 %r1, 1, 0, %p1;1276; SM90-NEXT: cvt.rn.bf16.u32 %rs3, %r1;1277; SM90-NEXT: st.param.b16 [func_retval0], %rs3;1278; SM90-NEXT: ret;1279 %r = uitofp i1 %a to bfloat1280 ret bfloat %r1281}1282 1283define bfloat @test_uitofp_i16(i16 %a) {1284; SM70-LABEL: test_uitofp_i16(1285; SM70: {1286; SM70-NEXT: .reg .pred %p<2>;1287; SM70-NEXT: .reg .b16 %rs<2>;1288; SM70-NEXT: .reg .b32 %r<8>;1289; SM70-EMPTY:1290; SM70-NEXT: // %bb.0:1291; SM70-NEXT: ld.param.b16 %rs1, [test_uitofp_i16_param_0];1292; SM70-NEXT: cvt.rn.f32.u16 %r1, %rs1;1293; SM70-NEXT: bfe.u32 %r2, %r1, 16, 1;1294; SM70-NEXT: add.s32 %r3, %r2, %r1;1295; SM70-NEXT: add.s32 %r4, %r3, 32767;1296; SM70-NEXT: setp.nan.f32 %p1, %r1, %r1;1297; SM70-NEXT: or.b32 %r5, %r1, 4194304;1298; SM70-NEXT: selp.b32 %r6, %r5, %r4, %p1;1299; SM70-NEXT: shr.u32 %r7, %r6, 16;1300; SM70-NEXT: st.param.b16 [func_retval0], %r7;1301; SM70-NEXT: ret;1302;1303; SM80-LABEL: test_uitofp_i16(1304; SM80: {1305; SM80-NEXT: .reg .b16 %rs<3>;1306; SM80-NEXT: .reg .b32 %r<2>;1307; SM80-EMPTY:1308; SM80-NEXT: // %bb.0:1309; SM80-NEXT: ld.param.b16 %rs1, [test_uitofp_i16_param_0];1310; SM80-NEXT: cvt.rn.f32.u16 %r1, %rs1;1311; SM80-NEXT: cvt.rn.bf16.f32 %rs2, %r1;1312; SM80-NEXT: st.param.b16 [func_retval0], %rs2;1313; SM80-NEXT: ret;1314;1315; SM80-FTZ-LABEL: test_uitofp_i16(1316; SM80-FTZ: {1317; SM80-FTZ-NEXT: .reg .b16 %rs<3>;1318; SM80-FTZ-NEXT: .reg .b32 %r<2>;1319; SM80-FTZ-EMPTY:1320; SM80-FTZ-NEXT: // %bb.0:1321; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_uitofp_i16_param_0];1322; SM80-FTZ-NEXT: cvt.rn.f32.u16 %r1, %rs1;1323; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %r1;1324; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1325; SM80-FTZ-NEXT: ret;1326;1327; SM90-FTZ-LABEL: test_uitofp_i16(1328; SM90-FTZ: {1329; SM90-FTZ-NEXT: .reg .b16 %rs<3>;1330; SM90-FTZ-EMPTY:1331; SM90-FTZ-NEXT: // %bb.0:1332; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_uitofp_i16_param_0];1333; SM90-FTZ-NEXT: cvt.rn.bf16.u16 %rs2, %rs1;1334; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1335; SM90-FTZ-NEXT: ret;1336;1337; SM90-LABEL: test_uitofp_i16(1338; SM90: {1339; SM90-NEXT: .reg .b16 %rs<3>;1340; SM90-EMPTY:1341; SM90-NEXT: // %bb.0:1342; SM90-NEXT: ld.param.b16 %rs1, [test_uitofp_i16_param_0];1343; SM90-NEXT: cvt.rn.bf16.u16 %rs2, %rs1;1344; SM90-NEXT: st.param.b16 [func_retval0], %rs2;1345; SM90-NEXT: ret;1346 %r = uitofp i16 %a to bfloat1347 ret bfloat %r1348}1349 1350define bfloat @test_uitofp_i32(i32 %a) {1351; SM70-LABEL: test_uitofp_i32(1352; SM70: {1353; SM70-NEXT: .reg .pred %p<2>;1354; SM70-NEXT: .reg .b32 %r<9>;1355; SM70-EMPTY:1356; SM70-NEXT: // %bb.0:1357; SM70-NEXT: ld.param.b32 %r1, [test_uitofp_i32_param_0];1358; SM70-NEXT: cvt.rn.f32.u32 %r2, %r1;1359; SM70-NEXT: bfe.u32 %r3, %r2, 16, 1;1360; SM70-NEXT: add.s32 %r4, %r3, %r2;1361; SM70-NEXT: add.s32 %r5, %r4, 32767;1362; SM70-NEXT: setp.nan.f32 %p1, %r2, %r2;1363; SM70-NEXT: or.b32 %r6, %r2, 4194304;1364; SM70-NEXT: selp.b32 %r7, %r6, %r5, %p1;1365; SM70-NEXT: shr.u32 %r8, %r7, 16;1366; SM70-NEXT: st.param.b16 [func_retval0], %r8;1367; SM70-NEXT: ret;1368;1369; SM80-LABEL: test_uitofp_i32(1370; SM80: {1371; SM80-NEXT: .reg .b16 %rs<2>;1372; SM80-NEXT: .reg .b32 %r<3>;1373; SM80-EMPTY:1374; SM80-NEXT: // %bb.0:1375; SM80-NEXT: ld.param.b32 %r1, [test_uitofp_i32_param_0];1376; SM80-NEXT: cvt.rn.f32.u32 %r2, %r1;1377; SM80-NEXT: cvt.rn.bf16.f32 %rs1, %r2;1378; SM80-NEXT: st.param.b16 [func_retval0], %rs1;1379; SM80-NEXT: ret;1380;1381; SM80-FTZ-LABEL: test_uitofp_i32(1382; SM80-FTZ: {1383; SM80-FTZ-NEXT: .reg .b16 %rs<2>;1384; SM80-FTZ-NEXT: .reg .b32 %r<3>;1385; SM80-FTZ-EMPTY:1386; SM80-FTZ-NEXT: // %bb.0:1387; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_uitofp_i32_param_0];1388; SM80-FTZ-NEXT: cvt.rn.f32.u32 %r2, %r1;1389; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs1, %r2;1390; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs1;1391; SM80-FTZ-NEXT: ret;1392;1393; SM90-FTZ-LABEL: test_uitofp_i32(1394; SM90-FTZ: {1395; SM90-FTZ-NEXT: .reg .b16 %rs<2>;1396; SM90-FTZ-NEXT: .reg .b32 %r<2>;1397; SM90-FTZ-EMPTY:1398; SM90-FTZ-NEXT: // %bb.0:1399; SM90-FTZ-NEXT: ld.param.b32 %r1, [test_uitofp_i32_param_0];1400; SM90-FTZ-NEXT: cvt.rn.bf16.u32 %rs1, %r1;1401; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs1;1402; SM90-FTZ-NEXT: ret;1403;1404; SM90-LABEL: test_uitofp_i32(1405; SM90: {1406; SM90-NEXT: .reg .b16 %rs<2>;1407; SM90-NEXT: .reg .b32 %r<2>;1408; SM90-EMPTY:1409; SM90-NEXT: // %bb.0:1410; SM90-NEXT: ld.param.b32 %r1, [test_uitofp_i32_param_0];1411; SM90-NEXT: cvt.rn.bf16.u32 %rs1, %r1;1412; SM90-NEXT: st.param.b16 [func_retval0], %rs1;1413; SM90-NEXT: ret;1414 %r = uitofp i32 %a to bfloat1415 ret bfloat %r1416}1417 1418define bfloat @test_uitofp_i64(i64 %a) {1419; SM70-LABEL: test_uitofp_i64(1420; SM70: {1421; SM70-NEXT: .reg .pred %p<2>;1422; SM70-NEXT: .reg .b32 %r<8>;1423; SM70-NEXT: .reg .b64 %rd<2>;1424; SM70-EMPTY:1425; SM70-NEXT: // %bb.0:1426; SM70-NEXT: ld.param.b64 %rd1, [test_uitofp_i64_param_0];1427; SM70-NEXT: cvt.rn.f32.u64 %r1, %rd1;1428; SM70-NEXT: bfe.u32 %r2, %r1, 16, 1;1429; SM70-NEXT: add.s32 %r3, %r2, %r1;1430; SM70-NEXT: add.s32 %r4, %r3, 32767;1431; SM70-NEXT: setp.nan.f32 %p1, %r1, %r1;1432; SM70-NEXT: or.b32 %r5, %r1, 4194304;1433; SM70-NEXT: selp.b32 %r6, %r5, %r4, %p1;1434; SM70-NEXT: shr.u32 %r7, %r6, 16;1435; SM70-NEXT: st.param.b16 [func_retval0], %r7;1436; SM70-NEXT: ret;1437;1438; SM80-LABEL: test_uitofp_i64(1439; SM80: {1440; SM80-NEXT: .reg .b16 %rs<2>;1441; SM80-NEXT: .reg .b32 %r<2>;1442; SM80-NEXT: .reg .b64 %rd<2>;1443; SM80-EMPTY:1444; SM80-NEXT: // %bb.0:1445; SM80-NEXT: ld.param.b64 %rd1, [test_uitofp_i64_param_0];1446; SM80-NEXT: cvt.rn.f32.u64 %r1, %rd1;1447; SM80-NEXT: cvt.rn.bf16.f32 %rs1, %r1;1448; SM80-NEXT: st.param.b16 [func_retval0], %rs1;1449; SM80-NEXT: ret;1450;1451; SM80-FTZ-LABEL: test_uitofp_i64(1452; SM80-FTZ: {1453; SM80-FTZ-NEXT: .reg .b16 %rs<2>;1454; SM80-FTZ-NEXT: .reg .b32 %r<2>;1455; SM80-FTZ-NEXT: .reg .b64 %rd<2>;1456; SM80-FTZ-EMPTY:1457; SM80-FTZ-NEXT: // %bb.0:1458; SM80-FTZ-NEXT: ld.param.b64 %rd1, [test_uitofp_i64_param_0];1459; SM80-FTZ-NEXT: cvt.rn.f32.u64 %r1, %rd1;1460; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs1, %r1;1461; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs1;1462; SM80-FTZ-NEXT: ret;1463;1464; SM90-FTZ-LABEL: test_uitofp_i64(1465; SM90-FTZ: {1466; SM90-FTZ-NEXT: .reg .b16 %rs<2>;1467; SM90-FTZ-NEXT: .reg .b64 %rd<2>;1468; SM90-FTZ-EMPTY:1469; SM90-FTZ-NEXT: // %bb.0:1470; SM90-FTZ-NEXT: ld.param.b64 %rd1, [test_uitofp_i64_param_0];1471; SM90-FTZ-NEXT: cvt.rn.bf16.u64 %rs1, %rd1;1472; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs1;1473; SM90-FTZ-NEXT: ret;1474;1475; SM90-LABEL: test_uitofp_i64(1476; SM90: {1477; SM90-NEXT: .reg .b16 %rs<2>;1478; SM90-NEXT: .reg .b64 %rd<2>;1479; SM90-EMPTY:1480; SM90-NEXT: // %bb.0:1481; SM90-NEXT: ld.param.b64 %rd1, [test_uitofp_i64_param_0];1482; SM90-NEXT: cvt.rn.bf16.u64 %rs1, %rd1;1483; SM90-NEXT: st.param.b16 [func_retval0], %rs1;1484; SM90-NEXT: ret;1485 %r = uitofp i64 %a to bfloat1486 ret bfloat %r1487}1488 1489define bfloat @test_roundeven(bfloat %a) {1490; SM70-LABEL: test_roundeven(1491; SM70: {1492; SM70-NEXT: .reg .pred %p<2>;1493; SM70-NEXT: .reg .b32 %r<10>;1494; SM70-EMPTY:1495; SM70-NEXT: // %bb.0:1496; SM70-NEXT: ld.param.b16 %r1, [test_roundeven_param_0];1497; SM70-NEXT: shl.b32 %r2, %r1, 16;1498; SM70-NEXT: cvt.rni.f32.f32 %r3, %r2;1499; SM70-NEXT: bfe.u32 %r4, %r3, 16, 1;1500; SM70-NEXT: add.s32 %r5, %r4, %r3;1501; SM70-NEXT: add.s32 %r6, %r5, 32767;1502; SM70-NEXT: setp.nan.f32 %p1, %r3, %r3;1503; SM70-NEXT: or.b32 %r7, %r3, 4194304;1504; SM70-NEXT: selp.b32 %r8, %r7, %r6, %p1;1505; SM70-NEXT: shr.u32 %r9, %r8, 16;1506; SM70-NEXT: st.param.b16 [func_retval0], %r9;1507; SM70-NEXT: ret;1508;1509; SM80-LABEL: test_roundeven(1510; SM80: {1511; SM80-NEXT: .reg .b16 %rs<3>;1512; SM80-NEXT: .reg .b32 %r<3>;1513; SM80-EMPTY:1514; SM80-NEXT: // %bb.0:1515; SM80-NEXT: ld.param.b16 %rs1, [test_roundeven_param_0];1516; SM80-NEXT: cvt.f32.bf16 %r1, %rs1;1517; SM80-NEXT: cvt.rni.f32.f32 %r2, %r1;1518; SM80-NEXT: cvt.rn.bf16.f32 %rs2, %r2;1519; SM80-NEXT: st.param.b16 [func_retval0], %rs2;1520; SM80-NEXT: ret;1521;1522; SM80-FTZ-LABEL: test_roundeven(1523; SM80-FTZ: {1524; SM80-FTZ-NEXT: .reg .b16 %rs<3>;1525; SM80-FTZ-NEXT: .reg .b32 %r<3>;1526; SM80-FTZ-EMPTY:1527; SM80-FTZ-NEXT: // %bb.0:1528; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_roundeven_param_0];1529; SM80-FTZ-NEXT: cvt.f32.bf16 %r1, %rs1;1530; SM80-FTZ-NEXT: cvt.rni.ftz.f32.f32 %r2, %r1;1531; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs2, %r2;1532; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1533; SM80-FTZ-NEXT: ret;1534;1535; SM90-FTZ-LABEL: test_roundeven(1536; SM90-FTZ: {1537; SM90-FTZ-NEXT: .reg .b16 %rs<3>;1538; SM90-FTZ-EMPTY:1539; SM90-FTZ-NEXT: // %bb.0:1540; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_roundeven_param_0];1541; SM90-FTZ-NEXT: cvt.rni.bf16.bf16 %rs2, %rs1;1542; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs2;1543; SM90-FTZ-NEXT: ret;1544;1545; SM90-LABEL: test_roundeven(1546; SM90: {1547; SM90-NEXT: .reg .b16 %rs<3>;1548; SM90-EMPTY:1549; SM90-NEXT: // %bb.0:1550; SM90-NEXT: ld.param.b16 %rs1, [test_roundeven_param_0];1551; SM90-NEXT: cvt.rni.bf16.bf16 %rs2, %rs1;1552; SM90-NEXT: st.param.b16 [func_retval0], %rs2;1553; SM90-NEXT: ret;1554 %r = call bfloat @llvm.roundeven.bf16(bfloat %a)1555 ret bfloat %r1556}1557 1558define bfloat @test_maximum(bfloat %a, bfloat %b) {1559; SM70-LABEL: test_maximum(1560; SM70: {1561; SM70-NEXT: .reg .pred %p<6>;1562; SM70-NEXT: .reg .b16 %rs<8>;1563; SM70-NEXT: .reg .b32 %r<7>;1564; SM70-EMPTY:1565; SM70-NEXT: // %bb.0:1566; SM70-NEXT: ld.param.b16 %rs1, [test_maximum_param_0];1567; SM70-NEXT: ld.param.b16 %rs2, [test_maximum_param_1];1568; SM70-NEXT: cvt.u32.u16 %r1, %rs2;1569; SM70-NEXT: shl.b32 %r2, %r1, 16;1570; SM70-NEXT: cvt.u32.u16 %r3, %rs1;1571; SM70-NEXT: shl.b32 %r4, %r3, 16;1572; SM70-NEXT: setp.gt.f32 %p1, %r4, %r2;1573; SM70-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;1574; SM70-NEXT: setp.nan.f32 %p2, %r4, %r2;1575; SM70-NEXT: selp.b16 %rs4, 0x7FC0, %rs3, %p2;1576; SM70-NEXT: setp.eq.b16 %p3, %rs1, 0;1577; SM70-NEXT: selp.b16 %rs5, %rs1, %rs4, %p3;1578; SM70-NEXT: setp.eq.b16 %p4, %rs2, 0;1579; SM70-NEXT: selp.b16 %rs6, %rs2, %rs5, %p4;1580; SM70-NEXT: cvt.u32.u16 %r5, %rs4;1581; SM70-NEXT: shl.b32 %r6, %r5, 16;1582; SM70-NEXT: setp.eq.f32 %p5, %r6, 0f00000000;1583; SM70-NEXT: selp.b16 %rs7, %rs6, %rs4, %p5;1584; SM70-NEXT: st.param.b16 [func_retval0], %rs7;1585; SM70-NEXT: ret;1586;1587; SM80-LABEL: test_maximum(1588; SM80: {1589; SM80-NEXT: .reg .b16 %rs<4>;1590; SM80-EMPTY:1591; SM80-NEXT: // %bb.0:1592; SM80-NEXT: ld.param.b16 %rs1, [test_maximum_param_0];1593; SM80-NEXT: ld.param.b16 %rs2, [test_maximum_param_1];1594; SM80-NEXT: max.NaN.bf16 %rs3, %rs1, %rs2;1595; SM80-NEXT: st.param.b16 [func_retval0], %rs3;1596; SM80-NEXT: ret;1597;1598; SM80-FTZ-LABEL: test_maximum(1599; SM80-FTZ: {1600; SM80-FTZ-NEXT: .reg .b16 %rs<4>;1601; SM80-FTZ-EMPTY:1602; SM80-FTZ-NEXT: // %bb.0:1603; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_maximum_param_0];1604; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_maximum_param_1];1605; SM80-FTZ-NEXT: max.NaN.bf16 %rs3, %rs1, %rs2;1606; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;1607; SM80-FTZ-NEXT: ret;1608;1609; SM90-FTZ-LABEL: test_maximum(1610; SM90-FTZ: {1611; SM90-FTZ-NEXT: .reg .b16 %rs<4>;1612; SM90-FTZ-EMPTY:1613; SM90-FTZ-NEXT: // %bb.0:1614; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_maximum_param_0];1615; SM90-FTZ-NEXT: ld.param.b16 %rs2, [test_maximum_param_1];1616; SM90-FTZ-NEXT: max.NaN.bf16 %rs3, %rs1, %rs2;1617; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;1618; SM90-FTZ-NEXT: ret;1619;1620; SM90-LABEL: test_maximum(1621; SM90: {1622; SM90-NEXT: .reg .b16 %rs<4>;1623; SM90-EMPTY:1624; SM90-NEXT: // %bb.0:1625; SM90-NEXT: ld.param.b16 %rs1, [test_maximum_param_0];1626; SM90-NEXT: ld.param.b16 %rs2, [test_maximum_param_1];1627; SM90-NEXT: max.NaN.bf16 %rs3, %rs1, %rs2;1628; SM90-NEXT: st.param.b16 [func_retval0], %rs3;1629; SM90-NEXT: ret;1630 %r = call bfloat @llvm.maximum.bf16(bfloat %a, bfloat %b)1631 ret bfloat %r1632}1633 1634define bfloat @test_maxnum(bfloat %a, bfloat %b) {1635; SM70-LABEL: test_maxnum(1636; SM70: {1637; SM70-NEXT: .reg .pred %p<2>;1638; SM70-NEXT: .reg .b32 %r<12>;1639; SM70-EMPTY:1640; SM70-NEXT: // %bb.0:1641; SM70-NEXT: ld.param.b16 %r1, [test_maxnum_param_1];1642; SM70-NEXT: shl.b32 %r2, %r1, 16;1643; SM70-NEXT: ld.param.b16 %r3, [test_maxnum_param_0];1644; SM70-NEXT: shl.b32 %r4, %r3, 16;1645; SM70-NEXT: max.f32 %r5, %r4, %r2;1646; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;1647; SM70-NEXT: add.s32 %r7, %r6, %r5;1648; SM70-NEXT: add.s32 %r8, %r7, 32767;1649; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;1650; SM70-NEXT: or.b32 %r9, %r5, 4194304;1651; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;1652; SM70-NEXT: shr.u32 %r11, %r10, 16;1653; SM70-NEXT: st.param.b16 [func_retval0], %r11;1654; SM70-NEXT: ret;1655;1656; SM80-LABEL: test_maxnum(1657; SM80: {1658; SM80-NEXT: .reg .b16 %rs<4>;1659; SM80-EMPTY:1660; SM80-NEXT: // %bb.0:1661; SM80-NEXT: ld.param.b16 %rs1, [test_maxnum_param_0];1662; SM80-NEXT: ld.param.b16 %rs2, [test_maxnum_param_1];1663; SM80-NEXT: max.bf16 %rs3, %rs1, %rs2;1664; SM80-NEXT: st.param.b16 [func_retval0], %rs3;1665; SM80-NEXT: ret;1666;1667; SM80-FTZ-LABEL: test_maxnum(1668; SM80-FTZ: {1669; SM80-FTZ-NEXT: .reg .b16 %rs<4>;1670; SM80-FTZ-EMPTY:1671; SM80-FTZ-NEXT: // %bb.0:1672; SM80-FTZ-NEXT: ld.param.b16 %rs1, [test_maxnum_param_0];1673; SM80-FTZ-NEXT: ld.param.b16 %rs2, [test_maxnum_param_1];1674; SM80-FTZ-NEXT: max.bf16 %rs3, %rs1, %rs2;1675; SM80-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;1676; SM80-FTZ-NEXT: ret;1677;1678; SM90-FTZ-LABEL: test_maxnum(1679; SM90-FTZ: {1680; SM90-FTZ-NEXT: .reg .b16 %rs<4>;1681; SM90-FTZ-EMPTY:1682; SM90-FTZ-NEXT: // %bb.0:1683; SM90-FTZ-NEXT: ld.param.b16 %rs1, [test_maxnum_param_0];1684; SM90-FTZ-NEXT: ld.param.b16 %rs2, [test_maxnum_param_1];1685; SM90-FTZ-NEXT: max.bf16 %rs3, %rs1, %rs2;1686; SM90-FTZ-NEXT: st.param.b16 [func_retval0], %rs3;1687; SM90-FTZ-NEXT: ret;1688;1689; SM90-LABEL: test_maxnum(1690; SM90: {1691; SM90-NEXT: .reg .b16 %rs<4>;1692; SM90-EMPTY:1693; SM90-NEXT: // %bb.0:1694; SM90-NEXT: ld.param.b16 %rs1, [test_maxnum_param_0];1695; SM90-NEXT: ld.param.b16 %rs2, [test_maxnum_param_1];1696; SM90-NEXT: max.bf16 %rs3, %rs1, %rs2;1697; SM90-NEXT: st.param.b16 [func_retval0], %rs3;1698; SM90-NEXT: ret;1699 %r = call bfloat @llvm.maxnum.bf16(bfloat %a, bfloat %b)1700 ret bfloat %r1701}1702 1703define <2 x bfloat> @test_maximum_v2(<2 x bfloat> %a, <2 x bfloat> %b) {1704; SM70-LABEL: test_maximum_v2(1705; SM70: {1706; SM70-NEXT: .reg .pred %p<11>;1707; SM70-NEXT: .reg .b16 %rs<15>;1708; SM70-NEXT: .reg .b32 %r<13>;1709; SM70-EMPTY:1710; SM70-NEXT: // %bb.0:1711; SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_maximum_v2_param_0];1712; SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_maximum_v2_param_1];1713; SM70-NEXT: cvt.u32.u16 %r1, %rs4;1714; SM70-NEXT: shl.b32 %r2, %r1, 16;1715; SM70-NEXT: cvt.u32.u16 %r3, %rs2;1716; SM70-NEXT: shl.b32 %r4, %r3, 16;1717; SM70-NEXT: setp.gt.f32 %p1, %r4, %r2;1718; SM70-NEXT: selp.b16 %rs5, %rs2, %rs4, %p1;1719; SM70-NEXT: setp.nan.f32 %p2, %r4, %r2;1720; SM70-NEXT: selp.b16 %rs6, 0x7FC0, %rs5, %p2;1721; SM70-NEXT: setp.eq.b16 %p3, %rs2, 0;1722; SM70-NEXT: selp.b16 %rs7, %rs2, %rs6, %p3;1723; SM70-NEXT: setp.eq.b16 %p4, %rs4, 0;1724; SM70-NEXT: selp.b16 %rs8, %rs4, %rs7, %p4;1725; SM70-NEXT: cvt.u32.u16 %r5, %rs6;1726; SM70-NEXT: shl.b32 %r6, %r5, 16;1727; SM70-NEXT: setp.eq.f32 %p5, %r6, 0f00000000;1728; SM70-NEXT: selp.b16 %rs9, %rs8, %rs6, %p5;1729; SM70-NEXT: cvt.u32.u16 %r7, %rs3;1730; SM70-NEXT: shl.b32 %r8, %r7, 16;1731; SM70-NEXT: cvt.u32.u16 %r9, %rs1;1732; SM70-NEXT: shl.b32 %r10, %r9, 16;1733; SM70-NEXT: setp.gt.f32 %p6, %r10, %r8;1734; SM70-NEXT: selp.b16 %rs10, %rs1, %rs3, %p6;1735; SM70-NEXT: setp.nan.f32 %p7, %r10, %r8;1736; SM70-NEXT: selp.b16 %rs11, 0x7FC0, %rs10, %p7;1737; SM70-NEXT: setp.eq.b16 %p8, %rs1, 0;1738; SM70-NEXT: selp.b16 %rs12, %rs1, %rs11, %p8;1739; SM70-NEXT: setp.eq.b16 %p9, %rs3, 0;1740; SM70-NEXT: selp.b16 %rs13, %rs3, %rs12, %p9;1741; SM70-NEXT: cvt.u32.u16 %r11, %rs11;1742; SM70-NEXT: shl.b32 %r12, %r11, 16;1743; SM70-NEXT: setp.eq.f32 %p10, %r12, 0f00000000;1744; SM70-NEXT: selp.b16 %rs14, %rs13, %rs11, %p10;1745; SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs14, %rs9};1746; SM70-NEXT: ret;1747;1748; SM80-LABEL: test_maximum_v2(1749; SM80: {1750; SM80-NEXT: .reg .b32 %r<4>;1751; SM80-EMPTY:1752; SM80-NEXT: // %bb.0:1753; SM80-NEXT: ld.param.b32 %r1, [test_maximum_v2_param_0];1754; SM80-NEXT: ld.param.b32 %r2, [test_maximum_v2_param_1];1755; SM80-NEXT: max.NaN.bf16x2 %r3, %r1, %r2;1756; SM80-NEXT: st.param.b32 [func_retval0], %r3;1757; SM80-NEXT: ret;1758;1759; SM80-FTZ-LABEL: test_maximum_v2(1760; SM80-FTZ: {1761; SM80-FTZ-NEXT: .reg .b32 %r<4>;1762; SM80-FTZ-EMPTY:1763; SM80-FTZ-NEXT: // %bb.0:1764; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_maximum_v2_param_0];1765; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_maximum_v2_param_1];1766; SM80-FTZ-NEXT: max.NaN.bf16x2 %r3, %r1, %r2;1767; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r3;1768; SM80-FTZ-NEXT: ret;1769;1770; SM90-FTZ-LABEL: test_maximum_v2(1771; SM90-FTZ: {1772; SM90-FTZ-NEXT: .reg .b32 %r<4>;1773; SM90-FTZ-EMPTY:1774; SM90-FTZ-NEXT: // %bb.0:1775; SM90-FTZ-NEXT: ld.param.b32 %r1, [test_maximum_v2_param_0];1776; SM90-FTZ-NEXT: ld.param.b32 %r2, [test_maximum_v2_param_1];1777; SM90-FTZ-NEXT: max.NaN.bf16x2 %r3, %r1, %r2;1778; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r3;1779; SM90-FTZ-NEXT: ret;1780;1781; SM90-LABEL: test_maximum_v2(1782; SM90: {1783; SM90-NEXT: .reg .b32 %r<4>;1784; SM90-EMPTY:1785; SM90-NEXT: // %bb.0:1786; SM90-NEXT: ld.param.b32 %r1, [test_maximum_v2_param_0];1787; SM90-NEXT: ld.param.b32 %r2, [test_maximum_v2_param_1];1788; SM90-NEXT: max.NaN.bf16x2 %r3, %r1, %r2;1789; SM90-NEXT: st.param.b32 [func_retval0], %r3;1790; SM90-NEXT: ret;1791 %r = call <2 x bfloat> @llvm.maximum.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b)1792 ret <2 x bfloat> %r1793}1794 1795define <2 x bfloat> @test_maxnum_v2(<2 x bfloat> %a, <2 x bfloat> %b) {1796; SM70-LABEL: test_maxnum_v2(1797; SM70: {1798; SM70-NEXT: .reg .pred %p<3>;1799; SM70-NEXT: .reg .b16 %rs<5>;1800; SM70-NEXT: .reg .b32 %r<22>;1801; SM70-EMPTY:1802; SM70-NEXT: // %bb.0:1803; SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_maxnum_v2_param_0];1804; SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_maxnum_v2_param_1];1805; SM70-NEXT: cvt.u32.u16 %r1, %rs4;1806; SM70-NEXT: shl.b32 %r2, %r1, 16;1807; SM70-NEXT: cvt.u32.u16 %r3, %rs2;1808; SM70-NEXT: shl.b32 %r4, %r3, 16;1809; SM70-NEXT: max.f32 %r5, %r4, %r2;1810; SM70-NEXT: bfe.u32 %r6, %r5, 16, 1;1811; SM70-NEXT: add.s32 %r7, %r6, %r5;1812; SM70-NEXT: add.s32 %r8, %r7, 32767;1813; SM70-NEXT: setp.nan.f32 %p1, %r5, %r5;1814; SM70-NEXT: or.b32 %r9, %r5, 4194304;1815; SM70-NEXT: selp.b32 %r10, %r9, %r8, %p1;1816; SM70-NEXT: cvt.u32.u16 %r11, %rs3;1817; SM70-NEXT: shl.b32 %r12, %r11, 16;1818; SM70-NEXT: cvt.u32.u16 %r13, %rs1;1819; SM70-NEXT: shl.b32 %r14, %r13, 16;1820; SM70-NEXT: max.f32 %r15, %r14, %r12;1821; SM70-NEXT: bfe.u32 %r16, %r15, 16, 1;1822; SM70-NEXT: add.s32 %r17, %r16, %r15;1823; SM70-NEXT: add.s32 %r18, %r17, 32767;1824; SM70-NEXT: setp.nan.f32 %p2, %r15, %r15;1825; SM70-NEXT: or.b32 %r19, %r15, 4194304;1826; SM70-NEXT: selp.b32 %r20, %r19, %r18, %p2;1827; SM70-NEXT: prmt.b32 %r21, %r20, %r10, 0x7632U;1828; SM70-NEXT: st.param.b32 [func_retval0], %r21;1829; SM70-NEXT: ret;1830;1831; SM80-LABEL: test_maxnum_v2(1832; SM80: {1833; SM80-NEXT: .reg .b32 %r<4>;1834; SM80-EMPTY:1835; SM80-NEXT: // %bb.0:1836; SM80-NEXT: ld.param.b32 %r1, [test_maxnum_v2_param_0];1837; SM80-NEXT: ld.param.b32 %r2, [test_maxnum_v2_param_1];1838; SM80-NEXT: max.bf16x2 %r3, %r1, %r2;1839; SM80-NEXT: st.param.b32 [func_retval0], %r3;1840; SM80-NEXT: ret;1841;1842; SM80-FTZ-LABEL: test_maxnum_v2(1843; SM80-FTZ: {1844; SM80-FTZ-NEXT: .reg .b32 %r<4>;1845; SM80-FTZ-EMPTY:1846; SM80-FTZ-NEXT: // %bb.0:1847; SM80-FTZ-NEXT: ld.param.b32 %r1, [test_maxnum_v2_param_0];1848; SM80-FTZ-NEXT: ld.param.b32 %r2, [test_maxnum_v2_param_1];1849; SM80-FTZ-NEXT: max.bf16x2 %r3, %r1, %r2;1850; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r3;1851; SM80-FTZ-NEXT: ret;1852;1853; SM90-FTZ-LABEL: test_maxnum_v2(1854; SM90-FTZ: {1855; SM90-FTZ-NEXT: .reg .b32 %r<4>;1856; SM90-FTZ-EMPTY:1857; SM90-FTZ-NEXT: // %bb.0:1858; SM90-FTZ-NEXT: ld.param.b32 %r1, [test_maxnum_v2_param_0];1859; SM90-FTZ-NEXT: ld.param.b32 %r2, [test_maxnum_v2_param_1];1860; SM90-FTZ-NEXT: max.bf16x2 %r3, %r1, %r2;1861; SM90-FTZ-NEXT: st.param.b32 [func_retval0], %r3;1862; SM90-FTZ-NEXT: ret;1863;1864; SM90-LABEL: test_maxnum_v2(1865; SM90: {1866; SM90-NEXT: .reg .b32 %r<4>;1867; SM90-EMPTY:1868; SM90-NEXT: // %bb.0:1869; SM90-NEXT: ld.param.b32 %r1, [test_maxnum_v2_param_0];1870; SM90-NEXT: ld.param.b32 %r2, [test_maxnum_v2_param_1];1871; SM90-NEXT: max.bf16x2 %r3, %r1, %r2;1872; SM90-NEXT: st.param.b32 [func_retval0], %r3;1873; SM90-NEXT: ret;1874 %r = call <2 x bfloat> @llvm.maxnum.v2bf16(<2 x bfloat> %a, <2 x bfloat> %b)1875 ret <2 x bfloat> %r1876}1877 1878define void @store_bf16(ptr %p1, ptr %p2, bfloat %v) {1879; CHECK-LABEL: store_bf16(1880; CHECK: {1881; CHECK-NEXT: .reg .b16 %rs<2>;1882; CHECK-NEXT: .reg .b64 %rd<3>;1883; CHECK-EMPTY:1884; CHECK-NEXT: // %bb.0:1885; CHECK-NEXT: ld.param.b64 %rd1, [store_bf16_param_0];1886; CHECK-NEXT: ld.param.b16 %rs1, [store_bf16_param_2];1887; CHECK-NEXT: st.b16 [%rd1], %rs1;1888; CHECK-NEXT: ld.param.b64 %rd2, [store_bf16_param_1];1889; CHECK-NEXT: st.b16 [%rd2], 0x3F80;1890; CHECK-NEXT: ret;1891 store bfloat %v, ptr %p11892 store bfloat 1.0, ptr %p21893 ret void1894}1895 1896declare bfloat @llvm.maximum.bf16(bfloat, bfloat)1897declare <2 x bfloat> @llvm.maximum.v2bf16(<2 x bfloat>, <2 x bfloat>)1898