[Mlir-commits] [mlir] [MLIR][NVVM] Update SM version requirements of Ops (PR #192257)
Srinivasa Ravi
llvmlistbot at llvm.org
Fri Apr 17 01:20:12 PDT 2026
https://github.com/Wolfram70 updated https://github.com/llvm/llvm-project/pull/192257
>From 54cadd73da00839c11f93937752a02bd705ce58f Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Wed, 15 Apr 2026 08:29:12 +0000
Subject: [PATCH 1/2] [MLIR][NVVM] Update SM version requirements of Ops
This change updates the SM version requirements of Ops with the
`NVVMRequiresSM` trait to include family-specific SM versions wherever
applicable.
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 48 +++++++++++++--------
1 file changed, 29 insertions(+), 19 deletions(-)
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 = [{
>From b7a5c821755b9fcfbbb93aebf86b5eba9930cf77 Mon Sep 17 00:00:00 2001
From: Srinivasa Ravi <srinivasar at nvidia.com>
Date: Fri, 17 Apr 2026 08:19:30 +0000
Subject: [PATCH 2/2] address comments
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 43 ++++++++-------------
1 file changed, 17 insertions(+), 26 deletions(-)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index ec3820b174995..c892ee18166f2 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -5075,8 +5075,7 @@ def Tcgen05WaitKindAttr :
let assemblyFormat = "`<` $value `>`";
}
-def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Tcgen05 alloc operation";
let description = [{
The `tcgen05.alloc` Op allocates tensor core memory for
@@ -5106,8 +5105,7 @@ def NVVM_Tcgen05AllocOp : NVVM_Op<"tcgen05.alloc",
}];
}
-def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Tcgen05 dealloc operation";
let description = [{
The `tcgen05.dealloc` Op de-allocates the tensor core memory
@@ -5136,7 +5134,7 @@ def NVVM_Tcgen05DeallocOp : NVVM_Op<"tcgen05.dealloc",
}
def NVVM_Tcgen05RelinquishAllocPermitOp : NVVM_Op<"tcgen05.relinquish_alloc_permit",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+ [NVVMRequiresSMf<[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
@@ -5159,8 +5157,7 @@ def NVVM_Tcgen05RelinquishAllocPermitOp : NVVM_Op<"tcgen05.relinquish_alloc_perm
}];
}
-def NVVM_Tcgen05FenceOp : NVVM_Op<"tcgen05.fence",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05FenceOp : NVVM_Op<"tcgen05.fence", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Tcgen05 fence operations";
let description = [{
The `tcgen05.fence<before>` orders all prior async tcgen05 operations
@@ -5182,8 +5179,7 @@ def NVVM_Tcgen05FenceOp : NVVM_Op<"tcgen05.fence",
}];
}
-def NVVM_Tcgen05WaitOp : NVVM_Op<"tcgen05.wait",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05WaitOp : NVVM_Op<"tcgen05.wait", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Tcgen05 wait operations";
let description = [{
The `tcgen05.wait<load>` causes the executing thread to block until
@@ -5205,8 +5201,7 @@ def NVVM_Tcgen05WaitOp : NVVM_Op<"tcgen05.wait",
}];
}
-def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05CommitOp : NVVM_Op<"tcgen05.commit", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Tcgen05 commit operations";
let description = [{
The `tcgen05.commit` makes the *mbarrier object*, specified by
@@ -5311,8 +5306,7 @@ def Tcgen05CpSrcFormatAttr : EnumAttr<NVVM_Dialect, Tcgen05CpSrcFormat, "tcgen05
let assemblyFormat = "`<` $value `>`";
}
-def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05CpOp : NVVM_Op<"tcgen05.cp", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Tcgen05 copy operation";
let description = [{
Instruction tcgen05.cp initiates an asynchronous copy operation from
@@ -5448,8 +5442,7 @@ def Tcgen05LdStShapeAttr: EnumAttr<NVVM_Dialect, Tcgen05LdStShape, "tcgen05_ldst
// NVVM tcgen05.ld Op
//===----------------------------------------------------------------------===//
-def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05LdOp : NVVM_Op<"tcgen05.ld", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "tensor memory load instructions";
let arguments = (ins
// Attributes
@@ -5631,8 +5624,7 @@ def NVVM_Tcgen05LdRedOp : NVVM_Op<"tcgen05.ld.red",
// NVVM tcgen05.st Op
//===----------------------------------------------------------------------===//
-def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+def NVVM_Tcgen05StOp : NVVM_Op<"tcgen05.st", [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "tensor memory store instructions";
let arguments = (ins
// Attributes
@@ -6011,8 +6003,7 @@ defvar Tcgen05MMABlockScaleKindAttr =
[EnumAttrIsOneOf<Tcgen05MMAKindAttr, Tcgen05MMABlockScaleKindList>]>;
def NVVM_Tcgen05MMAOp : NVVM_Op<"tcgen05.mma",
- [AttrSizedOperandSegments,
- NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+ [AttrSizedOperandSegments, NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Performs MMA operation on 5th-gen tensor cores";
let description = [{
@@ -6095,8 +6086,7 @@ def NVVM_Tcgen05MMAOp : NVVM_Op<"tcgen05.mma",
}
def NVVM_Tcgen05MMASparseOp : NVVM_Op<"tcgen05.mma.sp",
- [AttrSizedOperandSegments,
- NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+ [AttrSizedOperandSegments, NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Performs MMA operation with sparse A matrix on 5th-gen tensor cores";
let description = [{
@@ -6177,7 +6167,7 @@ def Tcgen05MMABlockScaleAttr : EnumAttr<NVVM_Dialect, Tcgen05MMABlockScale,
}
def NVVM_Tcgen05MMABlockScaleOp : NVVM_Op<"tcgen05.mma.block_scale",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+ [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Performs block scaled MMA operation on 5th-gen tensor cores";
let description = [{
@@ -6250,7 +6240,7 @@ def NVVM_Tcgen05MMABlockScaleOp : NVVM_Op<"tcgen05.mma.block_scale",
}
def NVVM_Tcgen05MMASparseBlockScaleOp : NVVM_Op<"tcgen05.mma.sp.block_scale",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+ [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Performs block scaled MMA operation with sparse A matrix on 5th-gen tensor cores";
let description = [{
@@ -6332,7 +6322,7 @@ def Tcgen05MMACollectorBBufferAttr : EnumAttr<NVVM_Dialect, Tcgen05MMACollectorB
}
def NVVM_Tcgen05MMAWsOp : NVVM_Op<"tcgen05.mma.ws",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+ [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Performs weight stationary convolution MMA operation on 5th-gen tensor cores";
let description = [{
@@ -6402,7 +6392,7 @@ def NVVM_Tcgen05MMAWsOp : NVVM_Op<"tcgen05.mma.ws",
}
def NVVM_Tcgen05MMAWsSparseOp : NVVM_Op<"tcgen05.mma.ws.sp",
- [NVVMRequiresSMaOrSMf<[100, 101, 110], [100, 101, 110]>]> {
+ [NVVMRequiresSMf<[100, 101, 110]>]> {
let summary = "Performs weight stationary convolution MMA with sparse A matrix on 5th-gen tensor cores";
let description = [{
@@ -6725,7 +6715,8 @@ def TensormapFieldValueAttr :
TensormapSwizzleModeAttr, TensormapSwizzleAtomicityAttr,
TensormapFillModeAttr]>;
-def NVVM_TensormapReplaceOp : NVVM_VoidIntrinsicOp<"tensormap.replace"> {
+def NVVM_TensormapReplaceOp : NVVM_VoidIntrinsicOp<"tensormap.replace",
+ [NVVMRequiresSMaOrSMf<[90, 103], [100, 101, 110, 120]>]> {
let summary = "Modifies a field of the tensor-map object";
let description = [{
The `nvvm.tensormap.replace` replaces the specified field of the tensor-map
More information about the Mlir-commits
mailing list