[mlir][nvvm] Introduce `mbarrier.inval`

Introduce support for PTX's `mbarrier.inval` .

Contiunation of D151334

Reviewed By: qcolombet

Differential Revision: https://reviews.llvm.org/D151338

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 118e7848515170029d31c57947d5f2cf6802de36..44aa4aa980ee2da594e1428341d417116697e437 100644 (file)
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -197,6 +197,22 @@ def NVVM_MBarrierInitSharedOp : NVVM_Op<"mbarrier.init.shared">,
   let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands)";
 }
 
+def NVVM_MBarrierInvalOp : NVVM_Op<"mbarrier.inval">,
+  Arguments<(ins LLVM_i64ptr_any:$addr)> {
+  string llvmBuilder = [{
+      createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_inval, {$addr});
+  }];
+  let assemblyFormat = "$addr attr-dict `:` type(operands)";
+}
+
+def NVVM_MBarrierInvalSharedOp : NVVM_Op<"mbarrier.inval.shared">,
+  Arguments<(ins LLVM_i64ptr_shared:$addr)> {
+  string llvmBuilder = [{
+      createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_inval_shared, {$addr});
+  }];
+  let assemblyFormat = "$addr attr-dict `:` type(operands)";
+}
+
 //===----------------------------------------------------------------------===//
 // NVVM synchronization op definitions
 //===----------------------------------------------------------------------===//

diff --git a/mlir/test/Dialect/LLVMIR/nvvm.mlir b/mlir/test/Dialect/LLVMIR/nvvm.mlir
index d08d02a04d8f81b9f299ebe9e1c639f7488b9e25..a32b33d410f59cd93c39e20c8ac15ab2f3c2b07a 100644 (file)
--- a/mlir/test/Dialect/LLVMIR/nvvm.mlir
+++ b/mlir/test/Dialect/LLVMIR/nvvm.mlir
@@ -353,3 +353,17 @@ llvm.func private @mbarrier_init_shared(%barrier: !llvm.ptr<3>) {
   nvvm.mbarrier.init.shared %barrier, %count : !llvm.ptr<3>, i32
   llvm.return
 }
+
+
+llvm.func private @mbarrier_inval_generic(%barrier: !llvm.ptr) {
+  // CHECK:   nvvm.mbarrier.inval %{{.*}} : !llvm.ptr
+  nvvm.mbarrier.inval %barrier : !llvm.ptr
+  llvm.return
+}
+
+
+llvm.func private @mbarrier_inval_shared(%barrier: !llvm.ptr<3>) {
+  // CHECK:   nvvm.mbarrier.inval.shared %{{.*}} : !llvm.ptr<3>
+  nvvm.mbarrier.inval.shared %barrier : !llvm.ptr<3>
+  llvm.return
+}