660 Commits

Author SHA1 Message Date
Twice
cb274ea176
[MLIR][Python] Support region in python-defined dialects (#179086)
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>
2026-02-02 22:11:22 +08:00
Twice
f992f9719f
[MLIR][Python] Support dialect conversion in python bindings (#177782)
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>
2026-01-31 12:37:49 +08:00
Pradeep Kumar
e5d6ed68fd
[mlir][NVVM] Add support for tcgen05.ld.red Op (#177330)
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
2026-01-29 11:53:59 +05:30
Ivan Butygin
77ae87ac07
[mlir][python] Add cluster_size to gpu.launch_func python binding (#177811) 2026-01-26 13:21:29 +03:00
Twice
2cc4d45715
[MLIR][Python] Add a DSL for defining dialects in Python bindings (#169045)
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>
2026-01-25 23:08:45 +08:00
Ryan Kim
ac88f7bcd4
[mlir][python] Support Arbitrary Precision Integers in MLIR C API and Python Bindings (#177733)
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.
2026-01-24 23:05:03 -08:00
Ryutaro Okada
4b066c7fff
[mlir][linalg] Extend linalg.pack and linalg.unpack to accept memref (#167675)
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>
2026-01-19 16:42:27 +01:00
Maksim Levental
2c16364d75
[MLIR][Python] add builtin module transform test (#176388)
See https://github.com/llvm/llvm-project/pull/176299
2026-01-16 15:50:58 +00:00
Maksim Levental
619a9c1e81
[mlir][Python] downcast Value to BlockArgument or OpResult (#175264)
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(...)`).
2026-01-12 17:20:53 +00:00
Twice
9bd910dae4
[MLIR][Python] Rename GreedyRewriteDriverConfig to GreedyRewriteConfig (#175409)
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.
2026-01-11 13:51:49 +08:00
Maksim Levental
0e4be262f4
[mlir][Python] fix dialect extensions which bind C types (#175405)
Fix some dialect bindings I missed in https://github.com/llvm/llvm-project/pull/174156 so they don't bind C structs (because that leads to multiple registration in the case when multiple packages are used simultaneously).
2026-01-10 21:24:55 -08:00
MaPePeR
524fde8a4d
[MLIR][Python] Register OpAttributeMap as Mapping for match compatibility (#174292)
This is a continuation of the idea from #174091 to add `match` support
for MLIR containers. In this PR the `OpAttributeMap` container is
registered as a `Mapping`, so be mapped as a "dictionary" in `match`
statements.

For this to work the `get(key, default=None)` method had to be
implemented. Those are pretty much copys of `dunderGetItemNamed` and
`dunderGetItemIndexed` with an added argument and `nb::object` as return
type, because they can now return other types than just `PyAttribute`.
Was unsure if I should refactor this to make `dunderGetItem...` use the
new `getWithDefault...` or if a separate method is preferred. Kept it as
a copy for simplicitys sake for now.

Even though the `OpAttributeMap` supports indexing by `int` and `str`,
Python does not allow to register it as a `Sequence` and a `Mapping` at
the same time. If it is registered as a Sequence it only returns the
attribute names as string, not as `NamedAttribute`. It is technically
possible to also use integer keys for the `dict`-like match, but it
doesn't provide any constraints on the number of attributes, etc., so
probably not recommended.

<details><summary>Example</summary>

```python
from mlir.ir import Context, Module, OpAttributeMap
from collections.abc import Sequence

ctx = Context()
ctx.allow_unregistered_dialects = True
module = Module.parse(
    r"""
"some.op"() { some.attribute = 1 : i8,
                other.attribute = 3.0,
                dependent = "text" } : () -> ()
""",
    ctx,
)
op = module.body.operations[0]

def test(attr):
    match attr:
        case [*args]:
            print("matched a Sequence", args)
        case _:
            print("Didn't match as Sequence")
    match attr:
        case {"some.attribute": a, "other.attribute": b, "dependent": c}:
            print("Matched as Mapping individually", a, b, c)
        case _:
            print("Didn't match a Mapping")
    match attr:
        case {0: a, 1: b}:
            print("Matched as Mapping with 2 int keys", a, b)
        case _:
            print("Didn't match as Mapping with 2 int keys")
print("Registered as Mapping only:")
test(op.attributes)
print("\nAfter additonally registering as Sequence:")
Sequence.register(OpAttributeMap)
test(op.attributes)
```
Output:
```
Registered as Mapping only:
Didn't match as Sequence
Matched as Mapping individually 1 : i8 3.000000e+00 : f64 "text"
Matched as Mapping with 2 int keys NamedAttribute(dependent="text") NamedAttribute(other.attribute=3.000000e+00 : f64)

After additonally registering as Sequence:
matched a Sequence ['dependent', 'other.attribute', 'some.attribute']
Didn't match a Mapping
Didn't match as Mapping with 2 int keys
```
</details>

makslevental Would be great if you could take a look again ❤️

---------

Co-authored-by: Maksim Levental <maksim.levental@gmail.com>
2026-01-10 09:09:05 -08:00
Twice
271a62f838
[MLIR][Python] Add attr_name for FloatAttr (#175306)
After #174756 I found that attribute name for `FloatAttr` is missing.
And this PR is to add it.

This is actually part of changes in #169045, but I think that we can
make it a separate PR to make #169045 easier to review.
2026-01-10 22:57:29 +08:00
Maksim Levental
94a95659c2
[MLIR][Python] Add GreedyRewriteDriverConfig parameter to apply_patterns_and_fold_greedily (#174913)
We already have `GreedyRewriteDriverConfig` on the Python side, but it
hasn’t yet been exposed as a parameter of
`apply_patterns_and_fold_greedily`. This PR does that.

Before:
```python
def apply_patterns_and_fold_greedily(module: ir.Module, set: FrozenRewritePatternSet) -> None
def apply_patterns_and_fold_greedily(op: ir._OperationBase, set: FrozenRewritePatternSet) -> None
```

After:
```python
def apply_patterns_and_fold_greedily(module: ir.Module, set: FrozenRewritePatternSet,
                                     config: GreedyRewriteDriverConfig | None = None) -> None
def apply_patterns_and_fold_greedily(op: ir._OperationBase, set: FrozenRewritePatternSet,
                                     config: GreedyRewriteDriverConfig | None = None) -> None
```

Note this PR is adapted from
https://github.com/llvm/llvm-project/pull/174785 but using
`std::optional` instead of `nb::object`. Note, this required refactoring
`PyGreedyRewriteDriverConfig` to have a `std::shared_ptr` so that it
could support a copy-ctor.

Co-authored-by: PragmaTwice <twice@apache.org>
2026-01-08 03:48:30 -08:00
Oleksandr "Alex" Zinenko
5958842aee
[mlir][py] ability to downcast AffineExpr after #172892 (#174808)
AffineExpr is a separate hierarchy of LLVM-style nested classes that
doesn't rely on TypeID and is not extensible. We need the ability to
downcast the Python equivalent of those to a specific subclass that was
seemingly lost in PR #172892. Bring it back by having an explicit cast.
We don't really need user-defined type casters here since AffineExpr is
entirely closed and not typed, unlike values.
2026-01-07 17:40:27 +00:00
Twice
ae4bbd0ec6
[MLIR][Python] Forward the name of MLIR attrs to Python side (#174756)
This PR is quite similiar to #174700.

In this PR, I added a C API for each (upstream) MLIR attributes to
retrieve its name (for example, `StringAttr -> mlirStringAttrGetName()
-> "builtin.string"`), and exposed a corresponding type_name class
attribute in the Python bindings (e.g., `StringAttr.attr_name ->
"builtin.string"`). This can be used in various places to avoid
hard-coded strings, such as eliminating the manual string in
`irdl.base("#builtin.string")`.

Note that parts of this PR (mainly mechanical changes) were produced via
GitHub Copilot and GPT-5.2. I have manually reviewed the changes and
verified them with tests to ensure correctness.
2026-01-07 22:57:14 +08:00
Twice
b919d62eae
[MLIR][Python] Forward the name of MLIR types to Python side (#174700)
In this PR, I added a C API for each (upstream) MLIR type to retrieve
its type name (for example, `IntegerType` -> `mlirIntegerTypeGetName()`
-> `"builtin.integer"`), and exposed a corresponding `type_name` class
attribute in the Python bindings (e.g., `IntegerType.type_name` ->
`"builtin.integer"`). This can be used in various places to avoid
hard-coded strings, such as eliminating the manual string in
`irdl.base("!builtin.integer")`.

Note that parts of this PR (mainly mechanical changes) were produced via
GitHub Copilot and GPT-5.2. I have manually reviewed the changes and
verified them with tests to ensure correctness.
2026-01-07 16:27:31 +08:00
Twice
c1de1543bf
[MLIR][Python] Add a .get method to IntegerType (#174406)
In this PR, I added a `.get` class method to `IntegerType`. The main
goal is to ensure that types from upstream dialects have a `.get` method
(at least for the builtin dialect). The benefit is that, for any MLIR
type, we can construct an instance directly without special-casing types
that don’t provide a `.get` method.

The design mirrors `mlir::IntegerType` in C++: it takes `width` and
`signedness` parameters, and `signedness` defaults to `signless`.

It is related to #169045.
2026-01-06 11:57:39 +08:00
Maksim Levental
6021ed572f
[mlir][Python] use maybeDowncast for PyType/PyAttribute returns after #174156 (#174489)
#174156 made all gettors return `Py*` but skipped downcasting where
possible. So restore it by calling `.maybeDowncast`.
2026-01-05 22:39:38 +00:00
Maksim Levental
fb8bbd4ed8
[mlir][Python] use canonical Python isinstance instead of Type.isinstance (#172892)
We've been able to do `isinstance(x, Type)` for a quite a while now
(since
bfb1ba7526)
so remove `Type.isinstance` and the the special-casing
(`_is_integer_type`, `_is_floating_point_type`, `_is_index_type`) in
some places (and therefore support various `fp8`, `fp6`, `fp4` types).
2026-01-05 21:07:24 +00:00
Maksim Levental
ee3338d135
[mlir][Python] port in-tree dialect extensions to use MLIRPythonSupport (#174156)
This PR ports all in-tree dialect extensions to use the
`PyConcreteType`, `PyConcreteAttribute` CRTPs instead of
`mlir_pure_subclass`. After this PR we can soft deprecate
`mlir_pure_subclass`. Also API signatures are updated to use `Py*`
instead of `Mlir*` so that type "inference" and hints are improved.
2026-01-05 10:23:22 -08:00
Maksim Levental
18fc908566
[mlir][Python] move IRTypes and IRAttributes to MLIRPythonSupport (#174118)
This PR continues the work of
https://github.com/llvm/llvm-project/pull/171775 by moving more useful
types/attributes into MLIRPythonSupport.

You can now do 

```c++
struct PyTestIntegerRankedTensorType
    : mlir::python::MLIR_BINDINGS_PYTHON_DOMAIN::PyConcreteType<
          PyTestIntegerRankedTensorType,
          mlir::python::MLIR_BINDINGS_PYTHON_DOMAIN::PyRankedTensorType>
struct PyTestTensorValue
    : mlir::python::MLIR_BINDINGS_PYTHON_DOMAIN::PyConcreteValue<
          PyTestTensorValue>
```
instead of `mlir_type_subclass` and `mlir_value_subclass`;
**specifically manual registration of the "value caster" via indirection
through the Python interpreter is no longer necessary** . You can also
now freely use all such types at the nanobind API level (e.g., overload
based on `FP*`):

```c++
using mlir::python::MLIR_BINDINGS_PYTHON_DOMAIN;
standaloneM.def("print_fp_type", [](PyF16Type &) { nb::print("this is a fp16 type"); });
standaloneM.def("print_fp_type", [](PyF32Type &) { nb::print("this is a fp32 type"); });
standaloneM.def("print_fp_type", [](PyF64Type &) { nb::print("this is a fp64 type"); });
```

Note, here we only port `PythonTestModuleNanobind` but there is a
follow-up PR that ports **all** in-tree dialect extensions
https://github.com/llvm/llvm-project/pull/174156 to use these. After
that one we can soft deprecate `mlir_pure_subclass`.

Note, depends on https://github.com/llvm/llvm-project/pull/171775
2026-01-05 09:34:58 -08:00
Maksim Levental
f0ef5dba6d
[mlir][Python] create MLIRPythonSupport (#171775)
# What

This PR adds a shared library `MLIRPythonSupport` which contains all of
the CRTP classes ike `PyConcreteValue`, `PyConcreteType`,
`PyConcreteAttribute`, as well as other useful code like `Defaulting*`
and etc enabling their reuse in downstream projects. Downstream projects
can now do

```c++
struct PyTestType : mlir::python::MLIR_BINDINGS_PYTHON_DOMAIN::PyConcreteType<PyTestType> {
  ...
};

class PyTestAttr : public mlir::python::MLIR_BINDINGS_PYTHON_DOMAIN::PyConcreteAttribute<PyTestAttr> {
  ...
}

NB_MODULE(_mlirPythonTestNanobind, m) {
  PyTestType::bind(m);
  PyTestAttr::bind(m);
}
```

instead of using the discordant alternative
`mlir_type_subclass`/`mlir_attr_subclass` (same goes for
`PyConcreteValue`/`mlir_value_subclass`).

# Why

This PR is mostly code motion (along with CMake) but before I describe
the changes I want to state the goals/benefits:

1. Currently upstream "core" extensions and "dialect" extensions ([all
of the `Dialect*` extensions
here](d7c734b5a1/mlir/lib/Bindings/Python))
are a two-tier system;
**a**. [core
extensions](https://github.com/llvm/llvm-project/blob/main/mlir/lib/Bindings/Python/IRTypes.cpp#L361)
enjoy first class support as far as type inference[^3], type stub
generation, and ease of implementation, while dialect extensions [have
poorer support](https://reviews.llvm.org/D150927), incorrect type stub
generation much more tedious (boilerplate) implementation;
**b**. Crucially, this two-tiered system is reflected in the fact that
**the two sets of types/attributes are not in the same Python object
hierarchy**. To wit: `isinstance(..., Type)` and `isinstance(...,
Attribute)` are not supported for the dialect extensions[^2];
**c**. Since these types are not exposed in public headers, downstream
users (dialect extensions or not) cannot write functions that overload
on e.g. `PyFloat8*Type` - that's quite a [useful
feature](fdbee98df8/cpp_ext/TorchOps.cpp (L29-L69))!
2. The dialect extensions incur a sizeable performance penalty relative
to the core extensions in that every single trip across the wire (either
`python->cpp` or `cpp->python`) requires work in addition to nanobind's
own casting/construction pipeline;
**a**. When going from `python->cpp`, [we extract the capsule object
from the Python
object](https://github.com/llvm/llvm-project/blob/main/mlir/include/mlir/Bindings/Python/NanobindAdaptors.h#L219C24-L219C46)
and then extract from the capsule the `Mlir*` opaque struct/ptr. This
side isn't so onerous;
**b**. When going from `cpp->python` we call long-hand call Python
`import` APIs and construct the Python object using `_CAPICreate`. Note,
there at least 2 `attr` calls incurred in addition to `_CAPICreate`;
this is already much more [efficiently handled by nanobind
itself](4ba51fcf79/src/nb_internals.h (L381-L382))!
3. This division blocks various features: in some configurations[^1] we
trigger a circular import bug because "dialect" types and attributes
perform an [import of the root `_mlir`
module](bd9651bf78/mlir/include/mlir/Bindings/Python/NanobindAdaptors.h (L585))
when they are created (the types themselves, not even instances of those
types). This blocks type stub generation for dialect extensions (i.e.,
the reason we currently only generate type stubs for `_mlir`).

# How

Prior this was not done/possible because of "ODR" issues but I have
resolved those issues; the basic idea for how we solve this is "move
things we want to share into shared libraries":

1. Move IRCore (stuff like `PyConcreteValue`, `PyConcreteType`,
`PyConcreteAttribute`) into `MLIRPythonSupport`;
- Note, we move the rest of the things in `IRModule.h` (renamed to
`IRCore.h`) because `PyConcreteValue`, `PyConcreteType`,
`PyConcreteAttribute` depend on them. This makes for a bigger PR than
one would hope for but ultimately I think we should give people access
to these classes to use as they see fit (specifically inherit from, but
also liberally use in bindings signatures instead of the opaque `Mlir*`
struct wrappers).
2. Put all of this code into a nested namespace
`MLIR_BINDINGS_PYTHON_DOMAIN` which is determined by a compile time
define (and tied to `MLIR_BINDINGS_PYTHON_NB_DOMAIN`). This is necessary
in order to prevent conflicts on both symbol name **and** typeid
(necessary for nanobind to not double register binded types) between
multiple bindings libraries (e.g., `torch-mlir`, and `jax`). Note
[nanobind doesn't support `module_local` like
pybind11](https://nanobind.readthedocs.io/en/latest/porting.html#removed-features).
It does support `NB_DOMAIN` but that is not sufficient for
disambiguating typeids across projects (to wit: we currently define
`NB_DOMAIN` and it was still necessary to move everything to a nested
namespace);
3. Build the [nanobind library itself as a shared
object](https://github.com/wjakob/nanobind/blob/master/cmake/nanobind-config.cmake#L127)
(and link it to both the extensions and `MLIRPythonSupport`).
4. CMake to make this work, in-tree, out-of-tree, downstream, upstream,
etc.

# Testing

Three tests are added here 

1. `PythonTestModuleNanobind` is ported to use
`PyConcreteType<PyTestType>` instead of `mlir_type_subclass` and
`PyConcreteAttribute<PyTestAttr>` instead of `mlir_atrr_subclass`,
verifying this works for non-core extensions in-tree;
2. `StandaloneExtensionNanobind` is ported to use `struct PyCustomType :
mlir::python::MLIR_BINDINGS_PYTHON_DOMAIN::PyConcreteType<PyCustomType>`
instead of `mlir_type_subclass` verifying this works for non-core
extensions out-of-tree;
3. `StandaloneExtensionNanobind`'s `smoketest` is extended to also load
another bindings package (namely `mlir`) verifying
`MLIR_BINDINGS_PYTHON_DOMAIN` successfully disambiguates symbols and
typeids.

I have also tested this downstream:
https://github.com/llvm/eudsl/pull/287 as well run the following builder
bots:

mlir-nvidia-gcc7:
https://lab.llvm.org/buildbot/#/buildrequests/6654424?redirect_to_build=true

I have also tested against IREE:
https://github.com/iree-org/iree/pull/21916

# Integration

It is highly recommended to set the CMake var
`MLIR_BINDINGS_PYTHON_NB_DOMAIN` (which will also determine
`MLIR_BINDINGS_PYTHON_DOMAIN`) to something unique for each downstream.
This can also be passed explicitly to `add_mlir_python_modules` if your
project builds multiple bindings packages. I added a `WARNING` to this
effect in `AddMLIRPython.cmake`.

[^3]: Python values being typed correctly when exiting from cpp;
[^1]: Specifically when the modules are imported using `importlib`,
which occurs with nanobind's
[stubgen](https://github.com/wjakob/nanobind/blob/master/src/stubgen.py#L965);
[^2]: The workaround we implemented was a class method for the dialect
bindings called `Class.isinstance(...)`;
2026-01-05 09:08:13 -08:00
MaPePeR
6d8dd3da4b
[MLIR][Python] Register Containers as Sequences for match compatibility (#174091)
This allows these containers to be used in `match` statements, which
allows extracting properties and asserting a shape at the same time.

It seems to be only possible, to match as _either_ a `Mapping` _or_ a
`Sequence`, so the `OpAttributeMap` is only a `Mapping`.

I couldn't find a way to make these C++ based types properly inherit
from `Sequence` or `Mapping`, so the Mixins are not provided (nanobind
only allows C++ parent classes, modifying `__base__` complains about
differing destructors).
`OpAttributeMap` was lacking the `get` method, so I simply copied it
from `collections.abc.Mapping`.

When writing the tests i ran into the error, that I wrote
`func.FuncOp(body=[Block(...)])` instead of
`func.FuncOp(body=Region(blocks=[Block(...)]))`. So maybe also turning
`Region` itself into a Sequence would be a good addition as well? Would
extend the Scope of this PR, though.

makslevental You suggested I make the PR, so i'm tagging you here as a
potential reviewer. I hope that is ok with you. :)

---------

Co-authored-by: Maksim Levental <maksim.levental@gmail.com>
2026-01-03 09:56:24 -08:00
Jacques Pienaar
654b3e844f
[mlir][c] Enable creating and setting greedy rewrite confing. (#162429)
Done very mechanically.

This changes that one cannot just pass null config to C API for config.
2026-01-02 06:12:30 +00:00
Akimasa Watanuki
ebb1c27198
[mlir][linalg] Reject unsigned pooling on non-integer element types (#166070)
Fixes: #164800 

Ensures unsigned pooling ops in Linalg stay in the integer domain: the
lowering now rejects floating/bool inputs with a clear diagnostic, new
regression tests lock in both the error path and a valid integer
example, and transform decompositions are updated to reflect the integer
typing.

Signed-off-by: Akimasa Watanuki <mencotton0410@gmail.com>
2026-01-01 13:04:41 +05:30
Twice
b6883607c0
[MLIR][Python] Refine the support of RewritePatternSet.add (#173874)
This patch includes the following changes:
- `RewritePatternSet.add` now accepts op name (e.g. `.add("arith.addi",
fn)`) besides op class (e.g. `.add(arith.AddIOp, fn)`)
- add a concrete signature and a more complete docstring to
`RewritePatternSet.add`.
2025-12-30 10:16:27 +08:00
Twice
3ed1e9c85d
[MLIR][Python] Add support of the walk pattern rewrite driver (#173562)
MLIR currently has three main pattern rewrite drivers (see
[https://mlir.llvm.org/docs/PatternRewriter/#common-pattern-drivers](https://mlir.llvm.org/docs/PatternRewriter/#common-pattern-drivers)):

* Dialect Conversion Driver
* Walk Pattern Rewrite Driver
* Greedy Pattern Rewrite Driver

Right now, we already support the greedy pattern rewrite driver in the C
API and Python bindings. This PR adds support for the walk pattern
rewrite driver. This lightweight driver, unlike the greedy driver, does
not repeatedly apply patterns; instead, it walks the IR once. API-wise,
the main change is adding the `walk_and_apply_patterns` function.

Note that the listener parameter is not supported now.
2025-12-26 16:11:06 +08:00
Hongzheng Chen
177072a763
[MLIR][Python] Update the scf.if interface to be consistent with affine.if (#173171)
This is a follow-up of #171957 that updates the argument names of
`scf.if` Python binding to be consistent with `affine.if`. Basically,
both operations should use `has_else` to determine whether the `if`
block is presented.

cc @makslevental
2025-12-20 21:33:37 -08:00
Maksim Levental
3d7018c70b
[MLIR][Python] remove pybind11 support (#172581)
This PR removes pybind which has been deprecated for over a year
(https://github.com/llvm/llvm-project/pull/117922).
2025-12-19 09:51:22 -08:00
Rolf Morel
79ed37ca92
[MLIR][Transform] Fix transform.smt.constrain_params's verifier (#172753)
Verifier was insisting on `!transform.param<...>` too early and hence
crashed on `!transform.any_param`.
2025-12-17 23:32:32 +00:00
Tim Gymnich
16f41cb1b8
[mlir][amdgpu] Add Python bindings for TDM types (#172309)
Add bindings for:
- `TDMBaseType`
- `TDMDescriptorType`
- `TDMGatherBaseType`
2025-12-16 10:43:08 +00:00
Rolf Morel
f12fcf030c
[MLIR][Transform][Python] transform.foreach wrapper and .owner OpViews (#172228)
Friendlier wrapper for transform.foreach.

To facilitate that friendliness, makes it so that OpResult.owner returns
the relevant OpView instead of Operation. For good measure, also changes
Value.owner to return OpView instead of Operation, thereby ensuring
consistency. That is, makes it is so that all op-returning .owner
accessors return OpView (and thereby give access to all goodies
available on registered OpViews.)

Reland of #171544 due to fixup for integration test.
2025-12-14 22:10:31 +00:00
Mehdi Amini
b9fe6532a7
Revert "[MLIR][Transform][Python] transform.foreach wrapper and .owner OpViews" (#172225)
Reverts llvm/llvm-project#171544 ; bots are broken.
2025-12-14 21:27:02 +00:00
Rolf Morel
4cdec92827
[MLIR][Transform][Python] transform.foreach wrapper and .owner OpViews (#171544)
Friendlier wrapper for `transform.foreach`.

To facilitate that friendliness, makes it so that `OpResult.owner`
returns the relevant `OpView` instead of `Operation`. For good measure,
also changes `Value.owner` to return `OpView` instead of `Operation`,
thereby ensuring consistency. That is, makes it is so that all
op-returning `.owner` accessors return `OpView` (and thereby give access
to all goodies available on registered `OpView`s.)
2025-12-14 20:44:15 +00:00
Hongzheng Chen
1335a05ab8
[MLIR][Python] Fix AffineIfOp insertion point (#171957)
This bug was introduced by #108323, where the loc and ip were not
properly set. It may lead to errors when the operations are not linearly
asserted to the IR.
2025-12-11 20:23:46 -08:00
Hongzheng Chen
86cc934b4a
[python] Expose replaceUsesOfWith C API (#171892)
This PR exposes the `replaceUsesOfWith` C API to Python
2025-12-11 16:09:18 -08:00
Oleksandr "Alex" Zinenko
bd9651bf78
[mlir][py] avoid crashing on None contexts in custom gets (#171140)
Following a series of refactorings, MLIR Python bindings would crash if
a
dialect object requiring a context defined using
mlir_attribute/type_subclass
was constructed outside of the `ir.Context` context manager. The type
caster
for `MlirContext` would try using `ir.Context.current` when the default
`None`
value was provided to the `get`, which would also just return `None`.
The
caster would then attempt to obtain the MLIR capsule for that `None`,
fail,
but access it anyway without checking, leading to a C++ assertion
failure or
segfault.

Guard against this case in nanobind adaptors. Also emit a warning to the
user
to clarify expectations, as the default message confusingly says that
`None` is
accepted as context and then fails with a type error. Using Python C API
is
currently recommended by nanobind in this case since the surrounding
function
must be marked `noexcept`.

The corresponding test is in the PDL dialect since it is where I first
observed
the behavior. Core types are not using the `mlir_type_subclass`
mechanism and
are immune to the problem, so cannot be used for checking.
2025-12-10 08:37:45 +01:00
Tianqi Chen
11fd760e3a
[MLIR][ExecutionEngine] Enable PIC option (#170995)
This PR enables the MLIR execution engine to dump object file as PIC
code, which is needed when the object file is later bundled into a dynamic
shared library.

---------

Co-authored-by: Mehdi Amini <joker.eph@gmail.com>
2025-12-07 17:07:17 +00:00
Maksim Levental
e6d1deae3e
[MLIR][Python] make Sliceable inherit from Sequence (#170551)
Generates type stubs like

```python
class RegionSequence(Sequence[Region]):
    def __add__(self, arg: RegionSequence, /) -> list[Region]: ...

    def __iter__(self) -> RegionIterator:
        """Returns an iterator over the regions in the sequence."""
```
2025-12-05 05:41:55 -08:00
James Molloy
d17284ae29
[python] Fix loc_tracebacks() (#170831)
There were two bugs lurking in mlir.ir.loc_tracebacks():
  1) The default None parameter was not handled correctly (passed to a
     C++ function that expects ints.
  2) The `yield` was incorrectly indented meaning loc_tracebacks()
     could not be nested (a "generator didn't yield" exception would be
     raised).

Added testing of loc_tracebacks by replacing the custom contextmanager
in the auto_location.py test with the loc_tracebacks() API.

Had to harden the test to line number differences.

---------

Co-authored-by: James Molloy <jmolloy@google.com>
2025-12-05 11:58:40 +00:00
Maksim Levental
772ff0436d
[mlir][ExecutionEngine] propagate errors in mlirExecutionEngineCreate (#170592) 2025-12-04 23:14:05 +00:00
Jakub Kuderski
0bd2f12753
[mlir][linalg] Restrict fill initial value type to output element type (#169567)
Disallow implicit casting, which is surprising, and, IME, usually
indicative of copy-paste errors.

Because the initial value must be a scalar, I don't expect this to
affect any data movement.
2025-11-30 09:51:37 -05:00
Benjamin Chetioui
012721d320
[mlir][python] Propagate error diagnostics when an op couldn't be created. (#169499) 2025-11-25 17:41:01 +00:00
Maksim Levental
740d0bd385
[MLIR][Python] add GetTypeID for llvm.struct_type and llvm.ptr and enable downcasting (#169383) 2025-11-24 18:39:15 +00:00
Jacques Pienaar
5ab49edde2
[mlir][py][c] Enable setting block arg locations. (#169033)
This enables changing the location of a block argument. Follows the
approach for updating type of block arg.
2025-11-21 13:31:46 +00:00
Tuomas Kärnä
e9fc393a9e
[MLIR][XeGPU][TransformOps] Add slice_dims argument to set_op_layout_attr and set_desc_layout (#168929)
`set_op_layout_attr` and `set_desc_layout` transform ops wrap
`xegpu.layout` in an `xegpu.slice` attribute if `slice_dims` argument is
set.
2025-11-21 10:08:12 +02:00
Asher Mancinelli
47d9d735a7
[MLIR][Python] Add arg_attrs and res_attrs to gpu func (#168475)
I missed these attributes when I added the wrapper for GPUFuncOp in
fbdd98f74f0d.
2025-11-18 07:55:11 -08:00
Rolf Morel
eb9d56cb55
[MLIR][Transform][Python] Expose applying named_sequences as a method (#168223)
Makes it so that a NamedSequenceOp can be directly applied to a Module,
via a method `apply(...)`.
2025-11-15 18:31:17 +00:00
Asher Mancinelli
a4e7d150ea
[MLIR][Python] Add tests for nvvm barrier ops (#167976)
Found this issue #167958 when adding these tests, thanks for the quick
fix @clementval.
2025-11-15 08:57:41 -08:00