[llvm] 6f318d4 - [NVPTX] Make minimum/maximum work on older GPUs
David Majnemer via llvm-commits
llvm-commits at lists.llvm.org
Wed Jul 31 15:38:59 PDT 2024
Author: David Majnemer
Date: 2024-07-31T22:38:19Z
New Revision: 6f318d47bfba56df65394db6c20befe3ed5bf243
URL: https://github.com/llvm/llvm-project/commit/6f318d47bfba56df65394db6c20befe3ed5bf243
DIFF: https://github.com/llvm/llvm-project/commit/6f318d47bfba56df65394db6c20befe3ed5bf243.diff
LOG: [NVPTX] Make minimum/maximum work on older GPUs
We want to use newer instructions if we are targeting sufficiently new
SM and PTX versions. If we cannot use those newer instructions, let LLVM
synthesize the sequence from more fundamental instructions.
Added:
Modified:
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/test/CodeGen/NVPTX/math-intrins.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 91b239a52d17f..a5bdc6fac3ca6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -268,16 +268,12 @@ multiclass ADD_SUB_INT_CARRY<string OpcStr, SDNode OpNode> {
}
}
-// Template for instructions which take three fp64 or fp32 args. The
-// instructions are named "<OpcStr>.f<Width>" (e.g. "min.f64").
+// Template for minimum/maximum instructions.
//
// Also defines ftz (flush subnormal inputs and results to sign-preserving
// zero) variants for fp32 functions.
-//
-// This multiclass should be used for nodes that cannot be folded into FMAs.
-// For nodes that can be folded into FMAs (i.e. adds and muls), use
-// F3_fma_component.
-multiclass F3<string OpcStr, SDNode OpNode> {
+multiclass FMINIMUMMAXIMUM<string OpcStr, bit NaN, SDNode OpNode> {
+ if !not(NaN) then {
def f64rr :
NVPTXInst<(outs Float64Regs:$dst),
(ins Float64Regs:$a, Float64Regs:$b),
@@ -288,6 +284,7 @@ multiclass F3<string OpcStr, SDNode OpNode> {
(ins Float64Regs:$a, f64imm:$b),
!strconcat(OpcStr, ".f64 \t$dst, $a, $b;"),
[(set Float64Regs:$dst, (OpNode Float64Regs:$a, fpimm:$b))]>;
+ }
def f32rr_ftz :
NVPTXInst<(outs Float32Regs:$dst),
(ins Float32Regs:$a, Float32Regs:$b),
@@ -322,45 +319,45 @@ multiclass F3<string OpcStr, SDNode OpNode> {
(ins Int16Regs:$a, Int16Regs:$b),
!strconcat(OpcStr, ".f16 \t$dst, $a, $b;"),
[(set Int16Regs:$dst, (OpNode (f16 Int16Regs:$a), (f16 Int16Regs:$b)))]>,
- Requires<[useFP16Math]>;
+ Requires<[useFP16Math, hasSM<80>, hasPTX<70>]>;
def f16x2rr_ftz :
NVPTXInst<(outs Int32Regs:$dst),
(ins Int32Regs:$a, Int32Regs:$b),
!strconcat(OpcStr, ".ftz.f16x2 \t$dst, $a, $b;"),
[(set Int32Regs:$dst, (OpNode (v2f16 Int32Regs:$a), (v2f16 Int32Regs:$b)))]>,
- Requires<[useFP16Math, doF32FTZ]>;
+ Requires<[useFP16Math, hasSM<80>, hasPTX<70>, doF32FTZ]>;
def f16x2rr :
NVPTXInst<(outs Int32Regs:$dst),
(ins Int32Regs:$a, Int32Regs:$b),
!strconcat(OpcStr, ".f16x2 \t$dst, $a, $b;"),
[(set Int32Regs:$dst, (OpNode (v2f16 Int32Regs:$a), (v2f16 Int32Regs:$b)))]>,
- Requires<[useFP16Math]>;
+ Requires<[useFP16Math, hasSM<80>, hasPTX<70>]>;
def bf16rr_ftz :
NVPTXInst<(outs Int16Regs:$dst),
(ins Int16Regs:$a, Int16Regs:$b),
!strconcat(OpcStr, ".ftz.bf16 \t$dst, $a, $b;"),
[(set Int16Regs:$dst, (OpNode (bf16 Int16Regs:$a), (bf16 Int16Regs:$b)))]>,
- Requires<[hasBF16Math, doF32FTZ]>;
+ Requires<[hasBF16Math, doF32FTZ, hasSM<80>, hasPTX<70>]>;
def bf16rr :
NVPTXInst<(outs Int16Regs:$dst),
(ins Int16Regs:$a, Int16Regs:$b),
!strconcat(OpcStr, ".bf16 \t$dst, $a, $b;"),
[(set Int16Regs:$dst, (OpNode (bf16 Int16Regs:$a), (bf16 Int16Regs:$b)))]>,
- Requires<[hasBF16Math]>;
+ Requires<[hasBF16Math, hasSM<80>, hasPTX<70>]>;
def bf16x2rr_ftz :
NVPTXInst<(outs Int32Regs:$dst),
(ins Int32Regs:$a, Int32Regs:$b),
!strconcat(OpcStr, ".ftz.bf16x2 \t$dst, $a, $b;"),
[(set Int32Regs:$dst, (OpNode (v2bf16 Int32Regs:$a), (v2bf16 Int32Regs:$b)))]>,
- Requires<[hasBF16Math, doF32FTZ]>;
+ Requires<[hasBF16Math, hasSM<80>, hasPTX<70>, doF32FTZ]>;
def bf16x2rr :
NVPTXInst<(outs Int32Regs:$dst),
(ins Int32Regs:$a, Int32Regs:$b),
!strconcat(OpcStr, ".bf16x2 \t$dst, $a, $b;"),
[(set Int32Regs:$dst, (OpNode (v2bf16 Int32Regs:$a), (v2bf16 Int32Regs:$b)))]>,
- Requires<[hasBF16Math]>;
+ Requires<[hasBF16Math, hasSM<80>, hasPTX<70>]>;
}
// Template for instructions which take three FP args. The
@@ -1178,11 +1175,10 @@ defm FADD : F3_fma_component<"add", fadd>;
defm FSUB : F3_fma_component<"sub", fsub>;
defm FMUL : F3_fma_component<"mul", fmul>;
-defm FMIN : F3<"min", fminnum>;
-defm FMAX : F3<"max", fmaxnum>;
-// Note: min.NaN.f64 and max.NaN.f64 do not actually exist.
-defm FMINNAN : F3<"min.NaN", fminimum>;
-defm FMAXNAN : F3<"max.NaN", fmaximum>;
+defm FMIN : FMINIMUMMAXIMUM<"min", /* NaN */ false, fminnum>;
+defm FMAX : FMINIMUMMAXIMUM<"max", /* NaN */ false, fmaxnum>;
+defm FMINNAN : FMINIMUMMAXIMUM<"min.NaN", /* NaN */ true, fminimum>;
+defm FMAXNAN : FMINIMUMMAXIMUM<"max.NaN", /* NaN */ true, fmaximum>;
defm FABS : F2<"abs", fabs>;
defm FNEG : F2<"neg", fneg>;
diff --git a/llvm/test/CodeGen/NVPTX/math-intrins.ll b/llvm/test/CodeGen/NVPTX/math-intrins.ll
index d31844549a322..fcc4ec6e4017f 100644
--- a/llvm/test/CodeGen/NVPTX/math-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/math-intrins.ll
@@ -1,6 +1,7 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s | FileCheck %s --check-prefixes=CHECK,CHECK-NOF16
-; RUN: llc < %s -mcpu=sm_80 | FileCheck %s --check-prefixes=CHECK,CHECK-F16
-; RUN: llc < %s -mcpu=sm_80 --nvptx-no-f16-math | FileCheck %s --check-prefixes=CHECK,CHECK-NOF16
+; RUN: llc < %s -mcpu=sm_80 -mattr +ptx70 | FileCheck %s --check-prefixes=CHECK,CHECK-F16
+; RUN: llc < %s -mcpu=sm_80 -mattr +ptx70 --nvptx-no-f16-math | FileCheck %s --check-prefixes=CHECK,CHECK-SM80-NOF16
; RUN: %if ptxas %{ llc < %s | %ptxas-verify %}
; RUN: %if ptxas-11.0 %{ llc < %s -mcpu=sm_80 | %ptxas-verify -arch=sm_80 %}
; RUN: %if ptxas-11.0 %{ llc < %s -mcpu=sm_80 --nvptx-no-f16-math | %ptxas-verify -arch=sm_80 %}
@@ -29,330 +30,1571 @@ declare half @llvm.minnum.f16(half, half) #0
declare float @llvm.minnum.f32(float, float) #0
declare double @llvm.minnum.f64(double, double) #0
declare <2 x half> @llvm.minnum.v2f16(<2 x half>, <2 x half>) #0
+declare half @llvm.minimum.f16(half, half) #0
+declare float @llvm.minimum.f32(float, float) #0
+declare double @llvm.minimum.f64(double, double) #0
+declare <2 x half> @llvm.minimum.v2f16(<2 x half>, <2 x half>) #0
declare half @llvm.maxnum.f16(half, half) #0
declare float @llvm.maxnum.f32(float, float) #0
declare double @llvm.maxnum.f64(double, double) #0
declare <2 x half> @llvm.maxnum.v2f16(<2 x half>, <2 x half>) #0
+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 float @llvm.fma.f32(float, float, float) #0
declare double @llvm.fma.f64(double, double, double) #0
; ---- ceil ----
-; CHECK-LABEL: ceil_float
define float @ceil_float(float %a) {
- ; CHECK: cvt.rpi.f32.f32
+; CHECK-LABEL: ceil_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [ceil_float_param_0];
+; CHECK-NEXT: cvt.rpi.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.ceil.f32(float %a)
ret float %b
}
-; CHECK-LABEL: ceil_float_ftz
define float @ceil_float_ftz(float %a) #1 {
- ; CHECK: cvt.rpi.ftz.f32.f32
+; CHECK-LABEL: ceil_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [ceil_float_ftz_param_0];
+; CHECK-NEXT: cvt.rpi.ftz.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.ceil.f32(float %a)
ret float %b
}
-; CHECK-LABEL: ceil_double
define double @ceil_double(double %a) {
- ; CHECK: cvt.rpi.f64.f64
+; CHECK-LABEL: ceil_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [ceil_double_param_0];
+; CHECK-NEXT: cvt.rpi.f64.f64 %fd2, %fd1;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd2;
+; CHECK-NEXT: ret;
%b = call double @llvm.ceil.f64(double %a)
ret double %b
}
; ---- floor ----
-; CHECK-LABEL: floor_float
define float @floor_float(float %a) {
- ; CHECK: cvt.rmi.f32.f32
+; CHECK-LABEL: floor_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [floor_float_param_0];
+; CHECK-NEXT: cvt.rmi.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.floor.f32(float %a)
ret float %b
}
-; CHECK-LABEL: floor_float_ftz
define float @floor_float_ftz(float %a) #1 {
- ; CHECK: cvt.rmi.ftz.f32.f32
+; CHECK-LABEL: floor_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [floor_float_ftz_param_0];
+; CHECK-NEXT: cvt.rmi.ftz.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.floor.f32(float %a)
ret float %b
}
-; CHECK-LABEL: floor_double
define double @floor_double(double %a) {
- ; CHECK: cvt.rmi.f64.f64
+; CHECK-LABEL: floor_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [floor_double_param_0];
+; CHECK-NEXT: cvt.rmi.f64.f64 %fd2, %fd1;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd2;
+; CHECK-NEXT: ret;
%b = call double @llvm.floor.f64(double %a)
ret double %b
}
; ---- round ----
-; CHECK-LABEL: round_float
define float @round_float(float %a) {
; check the use of sign mask and 0.5 to implement round
-; CHECK: and.b32 [[R1:%r[0-9]+]], {{.*}}, -2147483648;
-; CHECK: or.b32 {{.*}}, [[R1]], 1056964608;
+; CHECK-LABEL: round_float(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [round_float_param_0];
+; CHECK-NEXT: mov.b32 %r1, %f1;
+; CHECK-NEXT: and.b32 %r2, %r1, -2147483648;
+; CHECK-NEXT: or.b32 %r3, %r2, 1056964608;
+; CHECK-NEXT: mov.b32 %f2, %r3;
+; CHECK-NEXT: add.rn.f32 %f3, %f1, %f2;
+; CHECK-NEXT: cvt.rzi.f32.f32 %f4, %f3;
+; CHECK-NEXT: abs.f32 %f5, %f1;
+; CHECK-NEXT: setp.gt.f32 %p1, %f5, 0f4B000000;
+; CHECK-NEXT: selp.f32 %f6, %f1, %f4, %p1;
+; CHECK-NEXT: cvt.rzi.f32.f32 %f7, %f1;
+; CHECK-NEXT: setp.lt.f32 %p2, %f5, 0f3F000000;
+; CHECK-NEXT: selp.f32 %f8, %f7, %f6, %p2;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f8;
+; CHECK-NEXT: ret;
%b = call float @llvm.round.f32(float %a)
ret float %b
}
-; CHECK-LABEL: round_float_ftz
define float @round_float_ftz(float %a) #1 {
; check the use of sign mask and 0.5 to implement round
-; CHECK: and.b32 [[R1:%r[0-9]+]], {{.*}}, -2147483648;
-; CHECK: or.b32 {{.*}}, [[R1]], 1056964608;
+; CHECK-LABEL: round_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<3>;
+; CHECK-NEXT: .reg .b32 %r<4>;
+; CHECK-NEXT: .reg .f32 %f<9>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [round_float_ftz_param_0];
+; CHECK-NEXT: mov.b32 %r1, %f1;
+; CHECK-NEXT: and.b32 %r2, %r1, -2147483648;
+; CHECK-NEXT: or.b32 %r3, %r2, 1056964608;
+; CHECK-NEXT: mov.b32 %f2, %r3;
+; CHECK-NEXT: add.rn.ftz.f32 %f3, %f1, %f2;
+; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %f4, %f3;
+; CHECK-NEXT: abs.ftz.f32 %f5, %f1;
+; CHECK-NEXT: setp.gt.ftz.f32 %p1, %f5, 0f4B000000;
+; CHECK-NEXT: selp.f32 %f6, %f1, %f4, %p1;
+; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %f7, %f1;
+; CHECK-NEXT: setp.lt.ftz.f32 %p2, %f5, 0f3F000000;
+; CHECK-NEXT: selp.f32 %f8, %f7, %f6, %p2;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f8;
+; CHECK-NEXT: ret;
%b = call float @llvm.round.f32(float %a)
ret float %b
}
-; CHECK-LABEL: round_double
define double @round_double(double %a) {
; check the use of 0.5 to implement round
-; CHECK: setp.lt.f64 {{.*}}, [[R:%fd[0-9]+]], 0d3FE0000000000000;
-; CHECK: add.rn.f64 {{.*}}, [[R]], 0d3FE0000000000000;
+; CHECK-LABEL: round_double(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-NEXT: .reg .f64 %fd<10>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [round_double_param_0];
+; CHECK-NEXT: abs.f64 %fd2, %fd1;
+; CHECK-NEXT: setp.lt.f64 %p1, %fd2, 0d3FE0000000000000;
+; CHECK-NEXT: add.rn.f64 %fd3, %fd2, 0d3FE0000000000000;
+; CHECK-NEXT: cvt.rzi.f64.f64 %fd4, %fd3;
+; CHECK-NEXT: selp.f64 %fd5, 0d0000000000000000, %fd4, %p1;
+; CHECK-NEXT: abs.f64 %fd6, %fd5;
+; CHECK-NEXT: neg.f64 %fd7, %fd6;
+; CHECK-NEXT: mov.b64 %rd1, %fd1;
+; CHECK-NEXT: shr.u64 %rd2, %rd1, 63;
+; CHECK-NEXT: and.b64 %rd3, %rd2, 1;
+; CHECK-NEXT: setp.eq.b64 %p2, %rd3, 1;
+; CHECK-NEXT: selp.f64 %fd8, %fd7, %fd6, %p2;
+; CHECK-NEXT: setp.gt.f64 %p3, %fd2, 0d4330000000000000;
+; CHECK-NEXT: selp.f64 %fd9, %fd1, %fd8, %p3;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd9;
+; CHECK-NEXT: ret;
%b = call double @llvm.round.f64(double %a)
ret double %b
}
; ---- nearbyint ----
-; CHECK-LABEL: nearbyint_float
define float @nearbyint_float(float %a) {
- ; CHECK: cvt.rni.f32.f32
+; CHECK-LABEL: nearbyint_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [nearbyint_float_param_0];
+; CHECK-NEXT: cvt.rni.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.nearbyint.f32(float %a)
ret float %b
}
-; CHECK-LABEL: nearbyint_float_ftz
define float @nearbyint_float_ftz(float %a) #1 {
- ; CHECK: cvt.rni.ftz.f32.f32
+; CHECK-LABEL: nearbyint_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [nearbyint_float_ftz_param_0];
+; CHECK-NEXT: cvt.rni.ftz.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.nearbyint.f32(float %a)
ret float %b
}
-; CHECK-LABEL: nearbyint_double
define double @nearbyint_double(double %a) {
- ; CHECK: cvt.rni.f64.f64
+; CHECK-LABEL: nearbyint_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [nearbyint_double_param_0];
+; CHECK-NEXT: cvt.rni.f64.f64 %fd2, %fd1;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd2;
+; CHECK-NEXT: ret;
%b = call double @llvm.nearbyint.f64(double %a)
ret double %b
}
; ---- rint ----
-; CHECK-LABEL: rint_float
define float @rint_float(float %a) {
- ; CHECK: cvt.rni.f32.f32
+; CHECK-LABEL: rint_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [rint_float_param_0];
+; CHECK-NEXT: cvt.rni.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.rint.f32(float %a)
ret float %b
}
-; CHECK-LABEL: rint_float_ftz
define float @rint_float_ftz(float %a) #1 {
- ; CHECK: cvt.rni.ftz.f32.f32
+; CHECK-LABEL: rint_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [rint_float_ftz_param_0];
+; CHECK-NEXT: cvt.rni.ftz.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.rint.f32(float %a)
ret float %b
}
-; CHECK-LABEL: rint_double
define double @rint_double(double %a) {
- ; CHECK: cvt.rni.f64.f64
+; CHECK-LABEL: rint_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [rint_double_param_0];
+; CHECK-NEXT: cvt.rni.f64.f64 %fd2, %fd1;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd2;
+; CHECK-NEXT: ret;
%b = call double @llvm.rint.f64(double %a)
ret double %b
}
; ---- roundeven ----
-; CHECK-LABEL: roundeven_float
define float @roundeven_float(float %a) {
- ; CHECK: cvt.rni.f32.f32
+; CHECK-LABEL: roundeven_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [roundeven_float_param_0];
+; CHECK-NEXT: cvt.rni.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.roundeven.f32(float %a)
ret float %b
}
-; CHECK-LABEL: roundeven_float_ftz
define float @roundeven_float_ftz(float %a) #1 {
- ; CHECK: cvt.rni.ftz.f32.f32
+; CHECK-LABEL: roundeven_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [roundeven_float_ftz_param_0];
+; CHECK-NEXT: cvt.rni.ftz.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.roundeven.f32(float %a)
ret float %b
}
-; CHECK-LABEL: roundeven_double
define double @roundeven_double(double %a) {
- ; CHECK: cvt.rni.f64.f64
+; CHECK-LABEL: roundeven_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [roundeven_double_param_0];
+; CHECK-NEXT: cvt.rni.f64.f64 %fd2, %fd1;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd2;
+; CHECK-NEXT: ret;
%b = call double @llvm.roundeven.f64(double %a)
ret double %b
}
; ---- trunc ----
-; CHECK-LABEL: trunc_float
define float @trunc_float(float %a) {
- ; CHECK: cvt.rzi.f32.f32
+; CHECK-LABEL: trunc_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [trunc_float_param_0];
+; CHECK-NEXT: cvt.rzi.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.trunc.f32(float %a)
ret float %b
}
-; CHECK-LABEL: trunc_float_ftz
define float @trunc_float_ftz(float %a) #1 {
- ; CHECK: cvt.rzi.ftz.f32.f32
+; CHECK-LABEL: trunc_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [trunc_float_ftz_param_0];
+; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.trunc.f32(float %a)
ret float %b
}
-; CHECK-LABEL: trunc_double
define double @trunc_double(double %a) {
- ; CHECK: cvt.rzi.f64.f64
+; CHECK-LABEL: trunc_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [trunc_double_param_0];
+; CHECK-NEXT: cvt.rzi.f64.f64 %fd2, %fd1;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd2;
+; CHECK-NEXT: ret;
%b = call double @llvm.trunc.f64(double %a)
ret double %b
}
; ---- abs ----
-; CHECK-LABEL: abs_float
define float @abs_float(float %a) {
- ; CHECK: abs.f32
+; CHECK-LABEL: abs_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [abs_float_param_0];
+; CHECK-NEXT: abs.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.fabs.f32(float %a)
ret float %b
}
-; CHECK-LABEL: abs_float_ftz
define float @abs_float_ftz(float %a) #1 {
- ; CHECK: abs.ftz.f32
+; CHECK-LABEL: abs_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [abs_float_ftz_param_0];
+; CHECK-NEXT: abs.ftz.f32 %f2, %f1;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%b = call float @llvm.fabs.f32(float %a)
ret float %b
}
-; CHECK-LABEL: abs_double
define double @abs_double(double %a) {
- ; CHECK: abs.f64
+; CHECK-LABEL: abs_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [abs_double_param_0];
+; CHECK-NEXT: abs.f64 %fd2, %fd1;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd2;
+; CHECK-NEXT: ret;
%b = call double @llvm.fabs.f64(double %a)
ret double %b
}
-; ---- min ----
+; ---- minnum ----
-; CHECK-LABEL: min_half
-define half @min_half(half %a, half %b) {
- ; CHECK-NOF16: min.f32
- ; CHECK-F16: min.f16
+define half @minnum_half(half %a, half %b) {
+; CHECK-NOF16-LABEL: minnum_half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minnum_half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [minnum_half_param_1];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-NOF16-NEXT: min.f32 %f3, %f2, %f1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %f3;
+; CHECK-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minnum_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, [minnum_half_param_0];
+; CHECK-F16-NEXT: ld.param.b16 %rs2, [minnum_half_param_1];
+; CHECK-F16-NEXT: min.f16 %rs3, %rs1, %rs2;
+; CHECK-F16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minnum_half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minnum_half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [minnum_half_param_1];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-SM80-NOF16-NEXT: min.f32 %f3, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %f3;
+; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-SM80-NOF16-NEXT: ret;
%x = call half @llvm.minnum.f16(half %a, half %b)
ret half %x
}
-; CHECK-LABEL: min_float
-define float @min_float(float %a, float %b) {
- ; CHECK: min.f32
+define float @minnum_float(float %a, float %b) {
+; CHECK-LABEL: minnum_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [minnum_float_param_0];
+; CHECK-NEXT: ld.param.f32 %f2, [minnum_float_param_1];
+; CHECK-NEXT: min.f32 %f3, %f1, %f2;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-NEXT: ret;
%x = call float @llvm.minnum.f32(float %a, float %b)
ret float %x
}
-; CHECK-LABEL: min_imm1
-define float @min_imm1(float %a) {
- ; CHECK: min.f32
+define float @minnum_imm1(float %a) {
+; CHECK-LABEL: minnum_imm1(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [minnum_imm1_param_0];
+; CHECK-NEXT: min.f32 %f2, %f1, 0f00000000;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%x = call float @llvm.minnum.f32(float %a, float 0.0)
ret float %x
}
-; CHECK-LABEL: min_imm2
-define float @min_imm2(float %a) {
- ; CHECK: min.f32
+define float @minnum_imm2(float %a) {
+; CHECK-LABEL: minnum_imm2(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [minnum_imm2_param_0];
+; CHECK-NEXT: min.f32 %f2, %f1, 0f00000000;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%x = call float @llvm.minnum.f32(float 0.0, float %a)
ret float %x
}
-; CHECK-LABEL: min_float_ftz
-define float @min_float_ftz(float %a, float %b) #1 {
- ; CHECK: min.ftz.f32
+define float @minnum_float_ftz(float %a, float %b) #1 {
+; CHECK-LABEL: minnum_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [minnum_float_ftz_param_0];
+; CHECK-NEXT: ld.param.f32 %f2, [minnum_float_ftz_param_1];
+; CHECK-NEXT: min.ftz.f32 %f3, %f1, %f2;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-NEXT: ret;
%x = call float @llvm.minnum.f32(float %a, float %b)
ret float %x
}
-; CHECK-LABEL: min_double
-define double @min_double(double %a, double %b) {
- ; CHECK: min.f64
+define double @minnum_double(double %a, double %b) {
+; CHECK-LABEL: minnum_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [minnum_double_param_0];
+; CHECK-NEXT: ld.param.f64 %fd2, [minnum_double_param_1];
+; CHECK-NEXT: min.f64 %fd3, %fd1, %fd2;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd3;
+; CHECK-NEXT: ret;
%x = call double @llvm.minnum.f64(double %a, double %b)
ret double %x
}
-; CHECK-LABEL: min_v2half
-define <2 x half> @min_v2half(<2 x half> %a, <2 x half> %b) {
- ; CHECK-NOF16: min.f32
- ; CHECK-NOF16: min.f32
- ; CHECK-F16: min.f16x2
+define <2 x half> @minnum_v2half(<2 x half> %a, <2 x half> %b) {
+; CHECK-NOF16-LABEL: minnum_v2half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minnum_v2half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minnum_v2half_param_1];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-NOF16-NEXT: min.f32 %f3, %f2, %f1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %f3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-NOF16-NEXT: min.f32 %f6, %f5, %f4;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %f6;
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minnum_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, [minnum_v2half_param_1];
+; CHECK-F16-NEXT: ld.param.b32 %r2, [minnum_v2half_param_0];
+; CHECK-F16-NEXT: min.f16x2 %r3, %r2, %r1;
+; CHECK-F16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minnum_v2half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minnum_v2half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [minnum_v2half_param_1];
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-SM80-NOF16-NEXT: min.f32 %f3, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %f3;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-SM80-NOF16-NEXT: min.f32 %f6, %f5, %f4;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %f6;
+; CHECK-SM80-NOF16-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-SM80-NOF16-NEXT: ret;
%x = call <2 x half> @llvm.minnum.v2f16(<2 x half> %a, <2 x half> %b)
ret <2 x half> %x
}
-; ---- max ----
+; ---- minimum ----
-; CHECK-LABEL: max_half
-define half @max_half(half %a, half %b) {
- ; CHECK-NOF16: max.f32
- ; CHECK-F16: max.f16
+define half @minimum_half(half %a, half %b) {
+; CHECK-NOF16-LABEL: minimum_half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<6>;
+; CHECK-NOF16-NEXT: .reg .b16 %rs<10>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [minimum_half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b16 %rs3, [minimum_half_param_1];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-NOF16-NEXT: setp.lt.f32 %p1, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs4, %rs1, %rs3, %p1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs5, 0x7E00, %rs4, %p2;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p3, %rs1, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs6, %rs1, %rs5, %p3;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p4, %rs3, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs3, %rs6, %p4;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f3, %rs5;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs5, %p5;
+; CHECK-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs9;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimum_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, [minimum_half_param_0];
+; CHECK-F16-NEXT: ld.param.b16 %rs2, [minimum_half_param_1];
+; CHECK-F16-NEXT: min.NaN.f16 %rs3, %rs1, %rs2;
+; CHECK-F16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimum_half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .pred %p<6>;
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<10>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [minimum_half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs3, [minimum_half_param_1];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs3;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p1, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs4, %rs1, %rs3, %p1;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, 0x7E00, %rs4, %p2;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p3, %rs1, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, %rs1, %rs5, %p3;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p4, %rs3, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs3, %rs6, %p4;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f3, %rs5;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs5, %p5;
+; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs9;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call half @llvm.minimum.f16(half %a, half %b)
+ ret half %x
+}
+
+define float @minimum_float(float %a, float %b) {
+; CHECK-NOF16-LABEL: minimum_float(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
+; CHECK-NOF16-NEXT: ld.param.f32 %f2, [minimum_float_param_1];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f2;
+; CHECK-NOF16-NEXT: min.f32 %f3, %f1, %f2;
+; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
+; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
+; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, -2147483648;
+; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %f4, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f7;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimum_float(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
+; CHECK-F16-NEXT: ld.param.f32 %f2, [minimum_float_param_1];
+; CHECK-F16-NEXT: min.NaN.f32 %f3, %f1, %f2;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimum_float(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f2, [minimum_float_param_1];
+; CHECK-SM80-NOF16-NEXT: min.NaN.f32 %f3, %f1, %f2;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.minimum.f32(float %a, float %b)
+ ret float %x
+}
+
+define float @minimum_imm1(float %a) {
+; CHECK-NOF16-LABEL: minimum_imm1(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<2>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<6>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
+; CHECK-NOF16-NEXT: min.f32 %f2, %f1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
+; CHECK-NOF16-NEXT: selp.f32 %f4, %f1, %f3, %p2;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f5, %f4, %f3, %p3;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f5;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimum_imm1(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
+; CHECK-F16-NEXT: min.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimum_imm1(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
+; CHECK-SM80-NOF16-NEXT: min.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.minimum.f32(float %a, float 0.0)
+ ret float %x
+}
+
+define float @minimum_imm2(float %a) {
+; CHECK-NOF16-LABEL: minimum_imm2(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<2>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<6>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
+; CHECK-NOF16-NEXT: min.f32 %f2, %f1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
+; CHECK-NOF16-NEXT: selp.f32 %f4, %f1, %f3, %p2;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f5, %f4, %f3, %p3;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f5;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimum_imm2(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
+; CHECK-F16-NEXT: min.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimum_imm2(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
+; CHECK-SM80-NOF16-NEXT: min.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.minimum.f32(float 0.0, float %a)
+ ret float %x
+}
+
+define float @minimum_float_ftz(float %a, float %b) #1 {
+; CHECK-NOF16-LABEL: minimum_float_ftz(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
+; CHECK-NOF16-NEXT: ld.param.f32 %f2, [minimum_float_ftz_param_1];
+; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %f1, %f2;
+; CHECK-NOF16-NEXT: min.ftz.f32 %f3, %f1, %f2;
+; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, -2147483648;
+; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
+; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, -2147483648;
+; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
+; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %f4, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f7;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimum_float_ftz(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
+; CHECK-F16-NEXT: ld.param.f32 %f2, [minimum_float_ftz_param_1];
+; CHECK-F16-NEXT: min.NaN.ftz.f32 %f3, %f1, %f2;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimum_float_ftz(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f2, [minimum_float_ftz_param_1];
+; CHECK-SM80-NOF16-NEXT: min.NaN.ftz.f32 %f3, %f1, %f2;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.minimum.f32(float %a, float %b)
+ ret float %x
+}
+
+define double @minimum_double(double %a, double %b) {
+; CHECK-LABEL: minimum_double(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<5>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .f64 %fd<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [minimum_double_param_0];
+; CHECK-NEXT: mov.b64 %rd1, %fd1;
+; CHECK-NEXT: ld.param.f64 %fd2, [minimum_double_param_1];
+; CHECK-NEXT: setp.nan.f64 %p1, %fd1, %fd2;
+; CHECK-NEXT: min.f64 %fd3, %fd1, %fd2;
+; CHECK-NEXT: selp.f64 %fd4, 0d7FF8000000000000, %fd3, %p1;
+; CHECK-NEXT: setp.eq.s64 %p2, %rd1, -9223372036854775808;
+; CHECK-NEXT: selp.f64 %fd5, %fd1, %fd4, %p2;
+; CHECK-NEXT: mov.b64 %rd2, %fd2;
+; CHECK-NEXT: setp.eq.s64 %p3, %rd2, -9223372036854775808;
+; CHECK-NEXT: selp.f64 %fd6, %fd2, %fd5, %p3;
+; CHECK-NEXT: setp.eq.f64 %p4, %fd4, 0d0000000000000000;
+; CHECK-NEXT: selp.f64 %fd7, %fd6, %fd4, %p4;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd7;
+; CHECK-NEXT: ret;
+ %x = call double @llvm.minimum.f64(double %a, double %b)
+ ret double %x
+}
+
+define <2 x half> @minimum_v2half(<2 x half> %a, <2 x half> %b) {
+; CHECK-NOF16-LABEL: minimum_v2half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<11>;
+; CHECK-NOF16-NEXT: .reg .b16 %rs<19>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_v2half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minimum_v2half_param_1];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-NOF16-NEXT: setp.lt.f32 %p1, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p3, %rs4, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs6, %p3;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p4, %rs2, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs2, %rs8, %p4;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f3, %rs6;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs11, %rs10, %rs6, %p5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-NOF16-NEXT: setp.lt.f32 %p6, %f5, %f4;
+; CHECK-NOF16-NEXT: selp.b16 %rs12, %rs3, %rs1, %p6;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p7, %f5, %f4;
+; CHECK-NOF16-NEXT: selp.b16 %rs13, 0x7E00, %rs12, %p7;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p8, %rs3, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs15, %rs3, %rs13, %p8;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p9, %rs1, -32768;
+; CHECK-NOF16-NEXT: selp.b16 %rs17, %rs1, %rs15, %p9;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f6, %rs13;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p10, %f6, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs18, %rs17, %rs13, %p10;
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs18, %rs11};
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: minimum_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, [minimum_v2half_param_1];
+; CHECK-F16-NEXT: ld.param.b32 %r2, [minimum_v2half_param_0];
+; CHECK-F16-NEXT: min.NaN.f16x2 %r3, %r2, %r1;
+; CHECK-F16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: minimum_v2half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .pred %p<11>;
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<19>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [minimum_v2half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [minimum_v2half_param_1];
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p1, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p3, %rs4, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs6, %p3;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p4, %rs2, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs10, %rs2, %rs8, %p4;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f3, %rs6;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs11, %rs10, %rs6, %p5;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-SM80-NOF16-NEXT: setp.lt.f32 %p6, %f5, %f4;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs12, %rs3, %rs1, %p6;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p7, %f5, %f4;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs13, 0x7E00, %rs12, %p7;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p8, %rs3, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs15, %rs3, %rs13, %p8;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p9, %rs1, -32768;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs17, %rs1, %rs15, %p9;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f6, %rs13;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p10, %f6, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs18, %rs17, %rs13, %p10;
+; CHECK-SM80-NOF16-NEXT: mov.b32 %r3, {%rs18, %rs11};
+; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call <2 x half> @llvm.minimum.v2f16(<2 x half> %a, <2 x half> %b)
+ ret <2 x half> %x
+}
+
+; ---- maxnum ----
+
+define half @maxnum_half(half %a, half %b) {
+; CHECK-NOF16-LABEL: maxnum_half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maxnum_half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b16 %rs2, [maxnum_half_param_1];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-NOF16-NEXT: max.f32 %f3, %f2, %f1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %f3;
+; CHECK-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maxnum_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, [maxnum_half_param_0];
+; CHECK-F16-NEXT: ld.param.b16 %rs2, [maxnum_half_param_1];
+; CHECK-F16-NEXT: max.f16 %rs3, %rs1, %rs2;
+; CHECK-F16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maxnum_half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maxnum_half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs2, [maxnum_half_param_1];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-SM80-NOF16-NEXT: max.f32 %f3, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %f3;
+; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-SM80-NOF16-NEXT: ret;
%x = call half @llvm.maxnum.f16(half %a, half %b)
ret half %x
}
-; CHECK-LABEL: max_imm1
-define float @max_imm1(float %a) {
- ; CHECK: max.f32
+define float @maxnum_imm1(float %a) {
+; CHECK-LABEL: maxnum_imm1(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [maxnum_imm1_param_0];
+; CHECK-NEXT: max.f32 %f2, %f1, 0f00000000;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%x = call float @llvm.maxnum.f32(float %a, float 0.0)
ret float %x
}
-; CHECK-LABEL: max_imm2
-define float @max_imm2(float %a) {
- ; CHECK: max.f32
+define float @maxnum_imm2(float %a) {
+; CHECK-LABEL: maxnum_imm2(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [maxnum_imm2_param_0];
+; CHECK-NEXT: max.f32 %f2, %f1, 0f00000000;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-NEXT: ret;
%x = call float @llvm.maxnum.f32(float 0.0, float %a)
ret float %x
}
-; CHECK-LABEL: max_float
-define float @max_float(float %a, float %b) {
- ; CHECK: max.f32
+define float @maxnum_float(float %a, float %b) {
+; CHECK-LABEL: maxnum_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [maxnum_float_param_0];
+; CHECK-NEXT: ld.param.f32 %f2, [maxnum_float_param_1];
+; CHECK-NEXT: max.f32 %f3, %f1, %f2;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-NEXT: ret;
%x = call float @llvm.maxnum.f32(float %a, float %b)
ret float %x
}
-; CHECK-LABEL: max_float_ftz
-define float @max_float_ftz(float %a, float %b) #1 {
- ; CHECK: max.ftz.f32
+define float @maxnum_float_ftz(float %a, float %b) #1 {
+; CHECK-LABEL: maxnum_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [maxnum_float_ftz_param_0];
+; CHECK-NEXT: ld.param.f32 %f2, [maxnum_float_ftz_param_1];
+; CHECK-NEXT: max.ftz.f32 %f3, %f1, %f2;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-NEXT: ret;
%x = call float @llvm.maxnum.f32(float %a, float %b)
ret float %x
}
-; CHECK-LABEL: max_double
-define double @max_double(double %a, double %b) {
- ; CHECK: max.f64
+define double @maxnum_double(double %a, double %b) {
+; CHECK-LABEL: maxnum_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [maxnum_double_param_0];
+; CHECK-NEXT: ld.param.f64 %fd2, [maxnum_double_param_1];
+; CHECK-NEXT: max.f64 %fd3, %fd1, %fd2;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd3;
+; CHECK-NEXT: ret;
%x = call double @llvm.maxnum.f64(double %a, double %b)
ret double %x
}
-; CHECK-LABEL: max_v2half
-define <2 x half> @max_v2half(<2 x half> %a, <2 x half> %b) {
- ; CHECK-NOF16: max.f32
- ; CHECK-NOF16: max.f32
- ; CHECK-F16: max.f16x2
+define <2 x half> @maxnum_v2half(<2 x half> %a, <2 x half> %b) {
+; CHECK-NOF16-LABEL: maxnum_v2half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .b16 %rs<7>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maxnum_v2half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maxnum_v2half_param_1];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-NOF16-NEXT: max.f32 %f3, %f2, %f1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %f3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-NOF16-NEXT: max.f32 %f6, %f5, %f4;
+; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %f6;
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maxnum_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, [maxnum_v2half_param_1];
+; CHECK-F16-NEXT: ld.param.b32 %r2, [maxnum_v2half_param_0];
+; CHECK-F16-NEXT: max.f16x2 %r3, %r2, %r1;
+; CHECK-F16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maxnum_v2half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<7>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maxnum_v2half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [maxnum_v2half_param_1];
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-SM80-NOF16-NEXT: max.f32 %f3, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %f3;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-SM80-NOF16-NEXT: max.f32 %f6, %f5, %f4;
+; CHECK-SM80-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %f6;
+; CHECK-SM80-NOF16-NEXT: mov.b32 %r3, {%rs6, %rs5};
+; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-SM80-NOF16-NEXT: ret;
%x = call <2 x half> @llvm.maxnum.v2f16(<2 x half> %a, <2 x half> %b)
ret <2 x half> %x
}
+; ---- maximum ----
+
+define half @maximum_half(half %a, half %b) {
+; CHECK-NOF16-LABEL: maximum_half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<6>;
+; CHECK-NOF16-NEXT: .reg .b16 %rs<10>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b16 %rs1, [maximum_half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b16 %rs3, [maximum_half_param_1];
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-NOF16-NEXT: setp.gt.f32 %p1, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs4, %rs1, %rs3, %p1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs5, 0x7E00, %rs4, %p2;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p3, %rs1, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs6, %rs1, %rs5, %p3;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p4, %rs3, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs3, %rs6, %p4;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f3, %rs5;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs5, %p5;
+; CHECK-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs9;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximum_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, [maximum_half_param_0];
+; CHECK-F16-NEXT: ld.param.b16 %rs2, [maximum_half_param_1];
+; CHECK-F16-NEXT: max.NaN.f16 %rs3, %rs1, %rs2;
+; CHECK-F16-NEXT: st.param.b16 [func_retval0+0], %rs3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximum_half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .pred %p<6>;
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<10>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs1, [maximum_half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b16 %rs3, [maximum_half_param_1];
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs3;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs1;
+; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p1, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs4, %rs1, %rs3, %p1;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, 0x7E00, %rs4, %p2;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p3, %rs1, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, %rs1, %rs5, %p3;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p4, %rs3, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs3, %rs6, %p4;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f3, %rs5;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs5, %p5;
+; CHECK-SM80-NOF16-NEXT: st.param.b16 [func_retval0+0], %rs9;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call half @llvm.maximum.f16(half %a, half %b)
+ ret half %x
+}
+
+define float @maximum_imm1(float %a) {
+; CHECK-NOF16-LABEL: maximum_imm1(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<3>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm1_param_0];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
+; CHECK-NOF16-NEXT: max.f32 %f2, %f1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f4, 0f00000000, %f3, %p2;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f4;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximum_imm1(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_imm1_param_0];
+; CHECK-F16-NEXT: max.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximum_imm1(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm1_param_0];
+; CHECK-SM80-NOF16-NEXT: max.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.maximum.f32(float %a, float 0.0)
+ ret float %x
+}
+
+define float @maximum_imm2(float %a) {
+; CHECK-NOF16-LABEL: maximum_imm2(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<3>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<5>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm2_param_0];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f1;
+; CHECK-NOF16-NEXT: max.f32 %f2, %f1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f3, 0f7FC00000, %f2, %p1;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f4, 0f00000000, %f3, %p2;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f4;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximum_imm2(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<3>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_imm2_param_0];
+; CHECK-F16-NEXT: max.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximum_imm2(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<3>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_imm2_param_0];
+; CHECK-SM80-NOF16-NEXT: max.NaN.f32 %f2, %f1, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f2;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.maximum.f32(float 0.0, float %a)
+ ret float %x
+}
+
+define float @maximum_float(float %a, float %b) {
+; CHECK-NOF16-LABEL: maximum_float(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
+; CHECK-NOF16-NEXT: ld.param.f32 %f2, [maximum_float_param_1];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %f1, %f2;
+; CHECK-NOF16-NEXT: max.f32 %f3, %f1, %f2;
+; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, 0;
+; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
+; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, 0;
+; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %f4, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f7;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximum_float(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
+; CHECK-F16-NEXT: ld.param.f32 %f2, [maximum_float_param_1];
+; CHECK-F16-NEXT: max.NaN.f32 %f3, %f1, %f2;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximum_float(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f2, [maximum_float_param_1];
+; CHECK-SM80-NOF16-NEXT: max.NaN.f32 %f3, %f1, %f2;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.maximum.f32(float %a, float %b)
+ ret float %x
+}
+
+define float @maximum_float_ftz(float %a, float %b) #1 {
+; CHECK-NOF16-LABEL: maximum_float_ftz(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<8>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, %f1;
+; CHECK-NOF16-NEXT: ld.param.f32 %f2, [maximum_float_ftz_param_1];
+; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %f1, %f2;
+; CHECK-NOF16-NEXT: max.ftz.f32 %f3, %f1, %f2;
+; CHECK-NOF16-NEXT: selp.f32 %f4, 0f7FC00000, %f3, %p1;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p2, %r1, 0;
+; CHECK-NOF16-NEXT: selp.f32 %f5, %f1, %f4, %p2;
+; CHECK-NOF16-NEXT: mov.b32 %r2, %f2;
+; CHECK-NOF16-NEXT: setp.eq.s32 %p3, %r2, 0;
+; CHECK-NOF16-NEXT: selp.f32 %f6, %f2, %f5, %p3;
+; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %f4, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %f7, %f6, %f4, %p4;
+; CHECK-NOF16-NEXT: st.param.f32 [func_retval0+0], %f7;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximum_float_ftz(
+; CHECK-F16: {
+; CHECK-F16-NEXT: .reg .f32 %f<4>;
+; CHECK-F16-EMPTY:
+; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
+; CHECK-F16-NEXT: ld.param.f32 %f2, [maximum_float_ftz_param_1];
+; CHECK-F16-NEXT: max.NaN.ftz.f32 %f3, %f1, %f2;
+; CHECK-F16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximum_float_ftz(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<4>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.f32 %f2, [maximum_float_ftz_param_1];
+; CHECK-SM80-NOF16-NEXT: max.NaN.ftz.f32 %f3, %f1, %f2;
+; CHECK-SM80-NOF16-NEXT: st.param.f32 [func_retval0+0], %f3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call float @llvm.maximum.f32(float %a, float %b)
+ ret float %x
+}
+
+define double @maximum_double(double %a, double %b) {
+; CHECK-LABEL: maximum_double(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<5>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-NEXT: .reg .f64 %fd<8>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [maximum_double_param_0];
+; CHECK-NEXT: mov.b64 %rd1, %fd1;
+; CHECK-NEXT: ld.param.f64 %fd2, [maximum_double_param_1];
+; CHECK-NEXT: setp.nan.f64 %p1, %fd1, %fd2;
+; CHECK-NEXT: max.f64 %fd3, %fd1, %fd2;
+; CHECK-NEXT: selp.f64 %fd4, 0d7FF8000000000000, %fd3, %p1;
+; CHECK-NEXT: setp.eq.s64 %p2, %rd1, 0;
+; CHECK-NEXT: selp.f64 %fd5, %fd1, %fd4, %p2;
+; CHECK-NEXT: mov.b64 %rd2, %fd2;
+; CHECK-NEXT: setp.eq.s64 %p3, %rd2, 0;
+; CHECK-NEXT: selp.f64 %fd6, %fd2, %fd5, %p3;
+; CHECK-NEXT: setp.eq.f64 %p4, %fd4, 0d0000000000000000;
+; CHECK-NEXT: selp.f64 %fd7, %fd6, %fd4, %p4;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd7;
+; CHECK-NEXT: ret;
+ %x = call double @llvm.maximum.f64(double %a, double %b)
+ ret double %x
+}
+
+define <2 x half> @maximum_v2half(<2 x half> %a, <2 x half> %b) {
+; CHECK-NOF16-LABEL: maximum_v2half(
+; CHECK-NOF16: {
+; CHECK-NOF16-NEXT: .reg .pred %p<11>;
+; CHECK-NOF16-NEXT: .reg .b16 %rs<19>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-NOF16-EMPTY:
+; CHECK-NOF16-NEXT: // %bb.0:
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_v2half_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maximum_v2half_param_1];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-NOF16-NEXT: setp.gt.f32 %p1, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p3, %rs4, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs6, %p3;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p4, %rs2, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs2, %rs8, %p4;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f3, %rs6;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs11, %rs10, %rs6, %p5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-NOF16-NEXT: setp.gt.f32 %p6, %f5, %f4;
+; CHECK-NOF16-NEXT: selp.b16 %rs12, %rs3, %rs1, %p6;
+; CHECK-NOF16-NEXT: setp.nan.f32 %p7, %f5, %f4;
+; CHECK-NOF16-NEXT: selp.b16 %rs13, 0x7E00, %rs12, %p7;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p8, %rs3, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs15, %rs3, %rs13, %p8;
+; CHECK-NOF16-NEXT: setp.eq.s16 %p9, %rs1, 0;
+; CHECK-NOF16-NEXT: selp.b16 %rs17, %rs1, %rs15, %p9;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %f6, %rs13;
+; CHECK-NOF16-NEXT: setp.eq.f32 %p10, %f6, 0f00000000;
+; CHECK-NOF16-NEXT: selp.b16 %rs18, %rs17, %rs13, %p10;
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs18, %rs11};
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-NOF16-NEXT: ret;
+;
+; CHECK-F16-LABEL: maximum_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, [maximum_v2half_param_1];
+; CHECK-F16-NEXT: ld.param.b32 %r2, [maximum_v2half_param_0];
+; CHECK-F16-NEXT: max.NaN.f16x2 %r3, %r2, %r1;
+; CHECK-F16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-F16-NEXT: ret;
+;
+; CHECK-SM80-NOF16-LABEL: maximum_v2half(
+; CHECK-SM80-NOF16: {
+; CHECK-SM80-NOF16-NEXT: .reg .pred %p<11>;
+; CHECK-SM80-NOF16-NEXT: .reg .b16 %rs<19>;
+; CHECK-SM80-NOF16-NEXT: .reg .b32 %r<4>;
+; CHECK-SM80-NOF16-NEXT: .reg .f32 %f<7>;
+; CHECK-SM80-NOF16-EMPTY:
+; CHECK-SM80-NOF16-NEXT: // %bb.0:
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r1, [maximum_v2half_param_0];
+; CHECK-SM80-NOF16-NEXT: ld.param.b32 %r2, [maximum_v2half_param_1];
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f1, %rs2;
+; CHECK-SM80-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f2, %rs4;
+; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p1, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p1;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p2, %f2, %f1;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs6, 0x7E00, %rs5, %p2;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p3, %rs4, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs8, %rs4, %rs6, %p3;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p4, %rs2, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs10, %rs2, %rs8, %p4;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f3, %rs6;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p5, %f3, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs11, %rs10, %rs6, %p5;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f4, %rs1;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f5, %rs3;
+; CHECK-SM80-NOF16-NEXT: setp.gt.f32 %p6, %f5, %f4;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs12, %rs3, %rs1, %p6;
+; CHECK-SM80-NOF16-NEXT: setp.nan.f32 %p7, %f5, %f4;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs13, 0x7E00, %rs12, %p7;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p8, %rs3, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs15, %rs3, %rs13, %p8;
+; CHECK-SM80-NOF16-NEXT: setp.eq.s16 %p9, %rs1, 0;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs17, %rs1, %rs15, %p9;
+; CHECK-SM80-NOF16-NEXT: cvt.f32.f16 %f6, %rs13;
+; CHECK-SM80-NOF16-NEXT: setp.eq.f32 %p10, %f6, 0f00000000;
+; CHECK-SM80-NOF16-NEXT: selp.b16 %rs18, %rs17, %rs13, %p10;
+; CHECK-SM80-NOF16-NEXT: mov.b32 %r3, {%rs18, %rs11};
+; CHECK-SM80-NOF16-NEXT: st.param.b32 [func_retval0+0], %r3;
+; CHECK-SM80-NOF16-NEXT: ret;
+ %x = call <2 x half> @llvm.maximum.v2f16(<2 x half> %a, <2 x half> %b)
+ ret <2 x half> %x
+}
+
; ---- fma ----
-; CHECK-LABEL: @fma_float
define float @fma_float(float %a, float %b, float %c) {
- ; CHECK: fma.rn.f32
+; CHECK-LABEL: fma_float(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [fma_float_param_0];
+; CHECK-NEXT: ld.param.f32 %f2, [fma_float_param_1];
+; CHECK-NEXT: ld.param.f32 %f3, [fma_float_param_2];
+; CHECK-NEXT: fma.rn.f32 %f4, %f1, %f2, %f3;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f4;
+; CHECK-NEXT: ret;
%x = call float @llvm.fma.f32(float %a, float %b, float %c)
ret float %x
}
-; CHECK-LABEL: @fma_float_ftz
define float @fma_float_ftz(float %a, float %b, float %c) #1 {
- ; CHECK: fma.rn.ftz.f32
+; CHECK-LABEL: fma_float_ftz(
+; CHECK: {
+; CHECK-NEXT: .reg .f32 %f<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f32 %f1, [fma_float_ftz_param_0];
+; CHECK-NEXT: ld.param.f32 %f2, [fma_float_ftz_param_1];
+; CHECK-NEXT: ld.param.f32 %f3, [fma_float_ftz_param_2];
+; CHECK-NEXT: fma.rn.ftz.f32 %f4, %f1, %f2, %f3;
+; CHECK-NEXT: st.param.f32 [func_retval0+0], %f4;
+; CHECK-NEXT: ret;
%x = call float @llvm.fma.f32(float %a, float %b, float %c)
ret float %x
}
-; CHECK-LABEL: @fma_double
define double @fma_double(double %a, double %b, double %c) {
- ; CHECK: fma.rn.f64
+; CHECK-LABEL: fma_double(
+; CHECK: {
+; CHECK-NEXT: .reg .f64 %fd<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.f64 %fd1, [fma_double_param_0];
+; CHECK-NEXT: ld.param.f64 %fd2, [fma_double_param_1];
+; CHECK-NEXT: ld.param.f64 %fd3, [fma_double_param_2];
+; CHECK-NEXT: fma.rn.f64 %fd4, %fd1, %fd2, %fd3;
+; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd4;
+; CHECK-NEXT: ret;
%x = call double @llvm.fma.f64(double %a, double %b, double %c)
ret double %x
}
More information about the llvm-commits
mailing list