Google Git
Sign in
llvm / llvm-project / llvm / 4065f8ffaff7fd23a176d4c968bba1ebac135a19 / . / test / CodeGen / NVPTX / fma-relu-instruction-flag.ll
blob: 60bfe3fa2cbf6cbbf4eba2fc7e68766a2d9820b9 [file] [log] [blame]
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
; Using FTZ should emit fma.ftz.relu for f16, not for bf16
; RUN: llc < %s -denormal-fp-math-f32=preserve-sign -mtriple=nvptx64 -mcpu=sm_80 | FileCheck %s --check-prefixes=CHECK-FTZ
; RUN: %if ptxas %{ llc < %s -denormal-fp-math-f32=preserve-sign -mtriple=nvptx64 -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
; SM < 80 or (which needs PTX version >= 70) should not emit fma{.ftz}.relu
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 | FileCheck %s --check-prefixes=CHECK-SM70
define half @fma_f16_expanded_no_nans(half %a, half %b, half %c) {
; CHECK-LABEL: fma_f16_expanded_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16_expanded_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16_expanded_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<2>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-SM70-NEXT: mov.b16 %rs5, 0x0000;
; CHECK-SM70-NEXT: setp.gt.f16 %p1, %rs4, %rs5;
; CHECK-SM70-NEXT: selp.b16 %rs6, %rs4, 0x0000, %p1;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs6;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast half %a, %b
%2 = fadd fast half %1, %c
%3 = fcmp nsz ogt half %2, 0.0
%4 = select nsz i1 %3, half %2, half 0.0
ret half %4
}
; FMA relu shouldn't be selected if the FMA operation has multiple uses
define half @fma_f16_expanded_no_nans_multiple_uses_of_fma(half %a, half %b, half %c) {
; CHECK-LABEL: fma_f16_expanded_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: mov.b16 %rs5, 0x0000;
; CHECK-NEXT: max.f16 %rs6, %rs4, %rs5;
; CHECK-NEXT: mov.b16 %rs7, 0x4700;
; CHECK-NEXT: add.rn.f16 %rs8, %rs4, %rs7;
; CHECK-NEXT: add.rn.f16 %rs9, %rs6, %rs8;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs9;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16_expanded_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<10>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x0000;
; CHECK-FTZ-NEXT: max.ftz.f16 %rs6, %rs4, %rs5;
; CHECK-FTZ-NEXT: mov.b16 %rs7, 0x4700;
; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs8, %rs4, %rs7;
; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs9, %rs6, %rs8;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs9;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16_expanded_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<2>;
; CHECK-SM70-NEXT: .reg .b16 %rs<10>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-SM70-NEXT: mov.b16 %rs5, 0x0000;
; CHECK-SM70-NEXT: setp.gt.f16 %p1, %rs4, %rs5;
; CHECK-SM70-NEXT: selp.b16 %rs6, %rs4, 0x0000, %p1;
; CHECK-SM70-NEXT: mov.b16 %rs7, 0x4700;
; CHECK-SM70-NEXT: add.rn.f16 %rs8, %rs4, %rs7;
; CHECK-SM70-NEXT: add.rn.f16 %rs9, %rs6, %rs8;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs9;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast half %a, %b
%2 = fadd fast half %1, %c
%3 = fcmp nsz ogt half %2, 0.0
%4 = select nsz i1 %3, half %2, half 0.0
%5 = fadd half %2, 7.0
%6 = fadd half %4, %5
ret half %6
}
define half @fma_f16_expanded_maxnum_no_nans(half %a, half %b, half %c) {
; CHECK-LABEL: fma_f16_expanded_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16_expanded_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16_expanded_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<6>;
; CHECK-SM70-NEXT: .reg .b32 %r<3>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_expanded_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_expanded_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_expanded_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-SM70-NEXT: cvt.f32.f16 %r1, %rs4;
; CHECK-SM70-NEXT: max.f32 %r2, %r1, 0f00000000;
; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs5, %r2;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs5;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast half %a, %b
%2 = fadd fast half %1, %c
%3 = call nsz half @llvm.maxnum.f16(half %2, half 0.0)
ret half %3
}
define bfloat @fma_bf16_expanded_no_nans(bfloat %a, bfloat %b, bfloat %c) {
; CHECK-LABEL: fma_bf16_expanded_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16_expanded_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16_expanded_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<3>;
; CHECK-SM70-NEXT: .reg .b32 %r<14>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_no_nans_param_2];
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_no_nans_param_1];
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_no_nans_param_0];
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }
; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;
; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast bfloat %a, %b
%2 = fadd fast bfloat %1, %c
%3 = fcmp nsz ogt bfloat %2, 0.0
%4 = select nsz i1 %3, bfloat %2, bfloat 0.0
ret bfloat %4
}
; FMA relu shouldn't be selected if the FMA operation has multiple uses
define bfloat @fma_bf16_expanded_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloat %c) {
; CHECK-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<11>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: mov.b16 %rs5, 0x0000;
; CHECK-NEXT: max.bf16 %rs6, %rs4, %rs5;
; CHECK-NEXT: mov.b16 %rs7, 0x3F80;
; CHECK-NEXT: mov.b16 %rs8, 0x40E0;
; CHECK-NEXT: fma.rn.bf16 %rs9, %rs4, %rs7, %rs8;
; CHECK-NEXT: fma.rn.bf16 %rs10, %rs6, %rs7, %rs9;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs10;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<9>;
; CHECK-FTZ-NEXT: .reg .b32 %r<9>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x0000;
; CHECK-FTZ-NEXT: max.bf16 %rs6, %rs4, %rs5;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-FTZ-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40E00000;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs7, %r3;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r4, %rs6;
; CHECK-FTZ-NEXT: shl.b32 %r5, %r4, 16;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r6, %rs7;
; CHECK-FTZ-NEXT: shl.b32 %r7, %r6, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r8, %r5, %r7;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs8, %r8;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs8;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16_expanded_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<4>;
; CHECK-SM70-NEXT: .reg .b32 %r<29>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }
; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;
; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;
; CHECK-SM70-NEXT: add.rn.f32 %r14, %r13, 0f40E00000;
; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;
; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r14, %r14;
; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p3;
; CHECK-SM70-NEXT: cvt.u32.u16 %r20, %rs2;
; CHECK-SM70-NEXT: shl.b32 %r21, %r20, 16;
; CHECK-SM70-NEXT: and.b32 %r22, %r19, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r23, %r21, %r22;
; CHECK-SM70-NEXT: bfe.u32 %r24, %r23, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r25, %r24, %r23;
; CHECK-SM70-NEXT: add.s32 %r26, %r25, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r23, %r23;
; CHECK-SM70-NEXT: or.b32 %r27, %r23, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r28, %r27, %r26, %p4;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs3}, %r28; }
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs3;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast bfloat %a, %b
%2 = fadd fast bfloat %1, %c
%3 = fcmp nsz ogt bfloat %2, 0.0
%4 = select nsz i1 %3, bfloat %2, bfloat 0.0
%5 = fadd bfloat %2, 7.0
%6 = fadd bfloat %4, %5
ret bfloat %6
}
define bfloat @fma_bf16_expanded_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) {
;
;
; CHECK-LABEL: fma_bf16_expanded_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16_expanded_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_expanded_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_expanded_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_expanded_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16_expanded_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<20>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_expanded_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_expanded_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_expanded_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;
; CHECK-SM70-NEXT: max.f32 %r14, %r13, 0f00000000;
; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;
; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;
; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r19; }
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs1;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast bfloat %a, %b
%2 = fadd fast bfloat %1, %c
%3 = call nsz bfloat @llvm.maxnum.bf16(bfloat %2, bfloat 0.0)
ret bfloat %3
}
define <2 x half> @fma_f16x2_expanded_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {
;
;
; CHECK-LABEL: fma_f16x2_expanded_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16x2_expanded_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16x2_expanded_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<6>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-SM70-NEXT: mov.b32 %r5, 0;
; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;
; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;
; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;
; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;
; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
; CHECK-SM70-NEXT: ret;
%1 = fmul fast <2 x half> %a, %b
%2 = fadd fast <2 x half> %1, %c
%3 = fcmp nsz ogt <2 x half> %2, <half 0.0, half 0.0>
%4 = select nsz <2 x i1> %3, <2 x half> %2, <2 x half> <half 0.0, half 0.0>
ret <2 x half> %4
}
; FMA relu shouldn't be selected if the FMA operation has multiple uses
define <2 x half> @fma_f16x2_expanded_no_nans_multiple_uses_of_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) {
;
;
; CHECK-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: mov.b32 %r5, 0;
; CHECK-NEXT: max.f16x2 %r6, %r4, %r5;
; CHECK-NEXT: mov.b32 %r7, 1191200512;
; CHECK-NEXT: add.rn.f16x2 %r8, %r4, %r7;
; CHECK-NEXT: add.rn.f16x2 %r9, %r6, %r8;
; CHECK-NEXT: st.param.b32 [func_retval0], %r9;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<10>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: mov.b32 %r5, 0;
; CHECK-FTZ-NEXT: max.ftz.f16x2 %r6, %r4, %r5;
; CHECK-FTZ-NEXT: mov.b32 %r7, 1191200512;
; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r8, %r4, %r7;
; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r9, %r6, %r8;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r9;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16x2_expanded_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<10>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-SM70-NEXT: mov.b32 %r5, 0;
; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;
; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;
; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;
; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;
; CHECK-SM70-NEXT: mov.b32 %r6, {%rs4, %rs3};
; CHECK-SM70-NEXT: mov.b32 %r7, 1191200512;
; CHECK-SM70-NEXT: add.rn.f16x2 %r8, %r4, %r7;
; CHECK-SM70-NEXT: add.rn.f16x2 %r9, %r6, %r8;
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast <2 x half> %a, %b
%2 = fadd fast <2 x half> %1, %c
%3 = fcmp nsz ogt <2 x half> %2, <half 0.0, half 0.0>
%4 = select nsz <2 x i1> %3, <2 x half> %2, <2 x half> <half 0.0, half 0.0>
%5 = fadd <2 x half> %2, <half 7.0, half 7.0>
%6 = fadd <2 x half> %4, %5
ret <2 x half> %6
}
define <2 x half> @fma_f16x2_expanded_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {
; CHECK-LABEL: fma_f16x2_expanded_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16x2_expanded_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16x2_expanded_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<10>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_expanded_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_expanded_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_expanded_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;
; CHECK-SM70-NEXT: cvt.f32.f16 %r5, %rs2;
; CHECK-SM70-NEXT: max.f32 %r6, %r5, 0f00000000;
; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs3, %r6;
; CHECK-SM70-NEXT: cvt.f32.f16 %r7, %rs1;
; CHECK-SM70-NEXT: max.f32 %r8, %r7, 0f00000000;
; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs4, %r8;
; CHECK-SM70-NEXT: mov.b32 %r9, {%rs4, %rs3};
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast <2 x half> %a, %b
%2 = fadd fast <2 x half> %1, %c
%3 = call nsz <2 x half> @llvm.maxnum.f16x2(<2 x half> %2, <2 x half> <half 0.0, half 0.0>)
ret <2 x half> %3
}
define <2 x bfloat> @fma_bf16x2_expanded_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {
; CHECK-LABEL: fma_bf16x2_expanded_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16x2_expanded_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16x2_expanded_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<27>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_no_nans_param_2];
; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs5;
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs3;
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }
; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs6;
; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs4;
; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;
; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;
; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;
; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;
; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }
; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;
; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;
; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;
; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;
; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};
; CHECK-SM70-NEXT: ret;
%1 = fmul fast <2 x bfloat> %a, %b
%2 = fadd fast <2 x bfloat> %1, %c
%3 = fcmp nsz ogt <2 x bfloat> %2, <bfloat 0.0, bfloat 0.0>
%4 = select nsz <2 x i1> %3, <2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>
ret <2 x bfloat> %4
}
; FMA relu shouldn't be selected if the FMA operation has multiple uses
define <2 x bfloat> @fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {
; CHECK-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<11>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: mov.b32 %r5, 0;
; CHECK-NEXT: max.bf16x2 %r6, %r4, %r5;
; CHECK-NEXT: mov.b32 %r7, 1065369472;
; CHECK-NEXT: mov.b32 %r8, 1088438496;
; CHECK-NEXT: fma.rn.bf16x2 %r9, %r4, %r7, %r8;
; CHECK-NEXT: fma.rn.bf16x2 %r10, %r6, %r7, %r9;
; CHECK-NEXT: st.param.b32 [func_retval0], %r10;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;
; CHECK-FTZ-NEXT: .reg .b32 %r<24>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: mov.b32 %r5, 0;
; CHECK-FTZ-NEXT: max.bf16x2 %r6, %r4, %r5;
; CHECK-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r4;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r7, %rs2;
; CHECK-FTZ-NEXT: shl.b32 %r8, %r7, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r9, %r8, 0f40E00000;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r9;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r10, %rs1;
; CHECK-FTZ-NEXT: shl.b32 %r11, %r10, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r12, %r11, 0f40E00000;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs4, %r12;
; CHECK-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r6;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r13, %rs5;
; CHECK-FTZ-NEXT: shl.b32 %r14, %r13, 16;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r15, %rs4;
; CHECK-FTZ-NEXT: shl.b32 %r16, %r15, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r17, %r14, %r16;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r18, %rs6;
; CHECK-FTZ-NEXT: shl.b32 %r19, %r18, 16;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r20, %rs3;
; CHECK-FTZ-NEXT: shl.b32 %r21, %r20, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r22, %r19, %r21;
; CHECK-FTZ-NEXT: cvt.rn.bf16x2.f32 %r23, %r22, %r17;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r23;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16x2_expanded_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<9>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<58>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs6;
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs4;
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs2;
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }
; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs5;
; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs3;
; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs1;
; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;
; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;
; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;
; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }
; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;
; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;
; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;
; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;
; CHECK-SM70-NEXT: add.rn.f32 %r27, %r26, 0f40E00000;
; CHECK-SM70-NEXT: bfe.u32 %r28, %r27, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r29, %r28, %r27;
; CHECK-SM70-NEXT: add.s32 %r30, %r29, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p5, %r27, %r27;
; CHECK-SM70-NEXT: or.b32 %r31, %r27, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r32, %r31, %r30, %p5;
; CHECK-SM70-NEXT: add.rn.f32 %r33, %r25, 0f40E00000;
; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;
; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p6, %r33, %r33;
; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p6;
; CHECK-SM70-NEXT: cvt.u32.u16 %r39, %rs10;
; CHECK-SM70-NEXT: shl.b32 %r40, %r39, 16;
; CHECK-SM70-NEXT: and.b32 %r41, %r38, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r42, %r40, %r41;
; CHECK-SM70-NEXT: bfe.u32 %r43, %r42, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r44, %r43, %r42;
; CHECK-SM70-NEXT: add.s32 %r45, %r44, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p7, %r42, %r42;
; CHECK-SM70-NEXT: or.b32 %r46, %r42, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r47, %r46, %r45, %p7;
; CHECK-SM70-NEXT: cvt.u32.u16 %r48, %rs9;
; CHECK-SM70-NEXT: shl.b32 %r49, %r48, 16;
; CHECK-SM70-NEXT: and.b32 %r50, %r32, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r51, %r49, %r50;
; CHECK-SM70-NEXT: bfe.u32 %r52, %r51, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r53, %r52, %r51;
; CHECK-SM70-NEXT: add.s32 %r54, %r53, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p8, %r51, %r51;
; CHECK-SM70-NEXT: or.b32 %r55, %r51, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r56, %r55, %r54, %p8;
; CHECK-SM70-NEXT: prmt.b32 %r57, %r56, %r47, 0x7632U;
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r57;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast <2 x bfloat> %a, %b
%2 = fadd fast <2 x bfloat> %1, %c
%3 = fcmp nsz ogt <2 x bfloat> %2, <bfloat 0.0, bfloat 0.0>
%4 = select nsz <2 x i1> %3, <2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>
%5 = fadd <2 x bfloat> %2, <bfloat 7.0, bfloat 7.0>
%6 = fadd <2 x bfloat> %4, %5
ret <2 x bfloat> %6
}
define <2 x bfloat> @fma_bf16x2_expanded_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {
; CHECK-LABEL: fma_bf16x2_expanded_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16x2_expanded_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_expanded_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_expanded_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_expanded_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16x2_expanded_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<40>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_expanded_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_expanded_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_expanded_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs5;
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs3;
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs6;
; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs4;
; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;
; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;
; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;
; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;
; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;
; CHECK-SM70-NEXT: max.f32 %r26, %r25, 0f00000000;
; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;
; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;
; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;
; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;
; CHECK-SM70-NEXT: max.f32 %r33, %r32, 0f00000000;
; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;
; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;
; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;
; CHECK-SM70-NEXT: prmt.b32 %r39, %r38, %r31, 0x7632U;
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r39;
; CHECK-SM70-NEXT: ret;
%1 = fmul fast <2 x bfloat> %a, %b
%2 = fadd fast <2 x bfloat> %1, %c
%3 = call nsz <2 x bfloat> @llvm.maxnum.bf16x2(<2 x bfloat> %2, <2 x bfloat> <bfloat 0.0, bfloat 0.0>)
ret <2 x bfloat> %3
}
define half @fma_f16_no_nans(half %a, half %b, half %c) {
; CHECK-LABEL: fma_f16_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<2>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-SM70-NEXT: mov.b16 %rs5, 0x0000;
; CHECK-SM70-NEXT: setp.gt.f16 %p1, %rs4, %rs5;
; CHECK-SM70-NEXT: selp.b16 %rs6, %rs4, 0x0000, %p1;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs6;
; CHECK-SM70-NEXT: ret;
%1 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)
%2 = fcmp nsz ogt half %1, 0.0
%3 = select nsz i1 %2, half %1, half 0.0
ret half %3
}
; FMA relu shouldn't be selected if the FMA operation has multiple uses
define half @fma_f16_no_nans_multiple_uses_of_fma(half %a, half %b, half %c) {
; CHECK-LABEL: fma_f16_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<8>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: mov.b16 %rs5, 0x4700;
; CHECK-NEXT: add.rn.f16 %rs6, %rs4, %rs5;
; CHECK-NEXT: add.rn.f16 %rs7, %rs6, %rs4;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs7;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<8>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: mov.b16 %rs5, 0x4700;
; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs6, %rs4, %rs5;
; CHECK-FTZ-NEXT: add.rn.ftz.f16 %rs7, %rs6, %rs4;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs7;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<8>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-SM70-NEXT: mov.b16 %rs5, 0x4700;
; CHECK-SM70-NEXT: add.rn.f16 %rs6, %rs4, %rs5;
; CHECK-SM70-NEXT: add.rn.f16 %rs7, %rs6, %rs4;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs7;
; CHECK-SM70-NEXT: ret;
%1 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)
%2 = fcmp nsz ogt half %1, 0.0
%3 = select nsz i1 %2, half %1, half 0.0
%4 = fadd half %1, 7.0
%5 = fadd half %4, %1
ret half %5
}
define half @fma_f16_maxnum_no_nans(half %a, half %b, half %c) {
; CHECK-LABEL: fma_f16_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<6>;
; CHECK-SM70-NEXT: .reg .b32 %r<3>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %rs1, [fma_f16_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b16 %rs2, [fma_f16_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b16 %rs3, [fma_f16_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16 %rs4, %rs1, %rs2, %rs3;
; CHECK-SM70-NEXT: cvt.f32.f16 %r1, %rs4;
; CHECK-SM70-NEXT: max.f32 %r2, %r1, 0f00000000;
; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs5, %r2;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs5;
; CHECK-SM70-NEXT: ret;
%1 = call nnan half @llvm.fma.f16(half %a, half %b, half %c)
%2 = call nsz half @llvm.maxnum.f16(half %1, half 0.0)
ret half %2
}
define bfloat @fma_bf16_no_nans(bfloat %a, bfloat %b, bfloat %c) {
; CHECK-LABEL: fma_bf16_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<3>;
; CHECK-SM70-NEXT: .reg .b32 %r<14>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_no_nans_param_2];
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_no_nans_param_1];
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_no_nans_param_0];
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r12; }
; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p2, %r13, 0f00000000;
; CHECK-SM70-NEXT: selp.b16 %rs2, %rs1, 0x0000, %p2;
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-SM70-NEXT: ret;
%1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)
%2 = fcmp nsz ogt bfloat %1, 0.0
%3 = select nsz i1 %2, bfloat %1, bfloat 0.0
ret bfloat %3
}
; FMA_relu shouldn't be selected if the FMA operation has multiple uses
define bfloat @fma_bf16_no_nans_multiple_uses_of_fma(bfloat %a, bfloat %b, bfloat %c) {
; CHECK-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: mov.b16 %rs5, 0x3F80;
; CHECK-NEXT: mov.b16 %rs6, 0x40E0;
; CHECK-NEXT: fma.rn.bf16 %rs7, %rs4, %rs5, %rs6;
; CHECK-NEXT: fma.rn.bf16 %rs8, %rs7, %rs5, %rs4;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs8;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<7>;
; CHECK-FTZ-NEXT: .reg .b32 %r<7>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-FTZ-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40E00000;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %r3;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r4, %rs5;
; CHECK-FTZ-NEXT: shl.b32 %r5, %r4, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r6, %r5, %r2;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %r6;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs6;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<4>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<27>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r14, %r13, 0f40E00000;
; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;
; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;
; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;
; CHECK-SM70-NEXT: and.b32 %r20, %r19, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r21, %r20, %r13;
; CHECK-SM70-NEXT: bfe.u32 %r22, %r21, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r23, %r22, %r21;
; CHECK-SM70-NEXT: add.s32 %r24, %r23, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r21, %r21;
; CHECK-SM70-NEXT: or.b32 %r25, %r21, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r26, %r25, %r24, %p3;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r26; }
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs1;
; CHECK-SM70-NEXT: ret;
%1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)
%2 = fcmp nsz ogt bfloat %1, 0.0
%3 = select nsz i1 %2, bfloat %1, bfloat 0.0
%4 = fadd bfloat %1, 7.0
%5 = fadd bfloat %4, %1
ret bfloat %5
}
define bfloat @fma_bf16_maxnum_no_nans(bfloat %a, bfloat %b, bfloat %c) {
; CHECK-LABEL: fma_bf16_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [fma_bf16_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b16 %rs2, [fma_bf16_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b16 %rs3, [fma_bf16_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b16 %rs1, [fma_bf16_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b16 %rs2, [fma_bf16_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b16 %rs3, [fma_bf16_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16 %rs4, %rs1, %rs2, %rs3;
; CHECK-FTZ-NEXT: st.param.b16 [func_retval0], %rs4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<2>;
; CHECK-SM70-NEXT: .reg .b32 %r<20>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b16 %r1, [fma_bf16_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r3, [fma_bf16_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: ld.param.b16 %r5, [fma_bf16_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: and.b32 %r13, %r12, -65536;
; CHECK-SM70-NEXT: max.f32 %r14, %r13, 0f00000000;
; CHECK-SM70-NEXT: bfe.u32 %r15, %r14, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r16, %r15, %r14;
; CHECK-SM70-NEXT: add.s32 %r17, %r16, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r14, %r14;
; CHECK-SM70-NEXT: or.b32 %r18, %r14, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r19, %r18, %r17, %p2;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r19; }
; CHECK-SM70-NEXT: st.param.b16 [func_retval0], %rs1;
; CHECK-SM70-NEXT: ret;
%1 = call nnan bfloat @llvm.fma.bf16(bfloat %a, bfloat %b, bfloat %c)
%2 = call nsz bfloat @llvm.maxnum.bf16(bfloat %1, bfloat 0.0)
ret bfloat %2
}
define <2 x half> @fma_f16x2_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {
; CHECK-LABEL: fma_f16x2_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];
; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: mov.b32 %r5, 0;
; CHECK-NEXT: max.f16x2 %r6, %r4, %r5;
; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16x2_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<7>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: mov.b32 %r5, 0;
; CHECK-FTZ-NEXT: max.ftz.f16x2 %r6, %r4, %r5;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r6;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16x2_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<3>;
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<6>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-SM70-NEXT: mov.b32 %r5, 0;
; CHECK-SM70-NEXT: setp.gt.f16x2 %p1|%p2, %r4, %r5;
; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;
; CHECK-SM70-NEXT: selp.b16 %rs3, %rs2, 0x0000, %p2;
; CHECK-SM70-NEXT: selp.b16 %rs4, %rs1, 0x0000, %p1;
; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
; CHECK-SM70-NEXT: ret;
%1 = call <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)
%2 = fcmp nsz ogt <2 x half> %1, <half 0.0, half 0.0>
%3 = select nsz <2 x i1> %2, <2 x half> %1, <2 x half> <half 0.0, half 0.0>
ret <2 x half> %3
}
; FMA relu shouldn't be selected if the FMA operation has multiple uses
define <2 x half> @fma_f16x2_no_nans_multiple_uses_of_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) {
; CHECK-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<8>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: mov.b32 %r5, 1191200512;
; CHECK-NEXT: add.rn.f16x2 %r6, %r4, %r5;
; CHECK-NEXT: add.rn.f16x2 %r7, %r6, %r4;
; CHECK-NEXT: st.param.b32 [func_retval0], %r7;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<8>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.f16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: mov.b32 %r5, 1191200512;
; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r6, %r4, %r5;
; CHECK-FTZ-NEXT: add.rn.ftz.f16x2 %r7, %r6, %r4;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r7;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16x2_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b32 %r<8>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-SM70-NEXT: mov.b32 %r5, 1191200512;
; CHECK-SM70-NEXT: add.rn.f16x2 %r6, %r4, %r5;
; CHECK-SM70-NEXT: add.rn.f16x2 %r7, %r6, %r4;
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r7;
; CHECK-SM70-NEXT: ret;
%1 = call nnan <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)
%2 = fcmp nsz ogt <2 x half> %1, <half 0.0, half 0.0>
%3 = select nsz <2 x i1> %2, <2 x half> %1, <2 x half> <half 0.0, half 0.0>
%4 = fadd <2 x half> %1, <half 7.0, half 7.0>
%5 = fadd <2 x half> %4, %1
ret <2 x half> %5
}
define <2 x half> @fma_f16x2_maxnum_no_nans(<2 x half> %a, <2 x half> %b, <2 x half> %c) {
; CHECK-LABEL: fma_f16x2_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.f16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_f16x2_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.ftz.relu.f16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_f16x2_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .b16 %rs<5>;
; CHECK-SM70-NEXT: .reg .b32 %r<10>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.b32 %r1, [fma_f16x2_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.b32 %r2, [fma_f16x2_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: ld.param.b32 %r3, [fma_f16x2_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: fma.rn.f16x2 %r4, %r1, %r2, %r3;
; CHECK-SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;
; CHECK-SM70-NEXT: cvt.f32.f16 %r5, %rs2;
; CHECK-SM70-NEXT: max.f32 %r6, %r5, 0f00000000;
; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs3, %r6;
; CHECK-SM70-NEXT: cvt.f32.f16 %r7, %rs1;
; CHECK-SM70-NEXT: max.f32 %r8, %r7, 0f00000000;
; CHECK-SM70-NEXT: cvt.rn.f16.f32 %rs4, %r8;
; CHECK-SM70-NEXT: mov.b32 %r9, {%rs4, %rs3};
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r9;
; CHECK-SM70-NEXT: ret;
%1 = call nnan <2 x half> @llvm.fma.f16x2(<2 x half> %a, <2 x half> %b, <2 x half> %c)
%2 = call nsz <2 x half> @llvm.maxnum.f16x2(<2 x half> %1, <2 x half> <half 0.0, half 0.0>)
ret <2 x half> %2
}
define <2 x bfloat> @fma_bf16x2_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {
; CHECK-LABEL: fma_bf16x2_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16x2_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16x2_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<11>;
; CHECK-SM70-NEXT: .reg .b32 %r<27>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_no_nans_param_2];
; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs3;
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_no_nans_param_1];
; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs5;
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs7}, %r12; }
; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs4;
; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs6;
; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;
; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;
; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;
; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;
; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-SM70-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs8}, %r24; }
; CHECK-SM70-NEXT: and.b32 %r25, %r12, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p3, %r25, 0f00000000;
; CHECK-SM70-NEXT: and.b32 %r26, %r24, -65536;
; CHECK-SM70-NEXT: setp.gt.f32 %p4, %r26, 0f00000000;
; CHECK-SM70-NEXT: selp.b16 %rs9, %rs8, 0x0000, %p4;
; CHECK-SM70-NEXT: selp.b16 %rs10, %rs7, 0x0000, %p3;
; CHECK-SM70-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};
; CHECK-SM70-NEXT: ret;
%1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)
%2 = fcmp nsz ogt <2 x bfloat> %1, <bfloat 0.0, bfloat 0.0>
%3 = select nsz <2 x i1> %2, <2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>
ret <2 x bfloat> %3
}
; FMA_relu shouldn't be selected if the FMA operation has multiple uses
define <2 x bfloat> @fma_bf16x2_no_nans_multiple_uses_of_fma(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {
; CHECK-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<9>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];
; CHECK-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: mov.b32 %r5, 1065369472;
; CHECK-NEXT: mov.b32 %r6, 1088438496;
; CHECK-NEXT: fma.rn.bf16x2 %r7, %r4, %r5, %r6;
; CHECK-NEXT: fma.rn.bf16x2 %r8, %r7, %r5, %r4;
; CHECK-NEXT: st.param.b32 [func_retval0], %r8;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b16 %rs<5>;
; CHECK-FTZ-NEXT: .reg .b32 %r<18>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];
; CHECK-FTZ-NEXT: fma.rn.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r4;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r5, %rs2;
; CHECK-FTZ-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r7, %r6, 0f40E00000;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs3, %r7;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r8, %rs1;
; CHECK-FTZ-NEXT: shl.b32 %r9, %r8, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r10, %r9, 0f40E00000;
; CHECK-FTZ-NEXT: cvt.rn.bf16.f32 %rs4, %r10;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r11, %rs4;
; CHECK-FTZ-NEXT: shl.b32 %r12, %r11, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r13, %r12, %r9;
; CHECK-FTZ-NEXT: cvt.u32.u16 %r14, %rs3;
; CHECK-FTZ-NEXT: shl.b32 %r15, %r14, 16;
; CHECK-FTZ-NEXT: add.rn.ftz.f32 %r16, %r15, %r6;
; CHECK-FTZ-NEXT: cvt.rn.bf16x2.f32 %r17, %r16, %r13;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r17;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16x2_no_nans_multiple_uses_of_fma(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<7>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<54>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_0];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_2];
; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs4;
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_no_nans_multiple_uses_of_fma_param_1];
; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs6;
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs2;
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs3;
; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs5;
; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs1;
; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;
; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;
; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;
; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r26, %r25, 0f40E00000;
; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;
; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;
; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;
; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r33, %r32, 0f40E00000;
; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;
; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;
; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;
; CHECK-SM70-NEXT: and.b32 %r39, %r38, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r40, %r39, %r32;
; CHECK-SM70-NEXT: bfe.u32 %r41, %r40, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r42, %r41, %r40;
; CHECK-SM70-NEXT: add.s32 %r43, %r42, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p5, %r40, %r40;
; CHECK-SM70-NEXT: or.b32 %r44, %r40, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r45, %r44, %r43, %p5;
; CHECK-SM70-NEXT: and.b32 %r46, %r31, -65536;
; CHECK-SM70-NEXT: add.rn.f32 %r47, %r46, %r25;
; CHECK-SM70-NEXT: bfe.u32 %r48, %r47, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r49, %r48, %r47;
; CHECK-SM70-NEXT: add.s32 %r50, %r49, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p6, %r47, %r47;
; CHECK-SM70-NEXT: or.b32 %r51, %r47, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r52, %r51, %r50, %p6;
; CHECK-SM70-NEXT: prmt.b32 %r53, %r52, %r45, 0x7632U;
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r53;
; CHECK-SM70-NEXT: ret;
%1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)
%2 = fcmp nsz ogt <2 x bfloat> %1, <bfloat 0.0, bfloat 0.0>
%3 = select nsz <2 x i1> %2, <2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>
%4 = fadd <2 x bfloat> %1, <bfloat 7.0, bfloat 7.0>
%5 = fadd <2 x bfloat> %4, %1
ret <2 x bfloat> %5
}
define <2 x bfloat> @fma_bf16x2_maxnum_no_nans(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c) {
; CHECK-LABEL: fma_bf16x2_maxnum_no_nans(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];
; CHECK-NEXT: ld.param.b32 %r2, [fma_bf16x2_maxnum_no_nans_param_1];
; CHECK-NEXT: ld.param.b32 %r3, [fma_bf16x2_maxnum_no_nans_param_2];
; CHECK-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NEXT: ret;
;
; CHECK-FTZ-LABEL: fma_bf16x2_maxnum_no_nans(
; CHECK-FTZ: {
; CHECK-FTZ-NEXT: .reg .b32 %r<5>;
; CHECK-FTZ-EMPTY:
; CHECK-FTZ-NEXT: // %bb.0:
; CHECK-FTZ-NEXT: ld.param.b32 %r1, [fma_bf16x2_maxnum_no_nans_param_0];
; CHECK-FTZ-NEXT: ld.param.b32 %r2, [fma_bf16x2_maxnum_no_nans_param_1];
; CHECK-FTZ-NEXT: ld.param.b32 %r3, [fma_bf16x2_maxnum_no_nans_param_2];
; CHECK-FTZ-NEXT: fma.rn.relu.bf16x2 %r4, %r1, %r2, %r3;
; CHECK-FTZ-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-FTZ-NEXT: ret;
;
; CHECK-SM70-LABEL: fma_bf16x2_maxnum_no_nans(
; CHECK-SM70: {
; CHECK-SM70-NEXT: .reg .pred %p<5>;
; CHECK-SM70-NEXT: .reg .b16 %rs<7>;
; CHECK-SM70-NEXT: .reg .b32 %r<40>;
; CHECK-SM70-EMPTY:
; CHECK-SM70-NEXT: // %bb.0:
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [fma_bf16x2_maxnum_no_nans_param_0];
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [fma_bf16x2_maxnum_no_nans_param_2];
; CHECK-SM70-NEXT: cvt.u32.u16 %r1, %rs3;
; CHECK-SM70-NEXT: shl.b32 %r2, %r1, 16;
; CHECK-SM70-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [fma_bf16x2_maxnum_no_nans_param_1];
; CHECK-SM70-NEXT: cvt.u32.u16 %r3, %rs5;
; CHECK-SM70-NEXT: shl.b32 %r4, %r3, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r5, %rs1;
; CHECK-SM70-NEXT: shl.b32 %r6, %r5, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
; CHECK-SM70-NEXT: bfe.u32 %r8, %r7, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r9, %r8, %r7;
; CHECK-SM70-NEXT: add.s32 %r10, %r9, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p1, %r7, %r7;
; CHECK-SM70-NEXT: or.b32 %r11, %r7, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r12, %r11, %r10, %p1;
; CHECK-SM70-NEXT: cvt.u32.u16 %r13, %rs4;
; CHECK-SM70-NEXT: shl.b32 %r14, %r13, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r15, %rs6;
; CHECK-SM70-NEXT: shl.b32 %r16, %r15, 16;
; CHECK-SM70-NEXT: cvt.u32.u16 %r17, %rs2;
; CHECK-SM70-NEXT: shl.b32 %r18, %r17, 16;
; CHECK-SM70-NEXT: fma.rn.f32 %r19, %r18, %r16, %r14;
; CHECK-SM70-NEXT: bfe.u32 %r20, %r19, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r21, %r20, %r19;
; CHECK-SM70-NEXT: add.s32 %r22, %r21, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p2, %r19, %r19;
; CHECK-SM70-NEXT: or.b32 %r23, %r19, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r24, %r23, %r22, %p2;
; CHECK-SM70-NEXT: and.b32 %r25, %r24, -65536;
; CHECK-SM70-NEXT: max.f32 %r26, %r25, 0f00000000;
; CHECK-SM70-NEXT: bfe.u32 %r27, %r26, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r28, %r27, %r26;
; CHECK-SM70-NEXT: add.s32 %r29, %r28, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p3, %r26, %r26;
; CHECK-SM70-NEXT: or.b32 %r30, %r26, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r31, %r30, %r29, %p3;
; CHECK-SM70-NEXT: and.b32 %r32, %r12, -65536;
; CHECK-SM70-NEXT: max.f32 %r33, %r32, 0f00000000;
; CHECK-SM70-NEXT: bfe.u32 %r34, %r33, 16, 1;
; CHECK-SM70-NEXT: add.s32 %r35, %r34, %r33;
; CHECK-SM70-NEXT: add.s32 %r36, %r35, 32767;
; CHECK-SM70-NEXT: setp.nan.f32 %p4, %r33, %r33;
; CHECK-SM70-NEXT: or.b32 %r37, %r33, 4194304;
; CHECK-SM70-NEXT: selp.b32 %r38, %r37, %r36, %p4;
; CHECK-SM70-NEXT: prmt.b32 %r39, %r38, %r31, 0x7632U;
; CHECK-SM70-NEXT: st.param.b32 [func_retval0], %r39;
; CHECK-SM70-NEXT: ret;
%1 = call nnan <2 x bfloat> @llvm.fma.bf16x2(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bfloat> %c)
%2 = call nsz <2 x bfloat> @llvm.maxnum.bf16x2(<2 x bfloat> %1, <2 x bfloat> <bfloat 0.0, bfloat 0.0>)
ret <2 x bfloat> %2
}