799 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_no_nans(half %a, half %b, half %c) {13; CHECK-LABEL: fma_f16_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_no_nans_param_0];19; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];20; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_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_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_no_nans_param_0];31; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];32; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_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_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_no_nans_param_0];44; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];45; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_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 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)53 %2 = fcmp ogt half %1, 0.054 %3 = select nsz i1 %2, half %1, half 0.055 ret half %356}57 58; FMA relu shouldn't be selected if the FMA operation has multiple uses59define half @fma_f16_no_nans_multiple_uses_of_fma(half %a, half %b, half %c) {60; CHECK-LABEL: fma_f16_no_nans_multiple_uses_of_fma(61; CHECK: {62; CHECK-NEXT: .reg .b16 %rs<8>;63; CHECK-EMPTY:64; CHECK-NEXT: // %bb.0:65; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];66; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];67; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];68; CHECK-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;69; CHECK-NEXT: mov.b16 %rs5, 0x4700;70; CHECK-NEXT: add.f16 %rs6, %rs4, %rs5;71; CHECK-NEXT: add.f16 %rs7, %rs6, %rs4;72; CHECK-NEXT: st.param.b16 [func_retval0], %rs7;73; CHECK-NEXT: ret;74;75; CHECK-FTZ-LABEL: fma_f16_no_nans_multiple_uses_of_fma(76; CHECK-FTZ: {77; CHECK-FTZ-NEXT: .reg .b16 %rs<8>;78; CHECK-FTZ-EMPTY:79; CHECK-FTZ-NEXT: // %bb.0:80; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];81; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];82; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];83; CHECK-FTZ-NEXT: fma.rn.ftz.f16 %rs4, %rs1, %rs2, %rs3;84; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x4700;85; CHECK-FTZ-NEXT: add.ftz.f16 %rs6, %rs4, %rs5;86; CHECK-FTZ-NEXT: add.ftz.f16 %rs7, %rs6, %rs4;87; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs7;88; CHECK-FTZ-NEXT: ret;89;90; CHECK-SM70-LABEL: fma_f16_no_nans_multiple_uses_of_fma(91; CHECK-SM70: {92; CHECK-SM70-NEXT: .reg .b16 %rs<8>;93; CHECK-SM70-EMPTY:94; CHECK-SM70-NEXT: // %bb.0:95; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];96; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];97; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];98; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;99; CHECK-SM70-NEXT: mov.b16 %rs5, 0x4700;100; CHECK-SM70-NEXT: add.f16 %rs6, %rs4, %rs5;101; CHECK-SM70-NEXT: add.f16 %rs7, %rs6, %rs4;102; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs7;103; CHECK-SM70-NEXT: ret;104 %1 = call half @llvm.fma.f16(half %a, half %b, half %c)105 %2 = fcmp ogt half %1, 0.0106 %3 = select i1 %2, half %1, half 0.0107 %4 = fadd contract half %1, 7.0108 %5 = fadd contract half %4, %1109 ret half %5110}111 112define half @fma_f16_maxnum_no_nans(half %a, half %b, half %c) {113; CHECK-LABEL: fma_f16_maxnum_no_nans(114; CHECK: {115; CHECK-NEXT: .reg .b16 %rs<5>;116; CHECK-EMPTY:117; CHECK-NEXT: // %bb.0:118; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];119; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];120; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];121; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;122; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;123; CHECK-NEXT: ret;124;125; CHECK-FTZ-LABEL: fma_f16_maxnum_no_nans(126; CHECK-FTZ: {127; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;128; CHECK-FTZ-EMPTY:129; CHECK-FTZ-NEXT: // %bb.0:130; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];131; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];132; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];133; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;134; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;135; CHECK-FTZ-NEXT: ret;136;137; CHECK-SM70-LABEL: fma_f16_maxnum_no_nans(138; CHECK-SM70: {139; CHECK-SM70-NEXT: .reg .b16 %rs<6>;140; CHECK-SM70-NEXT: .reg .b32 %r<3>;141; CHECK-SM70-EMPTY:142; CHECK-SM70-NEXT: // %bb.0:143; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];144; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];145; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];146; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;147; CHECK-SM70-NEXT: cvt.f32.f16 %r1, %rs4;148; CHECK-SM70-NEXT: max.f32 %r2, %r1, 0f00000000;149; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs5, %r2;150; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs5;151; CHECK-SM70-NEXT: ret;152 %1 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)153 %2 = call nsz half @llvm.maxnum.f16(half %1, half 0.0)154 ret half %2155}156 157define bfloat @fma_bf16_no_nans(bfloat %a, bfloat %b, bfloat %c) {158; CHECK-LABEL: fma_bf16_no_nans(159; CHECK: {160; CHECK-NEXT: .reg .b16 %rs<5>;161; CHECK-EMPTY:162; CHECK-NEXT: // %bb.0:163; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_param_0];164; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_param_1];165; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_param_2];166; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;167; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;168; CHECK-NEXT: ret;169;170; CHECK-FTZ-LABEL: fma_bf16_no_nans(171; CHECK-FTZ: {172; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;173; CHECK-FTZ-EMPTY:174; CHECK-FTZ-NEXT: // %bb.0:175; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_param_0];176; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_param_1];177; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_param_2];178; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;179; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;180; CHECK-FTZ-NEXT: ret;181;182; CHECK-SM70-LABEL: fma_bf16_no_nans(183; CHECK-SM70: {184; CHECK-SM70-NEXT: .reg .pred %p<3>;185; CHECK-SM70-NEXT: .reg .b16 %rs<3>;186; CHECK-SM70-NEXT: .reg .b32 %r<14>;187; CHECK-SM70-EMPTY:188; CHECK-SM70-NEXT: // %bb.0:189; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_no_nans_param_2];190; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;191; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_no_nans_param_1];192; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;193; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_no_nans_param_0];194; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;195; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;196; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;197; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;198; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;199; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;200; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;201; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;202; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }203; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;204; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;205; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;206; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs2;207; CHECK-SM70-NEXT: ret;208 %1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)209 %2 = fcmp ogt bfloat %1, 0.0210 %3 = select nsz i1 %2, bfloat %1, bfloat 0.0211 ret bfloat %3212}213 214; FMA_relu shouldn't be selected if the FMA operation has multiple uses215define bfloat @fma_bf16_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloat %c) {216; CHECK-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(217; CHECK: {218; CHECK-NEXT: .reg .b16 %rs<9>;219; CHECK-EMPTY:220; CHECK-NEXT: // %bb.0:221; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];222; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];223; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];224; CHECK-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;225; CHECK-NEXT: mov.b16 %rs5, 0x3F80;226; CHECK-NEXT: mov.b16 %rs6, 0x40E0;227; CHECK-NEXT: fma.rn.bf16 %rs7, %rs4, %rs5, %rs6;228; CHECK-NEXT: fma.rn.bf16 %rs8, %rs7, %rs5, %rs4;229; CHECK-NEXT: st.param.b16 [func_retval0], %rs8;230; CHECK-NEXT: ret;231;232; CHECK-FTZ-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(233; CHECK-FTZ: {234; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;235; CHECK-FTZ-NEXT: .reg .b32 %r<7>;236; CHECK-FTZ-EMPTY:237; CHECK-FTZ-NEXT: // %bb.0:238; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];239; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];240; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];241; CHECK-FTZ-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;242; CHECK-FTZ-NEXT: cvt.u32.u16 %r1, %rs4;243; CHECK-FTZ-NEXT: shl.b32 %r2, %r1, 16;244; CHECK-FTZ-NEXT: add.ftz.f32 %r3, %r2, 0f40E00000;245; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %r3;246; CHECK-FTZ-NEXT: cvt.u32.u16 %r4, %rs5;247; CHECK-FTZ-NEXT: shl.b32 %r5, %r4, 16;248; CHECK-FTZ-NEXT: add.ftz.f32 %r6, %r5, %r2;249; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %r6;250; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs6;251; CHECK-FTZ-NEXT: ret;252;253; CHECK-SM70-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(254; CHECK-SM70: {255; CHECK-SM70-NEXT: .reg .pred %p<4>;256; CHECK-SM70-NEXT: .reg .b32 %r<28>;257; CHECK-SM70-EMPTY:258; CHECK-SM70-NEXT: // %bb.0:259; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];260; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;261; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];262; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;263; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];264; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;265; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;266; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;267; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;268; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;269; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;270; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;271; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;272; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;273; CHECK-SM70-NEXT: add.f32 %r14, %r13, 0f40E00000;274; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;275; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;276; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;277; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;278; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;279; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;280; CHECK-SM70-NEXT: and.b32 %r20, %r19, -65536;281; CHECK-SM70-NEXT: add.f32 %r21, %r20, %r13;282; CHECK-SM70-NEXT: bfe.u32 %r22, %r21, 16, 1;283; CHECK-SM70-NEXT: add.s32 %r23, %r22, %r21;284; CHECK-SM70-NEXT: add.s32 %r24, %r23, 32767;285; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r21, %r21;286; CHECK-SM70-NEXT: or.b32 %r25, %r21, 4194304;287; CHECK-SM70-NEXT: selp.b32 %r26, %r25, %r24, %p3;288; CHECK-SM70-NEXT: shr.u32 %r27, %r26, 16;289; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r27;290; CHECK-SM70-NEXT: ret;291 %1 = call bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)292 %2 = fcmp ogt bfloat %1, 0.0293 %3 = select i1 %2, bfloat %1, bfloat 0.0294 %4 = fadd contract bfloat %1, 7.0295 %5 = fadd contract bfloat %4, %1296 ret bfloat %5297}298 299define bfloat @fma_bf16_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) {300; CHECK-LABEL: fma_bf16_maxnum_no_nans(301; CHECK: {302; CHECK-NEXT: .reg .b16 %rs<5>;303; CHECK-EMPTY:304; CHECK-NEXT: // %bb.0:305; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_maxnum_no_nans_param_0];306; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_maxnum_no_nans_param_1];307; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_maxnum_no_nans_param_2];308; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;309; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;310; CHECK-NEXT: ret;311;312; CHECK-FTZ-LABEL: fma_bf16_maxnum_no_nans(313; CHECK-FTZ: {314; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;315; CHECK-FTZ-EMPTY:316; CHECK-FTZ-NEXT: // %bb.0:317; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_maxnum_no_nans_param_0];318; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_maxnum_no_nans_param_1];319; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_maxnum_no_nans_param_2];320; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;321; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;322; CHECK-FTZ-NEXT: ret;323;324; CHECK-SM70-LABEL: fma_bf16_maxnum_no_nans(325; CHECK-SM70: {326; CHECK-SM70-NEXT: .reg .pred %p<3>;327; CHECK-SM70-NEXT: .reg .b32 %r<21>;328; CHECK-SM70-EMPTY:329; CHECK-SM70-NEXT: // %bb.0:330; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_maxnum_no_nans_param_2];331; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;332; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_maxnum_no_nans_param_1];333; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;334; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_maxnum_no_nans_param_0];335; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;336; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;337; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;338; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;339; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;340; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;341; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;342; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;343; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;344; CHECK-SM70-NEXT: max.f32 %r14, %r13, 0f00000000;345; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;346; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;347; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;348; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;349; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;350; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;351; CHECK-SM70-NEXT: shr.u32 %r20, %r19, 16;352; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %r20;353; CHECK-SM70-NEXT: ret;354 %1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)355 %2 = call nsz bfloat @llvm.maxnum.bf16(bfloat %1, bfloat 0.0)356 ret bfloat %2357}358 359define <2 x half> @fma_f16x2_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {360; CHECK-LABEL: fma_f16x2_no_nans(361; CHECK: {362; CHECK-NEXT: .reg .b32 %r<5>;363; CHECK-EMPTY:364; CHECK-NEXT: // %bb.0:365; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];366; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];367; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];368; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;369; CHECK-NEXT: st.param.b32 [func_retval0], %r4;370; CHECK-NEXT: ret;371;372; CHECK-FTZ-LABEL: fma_f16x2_no_nans(373; CHECK-FTZ: {374; CHECK-FTZ-NEXT: .reg .b32 %r<5>;375; CHECK-FTZ-EMPTY:376; CHECK-FTZ-NEXT: // %bb.0:377; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];378; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];379; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];380; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;381; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;382; CHECK-FTZ-NEXT: ret;383;384; CHECK-SM70-LABEL: fma_f16x2_no_nans(385; CHECK-SM70: {386; CHECK-SM70-NEXT: .reg .pred %p<3>;387; CHECK-SM70-NEXT: .reg .b16 %rs<5>;388; CHECK-SM70-NEXT: .reg .b32 %r<6>;389; CHECK-SM70-EMPTY:390; CHECK-SM70-NEXT: // %bb.0:391; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];392; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];393; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];394; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;395; CHECK-SM70-NEXT: mov.b32 %r5, 0;396; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;397; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;398; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;399; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;400; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};401; CHECK-SM70-NEXT: ret;402 %1 = call nnan <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)403 %2 = fcmp ogt <2 x half> %1, <half 0.0, half 0.0>404 %3 = select nsz <2 x i1> %2, <2 x half> %1, <2 x half> <half 0.0, half 0.0>405 ret <2 x half> %3406}407 408; FMA relu shouldn't be selected if the FMA operation has multiple uses409define <2 x half> @fma_f16x2_no_nans_multiple_uses_of_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) {410; CHECK-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(411; CHECK: {412; CHECK-NEXT: .reg .b32 %r<8>;413; CHECK-EMPTY:414; CHECK-NEXT: // %bb.0:415; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];416; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];417; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];418; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;419; CHECK-NEXT: mov.b32 %r5, 1191200512;420; CHECK-NEXT: add.f16x2 %r6, %r4, %r5;421; CHECK-NEXT: add.f16x2 %r7, %r6, %r4;422; CHECK-NEXT: st.param.b32 [func_retval0], %r7;423; CHECK-NEXT: ret;424;425; CHECK-FTZ-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(426; CHECK-FTZ: {427; CHECK-FTZ-NEXT: .reg .b32 %r<8>;428; CHECK-FTZ-EMPTY:429; CHECK-FTZ-NEXT: // %bb.0:430; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];431; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];432; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];433; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;434; CHECK-FTZ-NEXT: mov.b32 %r5, 1191200512;435; CHECK-FTZ-NEXT: add.ftz.f16x2 %r6, %r4, %r5;436; CHECK-FTZ-NEXT: add.ftz.f16x2 %r7, %r6, %r4;437; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r7;438; CHECK-FTZ-NEXT: ret;439;440; CHECK-SM70-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(441; CHECK-SM70: {442; CHECK-SM70-NEXT: .reg .b32 %r<8>;443; CHECK-SM70-EMPTY:444; CHECK-SM70-NEXT: // %bb.0:445; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];446; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];447; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];448; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;449; CHECK-SM70-NEXT: mov.b32 %r5, 1191200512;450; CHECK-SM70-NEXT: add.f16x2 %r6, %r4, %r5;451; CHECK-SM70-NEXT: add.f16x2 %r7, %r6, %r4;452; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r7;453; CHECK-SM70-NEXT: ret;454 %1 = call <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)455 %2 = fcmp ogt <2 x half> %1, <half 0.0, half 0.0>456 %3 = select <2 x i1> %2, <2 x half> %1, <2 x half> <half 0.0, half 0.0>457 %4 = fadd contract <2 x half> %1, <half 7.0, half 7.0>458 %5 = fadd contract <2 x half> %4, %1459 ret <2 x half> %5460}461 462define <2 x half> @fma_f16x2_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {463; CHECK-LABEL: fma_f16x2_maxnum_no_nans(464; CHECK: {465; CHECK-NEXT: .reg .b32 %r<5>;466; CHECK-EMPTY:467; CHECK-NEXT: // %bb.0:468; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];469; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];470; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];471; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;472; CHECK-NEXT: st.param.b32 [func_retval0], %r4;473; CHECK-NEXT: ret;474;475; CHECK-FTZ-LABEL: fma_f16x2_maxnum_no_nans(476; CHECK-FTZ: {477; CHECK-FTZ-NEXT: .reg .b32 %r<5>;478; CHECK-FTZ-EMPTY:479; CHECK-FTZ-NEXT: // %bb.0:480; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];481; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];482; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];483; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;484; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;485; CHECK-FTZ-NEXT: ret;486;487; CHECK-SM70-LABEL: fma_f16x2_maxnum_no_nans(488; CHECK-SM70: {489; CHECK-SM70-NEXT: .reg .b16 %rs<5>;490; CHECK-SM70-NEXT: .reg .b32 %r<10>;491; CHECK-SM70-EMPTY:492; CHECK-SM70-NEXT: // %bb.0:493; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];494; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];495; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];496; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;497; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;498; CHECK-SM70-NEXT: cvt.f32.f16 %r5, %rs2;499; CHECK-SM70-NEXT: max.f32 %r6, %r5, 0f00000000;500; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs3, %r6;501; CHECK-SM70-NEXT: cvt.f32.f16 %r7, %rs1;502; CHECK-SM70-NEXT: max.f32 %r8, %r7, 0f00000000;503; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs4, %r8;504; CHECK-SM70-NEXT: mov.b32 %r9, {%rs4, %rs3};505; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;506; CHECK-SM70-NEXT: ret;507 %1 = call nnan <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)508 %2 = call nsz <2 x half> @llvm.maxnum.f16x2(<2 x half> %1, <2 x half> <half 0.0, half 0.0>)509 ret <2 x half> %2510}511 512define <2 x bfloat> @fma_bf16x2_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {513; CHECK-LABEL: fma_bf16x2_no_nans(514; CHECK: {515; CHECK-NEXT: .reg .b32 %r<5>;516; CHECK-EMPTY:517; CHECK-NEXT: // %bb.0:518; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];519; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_param_1];520; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_param_2];521; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;522; CHECK-NEXT: st.param.b32 [func_retval0], %r4;523; CHECK-NEXT: ret;524;525; CHECK-FTZ-LABEL: fma_bf16x2_no_nans(526; CHECK-FTZ: {527; CHECK-FTZ-NEXT: .reg .b32 %r<5>;528; CHECK-FTZ-EMPTY:529; CHECK-FTZ-NEXT: // %bb.0:530; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];531; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_param_1];532; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_param_2];533; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;534; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;535; CHECK-FTZ-NEXT: ret;536;537; CHECK-SM70-LABEL: fma_bf16x2_no_nans(538; CHECK-SM70: {539; CHECK-SM70-NEXT: .reg .pred %p<5>;540; CHECK-SM70-NEXT: .reg .b16 %rs<11>;541; CHECK-SM70-NEXT: .reg .b32 %r<27>;542; CHECK-SM70-EMPTY:543; CHECK-SM70-NEXT: // %bb.0:544; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_no_nans_param_0];545; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_no_nans_param_2];546; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs3;547; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;548; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_no_nans_param_1];549; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs5;550; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;551; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;552; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;553; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;554; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;555; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;556; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;557; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;558; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;559; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;560; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }561; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs4;562; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;563; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs6;564; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;565; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;566; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;567; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;568; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;569; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;570; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;571; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;572; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;573; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;574; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }575; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;576; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;577; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;578; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;579; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;580; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;581; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};582; CHECK-SM70-NEXT: ret;583 %1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)584 %2 = fcmp ogt <2 x bfloat> %1, <bfloat 0.0, bfloat 0.0>585 %3 = select nsz <2 x i1> %2, <2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>586 ret <2 x bfloat> %3587}588 589; FMA_relu shouldn't be selected if the FMA operation has multiple uses590define <2 x bfloat> @fma_bf16x2_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {591; CHECK-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(592; CHECK: {593; CHECK-NEXT: .reg .b32 %r<9>;594; CHECK-EMPTY:595; CHECK-NEXT: // %bb.0:596; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];597; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];598; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];599; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;600; CHECK-NEXT: mov.b32 %r5, 1065369472;601; CHECK-NEXT: mov.b32 %r6, 1088438496;602; CHECK-NEXT: fma.rn.bf16x2 %r7, %r4, %r5, %r6;603; CHECK-NEXT: fma.rn.bf16x2 %r8, %r7, %r5, %r4;604; CHECK-NEXT: st.param.b32 [func_retval0], %r8;605; CHECK-NEXT: ret;606;607; CHECK-FTZ-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(608; CHECK-FTZ: {609; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;610; CHECK-FTZ-NEXT: .reg .b32 %r<18>;611; CHECK-FTZ-EMPTY:612; CHECK-FTZ-NEXT: // %bb.0:613; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];614; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];615; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];616; CHECK-FTZ-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;617; CHECK-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r4;618; CHECK-FTZ-NEXT: cvt.u32.u16 %r5, %rs2;619; CHECK-FTZ-NEXT: shl.b32 %r6, %r5, 16;620; CHECK-FTZ-NEXT: add.ftz.f32 %r7, %r6, 0f40E00000;621; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r7;622; CHECK-FTZ-NEXT: cvt.u32.u16 %r8, %rs1;623; CHECK-FTZ-NEXT: shl.b32 %r9, %r8, 16;624; CHECK-FTZ-NEXT: add.ftz.f32 %r10, %r9, 0f40E00000;625; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs4, %r10;626; CHECK-FTZ-NEXT: cvt.u32.u16 %r11, %rs4;627; CHECK-FTZ-NEXT: shl.b32 %r12, %r11, 16;628; CHECK-FTZ-NEXT: add.ftz.f32 %r13, %r12, %r9;629; CHECK-FTZ-NEXT: cvt.u32.u16 %r14, %rs3;630; CHECK-FTZ-NEXT: shl.b32 %r15, %r14, 16;631; CHECK-FTZ-NEXT: add.ftz.f32 %r16, %r15, %r6;632; CHECK-FTZ-NEXT: cvt.rn.bf16x2.f32 %r17, %r16, %r13;633; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r17;634; CHECK-FTZ-NEXT: ret;635;636; CHECK-SM70-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(637; CHECK-SM70: {638; CHECK-SM70-NEXT: .reg .pred %p<7>;639; CHECK-SM70-NEXT: .reg .b16 %rs<7>;640; CHECK-SM70-NEXT: .reg .b32 %r<54>;641; CHECK-SM70-EMPTY:642; CHECK-SM70-NEXT: // %bb.0:643; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];644; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];645; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs4;646; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;647; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];648; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs6;649; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;650; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs2;651; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;652; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;653; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;654; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;655; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;656; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;657; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;658; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;659; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs3;660; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;661; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs5;662; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;663; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs1;664; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;665; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;666; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;667; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;668; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;669; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;670; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;671; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;672; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;673; CHECK-SM70-NEXT: add.f32 %r26, %r25, 0f40E00000;674; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;675; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;676; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;677; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;678; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;679; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;680; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;681; CHECK-SM70-NEXT: add.f32 %r33, %r32, 0f40E00000;682; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;683; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;684; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;685; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;686; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;687; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;688; CHECK-SM70-NEXT: and.b32 %r39, %r38, -65536;689; CHECK-SM70-NEXT: add.f32 %r40, %r39, %r32;690; CHECK-SM70-NEXT: bfe.u32 %r41, %r40, 16, 1;691; CHECK-SM70-NEXT: add.s32 %r42, %r41, %r40;692; CHECK-SM70-NEXT: add.s32 %r43, %r42, 32767;693; CHECK-SM70-NEXT: setp.nan.f32 %p5, %r40, %r40;694; CHECK-SM70-NEXT: or.b32 %r44, %r40, 4194304;695; CHECK-SM70-NEXT: selp.b32 %r45, %r44, %r43, %p5;696; CHECK-SM70-NEXT: and.b32 %r46, %r31, -65536;697; CHECK-SM70-NEXT: add.f32 %r47, %r46, %r25;698; CHECK-SM70-NEXT: bfe.u32 %r48, %r47, 16, 1;699; CHECK-SM70-NEXT: add.s32 %r49, %r48, %r47;700; CHECK-SM70-NEXT: add.s32 %r50, %r49, 32767;701; CHECK-SM70-NEXT: setp.nan.f32 %p6, %r47, %r47;702; CHECK-SM70-NEXT: or.b32 %r51, %r47, 4194304;703; CHECK-SM70-NEXT: selp.b32 %r52, %r51, %r50, %p6;704; CHECK-SM70-NEXT: prmt.b32 %r53, %r52, %r45, 0x7632U;705; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r53;706; CHECK-SM70-NEXT: ret;707 %1 = call <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)708 %2 = fcmp ogt <2 x bfloat> %1, <bfloat 0.0, bfloat 0.0>709 %3 = select <2 x i1> %2, <2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>710 %4 = fadd contract <2 x bfloat> %1, <bfloat 7.0, bfloat 7.0>711 %5 = fadd contract <2 x bfloat> %4, %1712 ret <2 x bfloat> %5713}714 715define <2 x bfloat> @fma_bf16x2_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {716; CHECK-LABEL: fma_bf16x2_maxnum_no_nans(717; CHECK: {718; CHECK-NEXT: .reg .b32 %r<5>;719; CHECK-EMPTY:720; CHECK-NEXT: // %bb.0:721; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];722; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_maxnum_no_nans_param_1];723; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_maxnum_no_nans_param_2];724; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;725; CHECK-NEXT: st.param.b32 [func_retval0], %r4;726; CHECK-NEXT: ret;727;728; CHECK-FTZ-LABEL: fma_bf16x2_maxnum_no_nans(729; CHECK-FTZ: {730; CHECK-FTZ-NEXT: .reg .b32 %r<5>;731; CHECK-FTZ-EMPTY:732; CHECK-FTZ-NEXT: // %bb.0:733; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];734; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_maxnum_no_nans_param_1];735; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_maxnum_no_nans_param_2];736; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;737; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;738; CHECK-FTZ-NEXT: ret;739;740; CHECK-SM70-LABEL: fma_bf16x2_maxnum_no_nans(741; CHECK-SM70: {742; CHECK-SM70-NEXT: .reg .pred %p<5>;743; CHECK-SM70-NEXT: .reg .b16 %rs<7>;744; CHECK-SM70-NEXT: .reg .b32 %r<40>;745; CHECK-SM70-EMPTY:746; CHECK-SM70-NEXT: // %bb.0:747; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_maxnum_no_nans_param_0];748; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_maxnum_no_nans_param_2];749; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs3;750; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;751; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_maxnum_no_nans_param_1];752; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs5;753; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;754; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;755; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;756; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;757; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;758; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;759; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;760; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;761; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;762; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;763; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs4;764; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;765; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs6;766; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;767; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;768; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;769; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;770; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;771; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;772; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;773; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;774; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;775; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;776; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;777; CHECK-SM70-NEXT: max.f32 %r26, %r25, 0f00000000;778; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;779; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;780; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;781; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;782; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;783; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;784; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;785; CHECK-SM70-NEXT: max.f32 %r33, %r32, 0f00000000;786; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;787; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;788; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;789; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;790; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;791; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;792; CHECK-SM70-NEXT: prmt.b32 %r39, %r38, %r31, 0x7632U;793; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r39;794; CHECK-SM70-NEXT: ret;795 %1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)796 %2 = call nsz <2 x bfloat> @llvm.maxnum.bf16x2(<2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>)797 ret <2 x bfloat> %2798}799