[Mlir-commits] [mlir] [MLIR][NVVM] Add tcgen05.mma MLIR Ops (PR #164356)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Mon Nov 17 20:51:53 PST 2025
github-actions[bot] wrote:
<!--PREMERGE ADVISOR COMMENT: Linux-->
# :penguin: Linux x64 Test Results
* 7093 tests passed
* 594 tests skipped
* 1 test failed
## Failed Tests
(click on a test name to see its output)
### MLIR
<details>
<summary>MLIR.Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 1
/home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/mlir-translate --mlir-to-llvmir /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir
# executed command: /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/mlir-translate --mlir-to-llvmir /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir
# 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/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir:6:12: error: CHECK: expected string not found in input
# | // CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) {{%[0-9]+}}, ptr addrspace(6) {{%[0-9]+}}, i64 {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 {{%[0-9]+}}, i32 0, i32 1, i32 0)
# | ^
# | <stdin>:4:36: note: scanning from here
# | define void @nvvm_tcgen05_mma_cta_1(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %2, i64 %3, i32 %4, i1 %5) {
# | ^
# | <stdin>:13:2: note: possible intended match here
# | call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 0, i32 1, i32 0)
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir:108:12: error: CHECK: expected string not found in input
# | // CHECK: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) {{%[0-9]+}}, ptr addrspace(6) {{%[0-9]+}}, i64 {{%[0-9]+}}, i32 {{%[0-9]+}}, i1 {{%[0-9]+}}, i32 0, i32 2, i32 0)
# | ^
# | <stdin>:32:36: note: scanning from here
# | define void @nvvm_tcgen05_mma_cta_2(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %2, i64 %3, i32 %4, i1 %5) {
# | ^
# | <stdin>:41:2: note: possible intended match here
# | call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 0, i32 2, i32 0)
# | ^
# |
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/Target/LLVMIR/nvvm/tcgen05-mma-tensor.mlir
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | 1: ; ModuleID = 'LLVMDialectModule'
# | 2: source_filename = "LLVMDialectModule"
# | 3:
# | 4: define void @nvvm_tcgen05_mma_cta_1(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %2, i64 %3, i32 %4, i1 %5) {
# | check:6'0 X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
# | 5: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f16 */ i32 0, /* cta_group= */ i32 1, /* collector=discard */ i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 6: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=tf32 */ i32 1, /* cta_group= */ i32 1, /* collector=discard */ i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 7: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 1, /* collector=discard */ i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 8: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=i8 */ i32 3, /* cta_group= */ i32 1, /* collector=discard */ i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 9: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f16 */ i32 0, /* cta_group= */ i32 1, /* collector=lastuse */ i32 1)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 10: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=tf32 */ i32 1, /* cta_group= */ i32 1, /* collector=lastuse */ i32 1)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 11: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 1, /* collector=lastuse */ i32 1)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 12: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=i8 */ i32 3, /* cta_group= */ i32 1, /* collector=lastuse */ i32 1)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 13: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 0, i32 1, i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:6'1 ? possible intended match
# | 14: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 1, i32 1, i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 15: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 2, i32 1, i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 16: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 3, i32 1, i32 0)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 17: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 0, i32 1, i32 1)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 18: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 1, i32 1, i32 1)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | 27: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 1, /* collector=use */ i32 3)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 28: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=i8 */ i32 3, /* cta_group= */ i32 1, /* collector=use */ i32 3)
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 29: ret void
# | check:6'0 ~~~~~~~~~~
# | 30: }
# | check:6'0 ~~
# | 31:
# | check:6'0 ~
# | 32: define void @nvvm_tcgen05_mma_cta_2(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %2, i64 %3, i32 %4, i1 %5) {
# | check:6'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:108'0 X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
# | 33: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f16 */ i32 0, /* cta_group= */ i32 2, /* collector=discard */ i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 34: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=tf32 */ i32 1, /* cta_group= */ i32 2, /* collector=discard */ i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 35: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 2, /* collector=discard */ i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 36: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=i8 */ i32 3, /* cta_group= */ i32 2, /* collector=discard */ i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 37: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f16 */ i32 0, /* cta_group= */ i32 2, /* collector=lastuse */ i32 1)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 38: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=tf32 */ i32 1, /* cta_group= */ i32 2, /* collector=lastuse */ i32 1)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 39: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=f8f6f4 */ i32 2, /* cta_group= */ i32 2, /* collector=lastuse */ i32 1)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 40: call void @llvm.nvvm.tcgen05.mma.tensor(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, /* kind=i8 */ i32 3, /* cta_group= */ i32 2, /* collector=lastuse */ i32 1)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 41: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 0, i32 2, i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:108'1 ? possible intended match
# | 42: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 1, i32 2, i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 43: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 2, i32 2, i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 44: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 3, i32 2, i32 0)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 45: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 0, i32 2, i32 1)
# | check:108'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 46: call void @llvm.nvvm.tcgen05.mma.tensor.ashift(ptr addrspace(6) %0, ptr addrspace(6) %1, i64 %3, i32 %4, i1 %5, i32 1, i32 2, i32 1)
# | check:108'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/164356
More information about the Mlir-commits
mailing list