[Mlir-commits] [mlir] [MLIR][NVVM] Update SM version requirements of Ops (PR #192257)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Wed Apr 15 06:16:55 PDT 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir
Author: Srinivasa Ravi (Wolfram70)
<details>
<summary>Changes</summary>
This change updates the SM version requirements of Ops with the `NVVMRequiresSM` trait to include family-specific SM versions wherever applicable.
---
Full diff: https://github.com/llvm/llvm-project/pull/192257.diff
1 Files Affected:
- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+29-19)
``````````diff
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 9b2a8985a1a44..ec3820b174995 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -5075,7 +5075,8 @@ def Tcgen05WaitKindAttr :
let assemblyFormat = "`<` $value `>`";
}
-def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Tcgen05 alloc operation";
let description = [{
The `tcgen05.alloc` Op allocates tensor core memory for
@@ -5105,7 +5106,8 @@ def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMa<[100, 101]>]
}];
}
-def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Tcgen05 dealloc operation";
let description = [{
The `tcgen05.dealloc` Op de-allocates the tensor core memory
@@ -5133,7 +5135,8 @@ def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc", [NVVMRequiresSMa<[100, 10
}];
}
-def NVVM_Tcgen05RelinquishAllocPermitOp : NVVM_Op<"tcgen05.relinquish_alloc_permit", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05RelinquishAllocPermitOp : NVVM_Op<"tcgen05.relinquish_alloc_permit",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Tcgen05 Op to relinquish the right to allocate";
let description = [{
The `tcgen05.relinquish_alloc_permit` Op specifies that the CTA
@@ -5156,7 +5159,8 @@ def NVVM_Tcgen05RelinquishAllocPermitOp : NVVM_Op<"tcgen05.relinquish_alloc_perm
}];
}
-def NVVM_Tcgen05FenceOp : NVVM_Op<"tcgen05.fence", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05FenceOp : NVVM_Op<"tcgen05.fence",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Tcgen05 fence operations";
let description = [{
The `tcgen05.fence<before>` orders all prior async tcgen05 operations
@@ -5178,7 +5182,8 @@ def NVVM_Tcgen05FenceOp : NVVM_Op<"tcgen05.fence", [NVVMRequiresSMa<[100, 101]>]
}];
}
-def NVVM_Tcgen05WaitOp : NVVM_Op<"tcgen05.wait", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05WaitOp : NVVM_Op<"tcgen05.wait",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Tcgen05 wait operations";
let description = [{
The `tcgen05.wait<load>` causes the executing thread to block until
@@ -5200,7 +5205,8 @@ def NVVM_Tcgen05WaitOp : NVVM_Op<"tcgen05.wait", [NVVMRequiresSMa<[100, 101]>]>
}];
}
-def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Tcgen05 commit operations";
let description = [{
The `tcgen05.commit` makes the *mbarrier object*, specified by
@@ -5238,7 +5244,8 @@ def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit", [NVVMRequiresSMa<[100, 101]
}];
}
-def NVVM_Tcgen05ShiftOp : NVVM_Op<"tcgen05.shift", [NVVMRequiresSMa<[100, 101, 103]>]> {
+def NVVM_Tcgen05ShiftOp : NVVM_Op<"tcgen05.shift",
+ [NVVMRequiresSMa<[100, 101, 103, 110]>]> {
let summary = "Tcgen05 shift operation";
let description = [{
The `tcgen05.shift` is an asynchronous instruction which initiates
@@ -5304,7 +5311,8 @@ def Tcgen05CpSrcFormatAttr : EnumAttr<NVVM_Dialect, Tcgen05CpSrcFormat, "tcgen05
let assemblyFormat = "`<` $value `>`";
}
-def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Tcgen05 copy operation";
let description = [{
Instruction tcgen05.cp initiates an asynchronous copy operation from
@@ -5440,7 +5448,8 @@ def Tcgen05LdStShapeAttr: EnumAttr<NVVM_Dialect, Tcgen05LdStShape, "tcgen05_ldst
// NVVM tcgen05.ld Op
//===----------------------------------------------------------------------===//
-def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "tensor memory load instructions";
let arguments = (ins
// Attributes
@@ -5533,7 +5542,7 @@ def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMa<[100, 101]>]> {
//===----------------------------------------------------------------------===//
def NVVM_Tcgen05LdRedOp : NVVM_Op<"tcgen05.ld.red",
- [NVVMRequiresSMa<[101, 110]>]> {
+ [NVVMRequiresSMaOrSMf<[101, 110], [101, 103, 110]>]> {
let summary = "Tcgen05 tensor memory load and reduce instructions";
let arguments = (ins
Tcgen05LdStShapeAttr:$shape,
@@ -5622,7 +5631,8 @@ def NVVM_Tcgen05LdRedOp : NVVM_Op<"tcgen05.ld.red",
// NVVM tcgen05.st Op
//===----------------------------------------------------------------------===//
-def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMa<[100, 101]>]> {
+def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st",
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "tensor memory store instructions";
let arguments = (ins
// Attributes
@@ -6001,8 +6011,8 @@ defvar Tcgen05MMABlockScaleKindAttr =
[EnumAttrIsOneOf<Tcgen05MMAKindAttr, Tcgen05MMABlockScaleKindList>]>;
def NVVM_Tcgen05MMAOp : NVVM_Op<"tcgen05.mma",
- [AttrSizedOperandSegments,
- NVVMRequiresSMa<[100, 110]>]> {
+ [AttrSizedOperandSegments,
+ NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Performs MMA operation on 5th-gen tensor cores";
let description = [{
@@ -6085,8 +6095,8 @@ def NVVM_Tcgen05MMAOp : NVVM_Op<"tcgen05.mma",
}
def NVVM_Tcgen05MMASparseOp : NVVM_Op<"tcgen05.mma.sp",
- [AttrSizedOperandSegments,
- NVVMRequiresSMa<[100, 110]>]> {
+ [AttrSizedOperandSegments,
+ NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Performs MMA operation with sparse A matrix on 5th-gen tensor cores";
let description = [{
@@ -6167,7 +6177,7 @@ def Tcgen05MMABlockScaleAttr : EnumAttr<NVVM_Dialect, Tcgen05MMABlockScale,
}
def NVVM_Tcgen05MMABlockScaleOp : NVVM_Op<"tcgen05.mma.block_scale",
- [NVVMRequiresSMa<[100, 110]>]> {
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Performs block scaled MMA operation on 5th-gen tensor cores";
let description = [{
@@ -6240,7 +6250,7 @@ def NVVM_Tcgen05MMABlockScaleOp : NVVM_Op<"tcgen05.mma.block_scale",
}
def NVVM_Tcgen05MMASparseBlockScaleOp : NVVM_Op<"tcgen05.mma.sp.block_scale",
- [NVVMRequiresSMa<[100, 110]>]> {
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Performs block scaled MMA operation with sparse A matrix on 5th-gen tensor cores";
let description = [{
@@ -6322,7 +6332,7 @@ def Tcgen05MMACollectorBBufferAttr : EnumAttr<NVVM_Dialect, Tcgen05MMACollectorB
}
def NVVM_Tcgen05MMAWsOp : NVVM_Op<"tcgen05.mma.ws",
- [NVVMRequiresSMa<[100, 110]>]> {
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Performs weight stationary convolution MMA operation on 5th-gen tensor cores";
let description = [{
@@ -6392,7 +6402,7 @@ def NVVM_Tcgen05MMAWsOp : NVVM_Op<"tcgen05.mma.ws",
}
def NVVM_Tcgen05MMAWsSparseOp : NVVM_Op<"tcgen05.mma.ws.sp",
- [NVVMRequiresSMa<[100, 110]>]> {
+ [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
let summary = "Performs weight stationary convolution MMA with sparse A matrix on 5th-gen tensor cores";
let description = [{
``````````
</details>
https://github.com/llvm/llvm-project/pull/192257
More information about the Mlir-commits
mailing list