[llvm] acff049 - [NVPTX] Add max/minimumnum to ISel (#155804)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Aug 29 02:32:11 PDT 2025
Author: Lewis Crawford
Date: 2025-08-29T10:32:07+01:00
New Revision: acff04955defe719822287e9b854acaa07f1391c
URL: https://github.com/llvm/llvm-project/commit/acff04955defe719822287e9b854acaa07f1391c
DIFF: https://github.com/llvm/llvm-project/commit/acff04955defe719822287e9b854acaa07f1391c.diff
LOG: [NVPTX] Add max/minimumnum to ISel (#155804)
Add direct support for the LLVM `maximumnum` and `minimumnum`
intrinsics for NVPTX, rather than lowering them to a sequence of
compare + select instructions.
The `maximumnum` and `minimumnum` intrinsics map directly to PTX
`max`/`min` instructions.
In future, the LLVM `maxnum`/`minnum` intrinsics might need some fix-ups
for sNaN handling added, but currently both `llvm.maxnum` and
`llvm.maximumnum` will map directly to PTX `max` instructions.
Added:
Modified:
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/test/CodeGen/NVPTX/math-intrins.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 7d98bd824bf9d..d3fb657851fe2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -543,6 +543,8 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
case ISD::FMINNUM_IEEE:
case ISD::FMAXIMUM:
case ISD::FMINIMUM:
+ case ISD::FMAXIMUMNUM:
+ case ISD::FMINIMUMNUM:
IsOpSupported &= STI.getSmVersion() >= 80 && STI.getPTXVersion() >= 70;
break;
case ISD::FEXP2:
@@ -989,7 +991,8 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
if (getOperationAction(ISD::FABS, MVT::bf16) == Promote)
AddPromotedToType(ISD::FABS, MVT::bf16, MVT::f32);
- for (const auto &Op : {ISD::FMINNUM, ISD::FMAXNUM}) {
+ for (const auto &Op :
+ {ISD::FMINNUM, ISD::FMAXNUM, ISD::FMINIMUMNUM, ISD::FMAXIMUMNUM}) {
setOperationAction(Op, MVT::f32, Legal);
setOperationAction(Op, MVT::f64, Legal);
setFP16OperationAction(Op, MVT::f16, Legal, Promote);
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 7f29c3788d810..4e38e026e6bda 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -295,7 +295,7 @@ multiclass ADD_SUB_INT_CARRY<string op_str, SDNode op_node, bit commutative> {
//
// Also defines ftz (flush subnormal inputs and results to sign-preserving
// zero) variants for fp32 functions.
-multiclass FMINIMUMMAXIMUM<string OpcStr, bit NaN, SDNode OpNode> {
+multiclass FMINIMUMMAXIMUM<string OpcStr, bit NaN, SDPatternOperator OpNode> {
defvar nan_str = !if(NaN, ".NaN", "");
if !not(NaN) then {
def _f64_rr :
@@ -911,8 +911,15 @@ defm FADD : F3_fma_component<"add", fadd>;
defm FSUB : F3_fma_component<"sub", fsub>;
defm FMUL : F3_fma_component<"mul", fmul>;
-defm MIN : FMINIMUMMAXIMUM<"min", /* NaN */ false, fminnum>;
-defm MAX : FMINIMUMMAXIMUM<"max", /* NaN */ false, fmaxnum>;
+def fminnum_or_fminimumnum : PatFrags<(ops node:$a, node:$b),
+ [(fminnum node:$a, node:$b),
+ (fminimumnum node:$a, node:$b)]>;
+def fmaxnum_or_fmaximumnum : PatFrags<(ops node:$a, node:$b),
+ [(fmaxnum node:$a, node:$b),
+ (fmaximumnum node:$a, node:$b)]>;
+
+defm MIN : FMINIMUMMAXIMUM<"min", /* NaN */ false, fminnum_or_fminimumnum>;
+defm MAX : FMINIMUMMAXIMUM<"max", /* NaN */ false, fmaxnum_or_fmaximumnum>;
defm MIN_NAN : FMINIMUMMAXIMUM<"min", /* NaN */ true, fminimum>;
defm MAX_NAN : FMINIMUMMAXIMUM<"max", /* NaN */ true, fmaximum>;
diff --git a/llvm/test/CodeGen/NVPTX/math-intrins.ll b/llvm/test/CodeGen/NVPTX/math-intrins.ll
index e9635e9393984..3cbb6b680fc05 100644
--- a/llvm/test/CodeGen/NVPTX/math-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/math-intrins.ll
@@ -42,6 +42,14 @@ declare half @llvm.maximum.f16(half, half) #0
declare float @llvm.maximum.f32(float, float) #0
declare double @llvm.maximum.f64(double, double) #0
declare <2 x half> @llvm.maximum.v2f16(<2 x half>, <2 x half>) #0
+declare half @llvm.minimumnum.f16(half, half) #0
+declare float @llvm.minimumnum.f32(float, float) #0
+declare double @llvm.minimumnum.f64(double, double) #0
+declare <2 x half> @llvm.minimumnum.v2f16(<2 x half>, <2 x half>) #0
+declare half @llvm.maximumnum.f16(half, half) #0
+declare float @llvm.maximumnum.f32(float, float) #0
+declare double @llvm.maximumnum.f64(double, double) #0
+declare <2 x half> @llvm.maximumnum.v2f16(<2 x half>, <2 x half>) #0
declare float @llvm.fma.f32(float, float, float) #0
declare double @llvm.fma.f64(double, double, double) #0
@@ -1486,6 +1494,410 @@ define <2 x half> @maximum_v2half(<2 x half> %a, <2 x half> %b) {
ret <2 x half> %x
}
+; ---- minimumnum ----
+
+define half @minimumnum_half(half %a, half %b) {
+; CHECK-NOF16-LABEL: minimumnum_half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minimumnum_half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [minimumnum_half_param_1];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;
+; CHECK-NOF16-NEXT: min.f32 %r3, %r2, %r1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
+; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimumnum_half(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .b16 %rs<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.b16 %rs1, [minimumnum_half_param_0];
+; CHECK-F16-NEXT: ld.param.b16 %rs2, [minimumnum_half_param_1];
+; CHECK-F16-NEXT: min.f16 %rs3, %rs1, %rs2;
+; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimumnum_half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minimumnum_half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [minimumnum_half_param_1];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;
+; CHECK-SM80-NOF16-NEXT: min.f32 %r3, %r2, %r1;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
+; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call half @llvm.minimumnum.f16(half %a, half %b)
+ ret half %x
+}
+
+define float @minimumnum_float(float %a, float %b) {
+; CHECK-LABEL: minimumnum_float(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [minimumnum_float_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [minimumnum_float_param_1];
+; CHECK-NEXT: min.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %x = call float @llvm.minimumnum.f32(float %a, float %b)
+ ret float %x
+}
+
+define float @minimumnum_float_ftz(float %a, float %b) #1 {
+; CHECK-LABEL: minimumnum_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [minimumnum_float_ftz_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [minimumnum_float_ftz_param_1];
+; CHECK-NEXT: min.ftz.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %x = call float @llvm.minimumnum.f32(float %a, float %b)
+ ret float %x
+}
+
+define double @minimumnum_double(double %a, double %b) {
+; CHECK-LABEL: minimumnum_double(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [minimumnum_double_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [minimumnum_double_param_1];
+; CHECK-NEXT: min.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %x = call double @llvm.minimumnum.f64(double %a, double %b)
+ ret double %x
+}
+
+; TODO Improve the "Expand" path for minimumnum vectors on targets where
+; f16 is not supported. Ideally it should use two f32 minimumnums first instead of
+; fully expanding the minimumnum instruction into compare/select instructions.
+define <2 x half> @minimumnum_v2half(<2 x half> %a, <2 x half> %b) {
+; CHECK-NOF16-LABEL: minimumnum_v2half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<13>;
+; CHECK-NOF16-NEXT: .reg .b16 %rs<17>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<11>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minimumnum_v2half_param_0];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minimumnum_v2half_param_1];
+; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r3, %r3;
+; CHECK-NOF16-NEXT: selp.b16 %rs6, %rs5, %rs4, %p2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs6;
+; CHECK-NOF16-NEXT: setp.lt.f32 %p3, %r2, %r4;
+; CHECK-NOF16-NEXT: selp.b16 %rs7, %rs5, %rs6, %p3;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p4, %rs5, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs5, %rs7, %p4;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p5, %rs6, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs6, %rs8, %p5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs7;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p6, %r5, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs9, %rs7, %p6;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p7, %r6, %r6;
+; CHECK-NOF16-NEXT: selp.b16 %rs11, %rs3, %rs1, %p7;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs11;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p8, %r8, %r8;
+; CHECK-NOF16-NEXT: selp.b16 %rs12, %rs11, %rs3, %p8;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs12;
+; CHECK-NOF16-NEXT: setp.lt.f32 %p9, %r7, %r9;
+; CHECK-NOF16-NEXT: selp.b16 %rs13, %rs11, %rs12, %p9;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p10, %rs11, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs14, %rs11, %rs13, %p10;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p11, %rs12, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs15, %rs12, %rs14, %p11;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs13;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p12, %r10, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs16, %rs15, %rs13, %p12;
+; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs16, %rs10};
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimumnum_v2half(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .b32 %r<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.b32 %r1, [minimumnum_v2half_param_0];
+; CHECK-F16-NEXT: ld.param.b32 %r2, [minimumnum_v2half_param_1];
+; CHECK-F16-NEXT: min.f16x2 %r3, %r1, %r2;
+; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimumnum_v2half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .pred %p<13>;
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<17>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<11>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [minimumnum_v2half_param_0];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [minimumnum_v2half_param_1];
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs5;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %r3, %r3;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, %rs5, %rs4, %p2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs6;
+; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p3, %r2, %r4;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs7, %rs5, %rs6, %p3;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p4, %rs5, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs5, %rs7, %p4;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p5, %rs6, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs9, %rs6, %rs8, %p5;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs7;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p6, %r5, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs10, %rs9, %rs7, %p6;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p7, %r6, %r6;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs11, %rs3, %rs1, %p7;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r7, %rs11;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p8, %r8, %r8;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs12, %rs11, %rs3, %p8;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r9, %rs12;
+; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p9, %r7, %r9;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs13, %rs11, %rs12, %p9;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p10, %rs11, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs14, %rs11, %rs13, %p10;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p11, %rs12, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs15, %rs12, %rs14, %p11;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r10, %rs13;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p12, %r10, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs16, %rs15, %rs13, %p12;
+; CHECK-SM80-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs16, %rs10};
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call <2 x half> @llvm.minimumnum.v2f16(<2 x half> %a, <2 x half> %b)
+ ret <2 x half> %x
+}
+
+; ---- maximumnum ----
+
+define half @maximumnum_half(half %a, half %b) {
+; CHECK-NOF16-LABEL: maximumnum_half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maximumnum_half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [maximumnum_half_param_1];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;
+; CHECK-NOF16-NEXT: max.f32 %r3, %r2, %r1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
+; CHECK-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximumnum_half(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .b16 %rs<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.b16 %rs1, [maximumnum_half_param_0];
+; CHECK-F16-NEXT: ld.param.b16 %rs2, [maximumnum_half_param_1];
+; CHECK-F16-NEXT: max.f16 %rs3, %rs1, %rs2;
+; CHECK-F16-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximumnum_half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maximumnum_half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [maximumnum_half_param_1];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs1;
+; CHECK-SM80-NOF16-NEXT: max.f32 %r3, %r2, %r1;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
+; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0], %rs3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call half @llvm.maximumnum.f16(half %a, half %b)
+ ret half %x
+}
+
+define float @maximumnum_float(float %a, float %b) {
+; CHECK-LABEL: maximumnum_float(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [maximumnum_float_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [maximumnum_float_param_1];
+; CHECK-NEXT: max.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %x = call float @llvm.maximumnum.f32(float %a, float %b)
+ ret float %x
+}
+
+define float @maximumnum_float_ftz(float %a, float %b) #1 {
+; CHECK-LABEL: maximumnum_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [maximumnum_float_ftz_param_0];
+; CHECK-NEXT: ld.param.b32 %r2, [maximumnum_float_ftz_param_1];
+; CHECK-NEXT: max.ftz.f32 %r3, %r1, %r2;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %x = call float @llvm.maximumnum.f32(float %a, float %b)
+ ret float %x
+}
+
+define double @maximumnum_double(double %a, double %b) {
+; CHECK-LABEL: maximumnum_double(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b64 %rd1, [maximumnum_double_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [maximumnum_double_param_1];
+; CHECK-NEXT: max.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
+; CHECK-NEXT: ret;
+ %x = call double @llvm.maximumnum.f64(double %a, double %b)
+ ret double %x
+}
+
+; TODO Improve the "Expand" path for maximumnum vectors on targets where
+; f16 is not supported. Ideally it should use two f32 maximumnums first instead of
+; fully expanding the maximumnum instruction into compare/select instructions.
+define <2 x half> @maximumnum_v2half(<2 x half> %a, <2 x half> %b) {
+; CHECK-NOF16-LABEL: maximumnum_v2half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<13>;
+; CHECK-NOF16-NEXT: .reg .b16 %rs<17>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<11>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maximumnum_v2half_param_0];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maximumnum_v2half_param_1];
+; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r3, %r3;
+; CHECK-NOF16-NEXT: selp.b16 %rs6, %rs5, %rs4, %p2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs6;
+; CHECK-NOF16-NEXT: setp.gt.f32 %p3, %r2, %r4;
+; CHECK-NOF16-NEXT: selp.b16 %rs7, %rs5, %rs6, %p3;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p4, %rs5, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs5, %rs7, %p4;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p5, %rs6, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs6, %rs8, %p5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs7;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p6, %r5, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs9, %rs7, %p6;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p7, %r6, %r6;
+; CHECK-NOF16-NEXT: selp.b16 %rs11, %rs3, %rs1, %p7;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs11;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p8, %r8, %r8;
+; CHECK-NOF16-NEXT: selp.b16 %rs12, %rs11, %rs3, %p8;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs12;
+; CHECK-NOF16-NEXT: setp.gt.f32 %p9, %r7, %r9;
+; CHECK-NOF16-NEXT: selp.b16 %rs13, %rs11, %rs12, %p9;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p10, %rs11, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs14, %rs11, %rs13, %p10;
+; CHECK-NOF16-NEXT: setp.eq.b16 %p11, %rs12, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs15, %rs12, %rs14, %p11;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs13;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p12, %r10, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs16, %rs15, %rs13, %p12;
+; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs16, %rs10};
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximumnum_v2half(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .b32 %r<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.b32 %r1, [maximumnum_v2half_param_0];
+; CHECK-F16-NEXT: ld.param.b32 %r2, [maximumnum_v2half_param_1];
+; CHECK-F16-NEXT: max.f16x2 %r3, %r1, %r2;
+; CHECK-F16-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximumnum_v2half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .pred %p<13>;
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<17>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<11>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [maximumnum_v2half_param_0];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r1, %rs2;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-SM80-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [maximumnum_v2half_param_1];
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r2, %rs5;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %r3, %r3;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, %rs5, %rs4, %p2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r4, %rs6;
+; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p3, %r2, %r4;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs7, %rs5, %rs6, %p3;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p4, %rs5, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs5, %rs7, %p4;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p5, %rs6, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs9, %rs6, %rs8, %p5;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r5, %rs7;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p6, %r5, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs10, %rs9, %rs7, %p6;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p7, %r6, %r6;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs11, %rs3, %rs1, %p7;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r7, %rs11;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p8, %r8, %r8;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs12, %rs11, %rs3, %p8;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r9, %rs12;
+; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p9, %r7, %r9;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs13, %rs11, %rs12, %p9;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p10, %rs11, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs14, %rs11, %rs13, %p10;
+; CHECK-SM80-NOF16-NEXT: setp.eq.b16 %p11, %rs12, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs15, %rs12, %rs14, %p11;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %r10, %rs13;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p12, %r10, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs16, %rs15, %rs13, %p12;
+; CHECK-SM80-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs16, %rs10};
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call <2 x half> @llvm.maximumnum.v2f16(<2 x half> %a, <2 x half> %b)
+ ret <2 x half> %x
+}
+
; ---- fma ----
define float @fma_float(float %a, float %b, float %c) {
More information about the llvm-commits
mailing list