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.
Reject distinct nodes when converting generic LLVM metadata into LLVM
dialect attributes because `MDNodeAttr` cannot preserve node identity.
Apply the policy recursively so top-level and nested nodes use the
existing unsupported-metadata failure path while finite acyclic uniqued
nodes remain importable.
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
This PR fixes a crash in ConvertShardToMPI caused by ignoring a failed partial conversion and continuing to apply folding patterns to partially converted IR.
Assisted-by: codex
---------
Co-authored-by: Maksim Levental <maksim.levental@gmail.com>
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.
The im2col patterns for conv_2d_nhwc_hwcf, conv_2d_nchw_fchw and
conv_2d_nhwc_fhwc avoided non-unit dilations, but the restriction
doesn't seem to be fundamental: dilation only affects the gather step,
which can fold it into its indexing map. Generalize the convolved index
expression from `oh * stride + fh` to `oh * stride + fh * dilation` and
drop the match failures.
Verified by comparing the results of a dilated convolution lowered
directly to loops against the im2col decomposition with mlir-runner.
Assisted-By: Claude Code (for generating tests).
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.
Example:
```fortran
!$acc parallel loop gang reduction(+:a)
do i = 1, N
!$acc loop worker reduction(+:a)
do j = 1, M
a(i) = a(i) + b(j,i)
end do
end do
```
In this code, the worker accumulator is block-shared, so every worker
updates
the same element with a plain read-modify-write and all but one partial
is lost.
Fix: make in-place updates of a block-shared array accumulator atomic,
unless
the element index varies per thread, in which case each thread owns its
element
and the plain update is already race-free.
---------
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
This PR aims to fix the misleading implementation of `classof` which
accepted as a `QuantizedType` every type registered under the
`QuantDialect.`
The now updated implementation makes the checks more restrictive,
allowing only direct subtypes of `QuantizedType` to be recognized as a
`QuantizedType,` therefore preventing `QuantileType` and other types
that are registered under `QuantDialect` to be seen as a subtype of
`QuantizedType` and wrongfully use its methods which might lead to
potential crashes.
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 PR fixes an issue in getDpasSubgroupLayouts(): It failed to find
valid subgroup layouts for otherwise-legal DPAS ops. It treated the K
(contraction) dimension as distributed across subgroups — requiring
wgShape % sgLayout == 0 on every dim and gating on a
checkAlignedSgDataAB() equality — when K is actually broadcast: its full
extent stays in every subgroup. This rejected the only valid candidate,
causing layout propagation to bail out.
Update the two affected tests, which now propagate successfully instead
of failing.
assisted-by-claude
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 PR enhance `setupInsertStridedSliceResultLayout` to treat a size-1
source dim as a broadcast dim: keep its lane_data at 1 instead of
requiring divisibility.
Add regression tests for both the layout propagation setup and the
subgroup-to-lane distribution of such an op.
Assisted-by-claude
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
The TOSA specifications indicates that when computing the imaginary
output for locations (0,0), (0,W/2), (H/2,0), (H/2, W/2),
implementations may choose to skip the computation entirely and return
0, or compute each accumulation term as `-val_real * 0.0`.
The previous legalizations used `-val_real * sin(a)` but `sin(a)` may
not return exactly 0.0 due to representation error in the input to
`sin`.
This patch updates the legalization to instead check for indices where
`sin(a)` should be replaced by `0.0` and uses a `select` op to match the
expected behaviour according to the Specification when
tosa_extra_multiplies is set to `true`.
Signed-off-by: Ian Tayler Lessa <ian.taylerlessa@arm.com>
`[IntrNoMem, IntrHasSideEffects]` is a perfectly valid intrinsic trait
combination, and MLIR would CSE/DCE those without checking. This mirrors
how LLVM models this, see getEffectiveME in IntrinsicEmitter.cpp
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 patch fixes a crash in `transform.apply_patterns` when `apply_cse`
is enabled and the target operation does not have the
`IsolatedFromAbove` trait.
For non-`IsolatedFromAbove` targets, `transform.apply_patterns` cannot
apply patterns directly to the target operation. Instead, it collects
the operations within the target into a worklist and passes that
worklist to `applyOpPatternsGreedily`. When `apply_cse` is enabled,
pattern application and CSE are performed repeatedly until a fixpoint is
reached.
Previously, the worklist was collected only once before entering the
fixpoint loop. However, pattern application may erase operations from
the IR, leaving stale pointers in the cached worklist. Reusing this
worklist in subsequent iterations may therefore result in a crash.
This patch rebuilds the worklist at the beginning of each iteration,
ensuring that it always reflects the current IR state.
A representative reproducer is shown below:
```llvm
module attributes {transform.with_named_sequence} {
func.func @test(%A: tensor<128x256xf32>,
%B: tensor<256x512xf32>,
%C: tensor<128x512xf32>) -> tensor<128x512xf32> {
%res = linalg.matmul
ins(%A, %B : tensor<128x256xf32>, tensor<256x512xf32>)
outs(%C : tensor<128x512xf32>) -> tensor<128x512xf32>
return %res : tensor<128x512xf32>
}
transform.named_sequence @__transform_main(%root: !transform.any_op) {
%matmul = transform.structured.match
ops{["linalg.matmul"]}
in %root
: (!transform.any_op) -> !transform.any_op
%tiled, %loops:3 =
transform.structured.tile_using_for %matmul
tile_sizes [4, 4, 4]
: (!transform.any_op)
-> (!transform.any_op,
!transform.any_op,
!transform.any_op,
!transform.any_op)
transform.structured.vectorize %tiled : !transform.any_op
transform.apply_patterns to %loops#0 {
transform.apply_patterns.canonicalization
} {apply_cse} : !transform.any_op
transform.yield
}
}
```
Here, `transform.apply_patterns` is applied to the `scf.for` loop
produced by tiling, which is not `IsolatedFromAbove`. Running `mlir-opt
--transform-interpreter` on this example crashes before this change and
completes successfully with this patch.
Signed-off-by: Mikhail Romanov <mmromanov@ispras.ru>
The CancellingBlockScaledCastsOptimization optimises away sequences
where a block-scaled tensor is cast to fp32 and then cast back to its
original block-scaled type.
This patch extends it to also cover cases where the intermediate type is
bf16, since the 7 bits of mantissa and 8 bits in the exponent are
sufficient to encode any `element * scale` product with valid
block-scaled type combinations.
---------
Signed-off-by: Ian Tayler Lessa <ian.taylerlessa@arm.com>
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>
Add missing `| FileCheck %s` to `RUN` lines so `CHECK` lines can be
verified. Also correct the expected result from `30` to `36`. The result
should be `6*6=36`.
The example in the tutorial is correct (`3.600000e+01`).
https://mlir.llvm.org/docs/Tutorials/Toy/Ch-6/
Signed-off-by: Felix Ye <felixyjs@gmail.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
Commit 6c51938067 ([MLIR] Guard optional operand resolution in
generated op parsers) guarded resolution of every optional operand to
avoid indexing an empty inferred type range. The guard also skipped
resolveOperands when an optional operand was absent but an independently
optional type directive was present. This accepted malformed assembly
and silently discarded the orphan type when printing the operation.
Keep the guard for inferred type resolutions, which may access an empty
optional type vector, but always resolve directly parsed type ranges so
resolveOperands validates operand/type cardinality. Add coverage for
both valid forms and a type-only negative reproducer.
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>
Follow the same pattern as all the other test suites where the PATHS
option to configure_lit_site_cfg and the path() function in the lit site
config are used to make site configs relocatable.
`setPropertiesFromParsedAttr` (generated by `OpFormatGen.cpp` for ops
using a custom `assemblyFormat`) rejects the trait-injected
`operandSegmentSizes` /. `resultSegmentSizes` properties when they
appear in a `prop-dict`, even though the printer emits them there for ops
whose format uses a bulk operand/result type directive
(`type(operands)`, `type(results)`, `functional-type(operands, results)`).
This breaks round-tripping: any op combining `AttrSizedOperandSegments`
/ `AttrSizedResultSegments` with such a bulk directive fails to re-parse
its own printed output, because printer-side elision of these keys
(introduced in #115930) is intentionally skipped in that case — the
sizes can't be reconstructed from individually-typed operand/result
groups, so they must survive in the text, but the custom parser was never
taught to read them back.
`setPropertiesFromAttr` (used by the generic-form parser, bytecode, and
C++ construction) already special-cases these two keys, so this change
makes `setPropertiesFromParsedAttr` mirror that behavior: keys are
accepted from the dictionary attr when the corresponding trait is present,
and required exactly when the format can't reconstruct the sizes itself
(i.e. when `fmt.allOperands` / `fmt.allResultTypes` is true). Otherwise the
key remains optional, since `genParserVariadicSegmentResolution`
overwrites it with the sizes inferred from the parsed operand/result
groups.
The contents of the PR as well as description has been made using
github-copilot. However I've reviewed the contents and tested locally
the change.
Fixes#211220
When using `ptxas` to do ptx->cubin, the options passed (via `cmd`) have
higher priority than the gpuModule target (such as `opt-level`). When
using `nvptxcompiler`, it is the opposite and we need to be consistent.
- 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.
Example:
```fortran
!$acc parallel loop vector reduction(+:b)
do i = 1, n
b(i) = i
end do
```
In this code the array reduction is thread(vector)-only with no
gang/block launch dim, so a within-block `all_reduce` is a complete
reduction. ACCCGToGPU still rejected it as NYI because the accumulate
had no block context.
Fix: keep the NYI only when the region actually launches block
dimensions; allow the existing per-element `all_reduce` path for
single-block thread-only array accumulates.
…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.
