[Mlir-commits] [flang] [mlir] [MLIR][NVVM] Split nvvm.barrier into nvvm.barrier and nvvm.barrier.reduction (PR #199404)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Sun May 24 01:48:34 PDT 2026
github-actions[bot] wrote:
<!--PREMERGE ADVISOR COMMENT: Linux-->
# :penguin: Linux x64 Test Results
* 8082 tests passed
* 619 tests skipped
* 2 tests failed
## Failed Tests
(click on a test name to see its output)
### Flang
<details>
<summary>Flang.Lower/CUDA/cuda-device-proc.cuf</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 1
bbc -emit-hlfir -fcuda /home/gha/actions-runner/_work/llvm-project/llvm-project/flang/test/Lower/CUDA/cuda-device-proc.cuf -o - | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/flang/test/Lower/CUDA/cuda-device-proc.cuf
# executed command: bbc -emit-hlfir -fcuda /home/gha/actions-runner/_work/llvm-project/llvm-project/flang/test/Lower/CUDA/cuda-device-proc.cuf -o -
# 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/flang/test/Lower/CUDA/cuda-device-proc.cuf
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/flang/test/Lower/CUDA/cuda-device-proc.cuf:106:10: error: CHECK: expected string not found in input
# | ! CHECK: %{{.*}} = nvvm.barrier #nvvm.reduction<and> %c1{{.*}} -> i32
# | ^
# | <stdin>:51:34: note: scanning from here
# | nvvm.bar.warp.sync %c1_i32 : i32
# | ^
# | <stdin>:53:4: note: possible intended match here
# | %43 = nvvm.barrier.reduction #nvvm.reduction<and> %c1_i32_1 -> i32
# | ^
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/flang/test/Lower/CUDA/cuda-device-proc.cuf:217:10: error: CHECK: expected string not found in input
# | ! CHECK: nvvm.barrier #nvvm.reduction<and> %c1{{.*}} -> i32
# | ^
# | <stdin>:357:36: note: scanning from here
# | nvvm.bar.warp.sync %c1_i32_2 : i32
# | ^
# | <stdin>:359:20: note: possible intended match here
# | %39 = nvvm.barrier.reduction #nvvm.reduction<and> %c1_i32_3 -> i32
# | ^
# |
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/flang/test/Lower/CUDA/cuda-device-proc.cuf
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 46: %41 = hlfir.designate %31#0{"x"} : (!fir.ref<!fir.type<_QM__fortran_builtinsT__builtin_dim3{x:i32,y:i32,z:i32}>>) -> !fir.ref<i32>
# | 47: %42 = fir.load %41 : !fir.ref<i32>
# | 48: hlfir.assign %42 to %33#0 : i32, !fir.ref<i32>
# | 49: nvvm.barrier
# | 50: %c1_i32 = arith.constant 1 : i32
# | 51: nvvm.bar.warp.sync %c1_i32 : i32
# | check:106'0 X error: no match found
# | 52: %c1_i32_1 = arith.constant 1 : i32
# | check:106'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 53: %43 = nvvm.barrier.reduction #nvvm.reduction<and> %c1_i32_1 -> i32
# | check:106'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:106'1 ? possible intended match
# | 54: hlfir.assign %43 to %27#0 : i32, !fir.ref<i32>
# | check:106'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 55: %44 = fir.load %33#0 : !fir.ref<i32>
# | check:106'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 56: %45 = fir.load %16#0 : !fir.ref<i32>
# | check:106'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 57: %46 = arith.cmpi sgt, %44, %45 : i32
# | check:106'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 58: %47 = fir.convert %46 : (i1) -> i32
# | check:106'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | 352: %36 = fir.load %5#0 : !fir.ref<i32>
# | 353: %37 = fir.convert %36 : (i32) -> i64
# | 354: %38 = hlfir.designate %3#0 (%37) : (!fir.ref<!fir.array<32xi32>>, i64) -> !fir.ref<i32>
# | 355: hlfir.assign %35 to %38 : i32, !fir.ref<i32>
# | 356: %c1_i32_2 = arith.constant 1 : i32
# | 357: nvvm.bar.warp.sync %c1_i32_2 : i32
# | check:217'0 X error: no match found
# | 358: %c1_i32_3 = arith.constant 1 : i32
# | check:217'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 359: %39 = nvvm.barrier.reduction #nvvm.reduction<and> %c1_i32_3 -> i32
# | check:217'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | check:217'1 ? possible intended match
# | 360: hlfir.assign %39 to %9#0 : i32, !fir.ref<i32>
# | check:217'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 361: %c1_i32_4 = arith.constant 1 : i32
# | check:217'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 362: %40 = nvvm.barrier.reduction #nvvm.reduction<popc> %c1_i32_4 -> i32
# | check:217'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 363: hlfir.assign %40 to %9#0 : i32, !fir.ref<i32>
# | check:217'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 364: %c1_i32_5 = arith.constant 1 : i32
# | check:217'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | >>>>>>
# `-----------------------------
# error: command failed with exit status: 1
--
```
</details>
### MLIR
<details>
<summary>MLIR.python/dialects/nvvm.py</summary>
```
Exit Code: 1
Command Output (stdout):
--
# RUN: at line 1
/usr/bin/python3 /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py | /home/gha/actions-runner/_work/llvm-project/llvm-project/build/bin/FileCheck /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py
# executed command: /usr/bin/python3 /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py
# .---command stderr------------
# | Traceback (most recent call last):
# | File "/home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py", line 126, in <module>
# | @constructAndPrintInModule
# | ^^^^^^^^^^^^^^^^^^^^^^^^^
# | File "/home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py", line 17, in constructAndPrintInModule
# | f()
# | File "/home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py", line 131, in test_barriers
# | @func.FuncOp.from_py_func(i32, i32, f32)
# | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
# | File "/home/gha/actions-runner/_work/llvm-project/llvm-project/build/tools/mlir/python_packages/mlir_core/mlir/dialects/func.py", line 197, in decorator
# | return_values = f(*func_args, **func_kwargs)
# | ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
# | File "/home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py", line 142, in barriers
# | pred = nvvm.barrier_reduction(
# | ^^^^^^^^^^^^^^^^^^^^^^^
# | TypeError: barrier_reduction() got an unexpected keyword argument 'res'
# `-----------------------------
# error: command failed with exit status: 1
# 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/python/dialects/nvvm.py
# .---command stderr------------
# | /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py:161:16: error: CHECK-LABEL: expected string not found in input
# | # CHECK-LABEL: func.func @barriers(
# | ^
# | <stdin>:27:33: note: scanning from here
# | func.func @mbarrier_arrive_ops(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<7>, %arg2: i32) {
# | ^
# | <stdin>:41:1: note: possible intended match here
# | TEST: test_barriers
# | ^
# |
# | Input file: <stdin>
# | Check file: /home/gha/actions-runner/_work/llvm-project/llvm-project/mlir/test/python/dialects/nvvm.py
# |
# | -dump-input=help explains the following input dump.
# |
# | Input was:
# | <<<<<<
# | .
# | .
# | .
# | 22: }
# | 23:
# | 24:
# | 25: TEST: test_mbarrier_arrive
# | 26: module {
# | 27: func.func @mbarrier_arrive_ops(%arg0: !llvm.ptr<3>, %arg1: !llvm.ptr<7>, %arg2: i32) {
# | label:161'0 X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found
# | 28: %0 = nvvm.mbarrier.arrive %arg0 : !llvm.ptr<3> -> i64
# | label:161'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 29: nvvm.mbarrier.arrive %arg1 : !llvm.ptr<7>
# | label:161'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 30: %1 = nvvm.mbarrier.arrive_drop %arg0 : !llvm.ptr<3> -> i64
# | label:161'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 31: nvvm.mbarrier.arrive_drop %arg1 : !llvm.ptr<7>
# | label:161'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | 32: %2 = nvvm.mbarrier.arrive.expect_tx %arg0, %arg2 : !llvm.ptr<3>, i32 -> i64
# | label:161'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
# | .
# | .
# | .
# | 36: return
# | label:161'0 ~~~~~~~~
# | 37: }
# | label:161'0 ~~~
# | 38: }
# | label:161'0 ~~
# | 39:
# | label:161'0 ~
# | 40:
# | label:161'0 ~
# | 41: TEST: test_barriers
# | label:161'0 ~~~~~~~~~~~~~~~~~~~~
# | label:161'1 ? possible intended match
# | >>>>>>
# `-----------------------------
# 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/199404
More information about the Mlir-commits
mailing list