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

Guray Ozen llvmlistbot at llvm.org
Mon Sep 25 05:37:17 PDT 2023


https://github.com/grypp created https://github.com/llvm/llvm-project/pull/67321

None

>From df28aebbaa79dc4f07429e6580b62815c5b90761 Mon Sep 17 00:00:00 2001
From: Guray Ozen <guray.ozen at gmail.com>
Date: Mon, 25 Sep 2023 14:34:26 +0200
Subject: [PATCH] [mlir][nvgpu] Better doc for `warpgroup.mma` (nfc)

---
 mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

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