def LLVM_i8Ptr_global : LLVM_IntPtrBase<8, 1>;
def LLVM_i8Ptr_shared : LLVM_IntPtrBase<8, 3>;
+def LLVM_i64ptr_any : LLVM_IntPtrBase<64>;
+def LLVM_i64ptr_shared : LLVM_IntPtrBase<64, 3>;
//===----------------------------------------------------------------------===//
// NVVM dialect definitions
}
//===----------------------------------------------------------------------===//
+// NVVM Split arrive/wait barrier
+//===----------------------------------------------------------------------===//
+
+/// mbarrier.init instruction with generic pointer type
+def NVVM_MBarrierInitOp : NVVM_Op<"mbarrier.init">,
+ Arguments<(ins LLVM_i64ptr_any:$addr, I32:$count)> {
+ string llvmBuilder = [{
+ createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_init, {$addr, $count});
+ }];
+ let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands)";
+}
+
+/// mbarrier.init instruction with shared pointer type
+def NVVM_MBarrierInitSharedOp : NVVM_Op<"mbarrier.init.shared">,
+ Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$count)> {
+ string llvmBuilder = [{
+ createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_init_shared, {$addr, $count});
+ }];
+ let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands)";
+}
+
+//===----------------------------------------------------------------------===//
// NVVM synchronization op definitions
//===----------------------------------------------------------------------===//
// expected-error@below {{attribute attached to unexpected op}}
func.func private @expected_llvm_func() attributes { nvvm.kernel }
+
+// -----
+llvm.func private @mbarrier_init_generic(%barrier: !llvm.ptr) {
+ %count = nvvm.read.ptx.sreg.ntid.x : i32
+ // CHECK: nvvm.mbarrier.init %{{.*}}, %{{.*}} : !llvm.ptr, i32
+ nvvm.mbarrier.init %barrier, %count : !llvm.ptr, i32
+ llvm.return
+}
+
+
+llvm.func private @mbarrier_init_shared(%barrier: !llvm.ptr<3>) {
+ %count = nvvm.read.ptx.sreg.ntid.x : i32
+ // CHECK: nvvm.mbarrier.init.shared %{{.*}}, %{{.*}} : !llvm.ptr<3>, i32
+ nvvm.mbarrier.init.shared %barrier, %count : !llvm.ptr<3>, i32
+ llvm.return
+}