[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