[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