[Mlir-commits] [mlir] 721e91c - [mlir][nvgpu] Better doc for `warpgroup.mma` (nfc) (#67321)

llvmlistbot at llvm.org llvmlistbot at llvm.org
Mon Sep 25 05:51:27 PDT 2023


Author: Guray Ozen
Date: 2023-09-25T14:51:22+02:00
New Revision: 721e91cd2e23f391d6e3f17a2a66e1110f5d4222

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

LOG: [mlir][nvgpu] Better doc for `warpgroup.mma` (nfc) (#67321)

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td

Removed: 
    


################################################################################
diff  --git a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
index 5fcf08c6d3e1d7a..31b137160545772 100644
--- a/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
+++ b/mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td
@@ -694,11 +694,11 @@ def NVGPU_WarpgroupMmaOp : NVGPU_Op<"warpgroup.mma"> {
     results are thread-level ownership to the warpgroup-level mma operation 
     shape. The shape is deduced from the descriptor types and output vector.
 
-    The Op corresponds multiple `nvvm.wgmma.mma_async` operations to complete the 
-    given shape. As the instruction `nvvm.wgmma.async` is an asynchronous, 
-    this Op groups the `nvvm.wgmma.async` and surrounds them between 
-    `wgmma.fence.aligned` and `wgmma.commit.group.sync.aligned`, 
-    `wgmma.wait.group.sync.aligned` Ops.
+    The Op encapsulates multiple `nvvm.wgmma.mma_async` operations to complete 
+    the given shape. As `nvvm.wgmma.async` Op, or its corresponding PTX 
+    instruction, is asynchronous, this Op groups the `nvvm.wgmma.async` and 
+    surrounds them between `wgmma.fence.aligned` and 
+    `wgmma.commit.group.sync.aligned`, `wgmma.wait.group.sync.aligned` Ops.
 
     Example:
     ```mlir


        


More information about the Mlir-commits mailing list