[MLIR][ODS][NVVM] Add EnumAttrIsOneOf/IsNoneOf constraints and use them in NVVM fence ops#182662
Merged
rajatbajpai merged 1 commit intoFeb 24, 2026
Merged
Conversation
Member
|
@llvm/pr-subscribers-mlir-ods @llvm/pr-subscribers-mlir-core Author: Rajat Bajpai (rajatbajpai) ChangesAdd Use it in the NVVM dialect to declaratively constrain the This eliminates the It can also be parameterized through base class hierarchies, allowing derived ops to specify their own allowed subsets: Full diff: https://github.com/llvm/llvm-project/pull/182662.diff 4 Files Affected:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 9725cdb6a4c4d..ddbf7c22d8c62 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -1307,8 +1307,13 @@ def NVVM_FenceScClusterOp : NVVM_Op<"fence.sc.cluster"> {
let assemblyFormat = "attr-dict";
}
+defvar MemOrderAcquireOrRelease =
+ ConfinedAttr<MemOrderKindAttr,
+ [EnumAttrIsOneOf<MemOrderKindAttr,
+ [MemOrderKindAcquire, MemOrderKindRelease]>]>;
+
def NVVM_FenceSyncRestrictOp : NVVM_Op<"fence.sync_restrict">,
- Arguments<(ins MemOrderKindAttr:$order)> {
+ Arguments<(ins MemOrderAcquireOrRelease:$order)> {
let summary = "Uni-directional thread fence operation";
let description = [{
The `nvvm.fence.sync_restrict` Op restricts the class of memory
@@ -1322,8 +1327,6 @@ def NVVM_FenceSyncRestrictOp : NVVM_Op<"fence.sync_restrict">,
let llvmBuilder = [{
createIntrinsicCall(builder, getFenceSyncRestrictID($order));
}];
-
- let hasVerifier = 1;
}
def NVVM_FenceMbarrierInitOp : NVVM_Op<"fence.mbarrier.init"> {
@@ -1439,7 +1442,7 @@ def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
}
def NVVM_FenceProxySyncRestrictOp : NVVM_Op<"fence.proxy.sync_restrict">,
- Arguments<(ins MemOrderKindAttr:$order,
+ Arguments<(ins MemOrderAcquireOrRelease:$order,
DefaultValuedAttr<ProxyKindAttr, "ProxyKind::GENERIC">:$fromProxy,
DefaultValuedAttr<ProxyKindAttr, "ProxyKind::async">:$toProxy)> {
let summary = "Uni-directional proxy fence operation with sync_restrict";
diff --git a/mlir/include/mlir/IR/EnumAttr.td b/mlir/include/mlir/IR/EnumAttr.td
index ec57626ebde65..97241f8a12e90 100644
--- a/mlir/include/mlir/IR/EnumAttr.td
+++ b/mlir/include/mlir/IR/EnumAttr.td
@@ -750,4 +750,21 @@ class ConstantEnumCase<Attr attribute, string case>
"attribute must be one of 'EnumAttr' or 'EnumInfo'";
}
+/// Attribute constraint restricting the enum attribute to a subset of allowed
+/// cases. `enumAttr` is the `EnumAttr` attribute and `allowedCases` is a list
+/// of `EnumCase`s that are permitted.
+///
+/// Example:
+/// ```
+/// ConfinedAttr<MyEnumAttr, [EnumAttrIsOneOf<MyEnumAttr, [CaseA, CaseB]>]>
+/// ```
+class EnumAttrIsOneOf<EnumAttr enumAttr, list<EnumCase> allowedCases>
+ : AttrConstraint<
+ Or<!foreach(case, allowedCases,
+ CPred<"::llvm::cast<" # enumAttr.cppNamespace # "::"
+ # enumAttr.cppClassName # ">($_self).getValue() == "
+ # enumAttr.enum.cppType # "::" # case.symbol>)>,
+ "whose value is one of {"
+ # !interleave(!foreach(case, allowedCases, case.str), ", ") # "}">;
+
#endif // ENUMATTR_TD
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index b51eb3408df00..dc0abae881465 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -2731,13 +2731,6 @@ bool NVVM::WgmmaMmaAsyncOp::getAsmValues(
return true; // Has manual mapping
}
-LogicalResult NVVM::FenceSyncRestrictOp::verify() {
- if (getOrder() != NVVM::MemOrderKind::ACQUIRE &&
- getOrder() != NVVM::MemOrderKind::RELEASE)
- return emitOpError("only acquire and release semantics are supported");
- return success();
-}
-
LogicalResult NVVM::FenceProxyOp::verify() {
if (getKind() == NVVM::ProxyKind::TENSORMAP)
return emitOpError() << "tensormap proxy is not a supported proxy kind";
@@ -2775,10 +2768,6 @@ LogicalResult NVVM::FenceProxyReleaseOp::verify() {
}
LogicalResult NVVM::FenceProxySyncRestrictOp::verify() {
- if (getOrder() != NVVM::MemOrderKind::ACQUIRE &&
- getOrder() != NVVM::MemOrderKind::RELEASE)
- return emitOpError("only acquire and release semantics are supported");
-
if (getFromProxy() != NVVM::ProxyKind::GENERIC)
return emitOpError("only generic is support for from_proxy attribute");
diff --git a/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
index 22578b5581da4..606b3fcffb272 100644
--- a/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-translate --mlir-to-llvmir -verify-diagnostics -split-input-file %s
llvm.func @fence_sync_restrict() {
- // expected-error @below {{only acquire and release semantics are supported}}
+ // expected-error @below {{attribute 'order' failed to satisfy constraint: NVVM Memory Ordering kind whose value is one of {acquire, release}}}
nvvm.fence.sync_restrict {order = #nvvm.mem_order<weak>}
llvm.return
}
@@ -9,7 +9,7 @@ llvm.func @fence_sync_restrict() {
// -----
llvm.func @fence_sync_restrict() {
- // expected-error @below {{only acquire and release semantics are supported}}
+ // expected-error @below {{attribute 'order' failed to satisfy constraint: NVVM Memory Ordering kind whose value is one of {acquire, release}}}
nvvm.fence.sync_restrict {order = #nvvm.mem_order<mmio>}
llvm.return
}
@@ -65,7 +65,7 @@ llvm.func @fence_proxy_release() {
// -----
llvm.func @fence_proxy_sync_restrict() {
- // expected-error @below {{only acquire and release semantics are supported}}
+ // expected-error @below {{attribute 'order' failed to satisfy constraint: NVVM Memory Ordering kind whose value is one of {acquire, release}}}
nvvm.fence.proxy.sync_restrict {order = #nvvm.mem_order<mmio>}
llvm.return
}
|
Member
|
@llvm/pr-subscribers-mlir-llvm Author: Rajat Bajpai (rajatbajpai) ChangesAdd Use it in the NVVM dialect to declaratively constrain the This eliminates the It can also be parameterized through base class hierarchies, allowing derived ops to specify their own allowed subsets: Full diff: https://github.com/llvm/llvm-project/pull/182662.diff 4 Files Affected:
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 9725cdb6a4c4d..ddbf7c22d8c62 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -1307,8 +1307,13 @@ def NVVM_FenceScClusterOp : NVVM_Op<"fence.sc.cluster"> {
let assemblyFormat = "attr-dict";
}
+defvar MemOrderAcquireOrRelease =
+ ConfinedAttr<MemOrderKindAttr,
+ [EnumAttrIsOneOf<MemOrderKindAttr,
+ [MemOrderKindAcquire, MemOrderKindRelease]>]>;
+
def NVVM_FenceSyncRestrictOp : NVVM_Op<"fence.sync_restrict">,
- Arguments<(ins MemOrderKindAttr:$order)> {
+ Arguments<(ins MemOrderAcquireOrRelease:$order)> {
let summary = "Uni-directional thread fence operation";
let description = [{
The `nvvm.fence.sync_restrict` Op restricts the class of memory
@@ -1322,8 +1327,6 @@ def NVVM_FenceSyncRestrictOp : NVVM_Op<"fence.sync_restrict">,
let llvmBuilder = [{
createIntrinsicCall(builder, getFenceSyncRestrictID($order));
}];
-
- let hasVerifier = 1;
}
def NVVM_FenceMbarrierInitOp : NVVM_Op<"fence.mbarrier.init"> {
@@ -1439,7 +1442,7 @@ def NVVM_FenceProxyReleaseOp : NVVM_Op<"fence.proxy.release">,
}
def NVVM_FenceProxySyncRestrictOp : NVVM_Op<"fence.proxy.sync_restrict">,
- Arguments<(ins MemOrderKindAttr:$order,
+ Arguments<(ins MemOrderAcquireOrRelease:$order,
DefaultValuedAttr<ProxyKindAttr, "ProxyKind::GENERIC">:$fromProxy,
DefaultValuedAttr<ProxyKindAttr, "ProxyKind::async">:$toProxy)> {
let summary = "Uni-directional proxy fence operation with sync_restrict";
diff --git a/mlir/include/mlir/IR/EnumAttr.td b/mlir/include/mlir/IR/EnumAttr.td
index ec57626ebde65..97241f8a12e90 100644
--- a/mlir/include/mlir/IR/EnumAttr.td
+++ b/mlir/include/mlir/IR/EnumAttr.td
@@ -750,4 +750,21 @@ class ConstantEnumCase<Attr attribute, string case>
"attribute must be one of 'EnumAttr' or 'EnumInfo'";
}
+/// Attribute constraint restricting the enum attribute to a subset of allowed
+/// cases. `enumAttr` is the `EnumAttr` attribute and `allowedCases` is a list
+/// of `EnumCase`s that are permitted.
+///
+/// Example:
+/// ```
+/// ConfinedAttr<MyEnumAttr, [EnumAttrIsOneOf<MyEnumAttr, [CaseA, CaseB]>]>
+/// ```
+class EnumAttrIsOneOf<EnumAttr enumAttr, list<EnumCase> allowedCases>
+ : AttrConstraint<
+ Or<!foreach(case, allowedCases,
+ CPred<"::llvm::cast<" # enumAttr.cppNamespace # "::"
+ # enumAttr.cppClassName # ">($_self).getValue() == "
+ # enumAttr.enum.cppType # "::" # case.symbol>)>,
+ "whose value is one of {"
+ # !interleave(!foreach(case, allowedCases, case.str), ", ") # "}">;
+
#endif // ENUMATTR_TD
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index b51eb3408df00..dc0abae881465 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -2731,13 +2731,6 @@ bool NVVM::WgmmaMmaAsyncOp::getAsmValues(
return true; // Has manual mapping
}
-LogicalResult NVVM::FenceSyncRestrictOp::verify() {
- if (getOrder() != NVVM::MemOrderKind::ACQUIRE &&
- getOrder() != NVVM::MemOrderKind::RELEASE)
- return emitOpError("only acquire and release semantics are supported");
- return success();
-}
-
LogicalResult NVVM::FenceProxyOp::verify() {
if (getKind() == NVVM::ProxyKind::TENSORMAP)
return emitOpError() << "tensormap proxy is not a supported proxy kind";
@@ -2775,10 +2768,6 @@ LogicalResult NVVM::FenceProxyReleaseOp::verify() {
}
LogicalResult NVVM::FenceProxySyncRestrictOp::verify() {
- if (getOrder() != NVVM::MemOrderKind::ACQUIRE &&
- getOrder() != NVVM::MemOrderKind::RELEASE)
- return emitOpError("only acquire and release semantics are supported");
-
if (getFromProxy() != NVVM::ProxyKind::GENERIC)
return emitOpError("only generic is support for from_proxy attribute");
diff --git a/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir b/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
index 22578b5581da4..606b3fcffb272 100644
--- a/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
@@ -1,7 +1,7 @@
// RUN: mlir-translate --mlir-to-llvmir -verify-diagnostics -split-input-file %s
llvm.func @fence_sync_restrict() {
- // expected-error @below {{only acquire and release semantics are supported}}
+ // expected-error @below {{attribute 'order' failed to satisfy constraint: NVVM Memory Ordering kind whose value is one of {acquire, release}}}
nvvm.fence.sync_restrict {order = #nvvm.mem_order<weak>}
llvm.return
}
@@ -9,7 +9,7 @@ llvm.func @fence_sync_restrict() {
// -----
llvm.func @fence_sync_restrict() {
- // expected-error @below {{only acquire and release semantics are supported}}
+ // expected-error @below {{attribute 'order' failed to satisfy constraint: NVVM Memory Ordering kind whose value is one of {acquire, release}}}
nvvm.fence.sync_restrict {order = #nvvm.mem_order<mmio>}
llvm.return
}
@@ -65,7 +65,7 @@ llvm.func @fence_proxy_release() {
// -----
llvm.func @fence_proxy_sync_restrict() {
- // expected-error @below {{only acquire and release semantics are supported}}
+ // expected-error @below {{attribute 'order' failed to satisfy constraint: NVVM Memory Ordering kind whose value is one of {acquire, release}}}
nvvm.fence.proxy.sync_restrict {order = #nvvm.mem_order<mmio>}
llvm.return
}
|
durga4github
approved these changes
Feb 24, 2026
Contributor
durga4github
left a comment
durga4github
left a comment
There was a problem hiding this comment.
Nice change!
LGTM, though I have a trivial follow-up ask.
Please wait for Guray to take a look before landing.
rajatbajpai
force-pushed
the
dev/rbajpai/enum-attr-ods-verifier
branch
from
55cf32c to
fa87158
Compare
rajatbajpai
changed the title
grypp
approved these changes
Feb 24, 2026
…em in NVVM fence ops
Add `EnumAttrIsOneOf` and `EnumAttrIsNoneOf` to `EnumAttr.td`, a generic
`AttrConstraint`s that restrict an `EnumAttr` to an allowlist or denylist
of cases respectively. These are the enum analogs of the existing
`IntIsOneOf` constraint for integer attributes.
`EnumAttrIsOneOf` restricts the attribute to a closed set of permitted
cases. `EnumAttrIsNoneOf` excludes specific cases while allowing all
others, which is more ergonomic when the disallowed set is smaller than
the allowed set and more future-proof when new enum cases are added.
This replaces hand-written verifier logic with ODS-generated
verification, eliminating the `FenceSyncRestrictOp` verifier entirely
and simplifying both `FenceProxyOp` and `FenceProxySyncRestrictOp`
verifiers.
It can also be parameterized through base class hierarchies, allowing
derived ops to specify their own allowed subsets:
```
class MyBase<string mnemonic, list<EnumCase> allowed>
: NVVM_Op<mnemonic>,
Arguments<(ins
ConfinedAttr<MyEnumAttr,
[EnumAttrIsOneOf<MyEnumAttr, allowed>]>:$kind)>;
def Op1 : MyBase<"op1", [CaseA, CaseB]>;
def Op2 : MyBase<"op2", [CaseB, CaseC, CaseD]>;
```
rajatbajpai
force-pushed
the
dev/rbajpai/enum-attr-ods-verifier
branch
from
fa87158 to
8302475
Compare
schwarzschild-radius
approved these changes
Feb 24, 2026
Contributor
schwarzschild-radius
left a comment
schwarzschild-radius
left a comment
There was a problem hiding this comment.
Thanks @rajatbajpai! LGTM!
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Add
EnumAttrIsOneOfandEnumAttrIsNoneOftoEnumAttr.td, a genericAttrConstraints that restrict anEnumAttrto an allowlist or denylistof cases respectively. These are the enum analogs of the existing
IntIsOneOfconstraint for integer attributes.EnumAttrIsOneOfrestricts the attribute to a closed set of permittedcases.
EnumAttrIsNoneOfexcludes specific cases while allowing allothers, which is more ergonomic when the disallowed set is smaller than
the allowed set and more future-proof when new enum cases are added.
This replaces hand-written verifier logic with ODS-generated
verification, eliminating the
FenceSyncRestrictOpverifier entirelyand simplifying both
FenceProxyOpandFenceProxySyncRestrictOpverifiers.
It can also be parameterized through base class hierarchies, allowing
derived ops to specify their own allowed subsets: