Fixes#202374.
Preserve LLVM function `!prof` metadata for:
- `synthetic_function_entry_count`
- `function_entry_count` with import GUID operands
The importer represents the entry count, count kind, and trailing import
GUID operands in a single `#llvm.function_entry_count` attribute,
validates malformed operands before mutating the operation, and exports
the metadata through `llvm::Function::setEntryCount`.
Tests cover import, export, malformed metadata operands, synthetic
counts, synthetic counts with trailing operands, import GUID
round-tripping, and raw i64 bit-pattern preservation.
Checks:
- `git diff --check`
- `git clang-format --diff origin/main -- ...`
- `llvm-as -disable-output` smoke check
Codex was used in making this PR.
This change adds the `extf` and `truncf` ops to the NVGPU dialect to
support floating-point conversion operations.
These ops support scalar and vector inputs and lower to the corresponding
NVVM ops after padding and chunking the input into `i32` registers.
## Motivation
I'm adding NVIDIA FP8 mma.sync support in IREE for sm_89/sm_120, and
IREE's codegen lowers through `nvgpu.mma.sync` on this path, which in
turn lowers to `nvvm.mma.sync`. This MLIR-level gap is blocking that
work, so I'm upstreaming the fix here first. This is split into two PRs
per review feedback — this one covers `nvvm.mma.sync` only; a follow-up
PR will add the corresponding `nvgpu.mma.sync` support built on top of
this.
## Description:
`nvvm.mma.sync` (warp-level, dense) don't support FP8 (e4m3/e5m2)
operands, even though the underlying PTX/NVPTX intrinsics have supported
FP8 `mma.sync.m16n8k32` since sm_89. FP8 currently only exists on the
sparse mma path and on warpgroup-level `wgmma`; the warp-level dense
path was never updated.
This adds FP8 (e4m3/e5m2) support to `nvvm.mma.sync`, covering
`m16n8k16` and `m16n8k32`.
## Verification
`$ mlir-opt mlir/test/Dialect/LLVMIR/nvvm.mlir -split-input-file
-verify-diagnostics | FileCheck mlir/test/Dialect/LLVMIR/nvvm.mlir
`
```
%0 = nvvm.mma.sync A[...] B[...] C[...]
{multiplicandAPtxType = #nvvm.mma_type<e4m3>,
multiplicandBPtxType = #nvvm.mma_type<e4m3>,
shape = #nvvm.shape<m = 16, n = 8, k = 32>}
: (i32, i32, f32) -> !llvm.struct<(f32, f32, f32, f32)>
```
`$ mlir-translate -mlir-to-llvmir mlir/test/Target/LLVMIR/nvvmir.mlir
-split-input-file -verify-diagnostics | FileCheck
mlir/test/Target/LLVMIR/nvvmir.mlir
`
```
%8 = call { float, float, float, float } @llvm.nvvm.mma.m16n8k16.row.col.f32.e4m3.e4m3.f32(i32 %0, i32 %1, i32 %2, float %3, float %4, float %5, float %6)
%11 = call { float, float, float, float } @llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e5m2.f32(i32 %0, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, float %6, float %7, float %8, float %9)
```
All existing tests in nvvm.mlir (including -verify-diagnostics error
cases) still pass — no regression.
Co-authored with Claude Sonnet 5
Signed-off-by: weimin023 <tnwilly@gmail.com>
In preparation for the pass that converts `acc.compute_region` to GPU
dialect, this PR adds several utilities which are used in that pass.
Doing so to simplify review and to ensure that unit testing is added for
each of the utilities.
In preparation for the pass that converts `acc.compute_region` to GPU
dialect, this PR adds several utilities which are used in that pass
related to reductions. Doing so to simplify review and to ensure that
unit testing is added for the utilities.
---------
Co-authored-by: Scott Manley <rscottmanley@gmail.com>
Mirror the isEffectivelySerial() function in acc.compute_region to
acc.parallel and acc.kernel so this check is unified anywhere it needs
to be used, namely before ACCComputeLowering. Should be NFC
Extends `memref-elide-reinterpret-cast` from scalar-only copy
transformation to also handle copies of effectively 1D and static
multidimensional regions.
The rewrite is limited to cases where each non-unit `reinterpret_cast`
result dimension maps directly to a distinct identity-layout source
dimension with the same stride and sufficient static size. This lets the
rewrite use each loop IV
directly as a source memref index. Cases where result strides do not
match source identity strides are left untouched because they would
require linearization/delinearization across source dimensions.
Assisted-by: Codex (refine implementation + tests). I reviewed all code
and tests before submission.
## Examples
Scalar copy, unchanged(*) behavior:
```mlir
// BEFORE
%view = memref.reinterpret_cast %dst
to offset: [1], sizes: [1, 1], strides: [1, 1]
: memref<1x108xf32> to memref<1x1xf32, strided<[1, 1], offset: 1>>
memref.copy %src, %view : memref<1x1xf32> to memref<1x1xf32, strided<[1, 1], offset: 1>>
// AFTER
%v = memref.load %src[%c0, %c0] : memref<1x1xf32>
memref.store %v, %dst[%c0, %c1] : memref<1x108xf32>
```
Effectively 1D copy:
```mlir
// BEFORE
%view = memref.reinterpret_cast %dst
to offset: [0], sizes: [1, 33, 1], strides: [1386, 42, 1]
: memref<1x33x42xf32> to memref<1x33x1xf32, strided<[1386, 42, 1]>>
memref.copy %src, %view : memref<1x33x1xf32> to memref<1x33x1xf32, strided<[1386, 42, 1]>>
// AFTER
scf.for %i = %c0 to %c33 step %c1 {
%v = memref.load %src[%c0, %i, %c0] : memref<1x33x1xf32>
memref.store %v, %dst[%c0, %i, %c0] : memref<1x33x42xf32>
}
```
Static 2D copy with offset becomes a nested loop with the static offset
folded into the `reinterpret_cast` source store indices.
```mlir
// BEFORE
%view = memref.reinterpret_cast %dst
to offset: [16], sizes: [1, 33, 4], strides: [1386, 42, 1]
: memref<1x33x42xf32>
to memref<1x33x4xf32, strided<[1386, 42, 1], offset: 16>>
memref.copy %src, %view
// AFTER
%c0 = arith.constant 0 : index
%c1 = arith.constant 1 : index
%c4 = arith.constant 4 : index
%c16 = arith.constant 16 : index
%c33 = arith.constant 33 : index
scf.for %i = %c0 to %c33 step %c1 {
scf.for %j = %c0 to %c4 step %c1 {
%v = memref.load %src[%c0, %i, %j]
memref.store %v, %dst[%c0, %i, %c16 + %j]
}
}
```
## Scope
Supported:
- ranked, static-shape copy source and `reinterpret_cast` source/result
memrefs
- identity-layout `reinterpret_cast` sources
- rank-preserving `reinterpret_cast`s
- scalar copies, including dynamic result strides because all result
indices are zero
- non-scalar copies where every non-unit result dimension has a static
stride matching a distinct source dimension with sufficient static size
- static offsets, delinearized into source indices
- dynamic offsets only when the `reinterpret_cast` source has at most
one non-unit dimension and the result has at most one mapped non-unit
dimension
Not addressed:
- unranked or dynamic-shape copy/reinterpret_cast operands
- non-identity-layout `reinterpret_cast` sources
- rank-changing `reinterpret_cast`s
- dynamic strides on non-unit result dimensions
- result dimensions whose strides do not match source identity strides
- result dimensions that reuse a source dimension or cross source
dimension boundaries
- dynamic offsets requiring runtime div/rem delinearization
## Correctness
Result indices are translated to source indices by direct dimension
mapping: each non-unit result dimension must have the same stride as a
distinct source dimension.
Static offsets are delinearized into source indices before rewriting.
Dynamic offsets are only accepted when they can be used directly in the
single non-unit source dimension.
Cases assumed invalid by `reinterpret_cast`/`copy` semantics are guarded
with assertions.
(*) NFC to scalar copy case, we're now only creating zero constants for
indices once.
RFC
https://discourse.llvm.org/t/rfc-dwarfdebug-fix-and-improve-handling-imported-entities-types-and-static-local-in-subprogram-and-lexical-block-scopes/68544
Similar to imported declarations and types, the patch tracks
function-local static variables (globals in terms of LLVM IR) in
DISubprogram's 'retainedNodes'. DwarfDebug is adjusted in accordance
with the aforementioned metadata change and provided a support of static
locals scoped by a lexical block. CodeViewDebug is modified to collect
global variables from DISubprogram's 'retainedNodes' too.
The patch assumes that DICompileUnit's 'globals' no longer tracks static
locals and DwarfDebug would assert if any locally-scoped variables get
placed there.
This change has already been reviewed in
https://reviews.llvm.org/D144008, but I’m opening a new PR because it
has been a while since it was approved. Additional logic was added to
CloneFunction to deal with DISubprograms that have static locals: if
such a subprogram is cloned, DIGlobalVariableExpressions are removed
from retainedNodes of the clone. This is consistent with the current
behavior, in which DIGlobalVariableExpressions are not cloned.
Authored-by: Kristina Bessonova <kbessonova@accesssoftek.com>
Co-authored-by: Tobias Gysi <tobias.gysi@nextsilicon.com>
Extend negative stride checks to MaskedLoadOp, MaskedStoreOp, GatherOp,
and ScatterOp to match LoadOp and StoreOp behavior.
Depends on: https://github.com/llvm/llvm-project/pull/204611.
AI Disclaimer: I used AI for the tests.
@banach-space @dcaballe If you look at the latest commit, you can see
why the integration tests were failing in #204842. Those two calls were
added in this PR, I sincerely apologize to you!
---------
Signed-off-by: Federico Bruzzone <federico.bruzzone.i@gmail.com>
Add getStoreSubgroupLayouts() to derive the subgroup (sg_layout/sg_data)
anchor layout for store_scatter and store_matrix. It factorizes numSg
into the most balanced sg_layout that divides the workgroup tile with
sg_data a multiple of inst_data (lane_layout * lane_data).
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Inlining can leave a `fir.declare`'s dummy_scope operand (`!fir.dscope`)
defined outside an OpenACC compute region while the declare itself gets
rematerialized inside, since `fir.dummy_scope` wasn't registered as an
OutlineRematerializationOpInterface candidate.
This dangling, unmappable live-in is illegal for parallelization.
Fix: Added `OutlineIdentityOperandOpInterface`, implemented by
`fir.declare/hlfir.declare`, so `OffloadLiveInValueCanonicalization`
drops their dummy_scope operand (an identity token that must never be
duplicated) instead of cloning it when sinking/rematerializing them into
offload regions.
This commit adds support for block scaled tensors.
In particular, the block scaled type has been extended with the
`DenseElementTypeInterface` to allow block scaled data values to be
specified in dense element attributes.
The block scaled type has also been extended to allow optional scale
values to be specified by the type. For now, scale values are not
expected to be propagated beyond their use in the attribute input of a
constant operation. In the future, we may want to propagate these values
to allow certain optimizations.
The `tosa.const` operation has also been updated in the validation pass.
Clang 21.1.8 segfaults during SFINAE overload resolution for
T::getChecked
when matching non-pointer arguments against an unconstrained
MLIRContext*
parameter.
This patch avoids the crash by adding an explicit SFINAE constraint to
the second overload of StorageUserBase::getChecked, ensuring that it is
only considered when the second argument is convertible to MLIRContext*.
Assisted-by: Antigravity
MmaOp::getIntrinsicID triggers an uninitialized-read warning under
Valgrind. `intOverflowBehavior` is optional; in the generated selector
its `has_value()` guarded access can be speculatively loaded before the
guard, reading the empty std::optional's uninitialized payload.
The read is benign: the payload is correctly stack-allocated but
uninitialized while empty, and the loaded value is discarded by the
guard, so this is not incorrect / undefined behavior.
This MR keeps the attribute as optional, but the lowering now resolves
it to a concrete value (wrapped when omitted) before calling
getIntrinsicID, so the selector always reads an initialized value.
`wrapped` matches the PTX spec and the existing lowering, where an
omitted attribute already selected the non-satfinite intrinsic;
intrinsic selection and printed IR are unchanged. Documentation updated
to correct the stale default.
The MLIR pass that removes operations exclusively intended for the host
from OpenMP target offload modules is currently defined as part of
Flang. However, this is a feature that would benefit from being reusable
by other frontends, as removing such operations is a requirement for all
OpenMP target device modules prior to LLVM IR translation.
By moving the `omp-host-op-filtering` pass out of Flang, it had to be
updated to work on a lower-level LLVM dialect-based representation,
rather than FIR. This simplified some of the existing edge cases, such
as `fir.declare` ops and `fir.boxchar` type handling. In addition, new
function arguments are introduced as placeholders and return values from
host-only functions are removed, producing a cleaner result and
simplifying the pass as compared to previously.
As a result of a later execution of this pass, dynamic dispatch of host
functions via dispatch table using `fir.dispatch`, `fir.type_info` and
`fir.dt_entry` ops would break due to the removal of `fir.dt_entry`
operations pointing to deleted host functions, while `fir.dispatch` ops
pointing to the same functions would remain. The
`omp-function-filtering` pass is updated to prevent `fir.dt_entry` from
being deleted and let them point to an undefined symbol. This lets FIR
lowering to LLVM work well enough to get to the host op filtering stage,
where any resulting operations for this host dynamic dispatch are
cleaned up.
The nanobind type/attr/loc/value casters in IRCore.h return a by-value
PyType/PyAttribute/PyLocation/PyValue as their DerivedTy result.
Under pre-C++23 rules, this was a rvalue that would decay to a lvalue
when passed to the DerivedTy constructor.
Under C++23's P2266 (implicit move on return), that return operand is an
xvalue, which cannot bind the PyConcrete* constructors that take a
non-const lvalue reference.
This patch constructs the derived type explicitly (return
DerivedTy(arg);) using the lvalue constructor; this preserves the
pre-C++23 behavior on C++23 builds.
Assign each mlir::Block a stable ID within its parent region, mirroring
llvm::BasicBlock/llvm::Function (called ID, not number, since a block/op
number denotes position in MLIR): Block gains getBlockID() and reads -1u
while it has no parent region; Region gains nextBlockID with
getMaxBlockID() and getBlockIDEpoch(); the block ilist traits assign the
ID on add/transfer and invalidate it on removal; and
GraphTraits<mlir::Block*>/<mlir::Region*> expose
getNumber/getMaxNumber/getNumberEpoch. This makes
GraphHasNodeNumbers<mlir::Block*> true, moving MLIR's CFGLoopInfo and
dominator tree onto the number-indexed path. MLIR never renumbers
blocks,
so the epoch is a fixed 0.
Prerequisite for requiring GraphHasNodeNumbers in llvm::LoopInfoBase and
dropping its DenseMap fallback.
Aided by Claude Opus 4.8
Add an operation to represent the runtime call to acc_on_device.
This runtime call is important to fold early in the compilation pipeline
and having an operation allows us to easily recognize it when emitted by
frontends.
Introduce an attribute `OpenACCLoopLocAttr` that can be attached to an
`acc.loop` operation to record loop locations and directive location
separately.
Attach it to `acc.loop` when building the op in flang.
This commit enables custom legalization for v1 types for intrinsic results and operand by calling CustomLowerNode which gives the target a chance to custom handle the node before scalarizing
To demonstrate, I have used tcgen05.ld/st intrinsics which previously used a scalar i32 type while all wider variants used vector types
Assisted by: Claude Code
Introduces an algorithm selection option to the static memory planner
pass and adds a best-fit algorithm that reuses memory from expired allocations by
finding the smallest suitable gap.
I have tested the correctness of the best-fit algorithm with an
independent randomized verifier.
Signed-off-by: mabsar <mabsar@qti.qualcomm.com>
Add a general acc utility for computing the size and alignment of a
type. It works for simple scalar types as well as nested and aggregate
types like arrays, tuples, and structures.
Because some types come from other dialects, the utility can hand those
off to a specialized helper that understands them. This lets sizing work
seamlessly even for mixed types, such as an aggregate whose members come
from a different dialect.
Add a Fortran-specific helper so Fortran types are sized correctly,
falling back to the general utility for everything else.
Include unit tests covering a range of scenarios, including scalars,
arrays, aggregates, and mixed-dialect types.
Add a generated LLVM dialect op for llvm.arithmetic.fence, including
LLVM IR import and export coverage for scalar and vector floating-point
forms.
Assisted-by: Codex
There are three hooks that have a different, more complicated API than
the other ones: allocationFn, memCpyFn, castFn. All three are called via
"helper" functions that check whether the hooks are set to call them and
fall back to default implementations otherwise.
Other hooks (e.g. unknownTypeConverterFn) in the bufferization options
have a different "API": they are always set to some default
implementation and can be overwritten by the user. This is a simpler API
overall and seems it can be universal.
Note for LLVM integration: BufferizationOptions::{createAlloc,
createMemCpy, createCast} functions are "replaced" with direct calls to
the underlying std::function objects. At any call-site that relied on
any of the removed APIs, a call to the respective hook should now be
used instead. The exception is the `allocationFn` hook which needs an
explicit "bufferAlignment" parameter. Beforehand, it was implicitly
taken from BufferizationOptions::bufferAlignment field, so a
non-functional migration path is to go from
`options.createAlloc(builder, loc, destType, value)` to
`options.allocationFn(builder, loc, destType, value,
options.bufferAlignment)`.
This commit adds a new compound-block scaled tensor type and uses this
type to implement support for MXFP in the CAST operation, as per the
following specification changes:
https://github.com/arm/tosa-specification/pull/50,
https://github.com/arm/tosa-specification/pull/53.
The new block-scaled type is closely modelled after the `quant` dialect
type and supports the following parameters:
- block shape - The size and axis of each block. Support for specifying
constant scale values has not been added in this commit, but will be
added in a later one.
- scale type - The type of the scale value associated with each block.
- value type - The type of the data values in each block.
Example syntax for the new block-scaled type:
```
tensor<160x!tosa.block_scaled<BLOCK_SHAPE_32:f8E8M0FNU:f8E5M2>>
```
As a pre-requisite for supporting MXFP CAST, the commit also adds new
extensions intended to split the existing EXT-MXFP extension into
separate extensions for each block-scaled type. This allows for more
fine-grained control over which block-scaled types are supported by a
given target. See specification change
https://github.com/arm/tosa-specification/pull/26 for details.
Finally, support for casting to/from the new block-scaled type has been
added, aligning with the behaviour specified in the CAST operation
specification.
Currently, formatv erases types using a base class and calls the virtual
function format() to format the objects. To avoid these vtables,
refactor formatv to instead store function_refs to a functor (struct
with overloaded operator()). This saves ~5kiB in vtables.
Additionally, add a static assertion that a formatter is present instead
of relying on errors due to missing templates. This requires a little
change to MLIR's tgfmt to use a different name -- previously, the
substitution failure on ArrayRef<> would cause the variadic function to
be skipped, but a static assertion failure is not a substitution error.
Also, simplify the code in FormatVariadicDetails to use if constexpr
instead of template overloads with enable_if, making the code shorter
and easier to read.
Serialize IntegerSetAttr using VarInt for numDims and numSymbols and
arrays for contraints and eqFlags. Add ArrayWithKnownSize` an Array
variant that omits the length prefix when the size is known from a
previously serialized member. This avoids encoding redundant array
lengths when, for example, two arrays are always the same size.
Assisted-By: Gemini
# Inner tile alignment hints for scalable `linalg.pack`/`linalg.unpack`
tiling and fusion
## Overview
Tiling and fusing `linalg.pack`/`linalg.unpack` produces a clean result
only when the tiling implementation can tell how a loop tile size
relates to the op's inner tile size. When both are statically known this
is decided by comparing the constants. But with **scalable** (and, more
generally, dynamic) sizes, e.g. a loop tile of `8 * vscale` against an
inner tile of `8 * vscale`, that relationship is symbolic and cannot
cleanly be recovered from the IR, so the implementation conservatively
falls back to a dynamic, over-allocated tile. See #150185 for more
details.
This PR adds an optional **inner tile alignment hint**: a per-dimension
caller assertion about that relationship, threaded from the transform-op
surface through the SCF driver into the pack/unpack `TilingInterface`
implementations. The caller is already aware of both the loop tile
sizes, and the inner tile sizes, and therefore can set these hints
easily. With the hint, scalable pack/unpack tiling and fusion produce
the same tight IR they produce for the statically-aligned case.
## The inner tile alignment hint
`InnerTileAlignment` is a per-iteration-domain-dimension enum:
- `Unknown` (`0`) — no information; the implementation keeps its prior
behavior.
- `Multiple` (`1`) — the loop tile size is an integer multiple of the
inner tile.
- `Equal` (`2`) — the loop tile size equals the inner tile size.
It is a *caller assertion*, not a checked fact. It is consulted only
when the relationship cannot be decided from the IR; when the tile and
inner-tile sizes are both statically known, that static comparison wins
and a contradicting hint is ignored rather than allowed to corrupt the
result. Otherwise the hint is not verified, so an incorrect assertion
silently produces invalid tiling, its correctness is owned by the
caller. Only pack/unpack consult it; every other op ignores it.
## TilingInterface API
Four methods gain a second overload taking an extra
`ArrayRef<InnerTileAlignment>` hint:
- `getTiledImplementation`
- `generateResultTileValue`
- `getTiledImplementationFromOperandTiles`
- `getIterationDomainTileFromOperandTiles`
The hint-bearing overload defaults to forwarding to the hint-less one,
so the change is additive: existing implementers and callers that never
mention the hint are unaffected, and the rest of the interface is
untouched.
### Where it helps
The loop tiles the packed dimension by `8 * vscale`, so the size handed
to the pack/unpack is a **full-or-remainder** `affine.min` against that
step:
```mlir
// #map = affine_map<(d0)[s0] -> (-d0 + 4096, s0)> // min(4096 - iv, 8*vscale)
%ts = affine.min #map(%iv)[%c8_vscale]
// ... %ts feeds the extracted slice fed to the pack/unpack, whose inner tile
// on this dimension is also 8*vscale.
```
To fold the tiled op's outer (tile-count) dimension to a static `1`, the
current implementation asks
`ValueBoundsConstraintSet::computeConstantBound` for a *constant* upper
bound on `%ts` and compares it against the inner tile size, essentially
`UB(%ts) % innerTile == 0`. With a scalable `8 * vscale` inner tile,
neither the upper bound of `%ts` nor the inner tile is a constant, so
the comparison is undecidable and the tile conservatively falls back to
a dynamic outer dim (`tensor<?x...>`), over-allocating the tile.
`inner_tile_alignments = [Equal]` asserts `%ts == 8 * vscale`, so the
outer dim folds to a static `1` (`tensor<1x...>`); `[Multiple]` instead
takes the `ceilDiv` branch. The hint thus supplies exactly the
relationship the upper-bound comparison cannot recover for symbolic
sizes.
### How the hint is arranged
- One entry per **iteration-domain dimension**, in **pre-interchange**
order (`interchange` only reorders the generated loops, not the hint).
- Entries are **not** remapped through indexing maps or
`outer_dims_perm`: for a transposing pack they stay in source
(pre-permutation) order.
- For the consumer-fusion (`*FromOperandTiles`) methods, the indexing
follows the op's iteration domain: for a pack it coincides with the
unpacked source dims; for an unpack the entry for the i-th inner tile
sits at its destination dimension `inner_dims_pos[i]`.
The caller must pre-arrange the array to match this order. See the tests
for more examples.
### Out-of-tree implementers
Because the two overloads share a source name:
- An op that lists one of the four methods in
`DeclareOpInterfaceMethods` now has **both** arities declared on it, so
it must also define the hint-bearing overload (typically a one-line
forward to the hint-less one).
- An `ExternalModel` that defines the hint-less overload **hides** the
inherited defaulted hint-bearing overload (C++ name hiding); add `using
Base::<method>;` to keep it visible.
Implementers that do not care about the hint need no functional change
beyond the above visibility boilerplate.
## SCF tiling driver API
The producer- and consumer-fusion driver entry points
(`tileAndFuseProducerOfSlice`, `tileAndFuseConsumerOfSlices`,
`tileAndFuseConsumer`) and `SCFTilingOptions` accept an optional
**control function** rather than a single fixed array.
A single tile-and-fuse invocation can tile/fuse several ops with
**different iteration domains**, so no single array can be correct for
all of them. The driver therefore invokes the control function once per
op being tiled/fused and
expects the returned array to be in *that op's own* domain, it does not
remap it.
### Why a control function and not one array
A single tile-and-fuse can pull in more than one pack/unpack, and each
consults the hint in *its own* iteration domain:
```mlir
%unpacked = linalg.unpack %p0 inner_dims_pos = [0] inner_tiles = [%c8_vscale] ... // domain A
%mul = linalg.generic ins(%unpacked, ...) ... // tiled root
%packed = linalg.pack %mul inner_dims_pos = [1] inner_tiles = [%c4_vscale] ... // domain B
// One tile-and-fuse: tile %mul, fuse %unpacked (producer) and %packed (consumer).
```
Both `%unpacked` and `%packed` consume the hint, but their domains
differ, a different dimension is tiled and the inner-tile arrangement is
different, so the array indexed in `%unpacked`'s domain is simply wrong
for `%packed`. No single fixed array can be correct for both. The driver
therefore invokes the control function once per op being tiled/fused and
the caller returns each op's array in its own domain.
### Why the function gets slices (not tile sizes) for fusion
For pure tiling the driver knows the explicit tile sizes and passes
them. In fusion there are no caller-specified tile sizes for the fused
op, its tile is induced by the `tensor.extract_slice` (producer) or
`insert_slice` (consumer) being fused. The control function is therefore
handed those slices, from which the caller can derive the alignment in
the fused op's domain.
## Transform ops
`transform.structured.tile_using_for`, `transform.structured.fuse`, and
`transform.structured.fuse_into_containing_op` gain an optional
`inner_tile_alignments` array attribute (empty by default, preserving
prior behavior). It is wired to the driver as a **constant** control
function (the same array for every op). This is sufficient because each
of these ops drives tiling/fusion at single-op granularity; the per-op
control function remains available at the library/driver level for
callers that fuse multiple ops with differing domains. The test-only
`transform.test.fuse_consumer` is wired the same way.
## Testing
- **`linalg.unpack`**: standalone tiling, producer fusion
(`structured.fuse`, `fuse_into_containing_op`), and consumer fusion
(`test.fuse_consumer`), with both `Multiple` and `Equal` hints.
- **`linalg.pack`**: standalone tiling, producer fusion, and consumer
fusion (`test.fuse_consumer`), with both `Multiple` and `Equal` hints.
- Negative cases: missing hint falls back to the conservative tiling; a
hint that contradicts statically-known sizes is ignored; transposing
`outer_dims_perm`/consumers confirm source-order (un-permuted) indexing.
- Transform-op attribute verification (`inner_tile_alignments`
validity).
Assisted-by: Claude Code
---------
Signed-off-by: Ege Beysel <beyselege@gmail.com>
Fixes LeakSanitizer failure from #206161 (reverted in #207104);
`mlirFreezeRewritePattern` moves contents out of the `RewritePatternSet`
but does not free the container (passed by value in the C API), so the
allocation from `mlirRewritePatternSetCreate` was never freed (add
`mlirRewritePatternSetDestroy(patterns)` after freezing).
This patch modifies the handling of the
`-fopenmp-assume-teams-oversubscription` and
`-fopenmp-assume-threads-oversubscription` compiler options in Flang and
the corresponding `omp.flags` MLIR attribute to have this information
present on both host and target device MLIR modules. This is necessary
to ensure SPMD to SPMD-No-Loop kernel promotion happens uniformly and no
divergence between host and device representations of the same target
region is introduced.
The identification of the kernel type in Flang lowering is updated to
address some broken edge cases. In particular, no SPMD to SPMD-No-Loop
transformation is performed on `target parallel do [simd]` kernels. A
split `target + parallel loop` is now marked SPMD, which was a case
missed previously, and the trip count on `target teams ompx_bare`
kernels isn't host evaluated even if removing the `ompx_bare` clause
would result in a valid SPMD kernel.
While merging blocks, `transformToStructuredCFBranches` may try to erase
blocks that still have uses. This may happen when the block has
self-loop or when it has 1 successor, but the successor has multiple
predecessors. This probably happens only with unreachable blocks in the
function.
Note that running a DCE pass before CF to SCF transform would remove the
triggering assert in the lit test.
Fixes#206086.
This PR is the second part of the work that aims to allow customizations
in resolving mismatching buffer types.
(see 0b1fe74f07 for part one)
Add a new bufferization hook that lets downstream bufferization
implementations materialize an operation to resolve a pre-reconciled
buffer type mismatch.
The hook + API is used as a fallback mechanism in several upstream
operations, as well as in the general bufferization infrastructure.
Semantically, this is a "cast" creation akin to `createAlloc()` and
`createMemCpy()` but for the cases where a memref::CastOp-like operation
is acceptable.
A valid value is returned whenever a cast-like operation is created,
failure signifies a failed attempt to materialize a cast (in theory,
there could be a fall back to alloc + copy either externally, but this
right now is not considered in most contexts). By default, a
`memref.cast` operation is created that casts a bufferized value to a
reconciled type.
`vector.step` previously produced only `index` vectors. Allow its result
element type to also be a signless integer of at least 8 bits. If a
lane's value in the sequence is not representable in the element type it
wraps around (the value is truncated to the element bitwidth),
consistent with the
[`llvm.stepvector`](https://llvm.org/docs/LangRef.html#llvm-stepvector-intrinsic)
intrinsic that `vector.step` lowers to.
ConvertVectorToLLVM now lowers `vector.step` to `llvm.intr.stepvector`
for all vector types (fixed-width and scalable) rather than
materializing an `arith.constant` for fixed-width vectors; the constant
form does not make sense at the LLVM boundary, where the intrinsic
already carries the target index bitwidth.
The in-dialect `arith.constant` lowering remains for other pipelines and
is exposed through a new `transform.apply_patterns.vector.lower_step`
op. `populateVectorStepLoweringPatterns` takes an `indexBitwidth` used
to materialize `index`-typed steps (default 64, preserving existing
behavior; 0 leaves them untouched). Integer range inference and the
VectorToSPIRV lowering are likewise updated for the new element types
and wrap-around.
The wrap-around path has no SPIR-V test: SPIR-V vector types are limited
to at most a handful of elements, so no legal SPIR-V vector is large
enough to exceed an 8-bit (or wider) element range.
---
**Note**: AI used to generate the code in this PR.
Signed-off-by: Víctor Pérez Carrasco <victorperez@fb.com>
1sgtpepper
Scott Manley
Alexis Engelke
Srinivasa Ravi
WMC
Razvan Lupusoru
ioana ghiban
ShashwathiNavada
Vladislav DzhidzhoevandTobias Gysi
Federico Bruzzone
Longsheng Mou
Jianhui LiandClaude Opus 4.8
Luke Hutton
Susan Tan (ス-ザン タン)
Kazu Hirata
Stefan Mada
Sergio Afonso
Nikhil Kalra
Fangrui Song
Ivan R. Ivanov
Moazin K.
Pradeep Kumar
Vladimir Suvorov
Javed Absar
Tom Eccles
Andrei Golubev
Aliaksei Urbanski
Jacques Pienaar
Ege Beysel
Maksim Levental
Jeremy Kun
Martin P
Arseniy Obolenskiy
Victor Perez