Skip to content

SelectionDAG: Improve expandFMINIMUM_FMAXIMUM#137367

Draft
wzssyqa wants to merge 27 commits intollvm:mainfrom
wzssyqa:expandFMINIMUM_NO_IEEE
Draft

SelectionDAG: Improve expandFMINIMUM_FMAXIMUM#137367
wzssyqa wants to merge 27 commits intollvm:mainfrom
wzssyqa:expandFMINIMUM_NO_IEEE

Conversation

@wzssyqa
Copy link
Copy Markdown
Contributor

@wzssyqa wzssyqa commented Apr 25, 2025

  1. Use FMAXNUM_IEEE, FMAXNUM, FMAXIMUMNUM if Legal or FMAXNUM_IEEE, FMAXIMUMNUM if Custom.
    We try the Legal ones first and then Custom ones then.
  2. For NaNs:
    • If Max/Min Opc can be used, we set the result to qNaN if LHS is unordered with RHS.
    • If we use FCMP, we move all NaNs in RHS to LHS only, as we perfer LHS if LHS is unordered with RHS.
  3. For +0 vs -0
    • If Max/Min Opc can be used, as all of them process it correctly now, we can just skip it.
    • If we use FCMP, we prefer RHS if equal, and then we determine the Sign of LHS when the result equals 0:
      1. if IsMax and LHS is positive, we switch LHS.
      2. if IsMax and LHS is negitive, we switch MinMax.
      3. if not IsMax and LHS is positive, we switch MinMax.
      4. if not IsMax and LHS is negitive, we switch LHS.
  4. we introduce a new static function determineFloatSign, which return the SDValue to determine the Sign of a SDValue.

@wzssyqa wzssyqa marked this pull request as draft April 25, 2025 17:19
@llvmbot llvmbot added the llvm:SelectionDAG SelectionDAGISel as well label Apr 25, 2025
@llvmbot
Copy link
Copy Markdown
Member

llvmbot commented Apr 25, 2025

@llvm/pr-subscribers-backend-mips
@llvm/pr-subscribers-backend-amdgpu
@llvm/pr-subscribers-backend-nvptx
@llvm/pr-subscribers-backend-powerpc

@llvm/pr-subscribers-llvm-selectiondag

Author: YunQiang Su (wzssyqa)

Changes

ISD::FMAXNUM and ISD::FMINNUM treat +0.0>-0.0 now, so let's set MinMaxMustRespectOrderedZero for it.


Full diff: https://github.com/llvm/llvm-project/pull/137367.diff

1 Files Affected:

  • (modified) llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp (+1-2)
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index 6930b54ddb14a..7baed2d591514 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -8573,8 +8573,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)) {
@@ -8582,6 +8580,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);

@wzssyqa
Copy link
Copy Markdown
Contributor Author

wzssyqa commented Dec 2, 2025

See:#168838

@wzssyqa wzssyqa closed this Dec 2, 2025
@arsenm arsenm reopened this Dec 2, 2025
@arsenm
Copy link
Copy Markdown
Contributor

arsenm commented Dec 2, 2025

See:#168838

This was reverted, and the contentious part is signaling nan handling. One of the reasons simply reverting this is wrong is the signed zero behavior should remain regardless of that

wzssyqa and others added 2 commits December 4, 2025 10:45
ISD::FMAXNUM and ISD::FMINNUM treat +0.0>-0.0 now,
so let's set MinMaxMustRespectOrderedZero for it.
@wzssyqa wzssyqa force-pushed the expandFMINIMUM_NO_IEEE branch from c153bb8 to a0a2c79 Compare December 4, 2025 02:49
@wzssyqa wzssyqa marked this pull request as ready for review December 4, 2025 02:49
@github-actions
Copy link
Copy Markdown

github-actions bot commented Dec 4, 2025

🐧 Linux x64 Test Results

  • 169370 tests passed
  • 3047 tests skipped
  • 5 tests failed

Failed Tests

(click on a test name to see its output)

LLVM

LLVM.CodeGen/NVPTX/atomicrmw-sm60.ll
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50 | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll --check-prefix=SM60
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -march=nvptx64 -mcpu=sm_60 -mattr=+ptx50
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll --check-prefix=SM60
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:1839:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:1910:2: note: scanning from here
# | {
# |  ^
# | <stdin>:1911:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:1874:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:1940:2: note: scanning from here
# | {
# |  ^
# | <stdin>:1941:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:2002:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:2067:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2068:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:2036:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:2096:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2097:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:2244:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2303:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2304:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:2296:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2354:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2355:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:2548:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2609:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2610:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll:2604:14: error: SM60-NEXT: expected string not found in input
# | ; SM60-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2665:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2666:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |              .
# |              .
# |              .
# |           1905:  // .globl fminimum_acq_rel_float_global_cta // -- Begin function fminimum_acq_rel_float_global_cta 
# |           1906: .visible .func (.param .b32 func_retval0) fminimum_acq_rel_float_global_cta( 
# |           1907:  .param .b64 fminimum_acq_rel_float_global_cta_param_0, 
# |           1908:  .param .b32 fminimum_acq_rel_float_global_cta_param_1 
# |           1909: ) // @fminimum_acq_rel_float_global_cta 
# |           1910: { 
# | next:1839'0      X error: no match found
# |           1911:  .reg .pred %p<3>; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~
# | next:1839'1      ?                  possible intended match
# |           1912:  .reg .b32 %r<6>; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~
# |           1913:  .reg .b64 %rd<2>; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~
# |           1914:  
# | next:1839'0     ~
# |           1915: // %bb.0: 
# | next:1839'0     ~~~~~~~~~~
# |           1916:  ld.param.b32 %r2, [fminimum_acq_rel_float_global_cta_param_1]; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           1935:  // .globl fmaximum_acq_rel_float_global_cta // -- Begin function fmaximum_acq_rel_float_global_cta 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           1936: .visible .func (.param .b32 func_retval0) fmaximum_acq_rel_float_global_cta( 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           1937:  .param .b64 fmaximum_acq_rel_float_global_cta_param_0, 
# |           1938:  .param .b32 fmaximum_acq_rel_float_global_cta_param_1 
# |           1939: ) // @fmaximum_acq_rel_float_global_cta 
# |           1940: { 
# | next:1874'0      X error: no match found
# |           1941:  .reg .pred %p<3>; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~~
# | next:1874'1      ?                  possible intended match
# |           1942:  .reg .b32 %r<6>; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~
# |           1943:  .reg .b64 %rd<2>; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~~
# |           1944:  
# | next:1874'0     ~
# |           1945: // %bb.0: 
# | next:1874'0     ~~~~~~~~~~
# |           1946:  ld.param.b32 %r2, [fmaximum_acq_rel_float_global_cta_param_1]; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2062:  // .globl fminimum_acq_rel_double_global_cta // -- Begin function fminimum_acq_rel_double_global_cta 
# |           2063: .visible .func (.param .b64 func_retval0) fminimum_acq_rel_double_global_cta( 
# |           2064:  .param .b64 fminimum_acq_rel_double_global_cta_param_0, 
# |           2065:  .param .b64 fminimum_acq_rel_double_global_cta_param_1 
# |           2066: ) // @fminimum_acq_rel_double_global_cta 
# |           2067: { 
# | next:2002'0      X error: no match found
# |           2068:  .reg .pred %p<3>; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~
# | next:2002'1      ?                  possible intended match
# |           2069:  .reg .b64 %rd<7>; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~
# |           2070:  
# | next:2002'0     ~
# |           2071: // %bb.0: 
# | next:2002'0     ~~~~~~~~~~
# |           2072:  ld.param.b64 %rd3, [fminimum_acq_rel_double_global_cta_param_1]; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2073:  ld.param.b64 %rd2, [fminimum_acq_rel_double_global_cta_param_0]; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2091:  // .globl fmaximum_acq_rel_double_global_cta // -- Begin function fmaximum_acq_rel_double_global_cta 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2092: .visible .func (.param .b64 func_retval0) fmaximum_acq_rel_double_global_cta( 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2093:  .param .b64 fmaximum_acq_rel_double_global_cta_param_0, 
# |           2094:  .param .b64 fmaximum_acq_rel_double_global_cta_param_1 
# |           2095: ) // @fmaximum_acq_rel_double_global_cta 
# |           2096: { 
# | next:2036'0      X error: no match found
# |           2097:  .reg .pred %p<3>; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~
# | next:2036'1      ?                  possible intended match
# |           2098:  .reg .b64 %rd<7>; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~
# |           2099:  
# | next:2036'0     ~
# |           2100: // %bb.0: 
# | next:2036'0     ~~~~~~~~~~
# |           2101:  ld.param.b64 %rd3, [fmaximum_acq_rel_double_global_cta_param_1]; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2102:  ld.param.b64 %rd2, [fmaximum_acq_rel_double_global_cta_param_0]; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2298:  // .globl fminimum_acq_rel_half_global_cta // -- Begin function fminimum_acq_rel_half_global_cta 
# |           2299: .visible .func (.param .align 2 .b8 func_retval0[2]) fminimum_acq_rel_half_global_cta( 
# |           2300:  .param .b64 fminimum_acq_rel_half_global_cta_param_0, 
# |           2301:  .param .align 2 .b8 fminimum_acq_rel_half_global_cta_param_1[2] 
# |           2302: ) // @fminimum_acq_rel_half_global_cta 
# |           2303: { 
# | next:2244'0      X error: no match found
# |           2304:  .reg .pred %p<6>; 
# | next:2244'0     ~~~~~~~~~~~~~~~~~~~
# | next:2244'1      ?                  possible intended match
# |           2305:  .reg .b16 %rs<8>; 
# | next:2244'0     ~~~~~~~~~~~~~~~~~~~
# |           2306:  .reg .b32 %r<15>; 
# | next:2244'0     ~~~~~~~~~~~~~~~~~~~
# |           2307:  .reg .b64 %rd<3>; 
# | next:2244'0     ~~~~~~~~~~~~~~~~~~~
# |           2308:  
# | next:2244'0     ~
# |           2309: // %bb.0: 
# | next:2244'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2349:  // .globl fmaximum_acq_rel_half_global_cta // -- Begin function fmaximum_acq_rel_half_global_cta 
# | next:2244'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2350: .visible .func (.param .align 2 .b8 func_retval0[2]) fmaximum_acq_rel_half_global_cta( 
# | next:2244'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2351:  .param .b64 fmaximum_acq_rel_half_global_cta_param_0, 
# |           2352:  .param .align 2 .b8 fmaximum_acq_rel_half_global_cta_param_1[2] 
# |           2353: ) // @fmaximum_acq_rel_half_global_cta 
# |           2354: { 
# | next:2296'0      X error: no match found
# |           2355:  .reg .pred %p<6>; 
# | next:2296'0     ~~~~~~~~~~~~~~~~~~~
# | next:2296'1      ?                  possible intended match
# |           2356:  .reg .b16 %rs<8>; 
# | next:2296'0     ~~~~~~~~~~~~~~~~~~~
# |           2357:  .reg .b32 %r<15>; 
# | next:2296'0     ~~~~~~~~~~~~~~~~~~~
# |           2358:  .reg .b64 %rd<3>; 
# | next:2296'0     ~~~~~~~~~~~~~~~~~~~
# |           2359:  
# | next:2296'0     ~
# |           2360: // %bb.0: 
# | next:2296'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2604:  // .globl fminimum_acq_rel_bfloat_global_cta // -- Begin function fminimum_acq_rel_bfloat_global_cta 
# |           2605: .visible .func (.param .align 2 .b8 func_retval0[2]) fminimum_acq_rel_bfloat_global_cta( 
# |           2606:  .param .b64 fminimum_acq_rel_bfloat_global_cta_param_0, 
# |           2607:  .param .align 2 .b8 fminimum_acq_rel_bfloat_global_cta_param_1[2] 
# |           2608: ) // @fminimum_acq_rel_bfloat_global_cta 
# |           2609: { 
# | next:2548'0      X error: no match found
# |           2610:  .reg .pred %p<6>; 
# | next:2548'0     ~~~~~~~~~~~~~~~~~~~
# | next:2548'1      ?                  possible intended match
# |           2611:  .reg .b16 %rs<7>; 
# | next:2548'0     ~~~~~~~~~~~~~~~~~~~
# |           2612:  .reg .b32 %r<21>; 
# | next:2548'0     ~~~~~~~~~~~~~~~~~~~
# |           2613:  .reg .b64 %rd<3>; 
# | next:2548'0     ~~~~~~~~~~~~~~~~~~~
# |           2614:  
# | next:2548'0     ~
# |           2615: // %bb.0: 
# | next:2548'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2660:  // .globl fmaximum_acq_rel_bfloat_global_cta // -- Begin function fmaximum_acq_rel_bfloat_global_cta 
# | next:2548'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2661: .visible .func (.param .align 2 .b8 func_retval0[2]) fmaximum_acq_rel_bfloat_global_cta( 
# | next:2548'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2662:  .param .b64 fmaximum_acq_rel_bfloat_global_cta_param_0, 
# |           2663:  .param .align 2 .b8 fmaximum_acq_rel_bfloat_global_cta_param_1[2] 
# |           2664: ) // @fmaximum_acq_rel_bfloat_global_cta 
# |           2665: { 
# | next:2604'0      X error: no match found
# |           2666:  .reg .pred %p<6>; 
# | next:2604'0     ~~~~~~~~~~~~~~~~~~~
# | next:2604'1      ?                  possible intended match
# |           2667:  .reg .b16 %rs<7>; 
# | next:2604'0     ~~~~~~~~~~~~~~~~~~~
# |           2668:  .reg .b32 %r<21>; 
# | next:2604'0     ~~~~~~~~~~~~~~~~~~~
# |           2669:  .reg .b64 %rd<3>; 
# | next:2604'0     ~~~~~~~~~~~~~~~~~~~
# |           2670:  
# | next:2604'0     ~
# |           2671: // %bb.0: 
# | next:2604'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

LLVM.CodeGen/NVPTX/atomicrmw-sm70.ll
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63 | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll --check-prefix=SM70
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -march=nvptx64 -mcpu=sm_70 -mattr=+ptx63
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll --check-prefix=SM70
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:1839:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:1910:2: note: scanning from here
# | {
# |  ^
# | <stdin>:1911:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:1874:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:1940:2: note: scanning from here
# | {
# |  ^
# | <stdin>:1941:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:2002:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:2067:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2068:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:2036:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:2096:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2097:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:2218:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2277:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2278:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:2270:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2328:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2329:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:2522:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2583:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2584:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll:2578:14: error: SM70-NEXT: expected string not found in input
# | ; SM70-NEXT: .reg .pred %p<7>;
# |              ^
# | <stdin>:2639:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2640:2: note: possible intended match here
# |  .reg .pred %p<6>;
# |  ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |              .
# |              .
# |              .
# |           1905:  // .globl fminimum_acq_rel_float_global_cta // -- Begin function fminimum_acq_rel_float_global_cta 
# |           1906: .visible .func (.param .b32 func_retval0) fminimum_acq_rel_float_global_cta( 
# |           1907:  .param .b64 fminimum_acq_rel_float_global_cta_param_0, 
# |           1908:  .param .b32 fminimum_acq_rel_float_global_cta_param_1 
# |           1909: ) // @fminimum_acq_rel_float_global_cta 
# |           1910: { 
# | next:1839'0      X error: no match found
# |           1911:  .reg .pred %p<3>; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~
# | next:1839'1      ?                  possible intended match
# |           1912:  .reg .b32 %r<6>; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~
# |           1913:  .reg .b64 %rd<2>; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~
# |           1914:  
# | next:1839'0     ~
# |           1915: // %bb.0: 
# | next:1839'0     ~~~~~~~~~~
# |           1916:  ld.param.b32 %r2, [fminimum_acq_rel_float_global_cta_param_1]; 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           1935:  // .globl fmaximum_acq_rel_float_global_cta // -- Begin function fmaximum_acq_rel_float_global_cta 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           1936: .visible .func (.param .b32 func_retval0) fmaximum_acq_rel_float_global_cta( 
# | next:1839'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           1937:  .param .b64 fmaximum_acq_rel_float_global_cta_param_0, 
# |           1938:  .param .b32 fmaximum_acq_rel_float_global_cta_param_1 
# |           1939: ) // @fmaximum_acq_rel_float_global_cta 
# |           1940: { 
# | next:1874'0      X error: no match found
# |           1941:  .reg .pred %p<3>; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~~
# | next:1874'1      ?                  possible intended match
# |           1942:  .reg .b32 %r<6>; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~
# |           1943:  .reg .b64 %rd<2>; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~~
# |           1944:  
# | next:1874'0     ~
# |           1945: // %bb.0: 
# | next:1874'0     ~~~~~~~~~~
# |           1946:  ld.param.b32 %r2, [fmaximum_acq_rel_float_global_cta_param_1]; 
# | next:1874'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2062:  // .globl fminimum_acq_rel_double_global_cta // -- Begin function fminimum_acq_rel_double_global_cta 
# |           2063: .visible .func (.param .b64 func_retval0) fminimum_acq_rel_double_global_cta( 
# |           2064:  .param .b64 fminimum_acq_rel_double_global_cta_param_0, 
# |           2065:  .param .b64 fminimum_acq_rel_double_global_cta_param_1 
# |           2066: ) // @fminimum_acq_rel_double_global_cta 
# |           2067: { 
# | next:2002'0      X error: no match found
# |           2068:  .reg .pred %p<3>; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~
# | next:2002'1      ?                  possible intended match
# |           2069:  .reg .b64 %rd<7>; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~
# |           2070:  
# | next:2002'0     ~
# |           2071: // %bb.0: 
# | next:2002'0     ~~~~~~~~~~
# |           2072:  ld.param.b64 %rd3, [fminimum_acq_rel_double_global_cta_param_1]; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2073:  ld.param.b64 %rd2, [fminimum_acq_rel_double_global_cta_param_0]; 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2091:  // .globl fmaximum_acq_rel_double_global_cta // -- Begin function fmaximum_acq_rel_double_global_cta 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2092: .visible .func (.param .b64 func_retval0) fmaximum_acq_rel_double_global_cta( 
# | next:2002'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2093:  .param .b64 fmaximum_acq_rel_double_global_cta_param_0, 
# |           2094:  .param .b64 fmaximum_acq_rel_double_global_cta_param_1 
# |           2095: ) // @fmaximum_acq_rel_double_global_cta 
# |           2096: { 
# | next:2036'0      X error: no match found
# |           2097:  .reg .pred %p<3>; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~
# | next:2036'1      ?                  possible intended match
# |           2098:  .reg .b64 %rd<7>; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~
# |           2099:  
# | next:2036'0     ~
# |           2100: // %bb.0: 
# | next:2036'0     ~~~~~~~~~~
# |           2101:  ld.param.b64 %rd3, [fmaximum_acq_rel_double_global_cta_param_1]; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2102:  ld.param.b64 %rd2, [fmaximum_acq_rel_double_global_cta_param_0]; 
# | next:2036'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2272:  // .globl fminimum_acq_rel_half_global_cta // -- Begin function fminimum_acq_rel_half_global_cta 
# |           2273: .visible .func (.param .align 2 .b8 func_retval0[2]) fminimum_acq_rel_half_global_cta( 
# |           2274:  .param .b64 fminimum_acq_rel_half_global_cta_param_0, 
# |           2275:  .param .align 2 .b8 fminimum_acq_rel_half_global_cta_param_1[2] 
# |           2276: ) // @fminimum_acq_rel_half_global_cta 
# |           2277: { 
# | next:2218'0      X error: no match found
# |           2278:  .reg .pred %p<6>; 
# | next:2218'0     ~~~~~~~~~~~~~~~~~~~
# | next:2218'1      ?                  possible intended match
# |           2279:  .reg .b16 %rs<8>; 
# | next:2218'0     ~~~~~~~~~~~~~~~~~~~
# |           2280:  .reg .b32 %r<15>; 
# | next:2218'0     ~~~~~~~~~~~~~~~~~~~
# |           2281:  .reg .b64 %rd<3>; 
# | next:2218'0     ~~~~~~~~~~~~~~~~~~~
# |           2282:  
# | next:2218'0     ~
# |           2283: // %bb.0: 
# | next:2218'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2323:  // .globl fmaximum_acq_rel_half_global_cta // -- Begin function fmaximum_acq_rel_half_global_cta 
# | next:2218'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2324: .visible .func (.param .align 2 .b8 func_retval0[2]) fmaximum_acq_rel_half_global_cta( 
# | next:2218'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2325:  .param .b64 fmaximum_acq_rel_half_global_cta_param_0, 
# |           2326:  .param .align 2 .b8 fmaximum_acq_rel_half_global_cta_param_1[2] 
# |           2327: ) // @fmaximum_acq_rel_half_global_cta 
# |           2328: { 
# | next:2270'0      X error: no match found
# |           2329:  .reg .pred %p<6>; 
# | next:2270'0     ~~~~~~~~~~~~~~~~~~~
# | next:2270'1      ?                  possible intended match
# |           2330:  .reg .b16 %rs<8>; 
# | next:2270'0     ~~~~~~~~~~~~~~~~~~~
# |           2331:  .reg .b32 %r<15>; 
# | next:2270'0     ~~~~~~~~~~~~~~~~~~~
# |           2332:  .reg .b64 %rd<3>; 
# | next:2270'0     ~~~~~~~~~~~~~~~~~~~
# |           2333:  
# | next:2270'0     ~
# |           2334: // %bb.0: 
# | next:2270'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2578:  // .globl fminimum_acq_rel_bfloat_global_cta // -- Begin function fminimum_acq_rel_bfloat_global_cta 
# |           2579: .visible .func (.param .align 2 .b8 func_retval0[2]) fminimum_acq_rel_bfloat_global_cta( 
# |           2580:  .param .b64 fminimum_acq_rel_bfloat_global_cta_param_0, 
# |           2581:  .param .align 2 .b8 fminimum_acq_rel_bfloat_global_cta_param_1[2] 
# |           2582: ) // @fminimum_acq_rel_bfloat_global_cta 
# |           2583: { 
# | next:2522'0      X error: no match found
# |           2584:  .reg .pred %p<6>; 
# | next:2522'0     ~~~~~~~~~~~~~~~~~~~
# | next:2522'1      ?                  possible intended match
# |           2585:  .reg .b16 %rs<7>; 
# | next:2522'0     ~~~~~~~~~~~~~~~~~~~
# |           2586:  .reg .b32 %r<21>; 
# | next:2522'0     ~~~~~~~~~~~~~~~~~~~
# |           2587:  .reg .b64 %rd<3>; 
# | next:2522'0     ~~~~~~~~~~~~~~~~~~~
# |           2588:  
# | next:2522'0     ~
# |           2589: // %bb.0: 
# | next:2522'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2634:  // .globl fmaximum_acq_rel_bfloat_global_cta // -- Begin function fmaximum_acq_rel_bfloat_global_cta 
# | next:2522'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2635: .visible .func (.param .align 2 .b8 func_retval0[2]) fmaximum_acq_rel_bfloat_global_cta( 
# | next:2522'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2636:  .param .b64 fmaximum_acq_rel_bfloat_global_cta_param_0, 
# |           2637:  .param .align 2 .b8 fmaximum_acq_rel_bfloat_global_cta_param_1[2] 
# |           2638: ) // @fmaximum_acq_rel_bfloat_global_cta 
# |           2639: { 
# | next:2578'0      X error: no match found
# |           2640:  .reg .pred %p<6>; 
# | next:2578'0     ~~~~~~~~~~~~~~~~~~~
# | next:2578'1      ?                  possible intended match
# |           2641:  .reg .b16 %rs<7>; 
# | next:2578'0     ~~~~~~~~~~~~~~~~~~~
# |           2642:  .reg .b32 %r<21>; 
# | next:2578'0     ~~~~~~~~~~~~~~~~~~~
# |           2643:  .reg .b64 %rd<3>; 
# | next:2578'0     ~~~~~~~~~~~~~~~~~~~
# |           2644:  
# | next:2578'0     ~
# |           2645: // %bb.0: 
# | next:2578'0     ~~~~~~~~~~
# |              .
# |              .
# |              .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

LLVM.CodeGen/NVPTX/atomicrmw-sm90.ll
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll --check-prefix=SM90
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll --check-prefix=SM90
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll:1986:14: error: SM90-NEXT: expected string not found in input
# | ; SM90-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:2063:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2064:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll:2020:14: error: SM90-NEXT: expected string not found in input
# | ; SM90-NEXT: .reg .pred %p<6>;
# |              ^
# | <stdin>:2092:2: note: scanning from here
# | {
# |  ^
# | <stdin>:2093:2: note: possible intended match here
# |  .reg .pred %p<3>;
# |  ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |              .
# |              .
# |              .
# |           2058:  // .globl fminimum_acq_rel_double_global_cta // -- Begin function fminimum_acq_rel_double_global_cta 
# |           2059: .visible .func (.param .b64 func_retval0) fminimum_acq_rel_double_global_cta( 
# |           2060:  .param .b64 fminimum_acq_rel_double_global_cta_param_0, 
# |           2061:  .param .b64 fminimum_acq_rel_double_global_cta_param_1 
# |           2062: ) // @fminimum_acq_rel_double_global_cta 
# |           2063: { 
# | next:1986'0      X error: no match found
# |           2064:  .reg .pred %p<3>; 
# | next:1986'0     ~~~~~~~~~~~~~~~~~~~
# | next:1986'1      ?                  possible intended match
# |           2065:  .reg .b64 %rd<7>; 
# | next:1986'0     ~~~~~~~~~~~~~~~~~~~
# |           2066:  
# | next:1986'0     ~
# |           2067: // %bb.0: 
# | next:1986'0     ~~~~~~~~~~
# |           2068:  ld.param.b64 %rd3, [fminimum_acq_rel_double_global_cta_param_1]; 
# | next:1986'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2069:  ld.param.b64 %rd2, [fminimum_acq_rel_double_global_cta_param_0]; 
# | next:1986'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |           2087:  // .globl fmaximum_acq_rel_double_global_cta // -- Begin function fmaximum_acq_rel_double_global_cta 
# | next:1986'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2088: .visible .func (.param .b64 func_retval0) fmaximum_acq_rel_double_global_cta( 
# | next:1986'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2089:  .param .b64 fmaximum_acq_rel_double_global_cta_param_0, 
# |           2090:  .param .b64 fmaximum_acq_rel_double_global_cta_param_1 
# |           2091: ) // @fmaximum_acq_rel_double_global_cta 
# |           2092: { 
# | next:2020'0      X error: no match found
# |           2093:  .reg .pred %p<3>; 
# | next:2020'0     ~~~~~~~~~~~~~~~~~~~
# | next:2020'1      ?                  possible intended match
# |           2094:  .reg .b64 %rd<7>; 
# | next:2020'0     ~~~~~~~~~~~~~~~~~~~
# |           2095:  
# | next:2020'0     ~
# |           2096: // %bb.0: 
# | next:2020'0     ~~~~~~~~~~
# |           2097:  ld.param.b64 %rd3, [fmaximum_acq_rel_double_global_cta_param_1]; 
# | next:2020'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           2098:  ld.param.b64 %rd2, [fmaximum_acq_rel_double_global_cta_param_0]; 
# | next:2020'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

LLVM.CodeGen/PowerPC/fminimum-fmaximum.ll
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mtriple=powerpc64le-unknown-linux-gnu -mcpu=pwr8 -mattr=-vsx < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll --check-prefix=NOVSX
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mtriple=powerpc64le-unknown-linux-gnu -mcpu=pwr8 -mattr=-vsx
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll --check-prefix=NOVSX
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll:242:15: error: NOVSX-NEXT: expected string not found in input
# | ; NOVSX-NEXT: addi 3, 1, -48
# |               ^
# | <stdin>:206:18: note: scanning from here
# |  vcmpeqfp 4, 3, 3
# |                  ^
# | <stdin>:211:5: note: possible intended match here
# |  vsel 3, 3, 2, 4
# |     ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll:312:15: error: NOVSX-NEXT: expected string not found in input
# | ; NOVSX-NEXT: addi 3, 1, -48
# |               ^
# | <stdin>:231:18: note: scanning from here
# |  vcmpeqfp 4, 3, 3
# |                  ^
# | <stdin>:236:5: note: possible intended match here
# |  vsel 3, 3, 2, 4
# |     ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/PowerPC/fminimum-fmaximum.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |             .
# |             .
# |             .
# |           201:  .type v4f32_minimum,@function 
# |           202: v4f32_minimum: # @v4f32_minimum 
# |           203: .Lfunc_begin4: 
# |           204:  .cfi_startproc 
# |           205: # %bb.0: # %entry 
# |           206:  vcmpeqfp 4, 3, 3 
# | next:242'0                      X error: no match found
# |           207:  vnot 4, 4 
# | next:242'0     ~~~~~~~~~~~
# |           208:  vsel 2, 2, 3, 4 
# | next:242'0     ~~~~~~~~~~~~~~~~~
# |           209:  vcmpgefp 4, 2, 3 
# | next:242'0     ~~~~~~~~~~~~~~~~~~
# |           210:  vnot 4, 4 
# | next:242'0     ~~~~~~~~~~~
# |           211:  vsel 3, 3, 2, 4 
# | next:242'0     ~~~~~~~~~~~~~~~~~
# | next:242'1         ?             possible intended match
# |           212:  vxor 4, 4, 4 
# | next:242'0     ~~~~~~~~~~~~~~
# |           213:  vcmpgtsw 5, 4, 2 
# | next:242'0     ~~~~~~~~~~~~~~~~~~
# |           214:  vcmpeqfp 4, 3, 4 
# | next:242'0     ~~~~~~~~~~~~~~~~~~
# |           215:  vsel 2, 3, 2, 5 
# | next:242'0     ~~~~~~~~~~~~~~~~~
# |           216:  vsel 2, 3, 2, 4 
# | next:242'0     ~~~~~~~~~~~~~~~~~
# |             .
# |             .
# |             .
# |           226:  .type v4f32_maximum,@function 
# | next:242'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |           227: v4f32_maximum: # @v4f32_maximum 
# | next:242'0     ~~~~~~~~~~~~~~
# |           228: .Lfunc_begin5: 
# |           229:  .cfi_startproc 
# |           230: # %bb.0: # %entry 
# |           231:  vcmpeqfp 4, 3, 3 
# | next:312'0                      X error: no match found
# |           232:  vnot 4, 4 
# | next:312'0     ~~~~~~~~~~~
# |           233:  vsel 2, 2, 3, 4 
# | next:312'0     ~~~~~~~~~~~~~~~~~
# |           234:  vcmpgefp 4, 3, 2 
# | next:312'0     ~~~~~~~~~~~~~~~~~~
# |           235:  vnot 4, 4 
# | next:312'0     ~~~~~~~~~~~
# |           236:  vsel 3, 3, 2, 4 
# | next:312'0     ~~~~~~~~~~~~~~~~~
# | next:312'1         ?             possible intended match
# |           237:  vxor 4, 4, 4 
# | next:312'0     ~~~~~~~~~~~~~~
# |           238:  vcmpgtsw 5, 4, 2 
# | next:312'0     ~~~~~~~~~~~~~~~~~~
# |           239:  vcmpeqfp 4, 3, 4 
# | next:312'0     ~~~~~~~~~~~~~~~~~~
# |           240:  vnot 5, 5 
# | next:312'0     ~~~~~~~~~~~
# |           241:  vsel 2, 3, 2, 5 
# | next:312'0     ~~~~~~~~~~~~~~~~~
# |             .
# |             .
# |             .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

LLVM.CodeGen/X86/fminimum-fmaximum.ll
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 2
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll -mtriple=x86_64-unknown-unknown -mattr=+sse2     | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll --check-prefixes=SSE2
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mtriple=x86_64-unknown-unknown -mattr=+sse2
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll --check-prefixes=SSE2
# note: command had no output on stdout or stderr
# RUN: at line 3
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll -mtriple=x86_64-unknown-unknown -mattr=+avx      | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll --check-prefixes=AVX,AVX1
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mtriple=x86_64-unknown-unknown -mattr=+avx
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll --check-prefixes=AVX,AVX1
# note: command had no output on stdout or stderr
# RUN: at line 4
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc < /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll -mtriple=x86_64-unknown-unknown -mattr=+avx512f  | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll --check-prefixes=AVX,AVX512,AVX512F
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/llc -mtriple=x86_64-unknown-unknown -mattr=+avx512f
# note: command had no output on stdout or stderr
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll --check-prefixes=AVX,AVX512,AVX512F
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll:1789:17: error: AVX512F-NEXT: expected string not found in input
# | ; AVX512F-NEXT: vmovdqa %xmm1, %xmm2
# |                 ^
# | <stdin>:583:9: note: scanning from here
# | # %bb.0:
# |         ^
# | <stdin>:597:2: note: possible intended match here
# |  vmovshdup %xmm1, %xmm4 # xmm4 = xmm1[1,1,3,3]
# |  ^
# | 
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/llvm/test/CodeGen/X86/fminimum-fmaximum.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |              .
# |              .
# |              .
# |            578:  # -- End function 
# |            579:  .globl test_fmaximum_v4f16 # -- Begin function test_fmaximum_v4f16 
# |            580:  .p2align 4 
# |            581:  .type test_fmaximum_v4f16,@function 
# |            582: test_fmaximum_v4f16: # @test_fmaximum_v4f16 
# |            583: # %bb.0: 
# | next:1789'0             X error: no match found
# |            584:  vpsrld $16, %xmm1, %xmm2 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            585:  vcvtph2ps %xmm2, %xmm2 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~~
# |            586:  xorl %eax, %eax 
# | next:1789'0     ~~~~~~~~~~~~~~~~~
# |            587:  vucomiss %xmm2, %xmm2 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~
# |            588:  movl $65535, %ecx # imm = 0xFFFF 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |            592:  vucomiss %xmm3, %xmm3 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~
# |            593:  movl $0, %esi 
# | next:1789'0     ~~~~~~~~~~~~~~~
# |            594:  cmovpl %ecx, %esi 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~
# |            595:  vmovd %esi, %xmm4 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~
# |            596:  vpinsrw $1, %edx, %xmm4, %xmm5 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            597:  vmovshdup %xmm1, %xmm4 # xmm4 = xmm1[1,1,3,3] 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | next:1789'1      ?                                              possible intended match
# |            598:  vcvtph2ps %xmm4, %xmm4 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~~
# |            599:  vucomiss %xmm4, %xmm4 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~
# |            600:  movl $0, %edx 
# | next:1789'0     ~~~~~~~~~~~~~~~
# |            601:  cmovpl %ecx, %edx 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~
# |            602:  vpinsrw $2, %edx, %xmm5, %xmm6 
# | next:1789'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

If these failures are unrelated to your changes (for example tests are broken or flaky at HEAD), please open an issue at https://github.com/llvm/llvm-project/issues and add the infrastructure label.

@arsenm arsenm added the floating-point Floating-point math label Dec 4, 2025
Copy link
Copy Markdown
Contributor

@arsenm arsenm left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Globalisel presumably needs the same change

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@arsenm MinMaxMustRespectOrderedZero is not set here.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@arsenm MinMaxMustRespectOrderedZero is used to skip it if we have minmax OPCs.

@wzssyqa wzssyqa changed the title expandFMINIMUM_FMAXIMUM: FMAXNUM/FMINNUM treat +0>-0 SelectionDAG: Improve expandFMINIMUM_FMAXIMUM Dec 9, 2025
@wzssyqa wzssyqa requested review from nikic and phoebewang December 17, 2025 01:17
@wzssyqa wzssyqa requested a review from phoebewang December 17, 2025 04:18
@wzssyqa wzssyqa requested a review from nikic December 23, 2025 12:10
EVT IntVT = VT.changeTypeToInteger();
SDValue NTrunc = N;
if (!TLI.isTypeLegal(IntVT)) {
EVT FloatVT = VT.changeElementType(MVT::f32);
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hardcoding this to i32/f32 doesn't feel right. Can you find the next smallest FP legal FP type?

@wzssyqa
Copy link
Copy Markdown
Contributor Author

wzssyqa commented Feb 9, 2026

This PR introduces performance regression on some architectures. So Let's split into some small patches.

@wzssyqa wzssyqa marked this pull request as draft February 9, 2026 08:52
@wzssyqa
Copy link
Copy Markdown
Contributor Author

wzssyqa commented Feb 9, 2026

See: #180487

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants