This path transforms `vector.contract(A, B, Acc)` into
`vector.contract(A, B, 0) + Acc` to decouple the contraction computation
from the accumulator update for contractions over the loop.
Exposes `OperationEquivalence` through the MLIR C API so callers can compare operations structurally rather than by handle identity.
Assisted by: Claude
Add a lane-level XeGPU operation that re-distributes a subgroup's
fragments across its lanes without changing the element type.
The op takes a 1D vector — the fragment held by one lane — and returns a
fragment of the same type. A pack/unpack mode selects the direction:
viewing the subgroup as an S x N element grid (S = subgroup size, N =
elements per lane), pack moves element j of lane i from logical position
j*S + i to i*N + j, so a lane's elements end up at consecutive
positions; unpack is the reverse.
This implements xegpu.convert_layout semantics at the lane level when
lane_layout is unchanged but lane_data differs.
assisted-by-claude
---------
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Share LLVM dialect metadata materialization through `ModuleTranslation`
so named metadata and metadata-as-value lowering use one conversion
path. Resolve metadata symbol references to functions, globals, aliases,
and ifuncs, and diagnose malformed required metadata before lowering.
Rename `MDFuncAttr` to `MDGlobalValueAttr` to reflect its support for
symbol-backed global values.
Exposes `Value::replaceUsesWithIf` through the MLIR C API, allowing callers to replace only a subset of a value's uses based on a predicate.
Assisted by: Claude
Exposes `OpBuilder::saveInsertionPoint` / `restoreInsertionPoint` through the MLIR C API for `MlirRewriterBase`, continuing the buildout of the rewrite/conversion C bindings.
Assisted by: Claude
Adds initial infrastructure for converting the acc dialect to LLVM,
specifically around generating libacctarget runtime calls. The current
libacctarget APIs are not yet finalized, but the draft proposal can be
found at https://github.com/llvm/llvm-project/pull/197894. This PR adds
codegen for acc init, shutdown, set, and wait.
PTX with memory side effects (stores, atomics, mbarrier operations with
acquire/release semantics) emitted through `nvvm.inline_ptx` lowers to
`llvm.inline_asm` with register constraints only. Without a `~{memory}`
clobber, LLVM may reorder memory accesses across the inline assembly.
There was no way to express the clobber through this op or the
`BasicPtxBuilderInterface` machinery.
This patch adds, as discussed with @grypp:
- A `hasMemoryClobber` interface method (default `false`) on
`BasicPtxBuilderOpInterface`; when it returns `true`, `PtxBuilder`
appends `~{memory}` to the constraints of the generated inline assembly,
after all register constraints and tied indices.
- An opt-in `memory_clobber` boolean attribute (default `false`) on
`nvvm.inline_ptx` exposing this. Existing behavior is unchanged unless
the attribute is set to `true`.
For example:
```mlir
nvvm.inline_ptx "mbarrier.init.b64 [$0], $1;" ro(%p, %c : !llvm.ptr, i32) memory_clobber = true
```
now lowers to inline assembly with constraints `"l,r,~{memory}"`.
Tests: constraint emission with/without the attribute, combination with
predicate (`"l,r,b,~{memory}"`), read-write operands with tied indices
(`"=f,=f,r,r,0,1,~{memory}"`), and the no-operand case (`"~{memory}"`).
The op documentation's first example (mbarrier.init) is updated to show
correct usage.
Currently, only gpu.launch_func supports stream-based asynchronicity. This
patch adds the same mode to gpu.launch, and makes kernel outlining preserve
that when converting to gpu.launch_func.
This patch also documents the stream-based model in the operations description,
which was previously missing.
This patch changes the builders for the two launch operations:
To retain the current behavior for gpu.launch one needs to add a `nullptr` for
the new `Value asyncObject` argument after the `ValueRange asyncDependencies`,
and make sure to pass a valid type for the `Type asyncTokenType` parameter if
dependency-based async behavior is intended.
For gpu.launch_func, both stream-based and dependency-based async models now go
through the same builders, which take all three `Type asyncTokenType,
ValueRange asyncDependencies, Value asyncObject` arguments. To retain
stream-based behavior, pass `nullptr, {}` to the first two parameters, while to
retain dependency-based behavior, pass `nullptr` to the third parameter.
Assisted-by: Claude Code
`arith-expand` expanded `arith.maximumf`/`minimumf`/`maxnumf`/`minnumf`
and the signed/unsigned integer max/min ops into `cmpf`/`cmpi` +
`select` sequences. These ops also have a direct arith-to-llvm lowering
to the `llvm.intr.maximum`/`minimum`/... intrinsics, which are a single
hardware instruction on many targets. Pipelines that run arith-to-llvm
after arith-expand (e.g. the GPU-to-XeVM pipeline) therefore paid a
large, avoidable overhead.
Add `include-min-max-f` (the float ops
maximumf/minimumf/maxnumf/minnumf) and `include-min-max-i` (the
signed/unsigned integer ops maxsi/maxui/minsi/minui) option (default `false`, changing the existing
behavior) that controls whether these min/max ops are expanded. The
min/max converters are factored into new
`populateExpandMinMaxFPatterns` and
`populateExpandMinMaxIPatterns`; the ceil/floor-div and scaling ext/trunc
expansions (which have no LLVM lowering) always run.
---------
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
This commit adds a version of the existing s_buffer_load intrinsic that
more accurately models the memory semantics of the s_buffer_load
instruction, namely that it is, in fact, a memory load.
To preserve the existing behavior that the "nomem" s.buffer.load
intrinsic was using, Clang and MLIR add !invariant.load metadata when
constructing the intrinsic (matching documented requirements on
scalarazable buffer loads) and a late codegen pass adds the metadata
just to be safe.
Tests that were "about" s.buffer.load have been copied to create
versions that use the new intrinsic, as was done for the other
*.ptr.buffer.* operations.
Other tests have been upgraded to use the new intrinsic. This has mainly
resulted in minor instruction ordering changes in prologues, if any
change at all. However, CodeGen/AMDGPU/dagcombine-fma-fmad.ll has seen a
v_fma => v_mad pattern fail to match in one case, I don't know if this
is a real regression.
While I was here, the s_buffer_load => buffer_load fallback has been
updated to preserve cachepolity.
AI disclosure: Primarily AI-written code, but I have at least looked at
and tried to find the worst of the silliness.
---------
Co-authored-by: Codex <codex@openai.com>
This commit removes the "QueryProfileInterface" and
"QueryExtensionInterface" as they are superfluous in most cases, with
TosaComplianceData.h.inc already providing much of the data we need.
Removing these also helps prevent mistakes as observed in:
https://github.com/llvm/llvm-project/pull/206687
This commit removes these interfaces and updates the profile conformance
to extract supported profile/extension information from
TosaComplianceData.h.inc instead.
One factor that isn't duplicated is operations that rely on extensions,
but do not declare any data type support. This is a small set of
operations, such as tosa.custom, tosa.cond_if, tosa.while_loop and
tosa.const_shape, and we can add them to the TosaComplianceData.h.inc
file as needed.
This PR integrates constant bound inference into the existing
`AffineLoopNormalize` pass under a new `useExpensiveMath` option. When
`useExpensiveMath` is enabled (disabled by default due to potential
compilation time overhead), the pass leverages
`ValueBoundsConstraintSet` (uss presburger) analysis to refine dynamic
`affine.for` loop bounds into compile-time constant bounds. RFC:
https://discourse.llvm.org/t/rfc-mlir-enable-dynamic-and-tighter-affine-unrolling-via-valueboundsconstraintset/91055/2
---------
Co-authored-by: Oleksandr "Alex" Zinenko <azinenko@amd.com>
# Recategorize QuantileType as Storage Format, Not Quantization Scheme
## Summary
Reclassify `QuantileType` from inheriting `QuantizedType` to functioning
as a pure storage type.
## Motivation
### QuantileType is a Storage Format, Not a Quantization Scheme
After careful consideration, there was a conclusion reached that
**QuantileType should not be registered as a subtype of QuantizedType**.
The architectural distinction is critical:
- **QuantizedType** represents different quantization schemes (uniform,
per-axis, etc.), semantic operations that map values into a
lower-precision domain
- **QuantileType** is purely a **storage format** that specifies how to
interpret indices within a lookup table (LUT)
QuantileType was never meant to function as a quantization scheme on its
own. Instead, its purpose is to indicate **how to interpret indices
based on the given LUT**. It is an orthogonal concern to quantization
semantics; it describes data organization, not mathematical
transformation.
### Relaxed Type Conditions
The previous implementation restricted the LUT values to float types
only. However, float types are not the only types that can exist inside
a LUT, integer types are valid as well.
This PR removes the type restriction, allowing **both floating-point and
integer-based LUT interpretations** to coexist, better matching actual
use cases.
## Changes
- **Removed**: Incorrect inheritance of `QuantizedType` by
`QuantileType`
- **Relaxed**: Type constraints to accept both float and integer LUT
value types
This pass emits MLIR remarks for all OpenACC data mapping actions.
Intended to be used when all the data actions are committed (typically
right before codegen to call runtime).
The design of the ACC routine lowering supports multiple device side
specializations of an ACC routine function. Earlier in the pipeline, all
calls to the function are expected to refer the original host function
and not their device side specialized versions. Once there is enough
parallelism context, the calls can be changed to point to the
appropriate specialized routine.
The nohost handling in `ACCRoutineLowering` goes against this design.
This change removes the special handling of nohost routines and this
clean up should be done in later passes.
Assisted-by: Codex
SPIR-V has no size-1 vector type, so `convert-math-to-xevm` previously
rejected `math.exp fastmath<fast> : vector<1xf32>` (and the other native
OCL patterns), leaving them as unconverted `math.exp`. These degenerate
size-1 vectors are the result of the XeGPU pipeline distributing and
linearizing larger vectors down to a single element per lane, so in
practice fastmath exps in e.g. flash-attention softmax were never
lowered to `__spirv_ocl_native_exp`, hurting performance.
Handle size-1 vectors by unwrapping them to the scalar element type:
`vector.extract [0]` -> scalar native intrinsic -> `vector.broadcast`
back. `vector::ExtractOp`/`BroadcastOp` are marked legal in the
conversion target so the partial conversion does not roll back, and the
pass now depends on the Vector dialect.
---------
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
- Added FP8 and all six MXFP extension/type combinations
- Updated RESIZE availability and enforced NEAREST_NEIGHBOR for MXFP
- Added positive, negative, version, extension, mixed-type, and verifier
tests.
Signed-off-by: Peng Sun <peng.sun@arm.com>
Represent `default_func_attrs` on `LLVM::InvokeOp` and translate the
collection onto the generated LLVM invoke instruction. Preserve default
call-site attributes when throwing calls are lowered through the LLVM
dialect, including CIR `try_call` operations.
…raints
The IntegerSet constructor of FlatLinearValueConstraints asserts when
flattening fails, which happens for semi-affine integer sets. Since the
failure is signalled from inside a constructor, callers have no way to
detect the unsupported case ahead of time and no way to work around it;
they simply crash.
Add an optional `bool *error` parameter to the constructor. When it is
non-null, hitting the unimplemented case sets `*error` to true and
returns instead of asserting, letting callers bail out gracefully.
`*error` is set to false on success. When the pointer is null the
previous assert behavior is retained, so existing callers are
unaffected.
per
[#210738](https://github.com/llvm/llvm-project/pull/210738#discussion_r3622744153)
Currently, the following crashes on verification with a segfault:
```
func.func @broadcast_rank0_tensor(%arg0: tensor<i32>, %arg1: tensor<32x2xi32>) -> tensor<32x2xi32> {
%0 = linalg.broadcast ins(%arg0 : tensor<i32>) outs(%arg1 : tensor<32x2xi32>) dimensions = [0, 0]
return %0 : tensor<32x2xi32>
}
```
The verifier should reject cases like this where we have a duplicate
dimension. I've also added a couple of tests covering `linalg.broadcast`
on various input scenarios to confirm they are being handled properly.
Rename `insertDbgValueIntrinsic` to `insertDbgValue`, which is
consistent with the other insert helpers like `insertDbgAssign`,
`insertDeclare`, etc.
Also use debug record instead of intrinsic terminology in doc comments.
Adds two missing ops to the WasmSSA dialect.
- `wasmssa.global_set`
- `wasmssa.select`
---------
Co-authored-by: Ferdinand Lemaire <flscminecraft@gmail.com>
This PR adds dynamic `IsIsolatedFromAbove` trait support and exposes it
to Python as `IsIsolatedFromAboveTrait`.
The test verifies that a region with the trait cannot capture an SSA
value from above, while a region without the trait can.
Assisted by Codex/GPT 5.6 Sol (for writing test cases).
Make the `input_layout` attribute of `xegpu.convert_layout` optional.
When it is not set, the effective input layout defaults to
`target_layout`, i.e. the source is assumed to already be in the target
distribution.
assisted-by-Claude
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Block list construction appends each block to all enclosing loops
(O(depth) per block), and each loop owns a separately allocated vector.
Switch to an Euler tour representation: one BlockLayout array per
LoopInfo holding the in-loop blocks in a loop-contiguous reverse
postorder, each loop's list a [begin, end) slice of it, subloop slices
nested inside their parent's. Headers remain first;
`SubLoops`/`TopLevelLoops` orders are unchanged. A few tests observe the
order and are updated. (Previously a loop's list was the function's
reverse postorder restricted to its members, so a subloop's blocks were
interleaved with the parent's own blocks at their RPO positions.)
Mutations first copy a borrowed slice into private storage from
LoopInfo's allocator. `contains(BlockT *)` remains map-based: a
materialized loop's slice is not a membership test. This copy is rare:
across default<O2> over dagcombiner and sqlite3, only 12--14% of loops
built by analyze() are ever mutated; the rest keep the borrowed slice.
AllocateLoop() no longer forwards constructor arguments; analyze() uses
allocateLoop(Header), removing the header-taking loop constructors.
Aided by Claude Fable 5
Relax the create_mem_desc source operand constraint to accept a
statically shaped shared-memory memref of any rank, replacing the
1D/2D-only StaticShared{1,2}DMemRefOf classes with a rank-agnostic
StaticSharedMemRefOf.
Add a verifier requiring the source memref to be contiguous row-major,
update the op documentation, and add valid/invalid lit tests.
assisted-by-claude
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
## [flang][OpenMP] Implement the `error` directive
### Summary
Implements the OpenMP `error` directive (OpenMP 5.1, 2.5.4) in Flang.
The directive prints a message and, with `severity(fatal)`, stops.
Handling depends on the `at` clause:
- **`at(compilation)`** (the default): handled in semantics. Emits a
compile-time warning or error and produces no IR.
- **`at(execution)`**: lowered to a new `omp.error` op, which becomes a
call to `__kmpc_error`.
Before this patch the directive was unimplemented and hit a `not yet
implemented: OmpErrorDirective` error during lowering.
Fixes#204240.
### Details
Changes go through the compiler pipeline.
**Flang semantics**
- `Enter(OmpErrorDirective)` emits the compile-time diagnostic for the
`at(compilation)` form, using the directive's severity and the `MESSAGE`
text when present.
- `Enter(OmpClause::Message)` checks that the `MESSAGE` expression is a
scalar, default-kind `CHARACTER`, which is required to pass it to the
runtime as a `const char*`.
**Flang lowering**
- `genErrorDirective` lowers the `at(execution)` form only.
- A constant message is stored in the `message` attribute. A
non-constant message is copied to a null-terminated buffer
(`genNullTerminatedString`) and passed as the `message_expr` operand,
because `__kmpc_error` expects a NUL-terminated string.
**New `omp.error` op (MLIR OpenMP dialect)**
- Adds a `ClauseSeverity` enum (`warning = 1`, `fatal = 2`, matching the
`__kmpc_error` argument) and `ClauseSeverityAttr`.
- `omp.error` takes a required `severity` attribute and an optional
message, given either as:
- `message` — a constant string attribute, or
- `message_expr` — a pointer to a null-terminated string.
- A verifier rejects providing both `message` and `message_expr`.
**MLIR → LLVM IR translation**
- Translates `omp.error` to `__kmpc_error(ident_t *, i32 severity, ptr
message)`.
**OMPIRBuilder**
- Adds `OpenMPIRBuilder::createError(...)`, which builds the call, maps
`fatal`→`2` and `warning`→`1`, and passes a null pointer when there is
no message.
### Tests
- **Parser** (`error-unparse.f90`): unparse and parse-tree output.
`at(compilation) severity(fatal)` cases that would stop compilation were
moved to `at(execution)`.
- **Semantics** (`Semantics/OpenMP/error.f90`): compilation
warning/error/default-severity cases, plus the new `MESSAGE` type, rank,
and kind errors.
- **Lowering** (`Lower/OpenMP/error.f90`): constant message, no message,
empty message, runtime `message_expr`, and `at(compilation)` cases that
must not produce `omp.error`. The old `Lower/OpenMP/Todo/error.f90` test
is removed.
- **MLIR** (`ops.mlir`, `invalid.mlir`): round-trip of each `omp.error`
form and the both-messages verifier error.
- **LLVM IR** (`openmp-llvm.mlir`): `__kmpc_error` calls for warning and
fatal, constant and runtime messages, and the null-message cases.
Assisted-by: Github Copilot
Fixes#206628 crash by adding an earlier body argument verifier for
`affine.for`. This crash is caused because `LoopLikeOpInterface`
verifier would call `getRegionIterArgs()` and assumed induction var of
`affine.for` exists.
Python-defined dialect loading currently relies on
`Dialect._mlir_module` to infer whether a dialect has already been
loaded. This state belongs to the Python dialect class rather than an
MLIR context.
Consequently, loading the same dialect after switching contexts requires
`reload=True`, while reloading it in a context where it is already
present can hit the operation registration assertion reported in
#210053.
This patch adds `mlirContextGetLoadedDialect` (following
https://github.com/llvm/lighthouse/pull/228#discussion_r3589792891) to
the C API and exposes it as `Context.is_dialect_loaded`.
`Dialect.load()` now queries the active context:
- loading a dialect more than once in the same context raises a
`RuntimeError`;
- loading the same Python-defined dialect in another context succeeds
without a `reload` flag.
The `reload` parameter is semantically wrong (mentioned in
https://github.com/llvm/lighthouse/pull/228#discussion_r3586672116) thus
removed from the API as a small breaking change. Usually, users who are
using `reload` can just remove it and everything will work fine.
Fixes#210053.
Assisted by GPT 5.6 Sol (for writing test cases).
After #207274, libc++ fails __assume_aligned assertion in
__assume_valid_range.
https://lab.llvm.org/buildbot/#/builders/25/builds/18994
`DenseArrayAttrImpl<T>::operator ArrayRef<T>()` is just a cast of
`raw.data()`. `raw` is aligned copy of range from
from BytecodeReader done by
Builtin_DenseArrayRawDataParameter allocator.
However, if range is empty, aligned copying was
omitted, leaving unaligned `ArrayRef<char>`.
The fix is to replace unaligned `ArrayRef<>` with
aligned default constructed.
Added the `class-name-format` option that takes a format string used to
generate the wrapper class name.
This provides greater control over the generated `ClassOp` name. When
multiple identically structured IR files need to be programmatically
lowered with this pass and linked, there will be naming conflicts which
this patch provides a method of resolution for.
`ACCLoopTiling` used to lower an N-dimensional `tile()` clause by
*uncollapsing* the fused loop into a deep nest of single-IV `acc.loop`s
(`tile_1 → … → tile_N → elem_1 → … → elem_N`). This discards the natural
grouping of the tile clause: the 2N single-IV loops no longer express
"these are the tile iterations" and "these are the in-tile iterations"
as collapsible units, and gang/vector end up spread across a deep nest
where only the outermost loop of each group carries the attribute. This
could potentially lead to poor parallelism assignment.
With this MR:
`tileACCLoops` now rewrites the single fused `acc.loop` in-place into
exactly two multi-IV loops, each carrying all the tiled induction
variables:
- a **tile-group** loop whose steps are the original steps scaled by the
tile sizes, and
- a nested **element-group** loop that walks the iterations within one
tile, with upper bounds clamped to `min(origUB, tileStart +
tileExtent)`.
Keeping each group as a single multi-IV `acc.loop` preserves the
tile/element structure as two collapsible units instead of a 2N-deep
single-IV nest, which is a simpler canonical form. The
gang/worker/vector distribution across the two groups is unchanged from
before.
This change also diagnoses `tile` + `collapse` on the same loop because
these clauses give conflicting loop-association counts. The pass now
emits a NYI diagnostic instead of silently dropping the `collapse`
clause. Added a new invalid-input test to verify this.
With the pass no longer uncollapsing, `uncollapseLoops` and the
vector-of-loops `tileACCLoops(SmallVector<LoopOp>&, …)` overload (and
the now-orphaned `createInnerLoop` helper) are deleted, along with their
unit tests. New unit tests are added to cover the surviving fused-loop
overload.
Add convert-vector-to-xegpu (nested under gpu.module) to lower
kernel-side vector ops to XeGPU.
Move convert-vector-to-scf from the pre-GPU common pipeline into the
post-GPU pipeline, so host transfer ops (and any SCF the pass emits) are
lowered to loops before the vector-to-LLVM conversion.
---------
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Python implementations of `MemoryEffectsOpInterface` receive a
`MemoryEffectInstancesList` (added in #176920), but currently have no
generic way to populate it.
This adds the C API for constructing and appending memory effect
instances and exposes it in Python through `MemoryEffect`,
`SideEffectResource`, and `MemoryEffectInstancesList.append`. The tests
use CSE and DCE to verify that the declared effects are observed by MLIR
passes.
Assisted-by: GPT 5.6 Sol
Fixes#203860
In LoopLikeOpInterface::isDefinedOutsideOfLoop default implementation,
value.getParentRegion() can return null during signature conversion
rollbacks when blocks/ops are unlinked. Check for null region to avoid a
segmentation fault.
Also, add a regression test for convert-func-to-llvm with
index-bitwidth=32 on functions with affine.for loops.
Extends `vector.compressstore` + `vector.expandload` to support scalable
vectors and updates relevant tests.
An e2e test for `vector.compressstore` is added. For
`vector.expandload`, we need to wait for QEMU support:
https://github.com/llvm/llvm-project/issues/210942.
This commit refactors the defined TOSA types to combine unranked and
ranked tensor types into a single type `TosaTensorOf`. This helps
simplify the type definitions and allows all tensor types to support
both unranked and ranked tensors.
Builds on the source/target materialization C bindings (#208934) to
expose the 1:N dialect-conversion functionality through the MLIR C API.
Continues the buildout of the dialect-conversion C bindings (follows
#206146 and #206161).
Assisted by: Claude
Continues the buildout of the dialect-conversion C bindings (follows
#206146 and #206161).
- Exposes `TypeConverter::addSourceMaterialization` and
`addTargetMaterialization` through the MLIR C API.
- Introduces `MlirTypeConverterConversionStatus`
(`Success`/`Failure`/`Declined`) and switches
`MlirTypeConverterConversionCallback` to return it, replacing the old
`MlirLogicalResult` + `MlirType{NULL}` dual sentinel. The old convention
could not distinguish the C++ decline (`std::nullopt`, try the next
conversion) and hard-failure (`failure()`, stop) states; the enum maps
cleanly to all three. The Python binding and C API test are updated
accordingly.
Assisted by: Claude
This adds support for the OpenMP interop construct, which enables
interoperability with foreign runtime environments (e.g., CUDA, HIP).
The init, use, and destroy action clauses are lowered from the Fortran
parse tree to MLIR interop ops and translated to the appropriate runtime
calls. The omp_lib module is extended with interop query API interfaces,
property constants, and return code constants. A semantics fix relaxes
the depend/targetsync constraint for use/destroy directives where the
interop type was established at a separate init site. Unsupported depend
clauses and prefer_type selector-list forms emit explicit diagnostics
instead of being silently dropped.
Assisted-by: Claude Opus 4.6
Emit external device symbols for OpenACC declare create and copyin
globals in -gpu=mem:unified; preserve device definitions for
device_resident globals.
