[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