[clang] [llvm] [mlir] [NVPTX] Unify and extend barrier{.cta} intrinsic support (PR #140615)
Alex MacLean via cfe-commits
cfe-commits at lists.llvm.org
Mon May 19 16:07:49 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;
----------------
AlexMaclean wrote:
I just went ahead and removed this test completely. Both register and immediate cases are well tested in barriers.ll and we also test that these legacy intrinsics are auto-upgraded correctly, so keeping this test around at all seems confusing and pointless.
https://github.com/llvm/llvm-project/pull/140615
More information about the cfe-commits
mailing list