1660 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 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}4 5; Using FTZ should emit fma.ftz.relu for f16, not for bf166; RUN: llc < %s -denormal-fp-math-f32=preserve-sign -mtriple=nvptx64 -mcpu=sm_80 | FileCheck %s --check-prefixes=CHECK-FTZ7; RUN: %if ptxas %{ llc < %s -denormal-fp-math-f32=preserve-sign -mtriple=nvptx64 -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}8 9; SM < 80 or (which needs PTX version >= 70) should not emit fma{.ftz}.relu10; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 | FileCheck %s --check-prefixes=CHECK-SM7011 12define half @fma_f16_expanded_no_nans(half %a, half %b, half %c) {13; CHECK-LABEL: fma_f16_expanded_no_nans(14; CHECK: {15; CHECK-NEXT: .reg .b16 %rs<5>;16; CHECK-EMPTY:17; CHECK-NEXT: // %bb.0:18; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_param_0];19; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_param_1];20; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_param_2];21; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;22; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;23; CHECK-NEXT: ret;24;25; CHECK-FTZ-LABEL: fma_f16_expanded_no_nans(26; CHECK-FTZ: {27; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;28; CHECK-FTZ-EMPTY:29; CHECK-FTZ-NEXT: // %bb.0:30; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_param_0];31; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_param_1];32; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_param_2];33; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;34; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;35; CHECK-FTZ-NEXT: ret;36;37; CHECK-SM70-LABEL: fma_f16_expanded_no_nans(38; CHECK-SM70: {39; CHECK-SM70-NEXT: .reg .pred %p<2>;40; CHECK-SM70-NEXT: .reg .b16 %rs<7>;41; CHECK-SM70-EMPTY:42; CHECK-SM70-NEXT: // %bb.0:43; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_param_0];44; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_param_1];45; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_param_2];46; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;47; CHECK-SM70-NEXT: mov.b16 %rs5, 0x0000;48; CHECK-SM70-NEXT: setp.gt.f16 %p1, %rs4, %rs5;49; CHECK-SM70-NEXT: selp.b16 %rs6, %rs4, 0x0000, %p1;50; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs6;51; CHECK-SM70-NEXT: ret;52 %1 = fmul fast half %a, %b53 %2 = fadd fast half %1, %c54 %3 = fcmp nsz ogt half %2, 0.055 %4 = select nsz i1 %3, half %2, half 0.056 ret half %457}58 59; FMA relu shouldn't be selected if the FMA operation has multiple uses60define half @fma_f16_expanded_no_nans_multiple_uses_of_fma(half %a, half %b, half %c) {61; CHECK-LABEL: fma_f16_expanded_no_nans_multiple_uses_of_fma(62; CHECK: {63; CHECK-NEXT: .reg .b16 %rs<10>;64; CHECK-EMPTY:65; CHECK-NEXT: // %bb.0:66; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_0];67; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_1];68; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_2];69; CHECK-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;70; CHECK-NEXT: mov.b16 %rs5, 0x0000;71; CHECK-NEXT: max.f16 %rs6, %rs4, %rs5;72; CHECK-NEXT: mov.b16 %rs7, 0x4700;73; CHECK-NEXT: add.rn.f16 %rs8, %rs4, %rs7;74; CHECK-NEXT: add.rn.f16 %rs9, %rs6, %rs8;75; CHECK-NEXT: st.param.b16 [func_retval0], %rs9;76; CHECK-NEXT: ret;77;78; CHECK-FTZ-LABEL: fma_f16_expanded_no_nans_multiple_uses_of_fma(79; CHECK-FTZ: {80; CHECK-FTZ-NEXT: .reg .b16 %rs<10>;81; CHECK-FTZ-EMPTY:82; CHECK-FTZ-NEXT: // %bb.0:83; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_0];84; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_1];85; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_2];86; CHECK-FTZ-NEXT: fma.rn.ftz.f16 %rs4, %rs1, %rs2, %rs3;87; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x0000;88; CHECK-FTZ-NEXT: max.ftz.f16 %rs6, %rs4, %rs5;89; CHECK-FTZ-NEXT: mov.b16 %rs7, 0x4700;90; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs8, %rs4, %rs7;91; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs9, %rs6, %rs8;92; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs9;93; CHECK-FTZ-NEXT: ret;94;95; CHECK-SM70-LABEL: fma_f16_expanded_no_nans_multiple_uses_of_fma(96; CHECK-SM70: {97; CHECK-SM70-NEXT: .reg .pred %p<2>;98; CHECK-SM70-NEXT: .reg .b16 %rs<10>;99; CHECK-SM70-EMPTY:100; CHECK-SM70-NEXT: // %bb.0:101; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_0];102; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_1];103; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_2];104; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;105; CHECK-SM70-NEXT: mov.b16 %rs5, 0x0000;106; CHECK-SM70-NEXT: setp.gt.f16 %p1, %rs4, %rs5;107; CHECK-SM70-NEXT: selp.b16 %rs6, %rs4, 0x0000, %p1;108; CHECK-SM70-NEXT: mov.b16 %rs7, 0x4700;109; CHECK-SM70-NEXT: add.rn.f16 %rs8, %rs4, %rs7;110; CHECK-SM70-NEXT: add.rn.f16 %rs9, %rs6, %rs8;111; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs9;112; CHECK-SM70-NEXT: ret;113 %1 = fmul fast half %a, %b114 %2 = fadd fast half %1, %c115 %3 = fcmp nsz ogt half %2, 0.0116 %4 = select nsz i1 %3, half %2, half 0.0117 %5 = fadd half %2, 7.0118 %6 = fadd half %4, %5119 ret half %6120}121 122define half @fma_f16_expanded_maxnum_no_nans(half %a, half %b, half %c) {123; CHECK-LABEL: fma_f16_expanded_maxnum_no_nans(124; CHECK: {125; CHECK-NEXT: .reg .b16 %rs<5>;126; CHECK-EMPTY:127; CHECK-NEXT: // %bb.0:128; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];129; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];130; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];131; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;132; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;133; CHECK-NEXT: ret;134;135; CHECK-FTZ-LABEL: fma_f16_expanded_maxnum_no_nans(136; CHECK-FTZ: {137; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;138; CHECK-FTZ-EMPTY:139; CHECK-FTZ-NEXT: // %bb.0:140; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];141; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];142; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];143; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;144; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;145; CHECK-FTZ-NEXT: ret;146;147; CHECK-SM70-LABEL: fma_f16_expanded_maxnum_no_nans(148; CHECK-SM70: {149; CHECK-SM70-NEXT: .reg .b16 %rs<6>;150; CHECK-SM70-NEXT: .reg .b32 %r<3>;151; CHECK-SM70-EMPTY:152; CHECK-SM70-NEXT: // %bb.0:153; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];154; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];155; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];156; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;157; CHECK-SM70-NEXT: cvt.f32.f16 %r1, %rs4;158; CHECK-SM70-NEXT: max.f32 %r2, %r1, 0f00000000;159; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs5, %r2;160; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs5;161; CHECK-SM70-NEXT: ret;162 %1 = fmul fast half %a, %b163 %2 = fadd fast half %1, %c164 %3 = call nsz half @llvm.maxnum.f16(half %2, half 0.0)165 ret half %3166}167 168define bfloat @fma_bf16_expanded_no_nans(bfloat %a, bfloat %b, bfloat %c) {169; CHECK-LABEL: fma_bf16_expanded_no_nans(170; CHECK: {171; CHECK-NEXT: .reg .b16 %rs<5>;172; CHECK-EMPTY:173; CHECK-NEXT: // %bb.0:174; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_param_0];175; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_param_1];176; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_param_2];177; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;178; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;179; CHECK-NEXT: ret;180;181; CHECK-FTZ-LABEL: fma_bf16_expanded_no_nans(182; CHECK-FTZ: {183; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;184; CHECK-FTZ-EMPTY:185; CHECK-FTZ-NEXT: // %bb.0:186; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_param_0];187; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_param_1];188; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_param_2];189; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;190; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;191; CHECK-FTZ-NEXT: ret;192;193; CHECK-SM70-LABEL: fma_bf16_expanded_no_nans(194; CHECK-SM70: {195; CHECK-SM70-NEXT: .reg .pred %p<3>;196; CHECK-SM70-NEXT: .reg .b16 %rs<3>;197; CHECK-SM70-NEXT: .reg .b32 %r<14>;198; CHECK-SM70-EMPTY:199; CHECK-SM70-NEXT: // %bb.0:200; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_no_nans_param_2];201; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;202; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_no_nans_param_1];203; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;204; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_no_nans_param_0];205; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;206; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;207; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;208; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;209; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;210; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;211; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;212; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;213; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }214; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;215; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;216; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;217; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs2;218; CHECK-SM70-NEXT: ret;219 %1 = fmul fast bfloat %a, %b220 %2 = fadd fast bfloat %1, %c221 %3 = fcmp nsz ogt bfloat %2, 0.0222 %4 = select nsz i1 %3, bfloat %2, bfloat 0.0223 ret bfloat %4224}225 226; FMA relu shouldn't be selected if the FMA operation has multiple uses227define bfloat @fma_bf16_expanded_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloat %c) {228; CHECK-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(229; CHECK: {230; CHECK-NEXT: .reg .b16 %rs<11>;231; CHECK-EMPTY:232; CHECK-NEXT: // %bb.0:233; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];234; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];235; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];236; CHECK-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;237; CHECK-NEXT: mov.b16 %rs5, 0x0000;238; CHECK-NEXT: max.bf16 %rs6, %rs4, %rs5;239; CHECK-NEXT: mov.b16 %rs7, 0x3F80;240; CHECK-NEXT: mov.b16 %rs8, 0x40E0;241; CHECK-NEXT: fma.rn.bf16 %rs9, %rs4, %rs7, %rs8;242; CHECK-NEXT: fma.rn.bf16 %rs10, %rs6, %rs7, %rs9;243; CHECK-NEXT: st.param.b16 [func_retval0], %rs10;244; CHECK-NEXT: ret;245;246; CHECK-FTZ-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(247; CHECK-FTZ: {248; CHECK-FTZ-NEXT: .reg .b16 %rs<9>;249; CHECK-FTZ-NEXT: .reg .b32 %r<9>;250; CHECK-FTZ-EMPTY:251; CHECK-FTZ-NEXT: // %bb.0:252; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];253; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];254; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];255; CHECK-FTZ-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;256; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x0000;257; CHECK-FTZ-NEXT: max.bf16 %rs6, %rs4, %rs5;258; CHECK-FTZ-NEXT: cvt.u32.u16 %r1, %rs4;259; CHECK-FTZ-NEXT: shl.b32 %r2, %r1, 16;260; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40E00000;261; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs7, %r3;262; CHECK-FTZ-NEXT: cvt.u32.u16 %r4, %rs6;263; CHECK-FTZ-NEXT: shl.b32 %r5, %r4, 16;264; CHECK-FTZ-NEXT: cvt.u32.u16 %r6, %rs7;265; CHECK-FTZ-NEXT: shl.b32 %r7, %r6, 16;266; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r8, %r5, %r7;267; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs8, %r8;268; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs8;269; CHECK-FTZ-NEXT: ret;270;271; CHECK-SM70-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(272; CHECK-SM70: {273; CHECK-SM70-NEXT: .reg .pred %p<5>;274; CHECK-SM70-NEXT: .reg .b16 %rs<3>;275; CHECK-SM70-NEXT: .reg .b32 %r<30>;276; CHECK-SM70-EMPTY:277; CHECK-SM70-NEXT: // %bb.0:278; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];279; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;280; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];281; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;282; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];283; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;284; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;285; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;286; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;287; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;288; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;289; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;290; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;291; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }292; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;293; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;294; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;295; CHECK-SM70-NEXT: add.rn.f32 %r14, %r13, 0f40E00000;296; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;297; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;298; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;299; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r14, %r14;300; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;301; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p3;302; CHECK-SM70-NEXT: cvt.u32.u16 %r20, %rs2;303; CHECK-SM70-NEXT: shl.b32 %r21, %r20, 16;304; CHECK-SM70-NEXT: and.b32 %r22, %r19, -65536;305; CHECK-SM70-NEXT: add.rn.f32 %r23, %r21, %r22;306; CHECK-SM70-NEXT: bfe.u32 %r24, %r23, 16, 1;307; CHECK-SM70-NEXT: add.s32 %r25, %r24, %r23;308; CHECK-SM70-NEXT: add.s32 %r26, %r25, 32767;309; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r23, %r23;310; CHECK-SM70-NEXT: or.b32 %r27, %r23, 4194304;311; CHECK-SM70-NEXT: selp.b32 %r28, %r27, %r26, %p4;312; CHECK-SM70-NEXT: shr.u32 %r29, %r28, 16;313; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r29;314; CHECK-SM70-NEXT: ret;315 %1 = fmul fast bfloat %a, %b316 %2 = fadd fast bfloat %1, %c317 %3 = fcmp nsz ogt bfloat %2, 0.0318 %4 = select nsz i1 %3, bfloat %2, bfloat 0.0319 %5 = fadd bfloat %2, 7.0320 %6 = fadd bfloat %4, %5321 ret bfloat %6322}323 324define bfloat @fma_bf16_expanded_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) {325;326;327; CHECK-LABEL: fma_bf16_expanded_maxnum_no_nans(328; CHECK: {329; CHECK-NEXT: .reg .b16 %rs<5>;330; CHECK-EMPTY:331; CHECK-NEXT: // %bb.0:332; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_maxnum_no_nans_param_0];333; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_maxnum_no_nans_param_1];334; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_maxnum_no_nans_param_2];335; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;336; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;337; CHECK-NEXT: ret;338;339; CHECK-FTZ-LABEL: fma_bf16_expanded_maxnum_no_nans(340; CHECK-FTZ: {341; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;342; CHECK-FTZ-EMPTY:343; CHECK-FTZ-NEXT: // %bb.0:344; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_maxnum_no_nans_param_0];345; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_maxnum_no_nans_param_1];346; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_maxnum_no_nans_param_2];347; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;348; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;349; CHECK-FTZ-NEXT: ret;350;351; CHECK-SM70-LABEL: fma_bf16_expanded_maxnum_no_nans(352; CHECK-SM70: {353; CHECK-SM70-NEXT: .reg .pred %p<3>;354; CHECK-SM70-NEXT: .reg .b32 %r<21>;355; CHECK-SM70-EMPTY:356; CHECK-SM70-NEXT: // %bb.0:357; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_maxnum_no_nans_param_2];358; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;359; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_maxnum_no_nans_param_1];360; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;361; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_maxnum_no_nans_param_0];362; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;363; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;364; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;365; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;366; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;367; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;368; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;369; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;370; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;371; CHECK-SM70-NEXT: max.f32 %r14, %r13, 0f00000000;372; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;373; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;374; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;375; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;376; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;377; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;378; CHECK-SM70-NEXT: shr.u32 %r20, %r19, 16;379; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r20;380; CHECK-SM70-NEXT: ret;381 %1 = fmul fast bfloat %a, %b382 %2 = fadd fast bfloat %1, %c383 %3 = call nsz bfloat @llvm.maxnum.bf16(bfloat %2, bfloat 0.0)384 ret bfloat %3385}386 387define <2 x half> @fma_f16x2_expanded_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {388;389;390; CHECK-LABEL: fma_f16x2_expanded_no_nans(391; CHECK: {392; CHECK-NEXT: .reg .b32 %r<5>;393; CHECK-EMPTY:394; CHECK-NEXT: // %bb.0:395; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];396; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];397; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];398; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;399; CHECK-NEXT: st.param.b32 [func_retval0], %r4;400; CHECK-NEXT: ret;401;402; CHECK-FTZ-LABEL: fma_f16x2_expanded_no_nans(403; CHECK-FTZ: {404; CHECK-FTZ-NEXT: .reg .b32 %r<5>;405; CHECK-FTZ-EMPTY:406; CHECK-FTZ-NEXT: // %bb.0:407; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];408; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];409; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];410; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;411; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;412; CHECK-FTZ-NEXT: ret;413;414; CHECK-SM70-LABEL: fma_f16x2_expanded_no_nans(415; CHECK-SM70: {416; CHECK-SM70-NEXT: .reg .pred %p<3>;417; CHECK-SM70-NEXT: .reg .b16 %rs<5>;418; CHECK-SM70-NEXT: .reg .b32 %r<6>;419; CHECK-SM70-EMPTY:420; CHECK-SM70-NEXT: // %bb.0:421; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];422; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];423; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];424; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;425; CHECK-SM70-NEXT: mov.b32 %r5, 0;426; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;427; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;428; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;429; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;430; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};431; CHECK-SM70-NEXT: ret;432 %1 = fmul fast <2 x half> %a, %b433 %2 = fadd fast <2 x half> %1, %c434 %3 = fcmp nsz ogt <2 x half> %2, <half 0.0, half 0.0>435 %4 = select nsz <2 x i1> %3, <2 x half> %2, <2 x half> <half 0.0, half 0.0>436 ret <2 x half> %4437}438 439; FMA relu shouldn't be selected if the FMA operation has multiple uses440define <2 x half> @fma_f16x2_expanded_no_nans_multiple_uses_of_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) {441;442;443; CHECK-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(444; CHECK: {445; CHECK-NEXT: .reg .b32 %r<10>;446; CHECK-EMPTY:447; CHECK-NEXT: // %bb.0:448; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];449; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];450; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];451; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;452; CHECK-NEXT: mov.b32 %r5, 0;453; CHECK-NEXT: max.f16x2 %r6, %r4, %r5;454; CHECK-NEXT: mov.b32 %r7, 1191200512;455; CHECK-NEXT: add.rn.f16x2 %r8, %r4, %r7;456; CHECK-NEXT: add.rn.f16x2 %r9, %r6, %r8;457; CHECK-NEXT: st.param.b32 [func_retval0], %r9;458; CHECK-NEXT: ret;459;460; CHECK-FTZ-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(461; CHECK-FTZ: {462; CHECK-FTZ-NEXT: .reg .b32 %r<10>;463; CHECK-FTZ-EMPTY:464; CHECK-FTZ-NEXT: // %bb.0:465; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];466; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];467; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];468; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;469; CHECK-FTZ-NEXT: mov.b32 %r5, 0;470; CHECK-FTZ-NEXT: max.ftz.f16x2 %r6, %r4, %r5;471; CHECK-FTZ-NEXT: mov.b32 %r7, 1191200512;472; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r8, %r4, %r7;473; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r9, %r6, %r8;474; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r9;475; CHECK-FTZ-NEXT: ret;476;477; CHECK-SM70-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(478; CHECK-SM70: {479; CHECK-SM70-NEXT: .reg .pred %p<3>;480; CHECK-SM70-NEXT: .reg .b16 %rs<5>;481; CHECK-SM70-NEXT: .reg .b32 %r<10>;482; CHECK-SM70-EMPTY:483; CHECK-SM70-NEXT: // %bb.0:484; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];485; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];486; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];487; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;488; CHECK-SM70-NEXT: mov.b32 %r5, 0;489; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;490; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;491; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;492; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;493; CHECK-SM70-NEXT: mov.b32 %r6, {%rs4, %rs3};494; CHECK-SM70-NEXT: mov.b32 %r7, 1191200512;495; CHECK-SM70-NEXT: add.rn.f16x2 %r8, %r4, %r7;496; CHECK-SM70-NEXT: add.rn.f16x2 %r9, %r6, %r8;497; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;498; CHECK-SM70-NEXT: ret;499 %1 = fmul fast <2 x half> %a, %b500 %2 = fadd fast <2 x half> %1, %c501 %3 = fcmp nsz ogt <2 x half> %2, <half 0.0, half 0.0>502 %4 = select nsz <2 x i1> %3, <2 x half> %2, <2 x half> <half 0.0, half 0.0>503 %5 = fadd <2 x half> %2, <half 7.0, half 7.0>504 %6 = fadd <2 x half> %4, %5505 ret <2 x half> %6506}507 508define <2 x half> @fma_f16x2_expanded_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {509; CHECK-LABEL: fma_f16x2_expanded_maxnum_no_nans(510; CHECK: {511; CHECK-NEXT: .reg .b32 %r<5>;512; CHECK-EMPTY:513; CHECK-NEXT: // %bb.0:514; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];515; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];516; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];517; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;518; CHECK-NEXT: st.param.b32 [func_retval0], %r4;519; CHECK-NEXT: ret;520;521; CHECK-FTZ-LABEL: fma_f16x2_expanded_maxnum_no_nans(522; CHECK-FTZ: {523; CHECK-FTZ-NEXT: .reg .b32 %r<5>;524; CHECK-FTZ-EMPTY:525; CHECK-FTZ-NEXT: // %bb.0:526; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];527; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];528; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];529; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;530; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;531; CHECK-FTZ-NEXT: ret;532;533; CHECK-SM70-LABEL: fma_f16x2_expanded_maxnum_no_nans(534; CHECK-SM70: {535; CHECK-SM70-NEXT: .reg .b16 %rs<5>;536; CHECK-SM70-NEXT: .reg .b32 %r<10>;537; CHECK-SM70-EMPTY:538; CHECK-SM70-NEXT: // %bb.0:539; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];540; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];541; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];542; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;543; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;544; CHECK-SM70-NEXT: cvt.f32.f16 %r5, %rs2;545; CHECK-SM70-NEXT: max.f32 %r6, %r5, 0f00000000;546; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs3, %r6;547; CHECK-SM70-NEXT: cvt.f32.f16 %r7, %rs1;548; CHECK-SM70-NEXT: max.f32 %r8, %r7, 0f00000000;549; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs4, %r8;550; CHECK-SM70-NEXT: mov.b32 %r9, {%rs4, %rs3};551; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;552; CHECK-SM70-NEXT: ret;553 %1 = fmul fast <2 x half> %a, %b554 %2 = fadd fast <2 x half> %1, %c555 %3 = call nsz <2 x half> @llvm.maxnum.f16x2(<2 x half> %2, <2 x half> <half 0.0, half 0.0>)556 ret <2 x half> %3557}558 559define <2 x bfloat> @fma_bf16x2_expanded_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {560; CHECK-LABEL: fma_bf16x2_expanded_no_nans(561; CHECK: {562; CHECK-NEXT: .reg .b32 %r<5>;563; CHECK-EMPTY:564; CHECK-NEXT: // %bb.0:565; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];566; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_param_1];567; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_param_2];568; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;569; CHECK-NEXT: st.param.b32 [func_retval0], %r4;570; CHECK-NEXT: ret;571;572; CHECK-FTZ-LABEL: fma_bf16x2_expanded_no_nans(573; CHECK-FTZ: {574; CHECK-FTZ-NEXT: .reg .b32 %r<5>;575; CHECK-FTZ-EMPTY:576; CHECK-FTZ-NEXT: // %bb.0:577; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];578; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_param_1];579; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_param_2];580; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;581; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;582; CHECK-FTZ-NEXT: ret;583;584; CHECK-SM70-LABEL: fma_bf16x2_expanded_no_nans(585; CHECK-SM70: {586; CHECK-SM70-NEXT: .reg .pred %p<5>;587; CHECK-SM70-NEXT: .reg .b16 %rs<11>;588; CHECK-SM70-NEXT: .reg .b32 %r<27>;589; CHECK-SM70-EMPTY:590; CHECK-SM70-NEXT: // %bb.0:591; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_no_nans_param_0];592; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_no_nans_param_1];593; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_no_nans_param_2];594; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs5;595; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;596; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs3;597; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;598; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;599; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;600; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;601; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;602; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;603; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;604; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;605; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;606; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;607; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }608; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs6;609; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;610; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs4;611; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;612; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;613; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;614; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;615; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;616; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;617; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;618; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;619; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;620; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;621; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }622; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;623; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;624; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;625; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;626; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;627; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;628; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};629; CHECK-SM70-NEXT: ret;630 %1 = fmul fast <2 x bfloat> %a, %b631 %2 = fadd fast <2 x bfloat> %1, %c632 %3 = fcmp nsz ogt <2 x bfloat> %2, <bfloat 0.0, bfloat 0.0>633 %4 = select nsz <2 x i1> %3, <2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>634 ret <2 x bfloat> %4635}636 637; FMA relu shouldn't be selected if the FMA operation has multiple uses638define <2 x bfloat> @fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {639; CHECK-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(640; CHECK: {641; CHECK-NEXT: .reg .b32 %r<11>;642; CHECK-EMPTY:643; CHECK-NEXT: // %bb.0:644; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];645; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];646; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];647; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;648; CHECK-NEXT: mov.b32 %r5, 0;649; CHECK-NEXT: max.bf16x2 %r6, %r4, %r5;650; CHECK-NEXT: mov.b32 %r7, 1065369472;651; CHECK-NEXT: mov.b32 %r8, 1088438496;652; CHECK-NEXT: fma.rn.bf16x2 %r9, %r4, %r7, %r8;653; CHECK-NEXT: fma.rn.bf16x2 %r10, %r6, %r7, %r9;654; CHECK-NEXT: st.param.b32 [func_retval0], %r10;655; CHECK-NEXT: ret;656;657; CHECK-FTZ-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(658; CHECK-FTZ: {659; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;660; CHECK-FTZ-NEXT: .reg .b32 %r<24>;661; CHECK-FTZ-EMPTY:662; CHECK-FTZ-NEXT: // %bb.0:663; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];664; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];665; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];666; CHECK-FTZ-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;667; CHECK-FTZ-NEXT: mov.b32 %r5, 0;668; CHECK-FTZ-NEXT: max.bf16x2 %r6, %r4, %r5;669; CHECK-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r4;670; CHECK-FTZ-NEXT: cvt.u32.u16 %r7, %rs2;671; CHECK-FTZ-NEXT: shl.b32 %r8, %r7, 16;672; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r9, %r8, 0f40E00000;673; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r9;674; CHECK-FTZ-NEXT: cvt.u32.u16 %r10, %rs1;675; CHECK-FTZ-NEXT: shl.b32 %r11, %r10, 16;676; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r12, %r11, 0f40E00000;677; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs4, %r12;678; CHECK-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r6;679; CHECK-FTZ-NEXT: cvt.u32.u16 %r13, %rs5;680; CHECK-FTZ-NEXT: shl.b32 %r14, %r13, 16;681; CHECK-FTZ-NEXT: cvt.u32.u16 %r15, %rs4;682; CHECK-FTZ-NEXT: shl.b32 %r16, %r15, 16;683; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r17, %r14, %r16;684; CHECK-FTZ-NEXT: cvt.u32.u16 %r18, %rs6;685; CHECK-FTZ-NEXT: shl.b32 %r19, %r18, 16;686; CHECK-FTZ-NEXT: cvt.u32.u16 %r20, %rs3;687; CHECK-FTZ-NEXT: shl.b32 %r21, %r20, 16;688; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r22, %r19, %r21;689; CHECK-FTZ-NEXT: cvt.rn.bf16x2.f32 %r23, %r22, %r17;690; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r23;691; CHECK-FTZ-NEXT: ret;692;693; CHECK-SM70-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(694; CHECK-SM70: {695; CHECK-SM70-NEXT: .reg .pred %p<9>;696; CHECK-SM70-NEXT: .reg .b16 %rs<11>;697; CHECK-SM70-NEXT: .reg .b32 %r<58>;698; CHECK-SM70-EMPTY:699; CHECK-SM70-NEXT: // %bb.0:700; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];701; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];702; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];703; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs6;704; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;705; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs4;706; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;707; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs2;708; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;709; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;710; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;711; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;712; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;713; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;714; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;715; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;716; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }717; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs5;718; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;719; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs3;720; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;721; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs1;722; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;723; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;724; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;725; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;726; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;727; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;728; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;729; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;730; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }731; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;732; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;733; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;734; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;735; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;736; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;737; CHECK-SM70-NEXT: add.rn.f32 %r27, %r26, 0f40E00000;738; CHECK-SM70-NEXT: bfe.u32 %r28, %r27, 16, 1;739; CHECK-SM70-NEXT: add.s32 %r29, %r28, %r27;740; CHECK-SM70-NEXT: add.s32 %r30, %r29, 32767;741; CHECK-SM70-NEXT: setp.nan.f32 %p5, %r27, %r27;742; CHECK-SM70-NEXT: or.b32 %r31, %r27, 4194304;743; CHECK-SM70-NEXT: selp.b32 %r32, %r31, %r30, %p5;744; CHECK-SM70-NEXT: add.rn.f32 %r33, %r25, 0f40E00000;745; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;746; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;747; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;748; CHECK-SM70-NEXT: setp.nan.f32 %p6, %r33, %r33;749; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;750; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p6;751; CHECK-SM70-NEXT: cvt.u32.u16 %r39, %rs10;752; CHECK-SM70-NEXT: shl.b32 %r40, %r39, 16;753; CHECK-SM70-NEXT: and.b32 %r41, %r38, -65536;754; CHECK-SM70-NEXT: add.rn.f32 %r42, %r40, %r41;755; CHECK-SM70-NEXT: bfe.u32 %r43, %r42, 16, 1;756; CHECK-SM70-NEXT: add.s32 %r44, %r43, %r42;757; CHECK-SM70-NEXT: add.s32 %r45, %r44, 32767;758; CHECK-SM70-NEXT: setp.nan.f32 %p7, %r42, %r42;759; CHECK-SM70-NEXT: or.b32 %r46, %r42, 4194304;760; CHECK-SM70-NEXT: selp.b32 %r47, %r46, %r45, %p7;761; CHECK-SM70-NEXT: cvt.u32.u16 %r48, %rs9;762; CHECK-SM70-NEXT: shl.b32 %r49, %r48, 16;763; CHECK-SM70-NEXT: and.b32 %r50, %r32, -65536;764; CHECK-SM70-NEXT: add.rn.f32 %r51, %r49, %r50;765; CHECK-SM70-NEXT: bfe.u32 %r52, %r51, 16, 1;766; CHECK-SM70-NEXT: add.s32 %r53, %r52, %r51;767; CHECK-SM70-NEXT: add.s32 %r54, %r53, 32767;768; CHECK-SM70-NEXT: setp.nan.f32 %p8, %r51, %r51;769; CHECK-SM70-NEXT: or.b32 %r55, %r51, 4194304;770; CHECK-SM70-NEXT: selp.b32 %r56, %r55, %r54, %p8;771; CHECK-SM70-NEXT: prmt.b32 %r57, %r56, %r47, 0x7632U;772; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r57;773; CHECK-SM70-NEXT: ret;774 %1 = fmul fast <2 x bfloat> %a, %b775 %2 = fadd fast <2 x bfloat> %1, %c776 %3 = fcmp nsz ogt <2 x bfloat> %2, <bfloat 0.0, bfloat 0.0>777 %4 = select nsz <2 x i1> %3, <2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>778 %5 = fadd <2 x bfloat> %2, <bfloat 7.0, bfloat 7.0>779 %6 = fadd <2 x bfloat> %4, %5780 ret <2 x bfloat> %6781}782 783define <2 x bfloat> @fma_bf16x2_expanded_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {784; CHECK-LABEL: fma_bf16x2_expanded_maxnum_no_nans(785; CHECK: {786; CHECK-NEXT: .reg .b32 %r<5>;787; CHECK-EMPTY:788; CHECK-NEXT: // %bb.0:789; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];790; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_maxnum_no_nans_param_1];791; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_maxnum_no_nans_param_2];792; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;793; CHECK-NEXT: st.param.b32 [func_retval0], %r4;794; CHECK-NEXT: ret;795;796; CHECK-FTZ-LABEL: fma_bf16x2_expanded_maxnum_no_nans(797; CHECK-FTZ: {798; CHECK-FTZ-NEXT: .reg .b32 %r<5>;799; CHECK-FTZ-EMPTY:800; CHECK-FTZ-NEXT: // %bb.0:801; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];802; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_maxnum_no_nans_param_1];803; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_maxnum_no_nans_param_2];804; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;805; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;806; CHECK-FTZ-NEXT: ret;807;808; CHECK-SM70-LABEL: fma_bf16x2_expanded_maxnum_no_nans(809; CHECK-SM70: {810; CHECK-SM70-NEXT: .reg .pred %p<5>;811; CHECK-SM70-NEXT: .reg .b16 %rs<7>;812; CHECK-SM70-NEXT: .reg .b32 %r<40>;813; CHECK-SM70-EMPTY:814; CHECK-SM70-NEXT: // %bb.0:815; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_maxnum_no_nans_param_0];816; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_maxnum_no_nans_param_1];817; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_maxnum_no_nans_param_2];818; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs5;819; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;820; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs3;821; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;822; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;823; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;824; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;825; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;826; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;827; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;828; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;829; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;830; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;831; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs6;832; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;833; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs4;834; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;835; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;836; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;837; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;838; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;839; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;840; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;841; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;842; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;843; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;844; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;845; CHECK-SM70-NEXT: max.f32 %r26, %r25, 0f00000000;846; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;847; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;848; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;849; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;850; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;851; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;852; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;853; CHECK-SM70-NEXT: max.f32 %r33, %r32, 0f00000000;854; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;855; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;856; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;857; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;858; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;859; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;860; CHECK-SM70-NEXT: prmt.b32 %r39, %r38, %r31, 0x7632U;861; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r39;862; CHECK-SM70-NEXT: ret;863 %1 = fmul fast <2 x bfloat> %a, %b864 %2 = fadd fast <2 x bfloat> %1, %c865 %3 = call nsz <2 x bfloat> @llvm.maxnum.bf16x2(<2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>)866 ret <2 x bfloat> %3867}868 869define half @fma_f16_no_nans(half %a, half %b, half %c) {870; CHECK-LABEL: fma_f16_no_nans(871; CHECK: {872; CHECK-NEXT: .reg .b16 %rs<5>;873; CHECK-EMPTY:874; CHECK-NEXT: // %bb.0:875; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_param_0];876; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];877; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_param_2];878; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;879; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;880; CHECK-NEXT: ret;881;882; CHECK-FTZ-LABEL: fma_f16_no_nans(883; CHECK-FTZ: {884; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;885; CHECK-FTZ-EMPTY:886; CHECK-FTZ-NEXT: // %bb.0:887; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_param_0];888; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];889; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_param_2];890; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;891; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;892; CHECK-FTZ-NEXT: ret;893;894; CHECK-SM70-LABEL: fma_f16_no_nans(895; CHECK-SM70: {896; CHECK-SM70-NEXT: .reg .pred %p<2>;897; CHECK-SM70-NEXT: .reg .b16 %rs<7>;898; CHECK-SM70-EMPTY:899; CHECK-SM70-NEXT: // %bb.0:900; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_param_0];901; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];902; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_param_2];903; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;904; CHECK-SM70-NEXT: mov.b16 %rs5, 0x0000;905; CHECK-SM70-NEXT: setp.gt.f16 %p1, %rs4, %rs5;906; CHECK-SM70-NEXT: selp.b16 %rs6, %rs4, 0x0000, %p1;907; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs6;908; CHECK-SM70-NEXT: ret;909 %1 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)910 %2 = fcmp nsz ogt half %1, 0.0911 %3 = select nsz i1 %2, half %1, half 0.0912 ret half %3913}914 915; FMA relu shouldn't be selected if the FMA operation has multiple uses916define half @fma_f16_no_nans_multiple_uses_of_fma(half %a, half %b, half %c) {917; CHECK-LABEL: fma_f16_no_nans_multiple_uses_of_fma(918; CHECK: {919; CHECK-NEXT: .reg .b16 %rs<8>;920; CHECK-EMPTY:921; CHECK-NEXT: // %bb.0:922; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];923; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];924; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];925; CHECK-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;926; CHECK-NEXT: mov.b16 %rs5, 0x4700;927; CHECK-NEXT: add.rn.f16 %rs6, %rs4, %rs5;928; CHECK-NEXT: add.rn.f16 %rs7, %rs6, %rs4;929; CHECK-NEXT: st.param.b16 [func_retval0], %rs7;930; CHECK-NEXT: ret;931;932; CHECK-FTZ-LABEL: fma_f16_no_nans_multiple_uses_of_fma(933; CHECK-FTZ: {934; CHECK-FTZ-NEXT: .reg .b16 %rs<8>;935; CHECK-FTZ-EMPTY:936; CHECK-FTZ-NEXT: // %bb.0:937; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];938; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];939; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];940; CHECK-FTZ-NEXT: fma.rn.ftz.f16 %rs4, %rs1, %rs2, %rs3;941; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x4700;942; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs6, %rs4, %rs5;943; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs7, %rs6, %rs4;944; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs7;945; CHECK-FTZ-NEXT: ret;946;947; CHECK-SM70-LABEL: fma_f16_no_nans_multiple_uses_of_fma(948; CHECK-SM70: {949; CHECK-SM70-NEXT: .reg .b16 %rs<8>;950; CHECK-SM70-EMPTY:951; CHECK-SM70-NEXT: // %bb.0:952; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];953; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];954; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];955; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;956; CHECK-SM70-NEXT: mov.b16 %rs5, 0x4700;957; CHECK-SM70-NEXT: add.rn.f16 %rs6, %rs4, %rs5;958; CHECK-SM70-NEXT: add.rn.f16 %rs7, %rs6, %rs4;959; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs7;960; CHECK-SM70-NEXT: ret;961 %1 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)962 %2 = fcmp nsz ogt half %1, 0.0963 %3 = select nsz i1 %2, half %1, half 0.0964 %4 = fadd half %1, 7.0965 %5 = fadd half %4, %1966 ret half %5967}968 969define half @fma_f16_maxnum_no_nans(half %a, half %b, half %c) {970; CHECK-LABEL: fma_f16_maxnum_no_nans(971; CHECK: {972; CHECK-NEXT: .reg .b16 %rs<5>;973; CHECK-EMPTY:974; CHECK-NEXT: // %bb.0:975; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];976; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];977; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];978; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;979; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;980; CHECK-NEXT: ret;981;982; CHECK-FTZ-LABEL: fma_f16_maxnum_no_nans(983; CHECK-FTZ: {984; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;985; CHECK-FTZ-EMPTY:986; CHECK-FTZ-NEXT: // %bb.0:987; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];988; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];989; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];990; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;991; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;992; CHECK-FTZ-NEXT: ret;993;994; CHECK-SM70-LABEL: fma_f16_maxnum_no_nans(995; CHECK-SM70: {996; CHECK-SM70-NEXT: .reg .b16 %rs<6>;997; CHECK-SM70-NEXT: .reg .b32 %r<3>;998; CHECK-SM70-EMPTY:999; CHECK-SM70-NEXT: // %bb.0:1000; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];1001; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];1002; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];1003; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;1004; CHECK-SM70-NEXT: cvt.f32.f16 %r1, %rs4;1005; CHECK-SM70-NEXT: max.f32 %r2, %r1, 0f00000000;1006; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs5, %r2;1007; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs5;1008; CHECK-SM70-NEXT: ret;1009 %1 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)1010 %2 = call nsz half @llvm.maxnum.f16(half %1, half 0.0)1011 ret half %21012}1013 1014define bfloat @fma_bf16_no_nans(bfloat %a, bfloat %b, bfloat %c) {1015; CHECK-LABEL: fma_bf16_no_nans(1016; CHECK: {1017; CHECK-NEXT: .reg .b16 %rs<5>;1018; CHECK-EMPTY:1019; CHECK-NEXT: // %bb.0:1020; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_param_0];1021; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_param_1];1022; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_param_2];1023; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;1024; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;1025; CHECK-NEXT: ret;1026;1027; CHECK-FTZ-LABEL: fma_bf16_no_nans(1028; CHECK-FTZ: {1029; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;1030; CHECK-FTZ-EMPTY:1031; CHECK-FTZ-NEXT: // %bb.0:1032; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_param_0];1033; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_param_1];1034; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_param_2];1035; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;1036; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;1037; CHECK-FTZ-NEXT: ret;1038;1039; CHECK-SM70-LABEL: fma_bf16_no_nans(1040; CHECK-SM70: {1041; CHECK-SM70-NEXT: .reg .pred %p<3>;1042; CHECK-SM70-NEXT: .reg .b16 %rs<3>;1043; CHECK-SM70-NEXT: .reg .b32 %r<14>;1044; CHECK-SM70-EMPTY:1045; CHECK-SM70-NEXT: // %bb.0:1046; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_no_nans_param_2];1047; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;1048; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_no_nans_param_1];1049; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;1050; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_no_nans_param_0];1051; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;1052; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;1053; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;1054; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;1055; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;1056; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;1057; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;1058; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;1059; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }1060; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;1061; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;1062; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;1063; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs2;1064; CHECK-SM70-NEXT: ret;1065 %1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)1066 %2 = fcmp nsz ogt bfloat %1, 0.01067 %3 = select nsz i1 %2, bfloat %1, bfloat 0.01068 ret bfloat %31069}1070 1071; FMA_relu shouldn't be selected if the FMA operation has multiple uses1072define bfloat @fma_bf16_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloat %c) {1073; CHECK-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(1074; CHECK: {1075; CHECK-NEXT: .reg .b16 %rs<9>;1076; CHECK-EMPTY:1077; CHECK-NEXT: // %bb.0:1078; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];1079; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];1080; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];1081; CHECK-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;1082; CHECK-NEXT: mov.b16 %rs5, 0x3F80;1083; CHECK-NEXT: mov.b16 %rs6, 0x40E0;1084; CHECK-NEXT: fma.rn.bf16 %rs7, %rs4, %rs5, %rs6;1085; CHECK-NEXT: fma.rn.bf16 %rs8, %rs7, %rs5, %rs4;1086; CHECK-NEXT: st.param.b16 [func_retval0], %rs8;1087; CHECK-NEXT: ret;1088;1089; CHECK-FTZ-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(1090; CHECK-FTZ: {1091; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;1092; CHECK-FTZ-NEXT: .reg .b32 %r<7>;1093; CHECK-FTZ-EMPTY:1094; CHECK-FTZ-NEXT: // %bb.0:1095; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];1096; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];1097; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];1098; CHECK-FTZ-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;1099; CHECK-FTZ-NEXT: cvt.u32.u16 %r1, %rs4;1100; CHECK-FTZ-NEXT: shl.b32 %r2, %r1, 16;1101; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40E00000;1102; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %r3;1103; CHECK-FTZ-NEXT: cvt.u32.u16 %r4, %rs5;1104; CHECK-FTZ-NEXT: shl.b32 %r5, %r4, 16;1105; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r6, %r5, %r2;1106; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %r6;1107; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs6;1108; CHECK-FTZ-NEXT: ret;1109;1110; CHECK-SM70-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(1111; CHECK-SM70: {1112; CHECK-SM70-NEXT: .reg .pred %p<4>;1113; CHECK-SM70-NEXT: .reg .b32 %r<28>;1114; CHECK-SM70-EMPTY:1115; CHECK-SM70-NEXT: // %bb.0:1116; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];1117; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;1118; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];1119; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;1120; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];1121; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;1122; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;1123; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;1124; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;1125; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;1126; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;1127; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;1128; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;1129; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;1130; CHECK-SM70-NEXT: add.rn.f32 %r14, %r13, 0f40E00000;1131; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;1132; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;1133; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;1134; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;1135; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;1136; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;1137; CHECK-SM70-NEXT: and.b32 %r20, %r19, -65536;1138; CHECK-SM70-NEXT: add.rn.f32 %r21, %r20, %r13;1139; CHECK-SM70-NEXT: bfe.u32 %r22, %r21, 16, 1;1140; CHECK-SM70-NEXT: add.s32 %r23, %r22, %r21;1141; CHECK-SM70-NEXT: add.s32 %r24, %r23, 32767;1142; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r21, %r21;1143; CHECK-SM70-NEXT: or.b32 %r25, %r21, 4194304;1144; CHECK-SM70-NEXT: selp.b32 %r26, %r25, %r24, %p3;1145; CHECK-SM70-NEXT: shr.u32 %r27, %r26, 16;1146; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r27;1147; CHECK-SM70-NEXT: ret;1148 %1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)1149 %2 = fcmp nsz ogt bfloat %1, 0.01150 %3 = select nsz i1 %2, bfloat %1, bfloat 0.01151 %4 = fadd bfloat %1, 7.01152 %5 = fadd bfloat %4, %11153 ret bfloat %51154}1155 1156define bfloat @fma_bf16_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) {1157; CHECK-LABEL: fma_bf16_maxnum_no_nans(1158; CHECK: {1159; CHECK-NEXT: .reg .b16 %rs<5>;1160; CHECK-EMPTY:1161; CHECK-NEXT: // %bb.0:1162; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_maxnum_no_nans_param_0];1163; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_maxnum_no_nans_param_1];1164; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_maxnum_no_nans_param_2];1165; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;1166; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;1167; CHECK-NEXT: ret;1168;1169; CHECK-FTZ-LABEL: fma_bf16_maxnum_no_nans(1170; CHECK-FTZ: {1171; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;1172; CHECK-FTZ-EMPTY:1173; CHECK-FTZ-NEXT: // %bb.0:1174; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_maxnum_no_nans_param_0];1175; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_maxnum_no_nans_param_1];1176; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_maxnum_no_nans_param_2];1177; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;1178; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;1179; CHECK-FTZ-NEXT: ret;1180;1181; CHECK-SM70-LABEL: fma_bf16_maxnum_no_nans(1182; CHECK-SM70: {1183; CHECK-SM70-NEXT: .reg .pred %p<3>;1184; CHECK-SM70-NEXT: .reg .b32 %r<21>;1185; CHECK-SM70-EMPTY:1186; CHECK-SM70-NEXT: // %bb.0:1187; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_maxnum_no_nans_param_2];1188; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;1189; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_maxnum_no_nans_param_1];1190; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;1191; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_maxnum_no_nans_param_0];1192; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;1193; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;1194; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;1195; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;1196; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;1197; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;1198; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;1199; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;1200; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;1201; CHECK-SM70-NEXT: max.f32 %r14, %r13, 0f00000000;1202; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;1203; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;1204; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;1205; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;1206; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;1207; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;1208; CHECK-SM70-NEXT: shr.u32 %r20, %r19, 16;1209; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r20;1210; CHECK-SM70-NEXT: ret;1211 %1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)1212 %2 = call nsz bfloat @llvm.maxnum.bf16(bfloat %1, bfloat 0.0)1213 ret bfloat %21214}1215 1216define <2 x half> @fma_f16x2_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {1217; CHECK-LABEL: fma_f16x2_no_nans(1218; CHECK: {1219; CHECK-NEXT: .reg .b32 %r<7>;1220; CHECK-EMPTY:1221; CHECK-NEXT: // %bb.0:1222; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];1223; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];1224; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];1225; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;1226; CHECK-NEXT: mov.b32 %r5, 0;1227; CHECK-NEXT: max.f16x2 %r6, %r4, %r5;1228; CHECK-NEXT: st.param.b32 [func_retval0], %r6;1229; CHECK-NEXT: ret;1230;1231; CHECK-FTZ-LABEL: fma_f16x2_no_nans(1232; CHECK-FTZ: {1233; CHECK-FTZ-NEXT: .reg .b32 %r<7>;1234; CHECK-FTZ-EMPTY:1235; CHECK-FTZ-NEXT: // %bb.0:1236; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];1237; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];1238; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];1239; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;1240; CHECK-FTZ-NEXT: mov.b32 %r5, 0;1241; CHECK-FTZ-NEXT: max.ftz.f16x2 %r6, %r4, %r5;1242; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r6;1243; CHECK-FTZ-NEXT: ret;1244;1245; CHECK-SM70-LABEL: fma_f16x2_no_nans(1246; CHECK-SM70: {1247; CHECK-SM70-NEXT: .reg .pred %p<3>;1248; CHECK-SM70-NEXT: .reg .b16 %rs<5>;1249; CHECK-SM70-NEXT: .reg .b32 %r<6>;1250; CHECK-SM70-EMPTY:1251; CHECK-SM70-NEXT: // %bb.0:1252; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];1253; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];1254; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];1255; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;1256; CHECK-SM70-NEXT: mov.b32 %r5, 0;1257; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;1258; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;1259; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;1260; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;1261; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};1262; CHECK-SM70-NEXT: ret;1263 %1 = call <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)1264 %2 = fcmp nsz ogt <2 x half> %1, <half 0.0, half 0.0>1265 %3 = select nsz <2 x i1> %2, <2 x half> %1, <2 x half> <half 0.0, half 0.0>1266 ret <2 x half> %31267}1268 1269; FMA relu shouldn't be selected if the FMA operation has multiple uses1270define <2 x half> @fma_f16x2_no_nans_multiple_uses_of_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) {1271; CHECK-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(1272; CHECK: {1273; CHECK-NEXT: .reg .b32 %r<8>;1274; CHECK-EMPTY:1275; CHECK-NEXT: // %bb.0:1276; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];1277; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];1278; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];1279; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;1280; CHECK-NEXT: mov.b32 %r5, 1191200512;1281; CHECK-NEXT: add.rn.f16x2 %r6, %r4, %r5;1282; CHECK-NEXT: add.rn.f16x2 %r7, %r6, %r4;1283; CHECK-NEXT: st.param.b32 [func_retval0], %r7;1284; CHECK-NEXT: ret;1285;1286; CHECK-FTZ-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(1287; CHECK-FTZ: {1288; CHECK-FTZ-NEXT: .reg .b32 %r<8>;1289; CHECK-FTZ-EMPTY:1290; CHECK-FTZ-NEXT: // %bb.0:1291; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];1292; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];1293; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];1294; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;1295; CHECK-FTZ-NEXT: mov.b32 %r5, 1191200512;1296; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r6, %r4, %r5;1297; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r7, %r6, %r4;1298; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r7;1299; CHECK-FTZ-NEXT: ret;1300;1301; CHECK-SM70-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(1302; CHECK-SM70: {1303; CHECK-SM70-NEXT: .reg .b32 %r<8>;1304; CHECK-SM70-EMPTY:1305; CHECK-SM70-NEXT: // %bb.0:1306; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];1307; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];1308; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];1309; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;1310; CHECK-SM70-NEXT: mov.b32 %r5, 1191200512;1311; CHECK-SM70-NEXT: add.rn.f16x2 %r6, %r4, %r5;1312; CHECK-SM70-NEXT: add.rn.f16x2 %r7, %r6, %r4;1313; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r7;1314; CHECK-SM70-NEXT: ret;1315 %1 = call nnan <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)1316 %2 = fcmp nsz ogt <2 x half> %1, <half 0.0, half 0.0>1317 %3 = select nsz <2 x i1> %2, <2 x half> %1, <2 x half> <half 0.0, half 0.0>1318 %4 = fadd <2 x half> %1, <half 7.0, half 7.0>1319 %5 = fadd <2 x half> %4, %11320 ret <2 x half> %51321}1322 1323define <2 x half> @fma_f16x2_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {1324; CHECK-LABEL: fma_f16x2_maxnum_no_nans(1325; CHECK: {1326; CHECK-NEXT: .reg .b32 %r<5>;1327; CHECK-EMPTY:1328; CHECK-NEXT: // %bb.0:1329; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];1330; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];1331; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];1332; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;1333; CHECK-NEXT: st.param.b32 [func_retval0], %r4;1334; CHECK-NEXT: ret;1335;1336; CHECK-FTZ-LABEL: fma_f16x2_maxnum_no_nans(1337; CHECK-FTZ: {1338; CHECK-FTZ-NEXT: .reg .b32 %r<5>;1339; CHECK-FTZ-EMPTY:1340; CHECK-FTZ-NEXT: // %bb.0:1341; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];1342; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];1343; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];1344; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;1345; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;1346; CHECK-FTZ-NEXT: ret;1347;1348; CHECK-SM70-LABEL: fma_f16x2_maxnum_no_nans(1349; CHECK-SM70: {1350; CHECK-SM70-NEXT: .reg .b16 %rs<5>;1351; CHECK-SM70-NEXT: .reg .b32 %r<10>;1352; CHECK-SM70-EMPTY:1353; CHECK-SM70-NEXT: // %bb.0:1354; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];1355; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];1356; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];1357; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;1358; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;1359; CHECK-SM70-NEXT: cvt.f32.f16 %r5, %rs2;1360; CHECK-SM70-NEXT: max.f32 %r6, %r5, 0f00000000;1361; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs3, %r6;1362; CHECK-SM70-NEXT: cvt.f32.f16 %r7, %rs1;1363; CHECK-SM70-NEXT: max.f32 %r8, %r7, 0f00000000;1364; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs4, %r8;1365; CHECK-SM70-NEXT: mov.b32 %r9, {%rs4, %rs3};1366; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;1367; CHECK-SM70-NEXT: ret;1368 %1 = call nnan <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)1369 %2 = call nsz <2 x half> @llvm.maxnum.f16x2(<2 x half> %1, <2 x half> <half 0.0, half 0.0>)1370 ret <2 x half> %21371}1372 1373define <2 x bfloat> @fma_bf16x2_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {1374; CHECK-LABEL: fma_bf16x2_no_nans(1375; CHECK: {1376; CHECK-NEXT: .reg .b32 %r<5>;1377; CHECK-EMPTY:1378; CHECK-NEXT: // %bb.0:1379; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];1380; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_param_1];1381; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_param_2];1382; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;1383; CHECK-NEXT: st.param.b32 [func_retval0], %r4;1384; CHECK-NEXT: ret;1385;1386; CHECK-FTZ-LABEL: fma_bf16x2_no_nans(1387; CHECK-FTZ: {1388; CHECK-FTZ-NEXT: .reg .b32 %r<5>;1389; CHECK-FTZ-EMPTY:1390; CHECK-FTZ-NEXT: // %bb.0:1391; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];1392; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_param_1];1393; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_param_2];1394; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;1395; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;1396; CHECK-FTZ-NEXT: ret;1397;1398; CHECK-SM70-LABEL: fma_bf16x2_no_nans(1399; CHECK-SM70: {1400; CHECK-SM70-NEXT: .reg .pred %p<5>;1401; CHECK-SM70-NEXT: .reg .b16 %rs<11>;1402; CHECK-SM70-NEXT: .reg .b32 %r<27>;1403; CHECK-SM70-EMPTY:1404; CHECK-SM70-NEXT: // %bb.0:1405; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_no_nans_param_0];1406; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_no_nans_param_2];1407; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs3;1408; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;1409; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_no_nans_param_1];1410; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs5;1411; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;1412; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;1413; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;1414; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;1415; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;1416; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;1417; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;1418; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;1419; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;1420; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;1421; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }1422; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs4;1423; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;1424; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs6;1425; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;1426; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;1427; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;1428; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;1429; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;1430; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;1431; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;1432; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;1433; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;1434; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;1435; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }1436; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;1437; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;1438; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;1439; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;1440; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;1441; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;1442; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};1443; CHECK-SM70-NEXT: ret;1444 %1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)1445 %2 = fcmp nsz ogt <2 x bfloat> %1, <bfloat 0.0, bfloat 0.0>1446 %3 = select nsz <2 x i1> %2, <2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>1447 ret <2 x bfloat> %31448}1449 1450; FMA_relu shouldn't be selected if the FMA operation has multiple uses1451define <2 x bfloat> @fma_bf16x2_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {1452; CHECK-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(1453; CHECK: {1454; CHECK-NEXT: .reg .b32 %r<9>;1455; CHECK-EMPTY:1456; CHECK-NEXT: // %bb.0:1457; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];1458; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];1459; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];1460; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;1461; CHECK-NEXT: mov.b32 %r5, 1065369472;1462; CHECK-NEXT: mov.b32 %r6, 1088438496;1463; CHECK-NEXT: fma.rn.bf16x2 %r7, %r4, %r5, %r6;1464; CHECK-NEXT: fma.rn.bf16x2 %r8, %r7, %r5, %r4;1465; CHECK-NEXT: st.param.b32 [func_retval0], %r8;1466; CHECK-NEXT: ret;1467;1468; CHECK-FTZ-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(1469; CHECK-FTZ: {1470; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;1471; CHECK-FTZ-NEXT: .reg .b32 %r<18>;1472; CHECK-FTZ-EMPTY:1473; CHECK-FTZ-NEXT: // %bb.0:1474; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];1475; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];1476; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];1477; CHECK-FTZ-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;1478; CHECK-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r4;1479; CHECK-FTZ-NEXT: cvt.u32.u16 %r5, %rs2;1480; CHECK-FTZ-NEXT: shl.b32 %r6, %r5, 16;1481; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r7, %r6, 0f40E00000;1482; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r7;1483; CHECK-FTZ-NEXT: cvt.u32.u16 %r8, %rs1;1484; CHECK-FTZ-NEXT: shl.b32 %r9, %r8, 16;1485; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r10, %r9, 0f40E00000;1486; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs4, %r10;1487; CHECK-FTZ-NEXT: cvt.u32.u16 %r11, %rs4;1488; CHECK-FTZ-NEXT: shl.b32 %r12, %r11, 16;1489; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r13, %r12, %r9;1490; CHECK-FTZ-NEXT: cvt.u32.u16 %r14, %rs3;1491; CHECK-FTZ-NEXT: shl.b32 %r15, %r14, 16;1492; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r16, %r15, %r6;1493; CHECK-FTZ-NEXT: cvt.rn.bf16x2.f32 %r17, %r16, %r13;1494; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r17;1495; CHECK-FTZ-NEXT: ret;1496;1497; CHECK-SM70-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(1498; CHECK-SM70: {1499; CHECK-SM70-NEXT: .reg .pred %p<7>;1500; CHECK-SM70-NEXT: .reg .b16 %rs<7>;1501; CHECK-SM70-NEXT: .reg .b32 %r<54>;1502; CHECK-SM70-EMPTY:1503; CHECK-SM70-NEXT: // %bb.0:1504; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];1505; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];1506; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs4;1507; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;1508; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];1509; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs6;1510; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;1511; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs2;1512; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;1513; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;1514; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;1515; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;1516; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;1517; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;1518; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;1519; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;1520; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs3;1521; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;1522; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs5;1523; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;1524; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs1;1525; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;1526; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;1527; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;1528; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;1529; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;1530; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;1531; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;1532; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;1533; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;1534; CHECK-SM70-NEXT: add.rn.f32 %r26, %r25, 0f40E00000;1535; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;1536; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;1537; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;1538; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;1539; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;1540; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;1541; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;1542; CHECK-SM70-NEXT: add.rn.f32 %r33, %r32, 0f40E00000;1543; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;1544; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;1545; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;1546; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;1547; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;1548; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;1549; CHECK-SM70-NEXT: and.b32 %r39, %r38, -65536;1550; CHECK-SM70-NEXT: add.rn.f32 %r40, %r39, %r32;1551; CHECK-SM70-NEXT: bfe.u32 %r41, %r40, 16, 1;1552; CHECK-SM70-NEXT: add.s32 %r42, %r41, %r40;1553; CHECK-SM70-NEXT: add.s32 %r43, %r42, 32767;1554; CHECK-SM70-NEXT: setp.nan.f32 %p5, %r40, %r40;1555; CHECK-SM70-NEXT: or.b32 %r44, %r40, 4194304;1556; CHECK-SM70-NEXT: selp.b32 %r45, %r44, %r43, %p5;1557; CHECK-SM70-NEXT: and.b32 %r46, %r31, -65536;1558; CHECK-SM70-NEXT: add.rn.f32 %r47, %r46, %r25;1559; CHECK-SM70-NEXT: bfe.u32 %r48, %r47, 16, 1;1560; CHECK-SM70-NEXT: add.s32 %r49, %r48, %r47;1561; CHECK-SM70-NEXT: add.s32 %r50, %r49, 32767;1562; CHECK-SM70-NEXT: setp.nan.f32 %p6, %r47, %r47;1563; CHECK-SM70-NEXT: or.b32 %r51, %r47, 4194304;1564; CHECK-SM70-NEXT: selp.b32 %r52, %r51, %r50, %p6;1565; CHECK-SM70-NEXT: prmt.b32 %r53, %r52, %r45, 0x7632U;1566; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r53;1567; CHECK-SM70-NEXT: ret;1568 %1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)1569 %2 = fcmp nsz ogt <2 x bfloat> %1, <bfloat 0.0, bfloat 0.0>1570 %3 = select nsz <2 x i1> %2, <2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>1571 %4 = fadd <2 x bfloat> %1, <bfloat 7.0, bfloat 7.0>1572 %5 = fadd <2 x bfloat> %4, %11573 ret <2 x bfloat> %51574}1575 1576define <2 x bfloat> @fma_bf16x2_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {1577; CHECK-LABEL: fma_bf16x2_maxnum_no_nans(1578; CHECK: {1579; CHECK-NEXT: .reg .b32 %r<5>;1580; CHECK-EMPTY:1581; CHECK-NEXT: // %bb.0:1582; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];1583; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_maxnum_no_nans_param_1];1584; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_maxnum_no_nans_param_2];1585; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;1586; CHECK-NEXT: st.param.b32 [func_retval0], %r4;1587; CHECK-NEXT: ret;1588;1589; CHECK-FTZ-LABEL: fma_bf16x2_maxnum_no_nans(1590; CHECK-FTZ: {1591; CHECK-FTZ-NEXT: .reg .b32 %r<5>;1592; CHECK-FTZ-EMPTY:1593; CHECK-FTZ-NEXT: // %bb.0:1594; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];1595; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_maxnum_no_nans_param_1];1596; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_maxnum_no_nans_param_2];1597; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;1598; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;1599; CHECK-FTZ-NEXT: ret;1600;1601; CHECK-SM70-LABEL: fma_bf16x2_maxnum_no_nans(1602; CHECK-SM70: {1603; CHECK-SM70-NEXT: .reg .pred %p<5>;1604; CHECK-SM70-NEXT: .reg .b16 %rs<7>;1605; CHECK-SM70-NEXT: .reg .b32 %r<40>;1606; CHECK-SM70-EMPTY:1607; CHECK-SM70-NEXT: // %bb.0:1608; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_maxnum_no_nans_param_0];1609; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_maxnum_no_nans_param_2];1610; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs3;1611; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;1612; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_maxnum_no_nans_param_1];1613; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs5;1614; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;1615; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;1616; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;1617; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;1618; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;1619; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;1620; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;1621; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;1622; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;1623; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;1624; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs4;1625; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;1626; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs6;1627; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;1628; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;1629; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;1630; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;1631; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;1632; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;1633; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;1634; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;1635; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;1636; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;1637; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;1638; CHECK-SM70-NEXT: max.f32 %r26, %r25, 0f00000000;1639; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;1640; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;1641; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;1642; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;1643; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;1644; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;1645; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;1646; CHECK-SM70-NEXT: max.f32 %r33, %r32, 0f00000000;1647; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;1648; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;1649; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;1650; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;1651; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;1652; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;1653; CHECK-SM70-NEXT: prmt.b32 %r39, %r38, %r31, 0x7632U;1654; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r39;1655; CHECK-SM70-NEXT: ret;1656 %1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)1657 %2 = call nsz <2 x bfloat> @llvm.maxnum.bf16x2(<2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>)1658 ret <2 x bfloat> %21659}1660