[Mlir-commits] [mlir] fff4a46 - [MLIR][ODS][NVVM] Add EnumAttrIsOneOf/IsNoneOf constraints and use them in NVVM fence ops (#182662)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Tue Feb 24 08:14:25 PST 2026
Author: Rajat Bajpai
Date: 2026-02-24T21:44:19+05:30
New Revision: fff4a4642b7fed78bed4223f5eaf7cf1839ff048
URL: https://github.com/llvm/llvm-project/commit/fff4a4642b7fed78bed4223f5eaf7cf1839ff048
DIFF: https://github.com/llvm/llvm-project/commit/fff4a4642b7fed78bed4223f5eaf7cf1839ff048.diff
LOG: [MLIR][ODS][NVVM] Add EnumAttrIsOneOf/IsNoneOf constraints and use them in NVVM fence ops (#182662)
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]>;
```
Added:
Modified:
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
mlir/include/mlir/IR/EnumAttr.td
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
mlir/test/Target/LLVMIR/nvvm/fence-invalid.mlir
Removed:
################################################################################
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 74c5a3fa7f1dd..308b002e4ded3 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"> {
@@ -1360,8 +1363,12 @@ def ProxyKindAttr : EnumAttr<NVVM_Dialect, ProxyKind, "proxy_kind"> {
let assemblyFormat = "`<` $value `>`";
}
+defvar ProxyKindNotTensormapOrGeneric =
+ ConfinedAttr<ProxyKindAttr,
+ [EnumAttrIsNoneOf<ProxyKindAttr, [ProxyTensorMap, ProxyGeneric]>]>;
+
def NVVM_FenceProxyOp : NVVM_Op<"fence.proxy">,
- Arguments<(ins ProxyKindAttr:$kind,
+ Arguments<(ins ProxyKindNotTensormapOrGeneric:$kind,
OptionalAttr<SharedSpaceAttr>:$space)> {
let description = [{
Fence operation with proxy to establish an ordering between memory accesses
@@ -1439,7 +1446,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..cfa2ac89350a1 100644
--- a/mlir/include/mlir/IR/EnumAttr.td
+++ b/mlir/include/mlir/IR/EnumAttr.td
@@ -750,4 +750,35 @@ 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), ", ") # "}">;
+
+/// Attribute constraint restricting the enum attribute to exclude a subset of
+/// cases. `enumAttr` is the `EnumAttr` attribute and `disallowedCases` is a
+/// list of `EnumCase`s that are not permitted.
+///
+/// Example:
+/// ```
+/// ConfinedAttr<MyEnumAttr, [EnumAttrIsNoneOf<MyEnumAttr, [CaseA, CaseB]>]>
+/// ```
+class EnumAttrIsNoneOf<EnumAttr enumAttr, list<EnumCase> disallowedCases>
+ : AttrConstraint<
+ Neg<EnumAttrIsOneOf<enumAttr, disallowedCases>.predicate>,
+ "whose value is none of {"
+ # !interleave(!foreach(case, disallowedCases, case.str), ", ") # "}">;
+
#endif // ENUMATTR_TD
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index ddab6ae409f82..fa085d407d6ec 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -2733,18 +2733,7 @@ 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";
- if (getKind() == NVVM::ProxyKind::GENERIC)
- return emitOpError() << "generic proxy not a supported proxy kind";
if (getKind() == NVVM::ProxyKind::async_shared && !getSpace().has_value()) {
return emitOpError() << "async_shared fence requires space attribute";
}
@@ -2777,10 +2766,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..d2e14637f094f 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
}
@@ -17,7 +17,7 @@ llvm.func @fence_sync_restrict() {
// -----
llvm.func @fence_proxy() {
- // expected-error @below {{tensormap proxy is not a supported proxy kind}}
+ // expected-error @below {{attribute 'kind' failed to satisfy constraint: Proxy kind whose value is none of {tensormap, generic}}}
nvvm.fence.proxy {kind = #nvvm.proxy_kind<tensormap>}
llvm.return
}
@@ -25,7 +25,7 @@ llvm.func @fence_proxy() {
// -----
llvm.func @fence_proxy() {
- // expected-error @below {{generic proxy not a supported proxy kind}}
+ // expected-error @below {{attribute 'kind' failed to satisfy constraint: Proxy kind whose value is none of {tensormap, generic}}}
nvvm.fence.proxy {kind = #nvvm.proxy_kind<generic>}
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
}
More information about the Mlir-commits
mailing list