[Mlir-commits] [mlir] 84e0145 - [mlir][nvvm] Introduce `fence.mbarrier.init` (#74058)

[llvmlistbot at llvm.org]

Author: Guray Ozen
Date: 2023-12-06T12:03:20+01:00
New Revision: 84e01450a30eabab571103d2f1195565c0ab851a

URL: https://github.com/llvm/llvm-project/commit/84e01450a30eabab571103d2f1195565c0ab851a
DIFF: https://github.com/llvm/llvm-project/commit/84e01450a30eabab571103d2f1195565c0ab851a.diff

LOG: [mlir][nvvm] Introduce `fence.mbarrier.init` (#74058)

This PR introduce `fence.mbarrier.init` OP

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 6670d94f842e9..57986f291de74 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -471,6 +471,22 @@ def NVVM_SetMaxRegisterOp : NVVM_PTXBuilder_Op<"setmaxregister"> {
   let hasVerifier = 1;
 }
 
+def NVVM_FenceMbarrierInitOp : NVVM_PTXBuilder_Op<"fence.mbarrier.init"> {
+  let arguments = (ins );
+    let description = [{
+    Fence operation that applies on the prior nvvm.mbarrier.init
+    [For more information, see PTX ISA]
+    (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar)
+  }];
+  
+  let assemblyFormat = "attr-dict";
+  let extraClassDefinition = [{        
+    std::string $cppClass::getPtx() {
+      return std::string("fence.mbarrier_init.release.cluster;");
+    }
+  }];
+}
+
 def ShflKindBfly : I32EnumAttrCase<"bfly", 0>;
 def ShflKindUp   : I32EnumAttrCase<"up", 1>;
 def ShflKindDown : I32EnumAttrCase<"down", 2>;

diff  --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index 4d2d152845898..43de50f3dc8de 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -642,6 +642,15 @@ func.func @cp_bulk_commit() {
   nvvm.cp.async.bulk.commit.group
   func.return
 }
+
+
+// -----
+
+func.func @fence_mbarrier_init() {
+  //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.mbarrier_init.release.cluster;"
+  nvvm.fence.mbarrier.init
+  func.return 
+}
 // -----
 
 func.func @fence_proxy() {