[llvm] [AMDGPU] Add to RewriteMFMAForm coexec scheduler pipeline (PR #199050)

via llvm-commits llvm-commits at lists.llvm.org
Thu May 21 09:00:00 PDT 2026


github-actions[bot] wrote:

<!--PREMERGE ADVISOR COMMENT: Windows-->
# :window: Windows x64 Test Results

* 134960 tests passed
* 3305 tests skipped
* 2 tests failed

## Failed Tests
(click on a test name to see its output)

### LLVM
<details>
<summary>LLVM.CodeGen/AMDGPU/coexec-rewrite-mfma.ll</summary>

```
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 1
c:\_work\llvm-project\llvm-project\build\bin\llc.exe -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-sched-strategy=coexec --enable-post-misched=0 C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-rewrite-mfma.ll -o - 2>&1 | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-rewrite-mfma.ll
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\llc.exe' -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-sched-strategy=coexec --enable-post-misched=0 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-rewrite-mfma.ll' -o -
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-rewrite-mfma.ll'
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-rewrite-mfma.ll:13:19: error: CHECK-COUNT: expected string not found in input (1 out of 28)
# | ; CHECK-COUNT-28: v_mfma_f32_16x16x32_f16 a
# |                   ^
# | <stdin>:9:19: note: scanning from here
# | v5_local_prefetch: ; @v5_local_prefetch
# |                   ^
# | <stdin>:65:2: note: possible intended match here
# |  v_mfma_f32_16x16x32_f16 v[46:49], v[60:63], v[60:63], v[46:49]
# |  ^
# | 
# | Input file: <stdin>
# | Check file: C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-rewrite-mfma.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |             1: warning: <unknown>:0:0: in function v5_local_prefetch void (<8 x half>): 'amdgpu-sched-strategy'='coexec' is only supported for gfx1250 
# |             2:  
# |             3:  .amdgcn_target "amdgcn-amd-amdhsa--gfx950" 
# |             4:  .amdhsa_code_object_version 6 
# |             5:  .text 
# |             6:  .globl v5_local_prefetch ; -- Begin function v5_local_prefetch 
# |             7:  .p2align 8 
# |             8:  .type v5_local_prefetch, at function 
# |             9: v5_local_prefetch: ; @v5_local_prefetch 
# | count:13'0                       X~~~~~~~~~~~~~~~~~~~~~ error: no match found
# |            10: ; %bb.0: ; %.lr.ph 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~
# |            11:  s_load_dwordx4 s[0:3], s[4:5], 0x0 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            12:  v_mov_b32_e32 v2, 0 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~
# |            13:  v_mov_b32_e32 v9, v2 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~
# |            14:  v_mov_b32_e32 v47, v2 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~
# |             .
# |             .
# |             .
# |            60:  s_waitcnt lgkmcnt(0) 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~
# |            61:  v_mov_b32_e32 v0, v64 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~
# |            62:  v_mov_b32_e32 v1, v65 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~
# |            63:  .p2align 5, , 4 
# | count:13'0     ~~~~~~~~~~~~~~~~~
# |            64: .LBB0_1: ; =>This Inner Loop Header: Depth=1 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            65:  v_mfma_f32_16x16x32_f16 v[46:49], v[60:63], v[60:63], v[46:49] 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | count:13'1      ?                                                               possible intended match
# |            66:  s_nop 4 
# | count:13'0     ~~~~~~~~~
# |            67:  v_mov_b32_e32 v4, v12 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~
# |            68:  v_mov_b32_e32 v3, v2 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~
# |            69:  v_mov_b32_e32 v6, v2 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~
# |            70:  v_mfma_f32_16x16x32_f16 v[22:25], v[60:63], v[60:63], v[22:25] 
# | count:13'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |             .
# |             .
# |             .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

```
</details>
<details>
<summary>LLVM.CodeGen/PowerPC/fast-isel-cmp-imm.ll</summary>

```
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 1
c:\_work\llvm-project\llvm-project\build\bin\llc.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll -O0 -verify-machineinstrs -fast-isel-abort=1 -mtriple=powerpc64-unknown-linux-gnu -mcpu=pwr7 -mattr=-vsx | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll --check-prefix=ELF64
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\llc.exe' -O0 -verify-machineinstrs -fast-isel-abort=1 -mtriple=powerpc64-unknown-linux-gnu -mcpu=pwr7 -mattr=-vsx
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll' --check-prefix=ELF64
# note: command had no output on stdout or stderr
# RUN: at line 2
c:\_work\llvm-project\llvm-project\build\bin\llc.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll -O0 -verify-machineinstrs -fast-isel-abort=1 -mtriple=powerpc64le-unknown-linux-gnu -mattr=+vsx | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll --check-prefix=VSX
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\llc.exe' -O0 -verify-machineinstrs -fast-isel-abort=1 -mtriple=powerpc64le-unknown-linux-gnu -mattr=+vsx
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll' --check-prefix=VSX
# note: command had no output on stdout or stderr
# RUN: at line 3
c:\_work\llvm-project\llvm-project\build\bin\llc.exe < C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll -O0 -verify-machineinstrs -fast-isel-abort=1 -mtriple=powerpc-unknown-linux-gnu -mcpu=e500 -mattr=spe | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll --check-prefix=SPE
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\llc.exe' -O0 -verify-machineinstrs -fast-isel-abort=1 -mtriple=powerpc-unknown-linux-gnu -mcpu=e500 -mattr=spe
# note: command had no output on stdout or stderr
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll' --check-prefix=SPE
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll:19:8: error: SPE: expected string not found in input
# | ; SPE: efscmpeq
# |        ^
# | <stdin>:11:12: note: scanning from here
# | t1a: # @t1a
# |            ^
# | <stdin>:12:2: note: possible intended match here
# | .Lfunc_begin0:
# |  ^
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll:42:8: error: SPE: expected string not found in input
# | ; SPE: efscmpeq
# |        ^
# | <stdin>:43:12: note: scanning from here
# | t1b: # @t1b
# |            ^
# | <stdin>:44:2: note: possible intended match here
# | .Lfunc_begin1:
# |  ^
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll:65:8: error: SPE: expected string not found in input
# | ; SPE: efscmpeq
# |        ^
# | <stdin>:75:12: note: scanning from here
# | t1c: # @t1c
# |            ^
# | <stdin>:76:2: note: possible intended match here
# | .Lfunc_begin2:
# |  ^
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll:88:8: error: SPE: expected string not found in input
# | ; SPE: efdcmpeq
# |        ^
# | <stdin>:108:12: note: scanning from here
# | t2a: # @t2a
# |            ^
# | <stdin>:109:2: note: possible intended match here
# | .Lfunc_begin3:
# |  ^
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll:111:8: error: SPE: expected string not found in input
# | ; SPE: efdcmpeq
# |        ^
# | <stdin>:149:12: note: scanning from here
# | t2b: # @t2b
# |            ^
# | <stdin>:150:2: note: possible intended match here
# | .Lfunc_begin4:
# |  ^
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll:134:8: error: SPE: expected string not found in input
# | ; SPE: efdcmpeq
# |        ^
# | <stdin>:190:12: note: scanning from here
# | t2c: # @t2c
# |            ^
# | <stdin>:191:2: note: possible intended match here
# | .Lfunc_begin5:
# |  ^
# | 
# | Input file: <stdin>
# | Check file: C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\PowerPC\fast-isel-cmp-imm.ll
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |              .
# |              .
# |              .
# |              6:  .text 
# |              7:  .globl t1a 
# |              8:  .p2align 2 
# |              9:  .prefalign 4, .Lfunc_end0, nop 
# |             10:  .type t1a, at function 
# |             11: t1a: # @t1a 
# | check:19'0                 X error: no match found
# |             12: .Lfunc_begin0: 
# | check:19'0      ~~~~~~~~~~~~~~~
# | check:19'1       ?              possible intended match
# |             13: # %bb.0: # %entry 
# | check:19'0      ~~~~~~~~~~~~~~~~~~
# |             14:  mflr 0 
# | check:19'0      ~~~~~~~~
# |             15:  stwu 1, -16(1) 
# | check:19'0      ~~~~~~~~~~~~~~~~
# |             16:  stw 0, 20(1) 
# | check:19'0      ~~~~~~~~~~~~~~
# |             17:  lis 4, .LCPI0_0 at ha 
# | check:19'0      ~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |             38:  .text 
# | check:19'0      ~~~~~~~
# |             39:  .globl t1b 
# | check:19'0      ~~~~~~~~~~~~
# |             40:  .p2align 2 
# | check:19'0      ~~~~~~~~~~~~
# |             41:  .prefalign 4, .Lfunc_end1, nop 
# | check:19'0      ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |             42:  .type t1b, at function 
# | check:19'0      ~~~~~~~~~~~~~~~~~~~~~
# |             43: t1b: # @t1b 
# | check:19'0      ~~~~~~~~~~~
# | check:42'0                 X error: no match found
# |             44: .Lfunc_begin1: 
# | check:42'0      ~~~~~~~~~~~~~~~
# | check:42'1       ?              possible intended match
# |             45: # %bb.0: # %entry 
# | check:42'0      ~~~~~~~~~~~~~~~~~~
# |             46:  mflr 0 
# | check:42'0      ~~~~~~~~
# |             47:  stwu 1, -16(1) 
# | check:42'0      ~~~~~~~~~~~~~~~~
# |             48:  stw 0, 20(1) 
# | check:42'0      ~~~~~~~~~~~~~~
# |             49:  lis 4, .LCPI1_0 at ha 
# | check:42'0      ~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |             70:  .text 
# | check:42'0      ~~~~~~~
# |             71:  .globl t1c 
# | check:42'0      ~~~~~~~~~~~~
# |             72:  .p2align 2 
# | check:42'0      ~~~~~~~~~~~~
# |             73:  .prefalign 4, .Lfunc_end2, nop 
# | check:42'0      ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |             74:  .type t1c, at function 
# | check:42'0      ~~~~~~~~~~~~~~~~~~~~~
# |             75: t1c: # @t1c 
# | check:42'0      ~~~~~~~~~~~
# | check:65'0                 X error: no match found
# |             76: .Lfunc_begin2: 
# | check:65'0      ~~~~~~~~~~~~~~~
# | check:65'1       ?              possible intended match
# |             77: # %bb.0: # %entry 
# | check:65'0      ~~~~~~~~~~~~~~~~~~
# |             78:  mflr 0 
# | check:65'0      ~~~~~~~~
# |             79:  stwu 1, -16(1) 
# | check:65'0      ~~~~~~~~~~~~~~~~
# |             80:  stw 0, 20(1) 
# | check:65'0      ~~~~~~~~~~~~~~
# |             81:  lis 4, .LCPI2_0 at ha 
# | check:65'0      ~~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |            103:  .text 
# | check:65'0      ~~~~~~~
# |            104:  .globl t2a 
# | check:65'0      ~~~~~~~~~~~~
# |            105:  .p2align 2 
# | check:65'0      ~~~~~~~~~~~~
# |            106:  .prefalign 4, .Lfunc_end3, nop 
# | check:65'0      ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            107:  .type t2a, at function 
# | check:65'0      ~~~~~~~~~~~~~~~~~~~~~
# |            108: t2a: # @t2a 
# | check:65'0      ~~~~~~~~~~~
# | check:88'0                 X error: no match found
# |            109: .Lfunc_begin3: 
# | check:88'0      ~~~~~~~~~~~~~~~
# | check:88'1       ?              possible intended match
# |            110: # %bb.0: # %entry 
# | check:88'0      ~~~~~~~~~~~~~~~~~~
# |            111:  mflr 0 
# | check:88'0      ~~~~~~~~
# |            112:  stwu 1, -16(1) 
# | check:88'0      ~~~~~~~~~~~~~~~~
# |            113:  stw 0, 20(1) 
# | check:88'0      ~~~~~~~~~~~~~~
# |            114:  evmergelo 3, 3, 4 
# | check:88'0      ~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |            144:  .text 
# | check:88'0      ~~~~~~~
# |            145:  .globl t2b 
# | check:88'0      ~~~~~~~~~~~~
# |            146:  .p2align 2 
# | check:88'0      ~~~~~~~~~~~~
# |            147:  .prefalign 4, .Lfunc_end4, nop 
# | check:88'0      ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            148:  .type t2b, at function 
# | check:88'0      ~~~~~~~~~~~~~~~~~~~~~
# |            149: t2b: # @t2b 
# | check:88'0      ~~~~~~~~~~~
# | check:111'0                X error: no match found
# |            150: .Lfunc_begin4: 
# | check:111'0     ~~~~~~~~~~~~~~~
# | check:111'1      ?              possible intended match
# |            151: # %bb.0: # %entry 
# | check:111'0     ~~~~~~~~~~~~~~~~~~
# |            152:  mflr 0 
# | check:111'0     ~~~~~~~~
# |            153:  stwu 1, -16(1) 
# | check:111'0     ~~~~~~~~~~~~~~~~
# |            154:  stw 0, 20(1) 
# | check:111'0     ~~~~~~~~~~~~~~
# |            155:  evmergelo 3, 3, 4 
# | check:111'0     ~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# |            185:  .text 
# | check:111'0     ~~~~~~~
# |            186:  .globl t2c 
# | check:111'0     ~~~~~~~~~~~~
# |            187:  .p2align 2 
# | check:111'0     ~~~~~~~~~~~~
# |            188:  .prefalign 4, .Lfunc_end5, nop 
# | check:111'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# |            189:  .type t2c, at function 
# | check:111'0     ~~~~~~~~~~~~~~~~~~~~~
# |            190: t2c: # @t2c 
# | check:111'0     ~~~~~~~~~~~
# | check:134'0                X error: no match found
# |            191: .Lfunc_begin5: 
# | check:134'0     ~~~~~~~~~~~~~~~
# | check:134'1      ?              possible intended match
# |            192: # %bb.0: # %entry 
# | check:134'0     ~~~~~~~~~~~~~~~~~~
# |            193:  mflr 0 
# | check:134'0     ~~~~~~~~
# |            194:  stwu 1, -16(1) 
# | check:134'0     ~~~~~~~~~~~~~~~~
# |            195:  stw 0, 20(1) 
# | check:134'0     ~~~~~~~~~~~~~~
# |            196:  evmergelo 3, 3, 4 
# | check:134'0     ~~~~~~~~~~~~~~~~~~~
# |              .
# |              .
# |              .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1

--

```
</details>

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.

https://github.com/llvm/llvm-project/pull/199050


More information about the llvm-commits mailing list