[clang] [llvm] [mlir] [NVPTX] Unify and extend barrier{.cta} intrinsic support (PR #140615)
Guray Ozen via cfe-commits
cfe-commits at lists.llvm.org
Tue May 20 02:47: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)});
+ }];
}
def NVVM_BarrierOp : NVVM_Op<"barrier", [AttrSizedOperandSegments]> {
let arguments = (ins
Optional<I32>:$barrierId,
Optional<I32>:$numberOfThreads);
string llvmBuilder = [{
- if ($numberOfThreads && $barrierId) {
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier,
- {$barrierId, $numberOfThreads});
- } else if($barrierId) {
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier_n,
- {$barrierId});
- } else {
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier0);
- }
+ auto id = $barrierId ? $barrierId : builder.getInt32(0);
----------------
grypp wrote:
We don't use auto when the type isn't obvious.
https://github.com/llvm/llvm-project/pull/140615
More information about the cfe-commits
mailing list