[llvm] [NVPTX] Unify and extend barrier{.cta} intrinsic support (PR #140615)

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Mon May 19 14:53:08 PDT 2025


================
@@ -22,16 +25,15 @@ define ptx_device void @test_barrier_named_cta() {
 ; Use bar.sync to arrive at a pre-computed barrier number and
 ; wait for fixed number of cooperating threads to arrive:
 define ptx_device void @test_barrier_named() {
-; CHECK: mov.b32  %r[[REG0A:[0-9]+]], 32;
-; CHECK: mov.b32  %r[[REG0B:[0-9]+]], 0;
-; CHECK: bar.sync %r[[REG0B]], %r[[REG0A]];
-; CHECK: mov.b32  %r[[REG1A:[0-9]+]], 352;
-; CHECK: mov.b32  %r[[REG1B:[0-9]+]], 10;
-; CHECK: bar.sync %r[[REG1B]], %r[[REG1A]];
-; CHECK: mov.b32  %r[[REG2A:[0-9]+]], 992;
-; CHECK: mov.b32  %r[[REG2B:[0-9]+]], 15;
-; CHECK: bar.sync %r[[REG2B]], %r[[REG2A]];
-; CHECK: ret;
+; CHECK-LABEL: test_barrier_named(
+; CHECK:       {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    bar.sync 0, 32;
+; CHECK-NEXT:    bar.sync 10, 352;
+; CHECK-NEXT:    bar.sync 15, 992;
----------------
Artem-B wrote:

Ditto.

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


More information about the llvm-commits mailing list