1113 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 -fp-contract=fast | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -fp-contract=fast | %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 -fp-contract=fast | FileCheck %s --check-prefixes=CHECK-FTZ7; RUN: %if ptxas %{ llc < %s -denormal-fp-math-f32=preserve-sign -fp-contract=fast -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 -fp-contract=fast | FileCheck %s --check-prefixes=CHECK-SM7011 12define half @fma_f16_expanded_no_nans(half %a, half %b, half %c) #0 {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 half %a, %b53 %2 = fadd half %1, %c54 %3 = fcmp ogt half %2, 0.055 %4 = select 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) #0 {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.f16 %rs8, %rs4, %rs7;74; CHECK-NEXT: add.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.ftz.f16 %rs8, %rs4, %rs7;91; CHECK-FTZ-NEXT: add.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.f16 %rs8, %rs4, %rs7;110; CHECK-SM70-NEXT: add.f16 %rs9, %rs6, %rs8;111; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs9;112; CHECK-SM70-NEXT: ret;113 %1 = fmul half %a, %b114 %2 = fadd half %1, %c115 %3 = fcmp ogt half %2, 0.0116 %4 = select 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_unsafe_with_nans(half %a, half %b, half %c) {123; CHECK-LABEL: fma_f16_expanded_unsafe_with_nans(124; CHECK: {125; CHECK-NEXT: .reg .b16 %rs<7>;126; CHECK-EMPTY:127; CHECK-NEXT: // %bb.0:128; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_unsafe_with_nans_param_0];129; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_unsafe_with_nans_param_1];130; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_unsafe_with_nans_param_2];131; CHECK-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;132; CHECK-NEXT: mov.b16 %rs5, 0x0000;133; CHECK-NEXT: max.f16 %rs6, %rs4, %rs5;134; CHECK-NEXT: st.param.b16 [func_retval0], %rs6;135; CHECK-NEXT: ret;136;137; CHECK-FTZ-LABEL: fma_f16_expanded_unsafe_with_nans(138; CHECK-FTZ: {139; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;140; CHECK-FTZ-EMPTY:141; CHECK-FTZ-NEXT: // %bb.0:142; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_unsafe_with_nans_param_0];143; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_unsafe_with_nans_param_1];144; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_unsafe_with_nans_param_2];145; CHECK-FTZ-NEXT: fma.rn.ftz.f16 %rs4, %rs1, %rs2, %rs3;146; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x0000;147; CHECK-FTZ-NEXT: max.ftz.f16 %rs6, %rs4, %rs5;148; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs6;149; CHECK-FTZ-NEXT: ret;150;151; CHECK-SM70-LABEL: fma_f16_expanded_unsafe_with_nans(152; CHECK-SM70: {153; CHECK-SM70-NEXT: .reg .pred %p<2>;154; CHECK-SM70-NEXT: .reg .b16 %rs<7>;155; CHECK-SM70-EMPTY:156; CHECK-SM70-NEXT: // %bb.0:157; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_unsafe_with_nans_param_0];158; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_unsafe_with_nans_param_1];159; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_unsafe_with_nans_param_2];160; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;161; CHECK-SM70-NEXT: mov.b16 %rs5, 0x0000;162; CHECK-SM70-NEXT: setp.gt.f16 %p1, %rs4, %rs5;163; CHECK-SM70-NEXT: selp.b16 %rs6, %rs4, 0x0000, %p1;164; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs6;165; CHECK-SM70-NEXT: ret;166 %1 = fmul half %a, %b167 %2 = fadd half %1, %c168 %3 = fcmp ogt half %2, 0.0169 %4 = select i1 %3, half %2, half 0.0170 ret half %4171}172 173define half @fma_f16_expanded_maxnum_no_nans(half %a, half %b, half %c) #0 {174; CHECK-LABEL: fma_f16_expanded_maxnum_no_nans(175; CHECK: {176; CHECK-NEXT: .reg .b16 %rs<5>;177; CHECK-EMPTY:178; CHECK-NEXT: // %bb.0:179; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];180; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];181; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];182; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;183; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;184; CHECK-NEXT: ret;185;186; CHECK-FTZ-LABEL: fma_f16_expanded_maxnum_no_nans(187; CHECK-FTZ: {188; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;189; CHECK-FTZ-EMPTY:190; CHECK-FTZ-NEXT: // %bb.0:191; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];192; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];193; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];194; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;195; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;196; CHECK-FTZ-NEXT: ret;197;198; CHECK-SM70-LABEL: fma_f16_expanded_maxnum_no_nans(199; CHECK-SM70: {200; CHECK-SM70-NEXT: .reg .b16 %rs<6>;201; CHECK-SM70-NEXT: .reg .b32 %r<3>;202; CHECK-SM70-EMPTY:203; CHECK-SM70-NEXT: // %bb.0:204; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];205; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];206; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];207; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;208; CHECK-SM70-NEXT: cvt.f32.f16 %r1, %rs4;209; CHECK-SM70-NEXT: max.f32 %r2, %r1, 0f00000000;210; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs5, %r2;211; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs5;212; CHECK-SM70-NEXT: ret;213 %1 = fmul half %a, %b214 %2 = fadd half %1, %c215 %3 = call half @llvm.maxnum.f16(half %2, half 0.0)216 ret half %3217}218 219define bfloat @fma_bf16_expanded_unsafe_with_nans(bfloat %a, bfloat %b, bfloat %c) {220; CHECK-LABEL: fma_bf16_expanded_unsafe_with_nans(221; CHECK: {222; CHECK-NEXT: .reg .b16 %rs<7>;223; CHECK-EMPTY:224; CHECK-NEXT: // %bb.0:225; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_unsafe_with_nans_param_0];226; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_unsafe_with_nans_param_1];227; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_unsafe_with_nans_param_2];228; CHECK-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;229; CHECK-NEXT: mov.b16 %rs5, 0x0000;230; CHECK-NEXT: max.bf16 %rs6, %rs4, %rs5;231; CHECK-NEXT: st.param.b16 [func_retval0], %rs6;232; CHECK-NEXT: ret;233;234; CHECK-FTZ-LABEL: fma_bf16_expanded_unsafe_with_nans(235; CHECK-FTZ: {236; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;237; CHECK-FTZ-EMPTY:238; CHECK-FTZ-NEXT: // %bb.0:239; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_unsafe_with_nans_param_0];240; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_unsafe_with_nans_param_1];241; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_unsafe_with_nans_param_2];242; CHECK-FTZ-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;243; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x0000;244; CHECK-FTZ-NEXT: max.bf16 %rs6, %rs4, %rs5;245; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs6;246; CHECK-FTZ-NEXT: ret;247;248; CHECK-SM70-LABEL: fma_bf16_expanded_unsafe_with_nans(249; CHECK-SM70: {250; CHECK-SM70-NEXT: .reg .pred %p<3>;251; CHECK-SM70-NEXT: .reg .b16 %rs<3>;252; CHECK-SM70-NEXT: .reg .b32 %r<14>;253; CHECK-SM70-EMPTY:254; CHECK-SM70-NEXT: // %bb.0:255; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_unsafe_with_nans_param_2];256; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;257; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_unsafe_with_nans_param_1];258; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;259; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_unsafe_with_nans_param_0];260; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;261; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;262; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;263; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;264; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;265; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;266; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;267; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;268; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }269; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;270; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;271; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;272; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs2;273; CHECK-SM70-NEXT: ret;274 %1 = fmul bfloat %a, %b275 %2 = fadd bfloat %1, %c276 %3 = fcmp ogt bfloat %2, 0.0277 %4 = select i1 %3, bfloat %2, bfloat 0.0278 ret bfloat %4279}280 281define bfloat @fma_bf16_expanded_no_nans(bfloat %a, bfloat %b, bfloat %c) #0 {282; CHECK-LABEL: fma_bf16_expanded_no_nans(283; CHECK: {284; CHECK-NEXT: .reg .b16 %rs<5>;285; CHECK-EMPTY:286; CHECK-NEXT: // %bb.0:287; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_param_0];288; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_param_1];289; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_param_2];290; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;291; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;292; CHECK-NEXT: ret;293;294; CHECK-FTZ-LABEL: fma_bf16_expanded_no_nans(295; CHECK-FTZ: {296; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;297; CHECK-FTZ-EMPTY:298; CHECK-FTZ-NEXT: // %bb.0:299; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_param_0];300; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_param_1];301; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_param_2];302; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;303; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;304; CHECK-FTZ-NEXT: ret;305;306; CHECK-SM70-LABEL: fma_bf16_expanded_no_nans(307; CHECK-SM70: {308; CHECK-SM70-NEXT: .reg .pred %p<3>;309; CHECK-SM70-NEXT: .reg .b16 %rs<3>;310; CHECK-SM70-NEXT: .reg .b32 %r<14>;311; CHECK-SM70-EMPTY:312; CHECK-SM70-NEXT: // %bb.0:313; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_no_nans_param_2];314; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;315; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_no_nans_param_1];316; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;317; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_no_nans_param_0];318; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;319; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;320; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;321; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;322; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;323; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;324; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;325; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;326; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }327; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;328; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;329; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;330; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs2;331; CHECK-SM70-NEXT: ret;332 %1 = fmul bfloat %a, %b333 %2 = fadd bfloat %1, %c334 %3 = fcmp ogt bfloat %2, 0.0335 %4 = select i1 %3, bfloat %2, bfloat 0.0336 ret bfloat %4337}338 339; FMA relu shouldn't be selected if the FMA operation has multiple uses340define bfloat @fma_bf16_expanded_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloat %c) #0 {341; CHECK-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(342; CHECK: {343; CHECK-NEXT: .reg .b16 %rs<11>;344; CHECK-EMPTY:345; CHECK-NEXT: // %bb.0:346; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];347; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];348; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];349; CHECK-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;350; CHECK-NEXT: mov.b16 %rs5, 0x0000;351; CHECK-NEXT: max.bf16 %rs6, %rs4, %rs5;352; CHECK-NEXT: mov.b16 %rs7, 0x3F80;353; CHECK-NEXT: mov.b16 %rs8, 0x40E0;354; CHECK-NEXT: fma.rn.bf16 %rs9, %rs4, %rs7, %rs8;355; CHECK-NEXT: fma.rn.bf16 %rs10, %rs6, %rs7, %rs9;356; CHECK-NEXT: st.param.b16 [func_retval0], %rs10;357; CHECK-NEXT: ret;358;359; CHECK-FTZ-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(360; CHECK-FTZ: {361; CHECK-FTZ-NEXT: .reg .b16 %rs<9>;362; CHECK-FTZ-NEXT: .reg .b32 %r<9>;363; CHECK-FTZ-EMPTY:364; CHECK-FTZ-NEXT: // %bb.0:365; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];366; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];367; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];368; CHECK-FTZ-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;369; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x0000;370; CHECK-FTZ-NEXT: max.bf16 %rs6, %rs4, %rs5;371; CHECK-FTZ-NEXT: cvt.u32.u16 %r1, %rs4;372; CHECK-FTZ-NEXT: shl.b32 %r2, %r1, 16;373; CHECK-FTZ-NEXT: add.ftz.f32 %r3, %r2, 0f40E00000;374; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs7, %r3;375; CHECK-FTZ-NEXT: cvt.u32.u16 %r4, %rs6;376; CHECK-FTZ-NEXT: shl.b32 %r5, %r4, 16;377; CHECK-FTZ-NEXT: cvt.u32.u16 %r6, %rs7;378; CHECK-FTZ-NEXT: shl.b32 %r7, %r6, 16;379; CHECK-FTZ-NEXT: add.ftz.f32 %r8, %r5, %r7;380; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs8, %r8;381; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs8;382; CHECK-FTZ-NEXT: ret;383;384; CHECK-SM70-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(385; CHECK-SM70: {386; CHECK-SM70-NEXT: .reg .pred %p<5>;387; CHECK-SM70-NEXT: .reg .b16 %rs<3>;388; CHECK-SM70-NEXT: .reg .b32 %r<30>;389; CHECK-SM70-EMPTY:390; CHECK-SM70-NEXT: // %bb.0:391; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];392; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;393; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];394; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;395; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];396; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;397; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;398; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;399; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;400; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;401; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;402; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;403; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;404; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }405; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;406; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;407; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;408; CHECK-SM70-NEXT: add.f32 %r14, %r13, 0f40E00000;409; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;410; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;411; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;412; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r14, %r14;413; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;414; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p3;415; CHECK-SM70-NEXT: cvt.u32.u16 %r20, %rs2;416; CHECK-SM70-NEXT: shl.b32 %r21, %r20, 16;417; CHECK-SM70-NEXT: and.b32 %r22, %r19, -65536;418; CHECK-SM70-NEXT: add.f32 %r23, %r21, %r22;419; CHECK-SM70-NEXT: bfe.u32 %r24, %r23, 16, 1;420; CHECK-SM70-NEXT: add.s32 %r25, %r24, %r23;421; CHECK-SM70-NEXT: add.s32 %r26, %r25, 32767;422; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r23, %r23;423; CHECK-SM70-NEXT: or.b32 %r27, %r23, 4194304;424; CHECK-SM70-NEXT: selp.b32 %r28, %r27, %r26, %p4;425; CHECK-SM70-NEXT: shr.u32 %r29, %r28, 16;426; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r29;427; CHECK-SM70-NEXT: ret;428 %1 = fmul bfloat %a, %b429 %2 = fadd bfloat %1, %c430 %3 = fcmp ogt bfloat %2, 0.0431 %4 = select i1 %3, bfloat %2, bfloat 0.0432 %5 = fadd bfloat %2, 7.0433 %6 = fadd bfloat %4, %5434 ret bfloat %6435}436 437define bfloat @fma_bf16_expanded_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) #0 {438; CHECK-LABEL: fma_bf16_expanded_maxnum_no_nans(439; CHECK: {440; CHECK-NEXT: .reg .b16 %rs<5>;441; CHECK-EMPTY:442; CHECK-NEXT: // %bb.0:443; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_maxnum_no_nans_param_0];444; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_maxnum_no_nans_param_1];445; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_maxnum_no_nans_param_2];446; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;447; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;448; CHECK-NEXT: ret;449;450; CHECK-FTZ-LABEL: fma_bf16_expanded_maxnum_no_nans(451; CHECK-FTZ: {452; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;453; CHECK-FTZ-EMPTY:454; CHECK-FTZ-NEXT: // %bb.0:455; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_maxnum_no_nans_param_0];456; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_maxnum_no_nans_param_1];457; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_maxnum_no_nans_param_2];458; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;459; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;460; CHECK-FTZ-NEXT: ret;461;462; CHECK-SM70-LABEL: fma_bf16_expanded_maxnum_no_nans(463; CHECK-SM70: {464; CHECK-SM70-NEXT: .reg .pred %p<3>;465; CHECK-SM70-NEXT: .reg .b32 %r<21>;466; CHECK-SM70-EMPTY:467; CHECK-SM70-NEXT: // %bb.0:468; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_maxnum_no_nans_param_2];469; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;470; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_maxnum_no_nans_param_1];471; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;472; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_maxnum_no_nans_param_0];473; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;474; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;475; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;476; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;477; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;478; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;479; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;480; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;481; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;482; CHECK-SM70-NEXT: max.f32 %r14, %r13, 0f00000000;483; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;484; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;485; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;486; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;487; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;488; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;489; CHECK-SM70-NEXT: shr.u32 %r20, %r19, 16;490; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r20;491; CHECK-SM70-NEXT: ret;492 %1 = fmul bfloat %a, %b493 %2 = fadd bfloat %1, %c494 %3 = call bfloat @llvm.maxnum.bf16(bfloat %2, bfloat 0.0)495 ret bfloat %3496}497 498define <2 x half> @fma_f16x2_expanded_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 {499; CHECK-LABEL: fma_f16x2_expanded_no_nans(500; CHECK: {501; CHECK-NEXT: .reg .b32 %r<5>;502; CHECK-EMPTY:503; CHECK-NEXT: // %bb.0:504; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];505; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];506; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];507; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;508; CHECK-NEXT: st.param.b32 [func_retval0], %r4;509; CHECK-NEXT: ret;510;511; CHECK-FTZ-LABEL: fma_f16x2_expanded_no_nans(512; CHECK-FTZ: {513; CHECK-FTZ-NEXT: .reg .b32 %r<5>;514; CHECK-FTZ-EMPTY:515; CHECK-FTZ-NEXT: // %bb.0:516; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];517; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];518; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];519; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;520; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;521; CHECK-FTZ-NEXT: ret;522;523; CHECK-SM70-LABEL: fma_f16x2_expanded_no_nans(524; CHECK-SM70: {525; CHECK-SM70-NEXT: .reg .pred %p<3>;526; CHECK-SM70-NEXT: .reg .b16 %rs<5>;527; CHECK-SM70-NEXT: .reg .b32 %r<6>;528; CHECK-SM70-EMPTY:529; CHECK-SM70-NEXT: // %bb.0:530; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];531; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];532; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];533; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;534; CHECK-SM70-NEXT: mov.b32 %r5, 0;535; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;536; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;537; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;538; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;539; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};540; CHECK-SM70-NEXT: ret;541 %1 = fmul <2 x half> %a, %b542 %2 = fadd <2 x half> %1, %c543 %3 = fcmp ogt <2 x half> %2, <half 0.0, half 0.0>544 %4 = select <2 x i1> %3, <2 x half> %2, <2 x half> <half 0.0, half 0.0>545 ret <2 x half> %4546}547 548; FMA relu shouldn't be selected if the FMA operation has multiple uses549define <2 x half> @fma_f16x2_expanded_no_nans_multiple_uses_of_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 {550; CHECK-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(551; CHECK: {552; CHECK-NEXT: .reg .b32 %r<10>;553; CHECK-EMPTY:554; CHECK-NEXT: // %bb.0:555; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];556; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];557; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];558; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;559; CHECK-NEXT: mov.b32 %r5, 0;560; CHECK-NEXT: max.f16x2 %r6, %r4, %r5;561; CHECK-NEXT: mov.b32 %r7, 1191200512;562; CHECK-NEXT: add.f16x2 %r8, %r4, %r7;563; CHECK-NEXT: add.f16x2 %r9, %r6, %r8;564; CHECK-NEXT: st.param.b32 [func_retval0], %r9;565; CHECK-NEXT: ret;566;567; CHECK-FTZ-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(568; CHECK-FTZ: {569; CHECK-FTZ-NEXT: .reg .b32 %r<10>;570; CHECK-FTZ-EMPTY:571; CHECK-FTZ-NEXT: // %bb.0:572; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];573; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];574; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];575; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;576; CHECK-FTZ-NEXT: mov.b32 %r5, 0;577; CHECK-FTZ-NEXT: max.ftz.f16x2 %r6, %r4, %r5;578; CHECK-FTZ-NEXT: mov.b32 %r7, 1191200512;579; CHECK-FTZ-NEXT: add.ftz.f16x2 %r8, %r4, %r7;580; CHECK-FTZ-NEXT: add.ftz.f16x2 %r9, %r6, %r8;581; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r9;582; CHECK-FTZ-NEXT: ret;583;584; CHECK-SM70-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(585; CHECK-SM70: {586; CHECK-SM70-NEXT: .reg .pred %p<3>;587; CHECK-SM70-NEXT: .reg .b16 %rs<5>;588; CHECK-SM70-NEXT: .reg .b32 %r<10>;589; CHECK-SM70-EMPTY:590; CHECK-SM70-NEXT: // %bb.0:591; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];592; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];593; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];594; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;595; CHECK-SM70-NEXT: mov.b32 %r5, 0;596; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;597; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;598; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;599; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;600; CHECK-SM70-NEXT: mov.b32 %r6, {%rs4, %rs3};601; CHECK-SM70-NEXT: mov.b32 %r7, 1191200512;602; CHECK-SM70-NEXT: add.f16x2 %r8, %r4, %r7;603; CHECK-SM70-NEXT: add.f16x2 %r9, %r6, %r8;604; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;605; CHECK-SM70-NEXT: ret;606 %1 = fmul <2 x half> %a, %b607 %2 = fadd <2 x half> %1, %c608 %3 = fcmp ogt <2 x half> %2, <half 0.0, half 0.0>609 %4 = select <2 x i1> %3, <2 x half> %2, <2 x half> <half 0.0, half 0.0>610 %5 = fadd <2 x half> %2, <half 7.0, half 7.0>611 %6 = fadd <2 x half> %4, %5612 ret <2 x half> %6613}614 615define <2 x half> @fma_f16x2_expanded_unsafe_with_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {616; CHECK-LABEL: fma_f16x2_expanded_unsafe_with_nans(617; CHECK: {618; CHECK-NEXT: .reg .b32 %r<7>;619; CHECK-EMPTY:620; CHECK-NEXT: // %bb.0:621; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_unsafe_with_nans_param_0];622; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_unsafe_with_nans_param_1];623; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_unsafe_with_nans_param_2];624; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;625; CHECK-NEXT: mov.b32 %r5, 0;626; CHECK-NEXT: max.f16x2 %r6, %r4, %r5;627; CHECK-NEXT: st.param.b32 [func_retval0], %r6;628; CHECK-NEXT: ret;629;630; CHECK-FTZ-LABEL: fma_f16x2_expanded_unsafe_with_nans(631; CHECK-FTZ: {632; CHECK-FTZ-NEXT: .reg .b32 %r<7>;633; CHECK-FTZ-EMPTY:634; CHECK-FTZ-NEXT: // %bb.0:635; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_unsafe_with_nans_param_0];636; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_unsafe_with_nans_param_1];637; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_unsafe_with_nans_param_2];638; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;639; CHECK-FTZ-NEXT: mov.b32 %r5, 0;640; CHECK-FTZ-NEXT: max.ftz.f16x2 %r6, %r4, %r5;641; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r6;642; CHECK-FTZ-NEXT: ret;643;644; CHECK-SM70-LABEL: fma_f16x2_expanded_unsafe_with_nans(645; CHECK-SM70: {646; CHECK-SM70-NEXT: .reg .pred %p<3>;647; CHECK-SM70-NEXT: .reg .b16 %rs<5>;648; CHECK-SM70-NEXT: .reg .b32 %r<6>;649; CHECK-SM70-EMPTY:650; CHECK-SM70-NEXT: // %bb.0:651; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_unsafe_with_nans_param_0];652; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_unsafe_with_nans_param_1];653; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_unsafe_with_nans_param_2];654; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;655; CHECK-SM70-NEXT: mov.b32 %r5, 0;656; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;657; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;658; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;659; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;660; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};661; CHECK-SM70-NEXT: ret;662 %1 = fmul <2 x half> %a, %b663 %2 = fadd <2 x half> %1, %c664 %3 = fcmp ogt <2 x half> %2, <half 0.0, half 0.0>665 %4 = select <2 x i1> %3, <2 x half> %2, <2 x half> <half 0.0, half 0.0>666 ret <2 x half> %4667}668 669define <2 x half> @fma_f16x2_expanded_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 {670; CHECK-LABEL: fma_f16x2_expanded_maxnum_no_nans(671; CHECK: {672; CHECK-NEXT: .reg .b32 %r<5>;673; CHECK-EMPTY:674; CHECK-NEXT: // %bb.0:675; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];676; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];677; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];678; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;679; CHECK-NEXT: st.param.b32 [func_retval0], %r4;680; CHECK-NEXT: ret;681;682; CHECK-FTZ-LABEL: fma_f16x2_expanded_maxnum_no_nans(683; CHECK-FTZ: {684; CHECK-FTZ-NEXT: .reg .b32 %r<5>;685; CHECK-FTZ-EMPTY:686; CHECK-FTZ-NEXT: // %bb.0:687; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];688; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];689; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];690; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;691; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;692; CHECK-FTZ-NEXT: ret;693;694; CHECK-SM70-LABEL: fma_f16x2_expanded_maxnum_no_nans(695; CHECK-SM70: {696; CHECK-SM70-NEXT: .reg .b16 %rs<5>;697; CHECK-SM70-NEXT: .reg .b32 %r<10>;698; CHECK-SM70-EMPTY:699; CHECK-SM70-NEXT: // %bb.0:700; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];701; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];702; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];703; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;704; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;705; CHECK-SM70-NEXT: cvt.f32.f16 %r5, %rs2;706; CHECK-SM70-NEXT: max.f32 %r6, %r5, 0f00000000;707; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs3, %r6;708; CHECK-SM70-NEXT: cvt.f32.f16 %r7, %rs1;709; CHECK-SM70-NEXT: max.f32 %r8, %r7, 0f00000000;710; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs4, %r8;711; CHECK-SM70-NEXT: mov.b32 %r9, {%rs4, %rs3};712; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;713; CHECK-SM70-NEXT: ret;714 %1 = fmul <2 x half> %a, %b715 %2 = fadd <2 x half> %1, %c716 %3 = call <2 x half> @llvm.maxnum.f16x2(<2 x half> %2, <2 x half> <half 0.0, half 0.0>)717 ret <2 x half> %3718}719 720define <2 x bfloat> @fma_bf16x2_expanded_unsafe_with_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {721; CHECK-LABEL: fma_bf16x2_expanded_unsafe_with_nans(722; CHECK: {723; CHECK-NEXT: .reg .b32 %r<7>;724; CHECK-EMPTY:725; CHECK-NEXT: // %bb.0:726; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_unsafe_with_nans_param_0];727; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_unsafe_with_nans_param_1];728; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_unsafe_with_nans_param_2];729; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;730; CHECK-NEXT: mov.b32 %r5, 0;731; CHECK-NEXT: max.bf16x2 %r6, %r4, %r5;732; CHECK-NEXT: st.param.b32 [func_retval0], %r6;733; CHECK-NEXT: ret;734;735; CHECK-FTZ-LABEL: fma_bf16x2_expanded_unsafe_with_nans(736; CHECK-FTZ: {737; CHECK-FTZ-NEXT: .reg .b32 %r<7>;738; CHECK-FTZ-EMPTY:739; CHECK-FTZ-NEXT: // %bb.0:740; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_unsafe_with_nans_param_0];741; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_unsafe_with_nans_param_1];742; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_unsafe_with_nans_param_2];743; CHECK-FTZ-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;744; CHECK-FTZ-NEXT: mov.b32 %r5, 0;745; CHECK-FTZ-NEXT: max.bf16x2 %r6, %r4, %r5;746; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r6;747; CHECK-FTZ-NEXT: ret;748;749; CHECK-SM70-LABEL: fma_bf16x2_expanded_unsafe_with_nans(750; CHECK-SM70: {751; CHECK-SM70-NEXT: .reg .pred %p<5>;752; CHECK-SM70-NEXT: .reg .b16 %rs<11>;753; CHECK-SM70-NEXT: .reg .b32 %r<27>;754; CHECK-SM70-EMPTY:755; CHECK-SM70-NEXT: // %bb.0:756; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_unsafe_with_nans_param_0];757; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_unsafe_with_nans_param_1];758; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_unsafe_with_nans_param_2];759; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs5;760; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;761; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs3;762; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;763; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;764; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;765; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;766; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;767; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;768; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;769; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;770; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;771; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;772; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }773; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs6;774; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;775; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs4;776; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;777; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;778; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;779; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;780; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;781; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;782; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;783; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;784; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;785; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;786; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }787; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;788; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;789; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;790; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;791; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;792; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;793; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};794; CHECK-SM70-NEXT: ret;795 %1 = fmul <2 x bfloat> %a, %b796 %2 = fadd <2 x bfloat> %1, %c797 %3 = fcmp ogt <2 x bfloat> %2, <bfloat 0.0, bfloat 0.0>798 %4 = select <2 x i1> %3, <2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>799 ret <2 x bfloat> %4800}801 802define <2 x bfloat> @fma_bf16x2_expanded_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) #0 {803; CHECK-LABEL: fma_bf16x2_expanded_no_nans(804; CHECK: {805; CHECK-NEXT: .reg .b32 %r<5>;806; CHECK-EMPTY:807; CHECK-NEXT: // %bb.0:808; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];809; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_param_1];810; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_param_2];811; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;812; CHECK-NEXT: st.param.b32 [func_retval0], %r4;813; CHECK-NEXT: ret;814;815; CHECK-FTZ-LABEL: fma_bf16x2_expanded_no_nans(816; CHECK-FTZ: {817; CHECK-FTZ-NEXT: .reg .b32 %r<5>;818; CHECK-FTZ-EMPTY:819; CHECK-FTZ-NEXT: // %bb.0:820; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];821; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_param_1];822; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_param_2];823; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;824; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;825; CHECK-FTZ-NEXT: ret;826;827; CHECK-SM70-LABEL: fma_bf16x2_expanded_no_nans(828; CHECK-SM70: {829; CHECK-SM70-NEXT: .reg .pred %p<5>;830; CHECK-SM70-NEXT: .reg .b16 %rs<11>;831; CHECK-SM70-NEXT: .reg .b32 %r<27>;832; CHECK-SM70-EMPTY:833; CHECK-SM70-NEXT: // %bb.0:834; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_no_nans_param_0];835; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_no_nans_param_1];836; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_no_nans_param_2];837; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs5;838; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;839; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs3;840; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;841; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;842; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;843; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;844; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;845; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;846; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;847; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;848; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;849; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;850; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }851; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs6;852; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;853; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs4;854; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;855; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;856; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;857; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;858; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;859; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;860; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;861; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;862; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;863; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;864; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }865; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;866; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;867; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;868; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;869; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;870; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;871; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};872; CHECK-SM70-NEXT: ret;873 %1 = fmul <2 x bfloat> %a, %b874 %2 = fadd <2 x bfloat> %1, %c875 %3 = fcmp ogt <2 x bfloat> %2, <bfloat 0.0, bfloat 0.0>876 %4 = select <2 x i1> %3, <2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>877 ret <2 x bfloat> %4878}879 880; FMA relu shouldn't be selected if the FMA operation has multiple uses881define <2 x bfloat> @fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) #0 {882; CHECK-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(883; CHECK: {884; CHECK-NEXT: .reg .b32 %r<11>;885; CHECK-EMPTY:886; CHECK-NEXT: // %bb.0:887; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];888; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];889; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];890; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;891; CHECK-NEXT: mov.b32 %r5, 0;892; CHECK-NEXT: max.bf16x2 %r6, %r4, %r5;893; CHECK-NEXT: mov.b32 %r7, 1065369472;894; CHECK-NEXT: mov.b32 %r8, 1088438496;895; CHECK-NEXT: fma.rn.bf16x2 %r9, %r4, %r7, %r8;896; CHECK-NEXT: fma.rn.bf16x2 %r10, %r6, %r7, %r9;897; CHECK-NEXT: st.param.b32 [func_retval0], %r10;898; CHECK-NEXT: ret;899;900; CHECK-FTZ-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(901; CHECK-FTZ: {902; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;903; CHECK-FTZ-NEXT: .reg .b32 %r<24>;904; CHECK-FTZ-EMPTY:905; CHECK-FTZ-NEXT: // %bb.0:906; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];907; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];908; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];909; CHECK-FTZ-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;910; CHECK-FTZ-NEXT: mov.b32 %r5, 0;911; CHECK-FTZ-NEXT: max.bf16x2 %r6, %r4, %r5;912; CHECK-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r4;913; CHECK-FTZ-NEXT: cvt.u32.u16 %r7, %rs2;914; CHECK-FTZ-NEXT: shl.b32 %r8, %r7, 16;915; CHECK-FTZ-NEXT: add.ftz.f32 %r9, %r8, 0f40E00000;916; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r9;917; CHECK-FTZ-NEXT: cvt.u32.u16 %r10, %rs1;918; CHECK-FTZ-NEXT: shl.b32 %r11, %r10, 16;919; CHECK-FTZ-NEXT: add.ftz.f32 %r12, %r11, 0f40E00000;920; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs4, %r12;921; CHECK-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r6;922; CHECK-FTZ-NEXT: cvt.u32.u16 %r13, %rs5;923; CHECK-FTZ-NEXT: shl.b32 %r14, %r13, 16;924; CHECK-FTZ-NEXT: cvt.u32.u16 %r15, %rs4;925; CHECK-FTZ-NEXT: shl.b32 %r16, %r15, 16;926; CHECK-FTZ-NEXT: add.ftz.f32 %r17, %r14, %r16;927; CHECK-FTZ-NEXT: cvt.u32.u16 %r18, %rs6;928; CHECK-FTZ-NEXT: shl.b32 %r19, %r18, 16;929; CHECK-FTZ-NEXT: cvt.u32.u16 %r20, %rs3;930; CHECK-FTZ-NEXT: shl.b32 %r21, %r20, 16;931; CHECK-FTZ-NEXT: add.ftz.f32 %r22, %r19, %r21;932; CHECK-FTZ-NEXT: cvt.rn.bf16x2.f32 %r23, %r22, %r17;933; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r23;934; CHECK-FTZ-NEXT: ret;935;936; CHECK-SM70-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(937; CHECK-SM70: {938; CHECK-SM70-NEXT: .reg .pred %p<9>;939; CHECK-SM70-NEXT: .reg .b16 %rs<11>;940; CHECK-SM70-NEXT: .reg .b32 %r<58>;941; CHECK-SM70-EMPTY:942; CHECK-SM70-NEXT: // %bb.0:943; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];944; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];945; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];946; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs6;947; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;948; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs4;949; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;950; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs2;951; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;952; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;953; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;954; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;955; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;956; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;957; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;958; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;959; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }960; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs5;961; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;962; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs3;963; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;964; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs1;965; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;966; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;967; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;968; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;969; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;970; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;971; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;972; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;973; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }974; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;975; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;976; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;977; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;978; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;979; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;980; CHECK-SM70-NEXT: add.f32 %r27, %r26, 0f40E00000;981; CHECK-SM70-NEXT: bfe.u32 %r28, %r27, 16, 1;982; CHECK-SM70-NEXT: add.s32 %r29, %r28, %r27;983; CHECK-SM70-NEXT: add.s32 %r30, %r29, 32767;984; CHECK-SM70-NEXT: setp.nan.f32 %p5, %r27, %r27;985; CHECK-SM70-NEXT: or.b32 %r31, %r27, 4194304;986; CHECK-SM70-NEXT: selp.b32 %r32, %r31, %r30, %p5;987; CHECK-SM70-NEXT: add.f32 %r33, %r25, 0f40E00000;988; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;989; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;990; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;991; CHECK-SM70-NEXT: setp.nan.f32 %p6, %r33, %r33;992; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;993; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p6;994; CHECK-SM70-NEXT: cvt.u32.u16 %r39, %rs10;995; CHECK-SM70-NEXT: shl.b32 %r40, %r39, 16;996; CHECK-SM70-NEXT: and.b32 %r41, %r38, -65536;997; CHECK-SM70-NEXT: add.f32 %r42, %r40, %r41;998; CHECK-SM70-NEXT: bfe.u32 %r43, %r42, 16, 1;999; CHECK-SM70-NEXT: add.s32 %r44, %r43, %r42;1000; CHECK-SM70-NEXT: add.s32 %r45, %r44, 32767;1001; CHECK-SM70-NEXT: setp.nan.f32 %p7, %r42, %r42;1002; CHECK-SM70-NEXT: or.b32 %r46, %r42, 4194304;1003; CHECK-SM70-NEXT: selp.b32 %r47, %r46, %r45, %p7;1004; CHECK-SM70-NEXT: cvt.u32.u16 %r48, %rs9;1005; CHECK-SM70-NEXT: shl.b32 %r49, %r48, 16;1006; CHECK-SM70-NEXT: and.b32 %r50, %r32, -65536;1007; CHECK-SM70-NEXT: add.f32 %r51, %r49, %r50;1008; CHECK-SM70-NEXT: bfe.u32 %r52, %r51, 16, 1;1009; CHECK-SM70-NEXT: add.s32 %r53, %r52, %r51;1010; CHECK-SM70-NEXT: add.s32 %r54, %r53, 32767;1011; CHECK-SM70-NEXT: setp.nan.f32 %p8, %r51, %r51;1012; CHECK-SM70-NEXT: or.b32 %r55, %r51, 4194304;1013; CHECK-SM70-NEXT: selp.b32 %r56, %r55, %r54, %p8;1014; CHECK-SM70-NEXT: prmt.b32 %r57, %r56, %r47, 0x7632U;1015; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r57;1016; CHECK-SM70-NEXT: ret;1017 %1 = fmul <2 x bfloat> %a, %b1018 %2 = fadd <2 x bfloat> %1, %c1019 %3 = fcmp ogt <2 x bfloat> %2, <bfloat 0.0, bfloat 0.0>1020 %4 = select <2 x i1> %3, <2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>1021 %5 = fadd <2 x bfloat> %2, <bfloat 7.0, bfloat 7.0>1022 %6 = fadd <2 x bfloat> %4, %51023 ret <2 x bfloat> %61024}1025 1026define <2 x bfloat> @fma_bf16x2_expanded_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) #0 {1027; CHECK-LABEL: fma_bf16x2_expanded_maxnum_no_nans(1028; CHECK: {1029; CHECK-NEXT: .reg .b32 %r<5>;1030; CHECK-EMPTY:1031; CHECK-NEXT: // %bb.0:1032; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];1033; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_maxnum_no_nans_param_1];1034; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_maxnum_no_nans_param_2];1035; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;1036; CHECK-NEXT: st.param.b32 [func_retval0], %r4;1037; CHECK-NEXT: ret;1038;1039; CHECK-FTZ-LABEL: fma_bf16x2_expanded_maxnum_no_nans(1040; CHECK-FTZ: {1041; CHECK-FTZ-NEXT: .reg .b32 %r<5>;1042; CHECK-FTZ-EMPTY:1043; CHECK-FTZ-NEXT: // %bb.0:1044; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];1045; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_maxnum_no_nans_param_1];1046; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_maxnum_no_nans_param_2];1047; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;1048; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;1049; CHECK-FTZ-NEXT: ret;1050;1051; CHECK-SM70-LABEL: fma_bf16x2_expanded_maxnum_no_nans(1052; CHECK-SM70: {1053; CHECK-SM70-NEXT: .reg .pred %p<5>;1054; CHECK-SM70-NEXT: .reg .b16 %rs<7>;1055; CHECK-SM70-NEXT: .reg .b32 %r<40>;1056; CHECK-SM70-EMPTY:1057; CHECK-SM70-NEXT: // %bb.0:1058; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_maxnum_no_nans_param_0];1059; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_maxnum_no_nans_param_1];1060; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_maxnum_no_nans_param_2];1061; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs5;1062; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;1063; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs3;1064; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;1065; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;1066; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;1067; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;1068; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;1069; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;1070; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;1071; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;1072; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;1073; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;1074; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs6;1075; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;1076; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs4;1077; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;1078; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;1079; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;1080; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;1081; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;1082; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;1083; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;1084; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;1085; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;1086; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;1087; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;1088; CHECK-SM70-NEXT: max.f32 %r26, %r25, 0f00000000;1089; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;1090; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;1091; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;1092; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;1093; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;1094; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;1095; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;1096; CHECK-SM70-NEXT: max.f32 %r33, %r32, 0f00000000;1097; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;1098; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;1099; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;1100; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;1101; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;1102; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;1103; CHECK-SM70-NEXT: prmt.b32 %r39, %r38, %r31, 0x7632U;1104; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r39;1105; CHECK-SM70-NEXT: ret;1106 %1 = fmul <2 x bfloat> %a, %b1107 %2 = fadd <2 x bfloat> %1, %c1108 %3 = call <2 x bfloat> @llvm.maxnum.bf16x2(<2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>)1109 ret <2 x bfloat> %31110}1111 1112attributes #0 = { "no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true" }1113