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

Alex MacLean via llvm-commits llvm-commits at lists.llvm.org
Tue May 20 08:24:48 PDT 2025


================
@@ -462,24 +462,28 @@ def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">,
 // NVVM synchronization op definitions
 //===----------------------------------------------------------------------===//
 
-def NVVM_Barrier0Op : NVVM_IntrOp<"barrier0"> {
+def NVVM_Barrier0Op : NVVM_Op<"barrier0"> {
   let assemblyFormat = "attr-dict";
+  string llvmBuilder = [{
+      createIntrinsicCall(
+          builder, llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_all,
+          {builder.getInt32(0)});
+  }];
----------------
AlexMaclean wrote:

I agree this op can be completely removed but doing so is not trivial. I'd prefer to leave this for a subsequent change. CC @durga4github 

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


More information about the llvm-commits mailing list