404 Commits

Author SHA1 Message Date
Twice
eb687fb106
[MLIR][Python] Make location optional in Python-defined dialect loading (#186172)
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.
2026-03-13 10:17:26 +08:00
Charitha Saumya
ce227964cc
[mlir][xegpu] Add support for setting order in SetDescLayoutOp and SetOpLayoutAttrOp transform ops. (#184705)
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.
2026-03-09 08:37:28 -07:00
Jakub Kuderski
15e7177f08
[mlir][GPU] Fix double spaces in tests after ODS printer fix. NFC. (#185325)
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>
2026-03-08 18:46:54 -04:00
Jakub Kuderski
4df95b1a56
[mlir][NVVM] Fix double spaces in tests after ODS printer fix. NFC. (#185326)
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>
2026-03-08 18:46:46 -04:00
Erick Ochoa Lopez
613a5c555e
[mlir][vector] Replace OneDimMultiReductionToTwoDim with OneDimMultiReductionToReduction (#184241)
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
2026-03-04 16:13:11 +00:00
Rolf Morel
756d068ead
[MLIR][Python][Transform] Expose PatternDescriptorOpInterface to Python (#184331)
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.
2026-03-04 10:19:59 +00:00
Twice
e95dabef96
[MLIR][Python] Support attribute definitions in Python-defined dialects (#183907)
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.
2026-03-02 09:57:25 +08:00
Twice
361e235986
[MLIR][Python] Support op adaptor for Python-defined operations (#183528)
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.
2026-02-27 09:13:24 +08:00
Adam Siemieniuk
67ac275fee
[mlir][x86] Rename x86vector to x86 (#183311)
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.
2026-02-26 11:21:58 +01:00
Twice
3f6648422c
[MLIR][Python] Fix typeid support for DynamicType and DynamicAttr (#183076)
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.)
2026-02-25 21:58:13 +08:00
Srinivasa Ravi
43dfde4a07
[MLIR][NVVM] Enable result type inference (#181781)
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.
2026-02-25 10:27:25 +05:30
Twice
3b2c1db870
[MLIR][Python] Support type definitions in Python-defined dialects (#182805)
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]]]
```
2026-02-24 10:34:58 +08:00
Twice
de3cefe560
[MLIR][Python] Add C and Python API for mlir::DynamicAttr (#182820)
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.
2026-02-23 17:26:08 +08:00
Twice
059accca53
[MLIR][Python] Add Python and C API of mlir::DynamicType (#182751)
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>
2026-02-23 12:14:56 +08:00
Twice
8542514e5c
[MLIR][Python] Allow passing dialect as a class keyword argument (#182465)
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"):
  ...
```
2026-02-22 18:52:57 +08:00
Erick Ochoa Lopez
eeb6b394c5
[mlir][vector] remove lower_multi_reduction (#182332)
* Removes `ApplyLowerMultiReductionPatternsOp`
(`apply_patterns.vector.lower_multi_reduction`)
* Updates uses of `apply_patterns.vector.lower_multi_reduction` in tests
to use:
  *  reorder_and_expand_multi_reduction_dims
  * multi_reduction_flattening
   * multi_reduction_unrolling
* Removes `populateVectorMultiReductionLoweringPatterns` (unused)
2026-02-20 08:33:24 -05:00
Erick Ochoa Lopez
f33f9a0cf5
[mlir][vector] Add apply_patterns.vector.multi_reduction_unrolling. (#182113)
* 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
2026-02-19 16:19:58 -05:00
Erick Ochoa Lopez
6ec5c1e368
[mlir][vector] Add multi_reduction_flattening (#181244)
* Adds tests for `populateVectorMultiReductionFlatteningPatterns`
* Add apply_patterns.vector.multi_reduction_flattening transform op.

This follows PR #180977. 

Assisted-by: claude
2026-02-18 14:40:24 -05:00
Erick Ochoa Lopez
3cf4156198
[mlir][vector] add ApplyReorderMultiReductionDimsPatternsOp tests (#180977)
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
2026-02-17 16:14:59 -05:00
Tuomas Kärnä
48566b21a4
[MLIR][XeGPU][TransformOps] set_op_layout_attr supports setting anchor layout (#172542)
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.
2026-02-13 07:49:59 +02:00
Rolf Morel
a1d7cda1d7
[MLIR][Python] Impl XOpInterface(s) from Python, with X=Transform and X=MemoryEffects (#176920)
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).
2026-02-12 14:07:10 +00:00
Twice
972aa597de
[MLIR][Python] Make traits declarative in python-defined operations (#180748)
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>
2026-02-11 20:39:58 +08:00
Twice
fccbdcb15a
[MLIR][Python] Support dynamic traits in python-defined dialects (#179705)
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.
2026-02-09 22:26:56 +08:00
Arun Thangamani
fe91384a5b
[mlir][vector] Wrapping populateFlattenVectorTransferPatterns as a transform pass. (#178134)
This PR covers the `mlir::vector::populateFlattenVectorTransferPatterns`
as a transform pass.
2026-02-09 15:10:51 +05:30
Adam Siemieniuk
ba58225a0a
[mlir][x86vector] Python bindings for x86vector dialect (#179958)
Registers python bindings for x86vector dialect and transform ops.
2026-02-05 20:05:17 +01:00
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
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
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
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
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
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
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
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
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
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