[llvm] [mlir] [NFC][mlir][NVVM] Move op descriptions to NVVMOpsDoc.td (PR #225090)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Sep 21 07:01:01 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Rajat Bajpai (rajatbajpai)
<details>
<summary>Changes</summary>
NVVM op definitions in NVVMOps.td currently inline each op's summary and description. Keeping a definition and its description in one place is convenient, but long descriptions make the definitions hard to read. This patch addresses that by moving op descriptions to a separate TableGen file, NVVMOpsDoc.td. The summary stays inline so the definition still says what the op does at a glance, while NVVMOpsDoc.td covers the details.
We have already applied similar pattern elsewhere: the llvmBuilder logic moved from the op definitions into NVVMDialect.cpp.
This patch migrates two ops (sin and mma) to showcase the new approach, and shows that it coexists with the existing one. Ops with short descriptions may therefore stay inline; that choice is left to the op authors' judgment.
---
Full diff: https://github.com/llvm/llvm-project/pull/225090.diff
3 Files Affected:
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+4-75)
- (added) mlir/include/mlir/Dialect/LLVMIR/NVVMOpsDoc.td (+97)
- (modified) utils/bazel/llvm-project-overlay/mlir/BUILD.bazel (+1)
``````````diff
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 28846502267d9..da7484cb83a5e 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -19,6 +19,7 @@ include "mlir/Dialect/LLVMIR/LLVMOpBase.td"
include "mlir/Dialect/LLVMIR/LLVMTypes.td"
include "mlir/Dialect/LLVMIR/NVVMDialect.td"
include "mlir/Dialect/LLVMIR/NVVMEnums.td"
+include "mlir/Dialect/LLVMIR/NVVMOpsDoc.td"
include "mlir/Dialect/LLVMIR/NVVMRequiresSMTraits.td"
include "mlir/Dialect/Ptr/IR/MemorySpaceInterfaces.td"
include "mlir/IR/CommonAttrConstraints.td"
@@ -38,9 +39,11 @@ def LLVM_PointerSharedCluster : LLVM_PointerInAddressSpace<7>;
//===----------------------------------------------------------------------===//
// NVVM op definitions
//===----------------------------------------------------------------------===//
-
+// Ops' descriptions are defined in NVVMOpsDoc.td in `<DefName>_Doc` records.
class NVVM_Op<string mnemonic, list<Trait> traits = []> :
LLVM_OpBase<NVVM_Dialect, mnemonic, traits> {
+ let description = !if(!exists<NVVM_OpDoc>(NAME # "_Doc"),
+ !cast<NVVM_OpDoc>(NAME # "_Doc").text, "");
}
/// Base class that defines BasicPtxBuilderOpInterface.
@@ -466,14 +469,6 @@ def NVVM_RcpApproxFtzF32Op : NVVM_IntrOp<"rcp.approx.ftz.f", [Pure], 1> {
def NVVM_SinOp : NVVM_F32UnaryApproxOp<"sin"> {
let summary = "Sine (fast approximation)";
- let description = [{
- Computes a fast approximation of the sine of the input value (in radians).
- The `ftz` attribute, when set, flushes subnormal inputs and results to
- sign-preserving zero.
-
- For more information, see PTX ISA:
- [sin](https://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-sin)
- }];
}
def NVVM_CosOp : NVVM_F32UnaryApproxOp<"cos"> {
@@ -3548,72 +3543,6 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
let summary = "cooperative matrix-multiply and accumulate";
- let description = [{
- The `nvvm.mma.sync` operation collectively performs the operation
- `D = matmul(A, B) + C` using all threads in a warp.
-
- All the threads in the warp must execute the same `mma.sync` operation.
-
- For each possible multiplicand PTX data type, there are one or more possible
- instruction shapes given as "mMnNkK". The below table describes the posssibilities
- as well as the types required for the operands. Note that the data type for
- C (the accumulator) and D (the result) can vary independently when there are
- multiple possibilities in the "C/D Type" column.
-
- When an optional attribute cannot be immediately inferred from the types of
- the operands and the result during parsing or validation, an error will be
- raised.
-
- `b1Op` is only relevant when the binary (b1) type is given to
- `multiplicandDataType`. It specifies how the multiply-and-acumulate is
- performed and is either `xor_popc` or `and_poc`. The default is `xor_popc`.
-
- `intOverflowBehavior` is only relevant when the `multiplicandType` attribute
- is one of `u8, s8, u4, s4`, this attribute describes how overflow is handled
- in the accumulator. When the attribute is `satfinite`, the accumulator values
- are clamped in the int32 range on overflow. Alternatively, accumulator
- behavior `wrapped` can be specified (this is the default), in which case
- overflow wraps from one end of the range to the other.
-
- `layoutA` and `layoutB` are required and should generally be set to
- `#nvvm.mma_layout<row>` and `#nvvm.mma_layout<col>` respectively, but other
- combinations are possible for certain layouts according to the table below.
-
- ```
- | A/B Type | Shape | ALayout | BLayout | A Type | B Type | C/D Type |
- |----------|-----------|---------|---------|----------|----------|-------------------|
- | f64 | .m8n8k4 | row | col | 1x f64 | 1x f64 | 2x f64 |
- | f16 | .m8n8k4 | row/col | row/col | 2x f16x2 | 2x f16x2 | 4x f16x2 or 8xf32 |
- | | .m16n8k8 | row | col | 2x f16x2 | 1x f16x2 | 2x f16x2 or 4 f32 |
- | | .m16n8k16 | row | col | 4x f16x2 | 2x f16x2 | 2x f16x2 or 4 f32 |
- | bf16 | .m16n8k8 | row | col | 2x i32 | 1x i32 | 4x f32 |
- | | .m16n8k16 | row | col | 4x i32 | 2x i32 | 4x f32 |
- | tf32 | .m16n8k4 | row | col | 2x i32 | 1x i32 | 4x f32 |
- | | .m16n8k8 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
- | u8/s8 | .m8n8k16 | row | col | 1x i32 | 1x i32 | 2x i32 |
- | | .m16n8k16 | row | col | 2x i32 | 1x i32 | 4x i32 |
- | | .m16n8k32 | row | col | 4x i32 | 2x i32 | 4x i32 |
- | u4/s4 | .m8n8k32 | row | col | 1x i32 | 1x i32 | 2x i32 |
- | | m16n8k32 | row | col | 2x i32 | 1x i32 | 4x i32 |
- | | m16n8k64 | row | col | 4x i32 | 2x i32 | 4x i32 |
- | b1 | m8n8k128 | row | col | 1x i32 | 1x i32 | 2x i32 |
- | | m16n8k128 | row | col | 2x i32 | 1x i32 | 4x i32 |
- ```
-
-
- Example:
- ```mlir
-
- %128 = nvvm.mma.sync A[%120, %121, %122, %123]
- B[%124, %125]
- C[%126, %127]
- shape = <m = 16, n = 8, k = 16>,
- layout_a = row, layout_b = col
- : (vector<2xf16>, vector<2xf16>, vector<2xf16>)
- -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
- ```
- }];
-
let results = (outs LLVM_AnyStruct:$res);
let arguments = (ins NVVM_MMAShapeAttr:$shape,
OptionalAttr<MMAB1OpAttr>:$b1Op,
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsDoc.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsDoc.td
new file mode 100644
index 0000000000000..433629720107b
--- /dev/null
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOpsDoc.td
@@ -0,0 +1,97 @@
+//===-- NVVMOpsDoc.td - NVVM op documentation ---------*- tablegen -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Descriptions for NVVM ops. Each op that uses this file defines a record
+// named `<DefName>_Doc` deriving from `NVVM_OpDoc`; the `NVVM_Op` class
+// in NVVMOps.td uses its `text` field as the op description.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef NVVMIR_OPS_DOC
+#define NVVMIR_OPS_DOC
+
+class NVVM_OpDoc<string docText> {
+ string text = docText;
+}
+
+def NVVM_SinOp_Doc : NVVM_OpDoc<[{
+ Computes a fast approximation of the sine of the input value (in radians).
+ The `ftz` attribute, when set, flushes subnormal inputs and results to
+ sign-preserving zero.
+
+ For more information, see PTX ISA:
+ [sin](https://docs.nvidia.com/cuda/parallel-thread-execution/#floating-point-instructions-sin)
+ }]>;
+
+def NVVM_MmaOp_Doc : NVVM_OpDoc<[{
+ The `nvvm.mma.sync` operation collectively performs the operation
+ `D = matmul(A, B) + C` using all threads in a warp.
+
+ All the threads in the warp must execute the same `mma.sync` operation.
+
+ For each possible multiplicand PTX data type, there are one or more possible
+ instruction shapes given as "mMnNkK". The below table describes the posssibilities
+ as well as the types required for the operands. Note that the data type for
+ C (the accumulator) and D (the result) can vary independently when there are
+ multiple possibilities in the "C/D Type" column.
+
+ When an optional attribute cannot be immediately inferred from the types of
+ the operands and the result during parsing or validation, an error will be
+ raised.
+
+ `b1Op` is only relevant when the binary (b1) type is given to
+ `multiplicandDataType`. It specifies how the multiply-and-acumulate is
+ performed and is either `xor_popc` or `and_poc`. The default is `xor_popc`.
+
+ `intOverflowBehavior` is only relevant when the `multiplicandType` attribute
+ is one of `u8, s8, u4, s4`, this attribute describes how overflow is handled
+ in the accumulator. When the attribute is `satfinite`, the accumulator values
+ are clamped in the int32 range on overflow. Alternatively, accumulator
+ behavior `wrapped` can be specified (this is the default), in which case
+ overflow wraps from one end of the range to the other.
+
+ `layoutA` and `layoutB` are required and should generally be set to
+ `#nvvm.mma_layout<row>` and `#nvvm.mma_layout<col>` respectively, but other
+ combinations are possible for certain layouts according to the table below.
+
+ ```
+ | A/B Type | Shape | ALayout | BLayout | A Type | B Type | C/D Type |
+ |----------|-----------|---------|---------|----------|----------|-------------------|
+ | f64 | .m8n8k4 | row | col | 1x f64 | 1x f64 | 2x f64 |
+ | f16 | .m8n8k4 | row/col | row/col | 2x f16x2 | 2x f16x2 | 4x f16x2 or 8xf32 |
+ | | .m16n8k8 | row | col | 2x f16x2 | 1x f16x2 | 2x f16x2 or 4 f32 |
+ | | .m16n8k16 | row | col | 4x f16x2 | 2x f16x2 | 2x f16x2 or 4 f32 |
+ | bf16 | .m16n8k8 | row | col | 2x i32 | 1x i32 | 4x f32 |
+ | | .m16n8k16 | row | col | 4x i32 | 2x i32 | 4x f32 |
+ | tf32 | .m16n8k4 | row | col | 2x i32 | 1x i32 | 4x f32 |
+ | | .m16n8k8 | row | col | 4x i32 | 2x i32 | 2x f16x2 or 4 f32 |
+ | u8/s8 | .m8n8k16 | row | col | 1x i32 | 1x i32 | 2x i32 |
+ | | .m16n8k16 | row | col | 2x i32 | 1x i32 | 4x i32 |
+ | | .m16n8k32 | row | col | 4x i32 | 2x i32 | 4x i32 |
+ | u4/s4 | .m8n8k32 | row | col | 1x i32 | 1x i32 | 2x i32 |
+ | | m16n8k32 | row | col | 2x i32 | 1x i32 | 4x i32 |
+ | | m16n8k64 | row | col | 4x i32 | 2x i32 | 4x i32 |
+ | b1 | m8n8k128 | row | col | 1x i32 | 1x i32 | 2x i32 |
+ | | m16n8k128 | row | col | 2x i32 | 1x i32 | 4x i32 |
+ ```
+
+
+ Example:
+ ```mlir
+
+ %128 = nvvm.mma.sync A[%120, %121, %122, %123]
+ B[%124, %125]
+ C[%126, %127]
+ shape = <m = 16, n = 8, k = 16>,
+ layout_a = row, layout_b = col
+ : (vector<2xf16>, vector<2xf16>, vector<2xf16>)
+ -> !llvm.struct<(vector<2xf16>, vector<2xf16>)>
+ ```
+ }]>;
+
+#endif // NVVMIR_OPS_DOC
diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
index 4deef03e2955b..c018b94afea11 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -6622,6 +6622,7 @@ td_library(
srcs = [
"include/mlir/Dialect/LLVMIR/NVVMDialect.td",
"include/mlir/Dialect/LLVMIR/NVVMEnums.td",
+ "include/mlir/Dialect/LLVMIR/NVVMOpsDoc.td",
"include/mlir/Dialect/LLVMIR/NVVMOps.td",
],
includes = ["include"],
``````````
</details>
https://github.com/llvm/llvm-project/pull/225090
More information about the llvm-commits
mailing list