[Mlir-commits] [mlir] a224ddc - [mlir][nvvm] Introduce `cp.async.bulk.commit.group`

[Guray Ozen]

Author: Guray Ozen
Date: 2023-12-01T10:21:50+01:00
New Revision: a224ddc9b4458b1b9cf0a758c974a554f0f17dc4

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

LOG: [mlir][nvvm] Introduce `cp.async.bulk.commit.group`

This PR introduced `cp.async.bulk.commit.group` 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 54826f4196993d4..ecad1a16eb6c590 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -1420,6 +1420,15 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
 // NVVM TMA Ops
 //===----------------------------------------------------------------------===//
 
+def NVVM_CpAsyncBulkCommitGroupOp : NVVM_PTXBuilder_Op<"cp.async.bulk.commit.group">,
+  Arguments<(ins )> {
+  let assemblyFormat = "attr-dict";
+  let extraClassDefinition = [{
+    std::string $cppClass::getPtx() { return std::string("cp.async.bulk.commit_group;"); }
+  }];
+}
+
+
 def NVVM_CpAsyncBulkTensorGlobalToSharedClusterOp : 
   NVVM_Op<"cp.async.bulk.tensor.shared.cluster.global", 
   [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>, 

diff  --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index 7da4e98c40e54b4..5482cc194192ddb 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -621,3 +621,11 @@ func.func @set_max_register() {
   nvvm.setmaxregister decrease 40
   func.return
 }
+
+// -----
+
+func.func @cp_bulk_commit() {
+  //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "cp.async.bulk.commit_group;"
+  nvvm.cp.async.bulk.commit.group
+  func.return
+}