[llvm] [AMDGPU] Add initial version of gfx950 MFMA co-exec rules (PR #214682)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 01:35:46 PDT 2026
github-actions[bot] wrote:
<!--PREMERGE ADVISOR COMMENT: Windows-->
# :window: Windows x64 Test Results
* 140339 tests passed
* 3707 tests skipped
* 3 tests failed
## Failed Tests
(click on a test name to see its output)
### LLVM
<details>
<summary>LLVM.CodeGen/AMDGPU/coexec-mfma-interleave-gfx950.ll</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 2
c:\_work\llvm-project\llvm-project\build\bin\llc.exe -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-sched-strategy=coexec C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-mfma-interleave-gfx950.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-mfma-interleave-gfx950.ll
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\llc.exe' -mtriple=amdgcn-amd-amdhsa -mcpu=gfx950 -amdgpu-sched-strategy=coexec 'C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-mfma-interleave-gfx950.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-mfma-interleave-gfx950.ll'
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-mfma-interleave-gfx950.ll:26:15: error: CHECK-NEXT: is not on the line after the previous match
# | ; CHECK-NEXT: global_load_dwordx4 v[14:17], v[18:19], off
# | ^
# | <stdin>:23:2: note: 'next' match was here
# | global_load_dwordx4 v[14:17], v[18:19], off
# | ^
# | <stdin>:21:48: note: previous match ended here
# | v_lshl_add_u64 v[18:19], v[16:17], 0, v[14:15]
# | ^
# | <stdin>:22:1: note: non-matching line after previous match is here
# | global_load_dwordx2 v[30:31], v[18:19], off offset:16
# | ^
# |
# | Input file: <stdin>
# | Check file: C:\_work\llvm-project\llvm-project\llvm\test\CodeGen\AMDGPU\coexec-mfma-interleave-gfx950.ll
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 16: s_waitcnt lgkmcnt(0)
# | 17: v_lshl_add_u64 v[16:17], s[0:1], 0, v[0:1]
# | 18: global_load_dwordx4 v[2:5], v0, s[0:1]
# | 19: global_load_dwordx4 v[6:9], v0, s[0:1] offset:16
# | 20: global_load_dwordx4 v[10:13], v0, s[0:1] offset:32
# | 21: v_lshl_add_u64 v[18:19], v[16:17], 0, v[14:15]
# | next:26'0 { search range start (exclusive)
# | 22: global_load_dwordx2 v[30:31], v[18:19], off offset:16
# | 23: global_load_dwordx4 v[14:17], v[18:19], off
# | next:26'1 !~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: match on wrong line
# | 24: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
# | 25: ; sched_barrier mask(0x00000000)
# | 26: v_mfma_f32_16x16x32_f16 v[18:21], v[2:5], v[6:9], v[2:5]
# | 27: s_add_i32 s0, s8, s9
# | 28: v_add_u32_e32 v22, v14, v15
# | .
# | .
# | .
# | 252: amdhsa.version:
# | 253: - 1
# | 254: - 2
# | 255: ...
# | 256:
# | 257: .end_amdgpu_metadata
# | next:26'2 } search range end (exclusive)
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1
--
```
</details>
### Clang
<details>
<summary>Clang.CodeGen/AArch64/abi-classify-arg-types.c</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 1
c:\_work\llvm-project\llvm-project\build\bin\clang.exe -cc1 -internal-isystem C:\_work\llvm-project\llvm-project\build\lib\clang\24\include -nostdsysteminc -triple arm64-apple-ios7.0 -target-abi darwinpcs -fenable-matrix -fexperimental-max-bitint-width=1024 -emit-llvm -o - C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.c | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.c --check-prefixes=CHECK,DARWIN,LONG64
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\clang.exe' -cc1 -internal-isystem 'C:\_work\llvm-project\llvm-project\build\lib\clang\24\include' -nostdsysteminc -triple arm64-apple-ios7.0 -target-abi darwinpcs -fenable-matrix -fexperimental-max-bitint-width=1024 -emit-llvm -o - 'C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.c'
# 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\clang\test\CodeGen\AArch64\abi-classify-arg-types.c' --check-prefixes=CHECK,DARWIN,LONG64
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.c:112:11: error: CHECK: expected string not found in input
# | // CHECK: define{{.*}} void @arg_bitint129(ptr nofree noundef align 16 dead_on_return dereferenceable(32) %{{.*}})
# | ^
# | <stdin>:186:44: note: scanning from here
# | define void @arg_bitint128(i128 noundef %x) #0 {
# | ^
# | <stdin>:194:1: note: possible intended match here
# | define void @arg_bitint129(ptr nofreeobj noundef align 16 dead_on_return dereferenceable(32) %0) #0 {
# | ^
# |
# | Input file: <stdin>
# | Check file: C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.c
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 181: store i128 %storedv, ptr %x.addr, align 16
# | 182: ret void
# | 183: }
# | 184:
# | 185: ; Function Attrs: noinline nounwind optnone
# | 186: define void @arg_bitint128(i128 noundef %x) #0 {
# | check:112'0 { search range start (exclusive)
# | check:112'1 error: no match found in search range
# | 187: entry:
# | 188: %x.addr = alloca i128, align 16
# | 189: store i128 %x, ptr %x.addr, align 16
# | 190: ret void
# | 191: }
# | 192:
# | 193: ; Function Attrs: noinline nounwind optnone
# | 194: define void @arg_bitint129(ptr nofreeobj noundef align 16 dead_on_return dereferenceable(32) %0) #0 {
# | check:112'2 ? possible intended match
# | 195: entry:
# | 196: %x.addr = alloca i256, align 16
# | 197: %1 = load i256, ptr %0, align 16
# | 198: %x = trunc i256 %1 to i129
# | 199: %storedv = sext i129 %x to i256
# | 200: store i256 %storedv, ptr %x.addr, align 16
# | 201: ret void
# | 202: }
# | 203:
# | 204: attributes #0 = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
# | 205:
# | 206: !llvm.ident = !{!0}
# | 207:
# | 208: !0 = !{!"clang version 24.0.0git (https://github.com/llvm/llvm-project 6eeb935a0747cb32235ce9a4f24085ebf3f707ab)"}
# | check:112'3 } search range end (exclusive)
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1
--
```
</details>
<details>
<summary>Clang.CodeGen/AArch64/abi-classify-arg-types.cpp</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 1
c:\_work\llvm-project\llvm-project\build\bin\clang.exe -cc1 -internal-isystem C:\_work\llvm-project\llvm-project\build\lib\clang\24\include -nostdsysteminc -triple arm64-apple-ios7.0 -target-abi darwinpcs -emit-llvm -o - C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\clang.exe' -cc1 -internal-isystem 'C:\_work\llvm-project\llvm-project\build\lib\clang\24\include' -nostdsysteminc -triple arm64-apple-ios7.0 -target-abi darwinpcs -emit-llvm -o - 'C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp'
# 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\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp'
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp:39:11: error: CHECK: expected string not found in input
# | // CHECK: define{{.*}} void @arg_nontrivial_copy(ptr nofree noundef align 4 dead_on_return dereferenceable(4) %{{.*}})
# | ^
# | <stdin>:1:1: note: scanning from here
# | ; ModuleID = 'C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp'
# | ^
# | <stdin>:7:1: note: possible intended match here
# | define void @arg_nontrivial_copy(ptr nofreeobj noundef align 4 dead_on_return dereferenceable(4) %a) #0 {
# | ^
# |
# | Input file: <stdin>
# | Check file: C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | 1: ; ModuleID = 'C:\_work\llvm-project\llvm-project\clang\test\CodeGen\AArch64\abi-classify-arg-types.cpp'
# | check:39'0 { search range start (exclusive)
# | check:39'1 error: no match found in search range
# | 2: source_filename = "C:\\_work\\llvm-project\\llvm-project\\clang\\test\\CodeGen\\AArch64\\abi-classify-arg-types.cpp"
# | 3: target datalayout = "e-m:o-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-n32:64-S128-Fn32"
# | 4: target triple = "arm64-apple-ios7.0"
# | 5:
# | 6: ; Function Attrs: mustprogress noinline nounwind optnone
# | 7: define void @arg_nontrivial_copy(ptr nofreeobj noundef align 4 dead_on_return dereferenceable(4) %a) #0 {
# | check:39'2 ? possible intended match
# | 8: entry:
# | 9: %a.indirect_addr = alloca ptr, align 8
# | 10: store ptr %a, ptr %a.indirect_addr, align 8
# | 11: ret void
# | 12: }
# | .
# | .
# | .
# | 29:
# | 30: attributes #0 = { mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
# | 31:
# | 32: !llvm.ident = !{!0}
# | 33:
# | 34: !0 = !{!"clang version 24.0.0git (https://github.com/llvm/llvm-project 6eeb935a0747cb32235ce9a4f24085ebf3f707ab)"}
# | check:39'3 } search range end (exclusive)
# | >>>>>>
# `-----------------------------
# 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/214682
More information about the llvm-commits
mailing list