[llvm] expandFMINIMUM_FMAXIMUM: FMAXNUM/FMINNUM treat +0>-0 (PR #137367)
YunQiang Su via llvm-commits
llvm-commits at lists.llvm.org
Sun Dec 7 21:41:28 PST 2025
https://github.com/wzssyqa updated https://github.com/llvm/llvm-project/pull/137367
>From 924946cee5b062e4a5298d28e97ac1fe6df732da Mon Sep 17 00:00:00 2001
From: YunQiang Su <syq at debian.org>
Date: Sat, 26 Apr 2025 01:15:56 +0800
Subject: [PATCH 1/3] expandFMINIMUM_FMAXIMUM: FMAXNUM/FMINNUM treat +0>-0
ISD::FMAXNUM and ISD::FMINNUM treat +0.0>-0.0 now,
so let's set MinMaxMustRespectOrderedZero for it.
---
llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp | 3 +--
1 file changed, 1 insertion(+), 2 deletions(-)
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index 76278e99d4082..8f6e0011e5b01 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -8745,8 +8745,6 @@ SDValue TargetLowering::expandFMINIMUM_FMAXIMUM(SDNode *N,
unsigned CompOpcIeee = IsMax ? ISD::FMAXNUM_IEEE : ISD::FMINNUM_IEEE;
unsigned CompOpc = IsMax ? ISD::FMAXNUM : ISD::FMINNUM;
- // FIXME: We should probably define fminnum/fmaxnum variants with correct
- // signed zero behavior.
bool MinMaxMustRespectOrderedZero = false;
if (isOperationLegalOrCustom(CompOpcIeee, VT)) {
@@ -8754,6 +8752,7 @@ SDValue TargetLowering::expandFMINIMUM_FMAXIMUM(SDNode *N,
MinMaxMustRespectOrderedZero = true;
} else if (isOperationLegalOrCustom(CompOpc, VT)) {
MinMax = DAG.getNode(CompOpc, DL, VT, LHS, RHS, Flags);
+ MinMaxMustRespectOrderedZero = true;
} else {
if (VT.isVector() && !isOperationLegalOrCustom(ISD::VSELECT, VT))
return DAG.UnrollVectorOp(N);
>From a0a2c79ab0bc48a749c0697c85f49002cb08fc54 Mon Sep 17 00:00:00 2001
From: YunQiang Su <yunqiang at isrc.iscas.ac.cn>
Date: Sat, 26 Apr 2025 12:31:17 +0800
Subject: [PATCH 2/3] Fix testcase
---
llvm/test/CodeGen/NVPTX/math-intrins.ll | 200 +++++++-----------
.../test/CodeGen/PowerPC/fminimum-fmaximum.ll | 133 +++---------
2 files changed, 110 insertions(+), 223 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/math-intrins.ll b/llvm/test/CodeGen/NVPTX/math-intrins.ll
index 625c93c3f0a53..c11d0dae4c48c 100644
--- a/llvm/test/CodeGen/NVPTX/math-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/math-intrins.ll
@@ -681,22 +681,16 @@ define half @minimum_half(half %a, half %b) {
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<8>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_param_0];
-; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_param_1];
-; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r2;
-; CHECK-NOF16-NEXT: min.f32 %r3, %r1, %r2;
-; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;
-; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, -2147483648;
-; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;
-; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %r4, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
+; 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: st.param.f32 [func_retval0], %f4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_float(
@@ -727,19 +721,15 @@ define float @minimum_float(float %a, float %b) {
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<6>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm1_param_0];
-; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
-; CHECK-NOF16-NEXT: min.f32 %r2, %r1, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;
-; CHECK-NOF16-NEXT: selp.f32 %r4, %r1, %r3, %p2;
-; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %r3, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r5, %r4, %r3, %p3;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
+; 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: st.param.f32 [func_retval0], %f3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_imm1(
@@ -768,19 +758,15 @@ define float @minimum_imm1(float %a) {
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<6>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm2_param_0];
-; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
-; CHECK-NOF16-NEXT: min.f32 %r2, %r1, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;
-; CHECK-NOF16-NEXT: selp.f32 %r4, %r1, %r3, %p2;
-; CHECK-NOF16-NEXT: setp.eq.f32 %p3, %r3, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r5, %r4, %r3, %p3;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r5;
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
+; 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: st.param.f32 [func_retval0], %f3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_imm2(
@@ -809,22 +795,16 @@ define float @minimum_imm2(float %a) {
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<8>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_ftz_param_0];
-; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_ftz_param_1];
-; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %r1, %r2;
-; CHECK-NOF16-NEXT: min.ftz.f32 %r3, %r1, %r2;
-; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, -2147483648;
-; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, -2147483648;
-; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;
-; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %r4, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
+; 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: st.param.f32 [func_retval0], %f4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_float_ftz(
@@ -855,22 +835,16 @@ define float @minimum_float_ftz(float %a, float %b) #1 {
define double @minimum_double(double %a, double %b) {
; CHECK-LABEL: minimum_double(
; CHECK: {
-; CHECK-NEXT: .reg .pred %p<5>;
-; CHECK-NEXT: .reg .b64 %rd<8>;
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd1, [minimum_double_param_0];
-; CHECK-NEXT: ld.param.b64 %rd2, [minimum_double_param_1];
-; CHECK-NEXT: setp.nan.f64 %p1, %rd1, %rd2;
-; CHECK-NEXT: min.f64 %rd3, %rd1, %rd2;
-; CHECK-NEXT: selp.f64 %rd4, 0d7FF8000000000000, %rd3, %p1;
-; CHECK-NEXT: setp.eq.b64 %p2, %rd1, -9223372036854775808;
-; CHECK-NEXT: selp.f64 %rd5, %rd1, %rd4, %p2;
-; CHECK-NEXT: setp.eq.b64 %p3, %rd2, -9223372036854775808;
-; CHECK-NEXT: selp.f64 %rd6, %rd2, %rd5, %p3;
-; CHECK-NEXT: setp.eq.f64 %p4, %rd4, 0d0000000000000000;
-; CHECK-NEXT: selp.f64 %rd7, %rd6, %rd4, %p4;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;
+; CHECK-NEXT: ld.param.f64 %fd1, [minimum_double_param_0];
+; 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: st.param.f64 [func_retval0], %fd4;
; CHECK-NEXT: ret;
%x = call double @llvm.minimum.f64(double %a, double %b)
ret double %x
@@ -1212,17 +1186,15 @@ define half @maximum_half(half %a, half %b) {
define float @maximum_imm1(float %a) {
; CHECK-NOF16-LABEL: maximum_imm1(
; CHECK-NOF16: {
-; CHECK-NOF16-NEXT: .reg .pred %p<3>;
-; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm1_param_0];
-; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
-; CHECK-NOF16-NEXT: max.f32 %r2, %r1, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
-; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %r3, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r4, 0f00000000, %r3, %p2;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;
+; 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: st.param.f32 [func_retval0], %f3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_imm1(
@@ -1251,17 +1223,15 @@ define float @maximum_imm1(float %a) {
define float @maximum_imm2(float %a) {
; CHECK-NOF16-LABEL: maximum_imm2(
; CHECK-NOF16: {
-; CHECK-NOF16-NEXT: .reg .pred %p<3>;
-; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm2_param_0];
-; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
-; CHECK-NOF16-NEXT: max.f32 %r2, %r1, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
-; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %r3, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r4, 0f00000000, %r3, %p2;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;
+; 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: st.param.f32 [func_retval0], %f3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_imm2(
@@ -1290,22 +1260,16 @@ define float @maximum_imm2(float %a) {
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<8>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_param_0];
-; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_param_1];
-; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r2;
-; CHECK-NOF16-NEXT: max.f32 %r3, %r1, %r2;
-; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, 0;
-; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, 0;
-; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;
-; CHECK-NOF16-NEXT: setp.eq.f32 %p4, %r4, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
+; 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: st.param.f32 [func_retval0], %f4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_float(
@@ -1336,22 +1300,16 @@ define float @maximum_float(float %a, float %b) {
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<8>;
+; CHECK-NOF16-NEXT: .reg .pred %p<2>;
+; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_ftz_param_0];
-; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_ftz_param_1];
-; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %r1, %r2;
-; CHECK-NOF16-NEXT: max.ftz.f32 %r3, %r1, %r2;
-; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p2, %r1, 0;
-; CHECK-NOF16-NEXT: selp.f32 %r5, %r1, %r4, %p2;
-; CHECK-NOF16-NEXT: setp.eq.b32 %p3, %r2, 0;
-; CHECK-NOF16-NEXT: selp.f32 %r6, %r2, %r5, %p3;
-; CHECK-NOF16-NEXT: setp.eq.ftz.f32 %p4, %r4, 0f00000000;
-; CHECK-NOF16-NEXT: selp.f32 %r7, %r6, %r4, %p4;
-; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r7;
+; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
+; 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: st.param.f32 [func_retval0], %f4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_float_ftz(
@@ -1382,22 +1340,16 @@ define float @maximum_float_ftz(float %a, float %b) #1 {
define double @maximum_double(double %a, double %b) {
; CHECK-LABEL: maximum_double(
; CHECK: {
-; CHECK-NEXT: .reg .pred %p<5>;
-; CHECK-NEXT: .reg .b64 %rd<8>;
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b64 %fd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b64 %rd1, [maximum_double_param_0];
-; CHECK-NEXT: ld.param.b64 %rd2, [maximum_double_param_1];
-; CHECK-NEXT: setp.nan.f64 %p1, %rd1, %rd2;
-; CHECK-NEXT: max.f64 %rd3, %rd1, %rd2;
-; CHECK-NEXT: selp.f64 %rd4, 0d7FF8000000000000, %rd3, %p1;
-; CHECK-NEXT: setp.eq.b64 %p2, %rd1, 0;
-; CHECK-NEXT: selp.f64 %rd5, %rd1, %rd4, %p2;
-; CHECK-NEXT: setp.eq.b64 %p3, %rd2, 0;
-; CHECK-NEXT: selp.f64 %rd6, %rd2, %rd5, %p3;
-; CHECK-NEXT: setp.eq.f64 %p4, %rd4, 0d0000000000000000;
-; CHECK-NEXT: selp.f64 %rd7, %rd6, %rd4, %p4;
-; CHECK-NEXT: st.param.b64 [func_retval0], %rd7;
+; CHECK-NEXT: ld.param.f64 %fd1, [maximum_double_param_0];
+; 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: st.param.f64 [func_retval0], %fd4;
; CHECK-NEXT: ret;
%x = call double @llvm.maximum.f64(double %a, double %b)
ret double %x
diff --git a/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll b/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll
index a99c25a4e4479..39cf136e10d77 100644
--- a/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll
+++ b/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll
@@ -301,22 +301,13 @@ define <4 x float> @v4f32_minimum(<4 x float> %a, <4 x float> %b) {
; VSX-NEXT: xvcmpeqsp 1, 35, 35
; VSX-NEXT: xvcmpeqsp 2, 34, 34
; VSX-NEXT: addis 3, 2, .LCPI4_0 at toc@ha
-; VSX-NEXT: xxleqv 36, 36, 36
-; VSX-NEXT: xvminsp 0, 34, 35
-; VSX-NEXT: vslw 4, 4, 4
; VSX-NEXT: addi 3, 3, .LCPI4_0 at toc@l
; VSX-NEXT: xxlnor 1, 1, 1
; VSX-NEXT: xxlnor 2, 2, 2
-; VSX-NEXT: vcmpequw 5, 2, 4
+; VSX-NEXT: xvminsp 0, 34, 35
; VSX-NEXT: xxlor 1, 2, 1
; VSX-NEXT: lxvd2x 2, 0, 3
-; VSX-NEXT: xxsel 0, 0, 2, 1
-; VSX-NEXT: xxlxor 2, 2, 2
-; VSX-NEXT: xvcmpeqsp 2, 0, 2
-; VSX-NEXT: xxsel 1, 0, 34, 37
-; VSX-NEXT: vcmpequw 2, 3, 4
-; VSX-NEXT: xxsel 1, 1, 35, 34
-; VSX-NEXT: xxsel 34, 0, 1, 2
+; VSX-NEXT: xxsel 34, 0, 2, 1
; VSX-NEXT: blr
;
; AIX-LABEL: v4f32_minimum:
@@ -324,21 +315,12 @@ define <4 x float> @v4f32_minimum(<4 x float> %a, <4 x float> %b) {
; AIX-NEXT: xvcmpeqsp 1, 35, 35
; AIX-NEXT: xvcmpeqsp 2, 34, 34
; AIX-NEXT: ld 3, L..C4(2) # %const.0
-; AIX-NEXT: xxleqv 36, 36, 36
; AIX-NEXT: xvminsp 0, 34, 35
-; AIX-NEXT: vslw 4, 4, 4
; AIX-NEXT: xxlnor 1, 1, 1
; AIX-NEXT: xxlnor 2, 2, 2
-; AIX-NEXT: vcmpequw 5, 2, 4
; AIX-NEXT: xxlor 1, 2, 1
; AIX-NEXT: lxvw4x 2, 0, 3
-; AIX-NEXT: xxsel 0, 0, 2, 1
-; AIX-NEXT: xxlxor 2, 2, 2
-; AIX-NEXT: xvcmpeqsp 2, 0, 2
-; AIX-NEXT: xxsel 1, 0, 34, 37
-; AIX-NEXT: vcmpequw 2, 3, 4
-; AIX-NEXT: xxsel 1, 1, 35, 34
-; AIX-NEXT: xxsel 34, 0, 1, 2
+; AIX-NEXT: xxsel 34, 0, 2, 1
; AIX-NEXT: blr
entry:
%m = call <4 x float> @llvm.minimum.v4f32(<4 x float> %a, <4 x float> %b)
@@ -377,16 +359,9 @@ define <4 x float> @v4f32_maximum(<4 x float> %a, <4 x float> %b) {
; VSX-NEXT: xxlnor 1, 1, 1
; VSX-NEXT: xxlnor 2, 2, 2
; VSX-NEXT: xvmaxsp 0, 34, 35
-; VSX-NEXT: xxlxor 36, 36, 36
-; VSX-NEXT: vcmpequw 5, 2, 4
; VSX-NEXT: xxlor 1, 2, 1
; VSX-NEXT: lxvd2x 2, 0, 3
-; VSX-NEXT: xxsel 0, 0, 2, 1
-; VSX-NEXT: xvcmpeqsp 2, 0, 36
-; VSX-NEXT: xxsel 1, 0, 34, 37
-; VSX-NEXT: vcmpequw 2, 3, 4
-; VSX-NEXT: xxsel 1, 1, 35, 34
-; VSX-NEXT: xxsel 34, 0, 1, 2
+; VSX-NEXT: xxsel 34, 0, 2, 1
; VSX-NEXT: blr
;
; AIX-LABEL: v4f32_maximum:
@@ -395,18 +370,11 @@ define <4 x float> @v4f32_maximum(<4 x float> %a, <4 x float> %b) {
; AIX-NEXT: xvcmpeqsp 2, 34, 34
; AIX-NEXT: ld 3, L..C5(2) # %const.0
; AIX-NEXT: xvmaxsp 0, 34, 35
-; AIX-NEXT: xxlxor 36, 36, 36
; AIX-NEXT: xxlnor 1, 1, 1
; AIX-NEXT: xxlnor 2, 2, 2
-; AIX-NEXT: vcmpequw 5, 2, 4
; AIX-NEXT: xxlor 1, 2, 1
; AIX-NEXT: lxvw4x 2, 0, 3
-; AIX-NEXT: xxsel 0, 0, 2, 1
-; AIX-NEXT: xvcmpeqsp 2, 0, 36
-; AIX-NEXT: xxsel 1, 0, 34, 37
-; AIX-NEXT: vcmpequw 2, 3, 4
-; AIX-NEXT: xxsel 1, 1, 35, 34
-; AIX-NEXT: xxsel 34, 0, 1, 2
+; AIX-NEXT: xxsel 34, 0, 2, 1
; AIX-NEXT: blr
entry:
%m = call <4 x float> @llvm.maximum.v4f32(<4 x float> %a, <4 x float> %b)
@@ -493,47 +461,28 @@ define <2 x double> @v2f64_minimum(<2 x double> %a, <2 x double> %b) {
; VSX-LABEL: v2f64_minimum:
; VSX: # %bb.0: # %entry
; VSX-NEXT: addis 3, 2, .LCPI6_0 at toc@ha
-; VSX-NEXT: xvcmpeqdp 36, 35, 35
-; VSX-NEXT: xvcmpeqdp 37, 34, 34
-; VSX-NEXT: addi 3, 3, .LCPI6_0 at toc@l
-; VSX-NEXT: xxlnor 36, 36, 36
-; VSX-NEXT: xxlnor 37, 37, 37
; VSX-NEXT: xvmindp 0, 34, 35
+; VSX-NEXT: xvcmpeqdp 35, 35, 35
+; VSX-NEXT: addi 3, 3, .LCPI6_0 at toc@l
+; VSX-NEXT: xvcmpeqdp 34, 34, 34
+; VSX-NEXT: xxlnor 35, 35, 35
+; VSX-NEXT: xxlnor 34, 34, 34
; VSX-NEXT: lxvd2x 2, 0, 3
-; VSX-NEXT: addis 3, 2, .LCPI6_1 at toc@ha
-; VSX-NEXT: xxlor 1, 37, 36
-; VSX-NEXT: addi 3, 3, .LCPI6_1 at toc@l
-; VSX-NEXT: lxvd2x 36, 0, 3
-; VSX-NEXT: vcmpequd 5, 2, 4
-; VSX-NEXT: xxsel 0, 0, 2, 1
-; VSX-NEXT: xxlxor 2, 2, 2
-; VSX-NEXT: xxsel 1, 0, 34, 37
-; VSX-NEXT: vcmpequd 2, 3, 4
-; VSX-NEXT: xxsel 1, 1, 35, 34
-; VSX-NEXT: xvcmpeqdp 34, 0, 2
-; VSX-NEXT: xxsel 34, 0, 1, 34
+; VSX-NEXT: xxlor 1, 34, 35
+; VSX-NEXT: xxsel 34, 0, 2, 1
; VSX-NEXT: blr
;
; AIX-LABEL: v2f64_minimum:
; AIX: # %bb.0: # %entry
; AIX-NEXT: ld 3, L..C6(2) # %const.0
-; AIX-NEXT: xvcmpeqdp 36, 35, 35
-; AIX-NEXT: xvcmpeqdp 37, 34, 34
-; AIX-NEXT: lxvd2x 2, 0, 3
-; AIX-NEXT: ld 3, L..C7(2) # %const.1
-; AIX-NEXT: xxlnor 36, 36, 36
-; AIX-NEXT: xxlnor 37, 37, 37
; AIX-NEXT: xvmindp 0, 34, 35
-; AIX-NEXT: xxlor 1, 37, 36
-; AIX-NEXT: lxvd2x 36, 0, 3
-; AIX-NEXT: vcmpequd 5, 2, 4
-; AIX-NEXT: xxsel 0, 0, 2, 1
-; AIX-NEXT: xxlxor 2, 2, 2
-; AIX-NEXT: xxsel 1, 0, 34, 37
-; AIX-NEXT: vcmpequd 2, 3, 4
-; AIX-NEXT: xxsel 1, 1, 35, 34
-; AIX-NEXT: xvcmpeqdp 34, 0, 2
-; AIX-NEXT: xxsel 34, 0, 1, 34
+; AIX-NEXT: xvcmpeqdp 35, 35, 35
+; AIX-NEXT: lxvd2x 2, 0, 3
+; AIX-NEXT: xvcmpeqdp 34, 34, 34
+; AIX-NEXT: xxlnor 35, 35, 35
+; AIX-NEXT: xxlnor 34, 34, 34
+; AIX-NEXT: xxlor 1, 34, 35
+; AIX-NEXT: xxsel 34, 0, 2, 1
; AIX-NEXT: blr
entry:
%m = call <2 x double> @llvm.minimum.v2f64(<2 x double> %a, <2 x double> %b)
@@ -618,42 +567,28 @@ define <2 x double> @v2f64_maximum(<2 x double> %a, <2 x double> %b) {
; VSX-LABEL: v2f64_maximum:
; VSX: # %bb.0: # %entry
; VSX-NEXT: addis 3, 2, .LCPI7_0 at toc@ha
-; VSX-NEXT: xvcmpeqdp 36, 35, 35
-; VSX-NEXT: xvcmpeqdp 37, 34, 34
-; VSX-NEXT: addi 3, 3, .LCPI7_0 at toc@l
-; VSX-NEXT: xxlnor 36, 36, 36
-; VSX-NEXT: xxlnor 37, 37, 37
; VSX-NEXT: xvmaxdp 0, 34, 35
+; VSX-NEXT: xvcmpeqdp 35, 35, 35
+; VSX-NEXT: addi 3, 3, .LCPI7_0 at toc@l
+; VSX-NEXT: xvcmpeqdp 34, 34, 34
+; VSX-NEXT: xxlnor 35, 35, 35
+; VSX-NEXT: xxlnor 34, 34, 34
; VSX-NEXT: lxvd2x 2, 0, 3
-; VSX-NEXT: xxlor 1, 37, 36
-; VSX-NEXT: xxlxor 36, 36, 36
-; VSX-NEXT: vcmpequd 5, 2, 4
-; VSX-NEXT: xxsel 0, 0, 2, 1
-; VSX-NEXT: xxsel 1, 0, 34, 37
-; VSX-NEXT: vcmpequd 2, 3, 4
-; VSX-NEXT: xxsel 1, 1, 35, 34
-; VSX-NEXT: xvcmpeqdp 34, 0, 36
-; VSX-NEXT: xxsel 34, 0, 1, 34
+; VSX-NEXT: xxlor 1, 34, 35
+; VSX-NEXT: xxsel 34, 0, 2, 1
; VSX-NEXT: blr
;
; AIX-LABEL: v2f64_maximum:
; AIX: # %bb.0: # %entry
-; AIX-NEXT: ld 3, L..C8(2) # %const.0
-; AIX-NEXT: xvcmpeqdp 36, 35, 35
-; AIX-NEXT: xvcmpeqdp 37, 34, 34
-; AIX-NEXT: lxvd2x 2, 0, 3
-; AIX-NEXT: xxlnor 36, 36, 36
-; AIX-NEXT: xxlnor 37, 37, 37
+; AIX-NEXT: ld 3, L..C7(2) # %const.0
; AIX-NEXT: xvmaxdp 0, 34, 35
-; AIX-NEXT: xxlor 1, 37, 36
-; AIX-NEXT: xxlxor 36, 36, 36
-; AIX-NEXT: vcmpequd 5, 2, 4
-; AIX-NEXT: xxsel 0, 0, 2, 1
-; AIX-NEXT: xxsel 1, 0, 34, 37
-; AIX-NEXT: vcmpequd 2, 3, 4
-; AIX-NEXT: xxsel 1, 1, 35, 34
-; AIX-NEXT: xvcmpeqdp 34, 0, 36
-; AIX-NEXT: xxsel 34, 0, 1, 34
+; AIX-NEXT: xvcmpeqdp 35, 35, 35
+; AIX-NEXT: lxvd2x 2, 0, 3
+; AIX-NEXT: xvcmpeqdp 34, 34, 34
+; AIX-NEXT: xxlnor 35, 35, 35
+; AIX-NEXT: xxlnor 34, 34, 34
+; AIX-NEXT: xxlor 1, 34, 35
+; AIX-NEXT: xxsel 34, 0, 2, 1
; AIX-NEXT: blr
entry:
%m = call <2 x double> @llvm.maximum.v2f64(<2 x double> %a, <2 x double> %b)
>From 9018fb93fd7bbd5059c258b913fd465105025d93 Mon Sep 17 00:00:00 2001
From: YunQiang Su <yunqiang at isrc.iscas.ac.cn>
Date: Thu, 4 Dec 2025 16:57:26 +0800
Subject: [PATCH 3/3] Update nvptx tests
---
llvm/test/CodeGen/NVPTX/math-intrins.ll | 132 ++++++++++++------------
1 file changed, 66 insertions(+), 66 deletions(-)
diff --git a/llvm/test/CodeGen/NVPTX/math-intrins.ll b/llvm/test/CodeGen/NVPTX/math-intrins.ll
index c11d0dae4c48c..6fb0112631af7 100644
--- a/llvm/test/CodeGen/NVPTX/math-intrins.ll
+++ b/llvm/test/CodeGen/NVPTX/math-intrins.ll
@@ -682,15 +682,15 @@ define float @minimum_float(float %a, float %b) {
; CHECK-NOF16-LABEL: minimum_float(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_param_0];
-; 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: st.param.f32 [func_retval0], %f4;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_param_1];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r2;
+; CHECK-NOF16-NEXT: min.f32 %r3, %r1, %r2;
+; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_float(
@@ -722,14 +722,14 @@ define float @minimum_imm1(float %a) {
; CHECK-NOF16-LABEL: minimum_imm1(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm1_param_0];
-; 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: st.param.f32 [func_retval0], %f3;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm1_param_0];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-NOF16-NEXT: min.f32 %r2, %r1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_imm1(
@@ -759,14 +759,14 @@ define float @minimum_imm2(float %a) {
; CHECK-NOF16-LABEL: minimum_imm2(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_imm2_param_0];
-; 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: st.param.f32 [func_retval0], %f3;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_imm2_param_0];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-NOF16-NEXT: min.f32 %r2, %r1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_imm2(
@@ -796,15 +796,15 @@ define float @minimum_float_ftz(float %a, float %b) #1 {
; CHECK-NOF16-LABEL: minimum_float_ftz(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.f32 %f1, [minimum_float_ftz_param_0];
-; 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: st.param.f32 [func_retval0], %f4;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [minimum_float_ftz_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [minimum_float_ftz_param_1];
+; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %r1, %r2;
+; CHECK-NOF16-NEXT: min.ftz.f32 %r3, %r1, %r2;
+; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: minimum_float_ftz(
@@ -836,15 +836,15 @@ define double @minimum_double(double %a, double %b) {
; CHECK-LABEL: minimum_double(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
-; CHECK-NEXT: .reg .b64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.f64 %fd1, [minimum_double_param_0];
-; 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: st.param.f64 [func_retval0], %fd4;
+; CHECK-NEXT: ld.param.b64 %rd1, [minimum_double_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [minimum_double_param_1];
+; CHECK-NEXT: setp.nan.f64 %p1, %rd1, %rd2;
+; CHECK-NEXT: min.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: selp.f64 %rd4, 0d7FF8000000000000, %rd3, %p1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;
; CHECK-NEXT: ret;
%x = call double @llvm.minimum.f64(double %a, double %b)
ret double %x
@@ -1187,14 +1187,14 @@ define float @maximum_imm1(float %a) {
; CHECK-NOF16-LABEL: maximum_imm1(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
; 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: st.param.f32 [func_retval0], %f3;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm1_param_0];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-NOF16-NEXT: max.f32 %r2, %r1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_imm1(
@@ -1224,14 +1224,14 @@ define float @maximum_imm2(float %a) {
; CHECK-NOF16-LABEL: maximum_imm2(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<4>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<4>;
; 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: st.param.f32 [func_retval0], %f3;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_imm2_param_0];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r1;
+; CHECK-NOF16-NEXT: max.f32 %r2, %r1, 0f00000000;
+; CHECK-NOF16-NEXT: selp.f32 %r3, 0f7FC00000, %r2, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_imm2(
@@ -1261,15 +1261,15 @@ define float @maximum_float(float %a, float %b) {
; CHECK-NOF16-LABEL: maximum_float(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_param_0];
-; 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: st.param.f32 [func_retval0], %f4;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_param_1];
+; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r1, %r2;
+; CHECK-NOF16-NEXT: max.f32 %r3, %r1, %r2;
+; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_float(
@@ -1301,15 +1301,15 @@ define float @maximum_float_ftz(float %a, float %b) #1 {
; CHECK-NOF16-LABEL: maximum_float_ftz(
; CHECK-NOF16: {
; CHECK-NOF16-NEXT: .reg .pred %p<2>;
-; CHECK-NOF16-NEXT: .reg .b32 %f<5>;
+; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.f32 %f1, [maximum_float_ftz_param_0];
-; 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: st.param.f32 [func_retval0], %f4;
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [maximum_float_ftz_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [maximum_float_ftz_param_1];
+; CHECK-NOF16-NEXT: setp.nan.ftz.f32 %p1, %r1, %r2;
+; CHECK-NOF16-NEXT: max.ftz.f32 %r3, %r1, %r2;
+; CHECK-NOF16-NEXT: selp.f32 %r4, 0f7FC00000, %r3, %p1;
+; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r4;
; CHECK-NOF16-NEXT: ret;
;
; CHECK-F16-LABEL: maximum_float_ftz(
@@ -1341,15 +1341,15 @@ define double @maximum_double(double %a, double %b) {
; CHECK-LABEL: maximum_double(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
-; CHECK-NEXT: .reg .b64 %fd<5>;
+; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.f64 %fd1, [maximum_double_param_0];
-; 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: st.param.f64 [func_retval0], %fd4;
+; CHECK-NEXT: ld.param.b64 %rd1, [maximum_double_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [maximum_double_param_1];
+; CHECK-NEXT: setp.nan.f64 %p1, %rd1, %rd2;
+; CHECK-NEXT: max.f64 %rd3, %rd1, %rd2;
+; CHECK-NEXT: selp.f64 %rd4, 0d7FF8000000000000, %rd3, %p1;
+; CHECK-NEXT: st.param.b64 [func_retval0], %rd4;
; CHECK-NEXT: ret;
%x = call double @llvm.maximum.f64(double %a, double %b)
ret double %x
More information about the llvm-commits
mailing list