This patch adds support for `float8_e3m4` and `float8_e4m3` in
`np_to_memref.py` by adding the appropriate ctypes structures.
Additionally changes minimum numpy version to 2.1.0 and uses a single
ml_dtypes version of 0.5.0.
This PR adds some LLVM metadata attributes and an `llvm.named_metadata`
container op (similar to `llvm.module_flags`) for those attributes.
Summary:
- Add MLIR attributes modeling LLVM IR metadata: `#llvm.md_string`,
`#llvm.md_const`, `#llvm.md_func`, and `#llvm.md_node`;
- Add `llvm.named_metadata` container op for module-level named metadata
nodes;
- Add MLIR-to-LLVM-IR translation for the new attributes and op;
- Add C API functions (`mlirLLVMMDStringAttrGet`,
`mlirLLVMMDNodeAttrGet`, etc.);
- Add Python bindings (`llvm.MDStringAttr`, `llvm.MDConstantAttr`,
`llvm.MDFuncAttr`, `llvm.MDNodeAttr`, `llvm.FunctionType`).
This includes several changes:
- `Dialect.load(reload=False)` will fail if the dialect was already
loaded in a different context. To prevent the further program abortion.
- `Dialect.load(reload=True)` implies `replace=True` in
dialect/operation registering.
- `PyGlobals::registerDialectImpl` now has a parameter `replace`.
- `register_dialect` and `register_operation` is no longer exposed in
`mlir.dialects.ext`.
This should solve the registering problem found in writing transform
test cases by @rolfmorel.
MLIR's C++ `Operation::walk` supports type-filtered traversal (e.g.
`op->walk([](arith::AddIOp op) { ... })`), but the Python binding
`op.walk()` requires users to manually implement type filtering inside
the callback function.
This PR adds type filtering into the python binding `op.walk()`, if
users pass `op_class`, walk() will only apply callback to matching ops.
This PR also adds a common use helper in mlir/ir that collects all ops
of a given type into a list. Users can just call: `ops =
ir.get_ops_of_type(root, op_class)`.
Now we need to provide a location when call `load()`, e.g.
```python
with Context(), Location.unknown():
MyDialect.load()
```
But it's actually weird: IRDL is just an implementation details, so for
users they don't know why they need to provide a location for loading a
dialect, which is unrelated to constructing an IR module.
This PR made location optional for dialect loading.
The `op.parent` only returns the immediate parent, in which case
downstream users have to traverse the operation by themselves to find a
specific type op.
This PR adds a python function `get_parent_of_type()` to mlir.ir to
provide an API to do so.
The function mirrors the implementation here:
https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/IR/Operation.h#L257-L273.
Instead of creating a new binding, reimplement it in python using
`isinstance()` is simpler.
Currently XeGPU transform dialect does not allow the user to set the
`order` attribute of a layout in `SetDescLayoutOp` and
`SetOpLayoutAttrOp`. This PR adds `order` as an optional argument to
these transform ops.
Follow-up to #184253. The ODS attr/type printer fix removed the leading
space from generated print() methods. Update tests that checked for the
old double-space output of GPU ops using GPU_DimensionAttr and
GPU_MmaElementwiseOpAttr.
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
Follow-up to #184253. Update tests that checked for the old double-space
output of NVVM ops using ReductionKindAttr, ShflKindAttr, and
LoadCacheModifierAttr.
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
The `OneDimMultiReductionToTwoDim` pattern had some issues. For the
input program:
```mlir
func.func @rank1_multi_reduction(%arg0: vector<8xf32>, %acc: f32) -> f32 {
%0 = vector.multi_reduction <add>, %arg0, %acc [0] : vector<8xf32> to f32
return %0 : f32
}
```
* when lowering using the inner-parallel strategy, the compiler would
essentially produce scalar code:
```mlir
func.func @rank1_multi_reduction(%arg0: vector<8xf32>, %arg1: f32) -> f32 {
%0 = vector.shape_cast %arg0 : vector<8xf32> to vector<1x8xf32>
%1 = vector.broadcast %arg1 : f32 to vector<1xf32>
%2 = vector.transpose %0, [1, 0] : vector<1x8xf32> to vector<8x1xf32>
%3 = vector.extract %2[0] : vector<1xf32> from vector<8x1xf32>
%4 = arith.addf %3, %1 : vector<1xf32>
%5 = vector.extract %2[1] : vector<1xf32> from vector<8x1xf32>
%6 = arith.addf %5, %4 : vector<1xf32>
... (repeats for all 8 elements) ...
%17 = vector.extract %2[7] : vector<1xf32> from vector<8x1xf32>
%18 = arith.addf %17, %16 : vector<1xf32>
%19 = vector.extract %18[0] : f32 from vector<1xf32>
return %19 : f32
}
```
* when lowering using the inner-reduction strategy, the compiler would
first unnecessarily transform it into a 2-D multi_reduction operation
<1x8xf32> and then extract an <8xf32> vector and apply reduction. The
canonicalization and folding would lead to the following final result:
```mlir
func.func @rank1_multi_reduction(%arg0: vector<8xf32>, %arg1: f32) -> f32 {
%0 = vector.reduction <add>, %arg0, %arg1 : vector<8xf32> into f32
return %0 : f32
}
```
Now, after this change:
* when lowering the compiler now produces for both strategies in one
step.
```
func.func @rank1_multi_reduction(%arg0: vector<8xf32>, %arg1: f32) -> f32 {
%0 = vector.reduction <add>, %arg0, %arg1 : vector<8xf32> into f32
return %0 : f32
}
```
This pattern is also useful for an ongoing refactoring that is happening
in the multi_reduction patterns. It is the only pattern that increases
multi_reduction in rank and would lead to an infinite loop when
attempting to reach a fixed point once we generalize other unrolling
patterns.
Assisted-by: Claude
Makes it possible to include Python-defined rewrite patterns in
transform-dialect schedules, inside of `transform.apply_patterns`, which
upon execution of the schedule runs the pattern in a greedy rewriter.
With assistance of Claude.
This PR is quite similiar to
https://github.com/llvm/llvm-project/pull/182805.
We added basic support of attribute definitions in Python-defined
dialects, including:
- IRDL codegen for attribute definitions
- Attr builders like `MyAttr.get(..)` and attr parameter accessors (e.g.
`my_attr.param1`)
- Use Python-defined attrs in Python-defined operations
Assisted by GitHub Copilot.
Add `MLIR_ENABLE_PYTHON_STABLE_ABI` cmake flag to build bindings against
the Python limited/stable API (abi3 / PEP 384). This allow for
compatibility across different >=3.12 versions with a single .so /
wheel. We also require CMake >=3.26.
The stable ABI restricts usage to a subset of the CPython C API: frame
and code object structs are opaque, so introspection APIs like
`PyCode_Addr2Location`, `PyFrame_GetLasti`, and `PyFrame_GetCode` are
unavailable. The traceback-based auto-location logic is dropped because
we don’t have stable ABI to produce complete locations.
Assisted-by: claude
This PR adds the `convert_region_types` API to
`ConversionPatternRewriter` and introduces a new integration test,
`bf.py`, which demonstrates how to combine a Python-defined dialect, the
dialect conversion API, the pass manager, and the execution engine to
build a pure-Python JIT compilation pipeline.
Previously, in #177782, we added support for dialect conversion and
generated an `OpAdaptor` subtype for every ODS-defined operation. In
this PR, we will also generate `OpAdaptor` subtypes for Python-defined
operations, so that they can be applied in dialect conversion as well.
Renames 'x86vector' dialect to 'x86'.
This is the first PR in series of cleanups around dialects targeting x86
platforms.
The new naming scheme is shorter, cleaner, and opens possibility of
integrating other x86-specific operations not strictly fitting pure
vector representation. For example, the generalization will allow for
future merger of AMX dialect into the x86 dialect to create one-stop x86
operations collection and boost discoverability.
Before this, MLIR error capture in `apply_partial_conversion` and
`apply_full_conversion` wasn’t handled, which meant any `emitError`
would crash the entire program. This PR adds the handling.
Previously, we were using the static `typeid` of `DynamicType` for
checks, which is incorrect. We should instead check against the `typeid`
of `DynamicTypeDefinition` (which is a subclass of `SelfOwningTypeID`),
and register it via `register_type_caster` so that Python-defined types
can use `maybe_downcast`. (The attribute part is same.)
Includes `InferOpTypeInterface.td` in `NVVMOps.td` enabling result type
inference for NVVM operations.
Fixes a test for `nvvm.redux.sync` in `nvvm.py` due to a resulting
change in the python binding for the operation.
When `ndim == 0`, `view->strides[view->ndim - 1]` is an out-of-bounds
access (unsigned underflow to `SIZE_MAX`). Use `view->itemsize` for
alignment instead, since a scalar buffer is trivially aligned to its
element size.
Fixesiree-org/iree-turbine#1312.
In this PR, we added basic support of type definitions in Python-defined
dialects, including:
- IRDL codegen for type definitions
- Type builders like `MyType.get(..)` and type parameter accessors (e.g.
`my_type.param1`)
- Use Python-defined types in Python-defined oeprations
```python
class TestType(Dialect, name="ext_type"):
pass
class Array(TestType.Type, name="array"):
elem_type: IntegerType[32] | IntegerType[64]
length: IntegerAttr
class MakeArrayOp(TestType.Operation, name="make_array"):
arr: Result[Array]
class MakeArray3Op(TestType.Operation, name="make_array3"):
arr: Result[Array[IntegerType[32], IntegerAttr[IntegerType[32], 3]]]
```
This PR adds C and Python API support for `mlir::DynamicAttr`. It
primarily enables attributes in dialects that are dynamically generated
via IRDL to be constructed in Python, and allows retrieving the
parameters contained in a dynamic attribute from Python.
This PR is quite similiar to #182751, so I use tab to autocomplete some
code via github copilot, but manually verified.
This PR adds C and Python API support for `mlir::DynamicType`. It
primarily enables types in dialects that are dynamically generated via
IRDL to be constructed in Python, and allows retrieving the parameters
contained in a dynamic type from Python.
---------
Co-authored-by: Rolf Morel <rolfmorel@gmail.com>
Previously, we constructed new ops using the pattern `class
MyOp(MyInt.Operation)`.
Now we’ve added a new pattern: `class MyOp(Operation, dialect=MyInt)`,
which allows more flexible composition. For example:
```python
class BinOpBase(Operation): # it can be used in any dialect!
res: Result[Any]
lhs: Operand[Any]
rhs: Operand[Any]
class MyInt(Dialect, name="myint"):
pass
class AddOp(BinOpBase, dialect=MyInt, name="add"):
...
```
* Adds vector transform op
`apply_patterns.vector.multi_reduction_unrolling`
* Adds test for `populateVectorMultiReductionUnrollingPatterns`
* Deletes old test files `vector-multi-reduction-lowering.mlir` and
`vector-multi-reduction-lowering-outer.mlir`. Tests that exercise these
patterns exist in `vector-multi-reduction-flattening.mlir`,
`vector-multi-reduction-reorder-and-expand.mlir` and
`vector-multi-reduction-unrolling.mlir`
Assisted-by: claude
With the new finer grained populate methods introduced in
8dde3051504cb9ae42e654bbce39001f3946beea (#180750), there was a
discussion about refactoring tests such that only one of the patterns
applies at a time. This commit starts this process by adding the
structure for one of these populate methods. The goal is for the
populate methods to have their own file (each showing inner and outer
reduction); deprecating populateVectorMultiReductionLoweringPatterns and
ApplyLowerMultiReduction; and removing the test file for
mlir/test/Dialect/Vector/vector-multi-reduction-lowering.mlir
Essentially an NFC. It also adds a new transform op for testing the
dialect and which downstream projects may choose to use.
Assisted-By: claude-4.5-sonnet
Changes `transform.xegpu.set_op_layout_attr` to support xegpu anchor
layouts. By default, if `result` and `operand` bool arguments are unset,
this transform op sets the op's anchor layout, if the op supports it
(otherwise emits a silenceable failure).
In contrast to the earlier implementation, setting the operand layout
now requires setting the new `operand` argument.
Provides the infrastructure for implementing and late-binding
OpInterfaces from Python.
* On the mlir-c API declaration side, each `XOpInterface` has a callback
struct, with a callback for each method and a userdata member (provided
as an arg to each method), and a
`mlirXOpInterfaceAttachFallbackModel(ctx, op_name, callbacks)` func.
* This CAPI is implemented by defining a subclass of
`XOpInterface::FallbackModel` that holds the callback struct and has
each method call the corresponding callback (with userdata as an arg).
Given a callback struct, a new `FallbackModel` is created and attached,
i.e. late bound, to the named op. (MLIR's interface infrastructure is
such that the thus registered `FallbackModel` will be returned in case
the op gets cast to the `XOpInterface`.)
* On the Python side, we expose a stand-in `XOpInterface` base class
which has one (class)method: `XOpInterface.attach(cls, op_name, ctx)`.
Python users subclass this class (`class MyInterfaceImpl(XOpInterface):
...`) and implement the interface's methods (with the right names and
signatures). The user calls `attach` on the subclass
(`MyInterfaceImpl.attach("my_dialect.my_op", ctx)`) which prepares the
callbacks struct _with userdata set to the subclass_ (as we use it to
lookup methods). These callbacks (and userdata) are then registered as
an `XOpInterface::FallbackModel` by
`mlirXOpInterfaceAttachFallbackModel(...)`. From then on the Python
methods will be used to respond to calls to the interface methods
(originating in C++).
This PR enables implementing the TransformOpInterface and the
MemoryEffectsOpInterface, both of which are required for making an op
into a transform op.
Everything besides the above linked code is there to facilitate exposing
the interfaces: the right types for the arguments of the methods are
exposed as are functions/methods for manipulating these arguments (e.g.
specifying side effects on `OpOperand`s and `OpResult`s and being able
to access and set the transform handles associated with args and
results).
`DenseElementsAttr` stores elements in a `ArrayRef<char>` buffer, where
each element is padded to a full byte. Before this commit, there used to
be a special storage format for `i1` elements: they used to be densely
packed, i.e., 1 bit per element. This commit removes the dense packing
special case for `i1`.
This commit removes complexity from `DenseElementsAttr`. If dense
packing is needed in the future it could be implemented in a general way
that works for all element types (based on #179122).
Discussion:
https://discourse.llvm.org/t/denseelementsattr-i1-element-type/62525
This will support two syntax in python-defined dialects.
First is that traits can now be declared in class parameters, e.g.
```python
class ParentIsIfTrait(DynamicOpTrait): #define a python-side trait
@staticmethod
def verify_invariants(op) -> bool:
if not isinstance(op.parent.opview, IfOp):
op.location.emit_error(
f"{op.name} should be put inside {IfOp.OPERATION_NAME}"
)
return False
return True
class YieldOp( # attach two traits: IsTerminatorTrait, ParentIsIfTrait
TestRegion.Operation, name="yield", traits=[IsTerminatorTrait, ParentIsIfTrait]
):
...
```
Second is that users can directly define
`verify_invariants`/`verify_region_invariants` methods in the operation
to add additional custom verification logic. And this is implemented via
traits.
```python
class YieldOp(TestRegion.Operation, name="yield", ...):
value: Operand[Any]
def verify_invariants(self) -> bool: # define a method directly
if self.parent.results[0].type != self.value.type:
self.location.emit_error(
"result type mismatch between YieldOp and its parent IfOp"
)
return False
return True
```
Previously we use `verify`/`verify_region` as method names (in
yesterday's PR #179705), but in this PR they are renamed to
`verify_invariants`/`verify_region_invariants` because there are
conflicts between the newly-added `verify` method and `ir.OpView.verify`
method:
- `verify_invariants` is just to attach **additional** verification
logic. but `OpView.verify` is to construct an OperationVerifer and do
full verification for an operation, so the semantics is not same between
these two. We should not shadow the `OpView.verify` method by defining a
new semantically-different `verify` method.
- it will make users confuse between these two `verify` methods, since
they have different meaning.
- if users didn't define the `verify` method in their python-defined
operation, `DynamicOpTraits.attach(opname, MyOpCls)` still do the
attaching (because `hasattr("verify")` returns `True`) and seg fault
(because we cannot attach `OpView.verify`).
---------
Co-authored-by: Rolf Morel <rolfmorel@gmail.com>
This is a follow-up PR of #169045 and the second part of #179086.
In #179086, we added support for defining regions in Python-defined ops,
but its usefulness was quite limited because we still couldn’t mark an
op as a `Terminator` or `NoTerminator`. In this PR, we port the
`DynamicOpTrait` (introduced on the C++ side for `DynamicDialect` in
#177735) to Python, so we can dynamically attach traits to
Python-defined ops.
This PR adds basic support for defining regions in Python-defined
dialects. Example usage:
```python
class TestRegion(Dialect, name="ext_region"):
pass
class IfOp(TestRegion.Operation, name="if"):
cond: Operand[IntegerType[1]]
then: Region
else_: Region
```
Current limitations:
* We can’t specify region constraints yet (e.g., number of blocks or
block argument types). This will be addressed as a follow-up task.
* We can’t mark an op as a `Terminator` or `NoTerminator` yet. This
depends on `DynamicOpTraits` (#177735) and Python-side trait API
support, and will be implemented in a follow-up PR.
This is the first PR after splitting off #179032.
This is a follow-up PR of #169045.
---------
Co-authored-by: Rolf Morel <rolfmorel@gmail.com>
This PR adds dialect conversion support to the MLIR Python bindings.
Because it introduces a number of new APIs, it’s a fairly large PR. It
mainly includes the following parts:
* Add a set of types and APIs to the C API, including
`MlirConversionTarget`, `MlirConversionPattern`, `MlirTypeConverter`,
`MlirConversionPatternRewriter`, and others.
* Add the corresponding types and APIs to the Python bindings.
* Extend `mlir-tblgen` with codegen for Python adaptor classes, which
generates an adaptor class for each op.
Note that this PR only adds support for 1-to-1 conversions, 1-to-N
type/value conversions are not supported yet.
---------
Co-authored-by: Maksim Levental <maksim.levental@gmail.com>
The commit adds the following:
- Adds tcgen05.ld.red Op with tests under tcgen05-ld-red.mlir and
tcgen05-ld-red-invalid.mlir
- Renamed ReduxKind to ReductionKind and renamed it across NVVM and GPU
Dialects
- Replaced Tcgen05LdRedOperationAtr with ReductionKindAttr
- Updated tcgen05.ld.red and nvvm.redux.sync tests
Python bindings for the IRDL dialect were introduced in #158488. They
are currently usable—for constructing IR and dynamically loading modules
that contain `irdl.dialect` into MLIR. However, there are still several
pain points when working with them:
* The IRDL IR-building interface is not very intuitive and tends to be
quite verbose.
* We do not yet have the corresponding `OpView` classes for IRDL-defined
operations.
To address these issues, I propose creating a wrapper (effectively a
small “DSL”) on top of the existing IRDL Python bindings. This wrapper
aims to simplify IR construction and automatically generate the
corresponding `OpView` types. A simple example is shown below.
Currently, using the IRDL bindings looks like this:
```python
m = Module.create()
with InsertionPoint(m.body):
myint = irdl.dialect("myint")
with InsertionPoint(myint.body):
constant = irdl.operation_("constant")
with InsertionPoint(constant.body):
iattr = irdl.base(base_name="#builtin.integer")
i32 = irdl.is_(TypeAttr.get(IntegerType.get_signless(32)))
irdl.attributes_([iattr], ["value"])
irdl.results_([i32], ["cst"], [irdl.Variadicity.single])
add = irdl.operation_("add")
with InsertionPoint(add.body):
i32 = irdl.is_(TypeAttr.get(IntegerType.get_signless(32)))
irdl.operands_(
[i32, i32],
["lhs", "rhs"],
[irdl.Variadicity.single, irdl.Variadicity.single],
)
irdl.results_([i32], ["res"], [irdl.Variadicity.single])
irdl.load_dialects(m)
```
With the proposed DSL (module name `mlir.dialects.ext`), the equivalent
implementation becomes:
```python
class MyInt(Dialect, name="myint"):
pass
i32 = IntegerType[32]
class ConstantOp(MyInt.Operation, name="constant"):
value: IntegerAttr
cst: Result[i32]
class AddOp(MyInt.Operation, name="add"):
lhs: Operand[i32]
rhs: Operand[i32]
res: Result[i32]
MyInt.load()
```
Compared with the current IRDL Python bindings, this DSL mainly adds the
following:
* **A more intuitive interface** for constructing IRDL definitions (as
shown in the example).
* **Automatic generation of the corresponding `OpView`
classes**—including `__init__` methods and property getters for each
defined operation. Similar to TableGen’s `ins`, operands and attributes
can be interleaved in arbitrary order. Special handling is also
implemented for optional and variadic operands/results (such as
computing segment sizes) so that they feel as natural to use as native
operations.
* **Lazy insertion of ops**: all ops are created and inserted only when
`Dialect.load()` is called, which makes it unnecessary to specify an
MLIR context immediately when defining an IRDL dialect.
* **Basic type inference** in operation builders (i.e.
`OpViewCls.__init__`) for trivial result types.
The current DSL does not yet cover all IRDL operations. Several features
are not supported at the moment:
- Defining new types or attributes
- Parametric constraints
- Adding regions to operations
---------
Co-authored-by: Rolf Morel <rolfmorel@gmail.com>
This PR extends the MLIR C API and Python bindings to support
**arbitrary-precision integers (`APInt`)**, overcoming the previous
limitation where `IntegerAttr` values were restricted to 64 bits.
Cryptographic applications often require integer types much larger than
standard machine words (e.g., the 256-bit modulus for the BN254 curve).
Previously, attempting to bind these values resulted in truncation or
errors. This PR exposes the underlying word-based `APInt` structure via
the C API and updates the Python bindings to seamlessly handle Python's
arbitrary-precision integers.
Extend linalg.pack and linalg.unpack to accept memref operands in
addition to tensors. As part of this change, we now disable all
transformations when these ops have memref semantics.
Closes https://github.com/llvm/llvm-project/issues/129004
---------
Signed-off-by: Ryutaro Okada <1015ryu88@gmail.com>
Co-authored-by: Hyunsung Lee <ita9naiwa@gmail.com>
This PR adds "downcasting" of `ir.Value` to either `BlockArgument` or
`OpResult` (and then potentially further down if a user-registered
"value caster" exists). Also this PR changes `__str__` to return the
correct thing (`OpResult(...)` or `BlockArgument(...)` instead of
generic `Value(...)`).
This is mainly for two purposes:
1. to keep it consistent with the C++ class name
`mlir::GreedyRewriteConfig`,
2. to make it shorter.
Since this type was only added a few days ago
(654b3e844f21d3f64521e9cb028efdfebbf99bb4), it shouldn’t cause any
obvious compatibility issues.
