[Mlir-commits] [clang] [llvm] [mlir] [LLVM][NVPTX][MLIR] Add mbarrier layout support (PR #217252)

[llvmlistbot at llvm.org]

github-actions[bot] wrote:

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

* 140564 tests passed
* 3763 tests skipped
* 1 test failed

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

### MLIR
<details>
<summary>MLIR.Conversion/NVVMToLLVM/nvvm-to-llvm.mlir</summary>

```
Exit Code: 1

Command Output (stdout):
--
# RUN: at line 1
c:\_work\llvm-project\llvm-project\build\bin\mlir-opt.exe --convert-nvvm-to-llvm --convert-arith-to-llvm --split-input-file C:\_work\llvm-project\llvm-project\mlir\test\Conversion\NVVMToLLVM\nvvm-to-llvm.mlir | c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe C:\_work\llvm-project\llvm-project\mlir\test\Conversion\NVVMToLLVM\nvvm-to-llvm.mlir
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\mlir-opt.exe' --convert-nvvm-to-llvm --convert-arith-to-llvm --split-input-file 'C:\_work\llvm-project\llvm-project\mlir\test\Conversion\NVVMToLLVM\nvvm-to-llvm.mlir'
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\mlir\test\Conversion\NVVMToLLVM\nvvm-to-llvm.mlir:11:40: error: custom op 'nvvm.mbarrier.init' expected 'layout'
# |   nvvm.mbarrier.init %barrier, %count, predicate = %pred : !llvm.ptr<3>, i32, i1 
# |                                        ^
# `-----------------------------
# error: command failed with exit status: 1
# executed command: 'c:\_work\llvm-project\llvm-project\build\bin\filecheck.exe' 'C:\_work\llvm-project\llvm-project\mlir\test\Conversion\NVVMToLLVM\nvvm-to-llvm.mlir'
# .---command stderr------------
# | C:\_work\llvm-project\llvm-project\mlir\test\Conversion\NVVMToLLVM\nvvm-to-llvm.mlir:19:17: error: CHECK-LABEL: expected string not found in input
# | // CHECK-LABEL: @init_mbarrier_arrive_expect_tx
# |                 ^
# | <stdin>:391:26: note: scanning from here
# |  llvm.func @init_mbarrier(%arg0: !llvm.ptr, %arg1: !llvm.ptr<3>, %arg2: i32, %arg3: i1) {
# |                          ^
# | <stdin>:400:12: note: possible intended match here
# |  llvm.func @init_mbarrier_memory_clobber(%arg0: !llvm.ptr, %arg1: i32, %arg2: i1) {
# |            ^
# | 
# | Input file: <stdin>
# | Check file: C:\_work\llvm-project\llvm-project\mlir\test\Conversion\NVVMToLLVM\nvvm-to-llvm.mlir
# | 
# | -dump-input=help explains the following input dump.
# | 
# | Input was:
# | <<<<<<
# |             .
# |             .
# |             .
# |           386:  } 
# |           387: } 
# |           388:  
# |           389: // ----- 
# |           390: module { 
# |           391:  llvm.func @init_mbarrier(%arg0: !llvm.ptr, %arg1: !llvm.ptr<3>, %arg2: i32, %arg3: i1) { 
# | label:19'0                             {                                                                   search range start (exclusive)
# | label:19'1                                                                                                 error: no match found in search range
# |           392:  llvm.inline_asm has_side_effects asm_dialect = att "mbarrier.init.b64 [$0], $1;", "l,r" %arg0, %arg2 : (!llvm.ptr, i32) -> () 
# |           393:  llvm.inline_asm has_side_effects asm_dialect = att "@$2 mbarrier.init.b64 [$0], $1;", "l,r,b" %arg0, %arg2, %arg3 : (!llvm.ptr, i32, i1) -> () 
# |           394:  llvm.return 
# |           395:  } 
# |           396: } 
# |           397:  
# |           398: // ----- 
# |           399: module { 
# |           400:  llvm.func @init_mbarrier_memory_clobber(%arg0: !llvm.ptr, %arg1: i32, %arg2: i1) { 
# | label:19'2                ?                                                                          possible intended match
# |           401:  llvm.inline_asm has_side_effects asm_dialect = att "mbarrier.init.b64 [$0], $1;", "l,r,~{memory}" %arg0, %arg1 : (!llvm.ptr, i32) -> () 
# |           402:  llvm.inline_asm has_side_effects asm_dialect = att "@$2 mbarrier.init.b64 [$0], $1;", "l,r,b,~{memory}" %arg0, %arg1, %arg2 : (!llvm.ptr, i32, i1) -> () 
# |           403:  llvm.return 
# |           404:  } 
# |           405: } 
# |             .
# |             .
# |             .
# |           520:  llvm.func @inline_ptx_special_register_trailing_digit() -> i32 { 
# |           521:  %0 = llvm.inline_asm has_side_effects asm_dialect = att "mov.u32 $0, %pm0;", "=r" : () -> i32 
# |           522:  llvm.return %0 : i32 
# |           523:  } 
# |           524: } 
# |           525:  
# | label:19'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/217252