[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
================
@@ -6,13 +7,15 @@
; Use bar.sync to arrive at a pre-computed barrier number and
; wait for all threads in CTA to also arrive:
define ptx_device void @test_barrier_named_cta() {
-; CHECK: mov.b32 %r[[REG0:[0-9]+]], 0;
-; CHECK: bar.sync %r[[REG0]];
-; CHECK: mov.b32 %r[[REG1:[0-9]+]], 10;
-; CHECK: bar.sync %r[[REG1]];
-; CHECK: mov.b32 %r[[REG2:[0-9]+]], 15;
-; CHECK: bar.sync %r[[REG2]];
-; CHECK: ret;
+; CHECK-LABEL: test_barrier_named_cta(
+; CHECK: {
+; CHECK-EMPTY:
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: bar.sync 0;
+; CHECK-NEXT: bar.sync 10;
+; CHECK-NEXT: bar.sync 15;
----------------
Artem-B wrote:
I'd add a few cases that pass ID in a register.
https://github.com/llvm/llvm-project/pull/140615
More information about the llvm-commits
mailing list