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

[guraypp]

Summary

Introduce support for PTX's mbarrier.inval .

Contiunation of D151334

[guraypp]

update

[qcolombet]

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

201

Same comment as in https://reviews.llvm.org/D151334 on the pointer type.

[guraypp]

rebase

[guraypp]

add i64 pointer

[qcolombet]

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

201

You missed a spot :).

This should be LLVM_i64ptr, right?

With LLVM_i64ptr being:

def LLVM_i64ptr : LLVM_IntPtrBase<64>

I.e., no addrspace specifier.
To match:

def llvm_i64ptr_ty        : LLVMPointerType<llvm_i64_ty>;        // i64*

[guraypp]

rebase

[guraypp]

rebase

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

 #define NVVMIR_OPS
 
 include "mlir/IR/EnumAttr.td"
 include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
 def LLVM_i8Ptr_global : LLVM_IntPtrBase<8, 1>;
 def LLVM_i8Ptr_shared : LLVM_IntPtrBase<8, 3>;
+def LLVM_i64ptr_global : LLVM_IntPtrBase<64, 1>;
+def LLVM_i64ptr_shared : LLVM_IntPtrBase<64, 3>;
 
 //===----------------------------------------------------------------------===//
 // NVVM dialect definitions
 //===----------------------------------------------------------------------===//
 
 def NVVM_Dialect : Dialect {
   let name = "nvvm";
   let cppNamespace = "::mlir::NVVM";
def NVVM_MBarrierInitOp : NVVM_Op<"mbarrier.init">,
   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_i8Ptr_shared:$addr, I32:$count)> {
+  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)";
 }
 
+def NVVM_MBarrierInvalOp : NVVM_Op<"mbarrier.inval">,
+  Arguments<(ins LLVM_AnyPointer:$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
 //===----------------------------------------------------------------------===//
 
 def NVVM_Barrier0Op : NVVM_Op<"barrier0"> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier0);
   }];

mlir/test/Dialect/LLVMIR/nvvm.mlir

 
 
 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
 }
+
+
+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
+}