[AArch64][TTI] Cost extends folded into [US]SHLL as free (#233489)
A doubling NEON zext/sext feeding a shl by a splat immediate
is selected as one [US]SHLL/SHLL, so the extend is free.
[flang] Decouple HLFIR array expansion from optimization level (#233478)
Add optimizeArrayExpression option to the pass pipeline configuration,
set to false by default for now.
It will allow preserving runtime calls for array expressions when set to
false once array temporary
optimizations are default at O0.
Always attempt to inline scalar-to-array assignments, including host
code at O0.
Update pipeline expectations and add coverage for broadcasts, array
expressions, and reductions.
The goal is to make this option always true by default once we start
emitting some fir.assert operation.
Part of
https://discourse.llvm.org/t/rfc-flang-expanding-array-expressions-and-assignments-at-o0/91971
[mlir][vector][NFC] Rename getShapeForUnroll to getIterationShapeToUnroll (#232648)
The old name was ambiguous about whether it returned the original shape
or the target shape for unrolling.
The new name makes clear that it returns the shape of the iteration
space to be unrolled. This can differ from the result shape: for
example, vector.contract includes reduction dimensions in its iteration space
that are absent from its result shape.
Implements: https://github.com/llvm/llvm-project/issues/233408
[AArch64][TTI] Cost NEON Nots absorbed by And/Or users as free (#232788)
Cost a NEON Not as free when every user is an And/Or in the same block,
since the Not folds into its BIC, ORN or BSL. If a user negates both of
its operands, only one Not can fold, so it still counts.
---
<sub>Stack created with <a
href="https://github.com/github/gh-stack">GitHub Stacks CLI</a> • <a
href="https://gh.io/stacks-feedback">Give Feedback 💬</a></sub>
[AArch64][TTI] Cost the Ands of a NEON bit select as free (#232787)
Or(And(M, A), And(Not(M), B)) selects to a single BSL/BIT/BIF. Charge
the
Or and treat the single-use Ands in its block as free, using the matcher
shared with operand sinking.
---
<sub>Stack created with <a
href="https://github.com/github/gh-stack">GitHub Stacks CLI</a> • <a
href="https://gh.io/stacks-feedback">Give Feedback 💬</a></sub>
[AArch64][TTI] Drop single-use requirement on the mask from BSL sinking (#232786)
Drop the one-use requirement on the mask's Not. Sinking puts a copy of
the
Not next to the And, and the original stays for its other users.
ISel can then form bsl/bif even when the Not has other uses.
---
<sub>Stack created with <a
href="https://github.com/github/gh-stack">GitHub Stacks CLI</a> • <a
href="https://gh.io/stacks-feedback">Give Feedback 💬</a></sub>
[AMDGPU][NFC] Explicitly narrow conversions in libcalls, TTI and InstCombine (#215196)
Fixes 34 instances of MSVC warning C4244 and 3 of C4267 ("possible loss of data") across 4 files in llvm/lib/Target/AMDGPU.
Assisted-by: Claude <noreply at anthropic.com>
[Clang][AST][NFCI] Make Decl reference attributes directly (#219138)
This commit changes `Decl` attributes from being stored in a map to
being managed by an allocator and letting the individual `Decl` objects
reference the attributes directly. Doing so reduces the additional
complexity of maintaining a map and improves performance by reducing the
levels of directions used when accessing `Decl` attributes.
As an additional benefit, this removes the need for the `hasAttr` field
as the pointer can be used to indicate whether a `Decl` has attributes
or not. This coincidentally preserves the size of `Decl` on most
systems, as the single-bit `hasAttr` field caused additional padding for
the class.
Assisted-by: Claude Opus 5
---------
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
[IR][TableGen] Remove unused Attribute::isExistingAttribute (NFC) (#233682)
The last caller of Attribute::isExistingAttribute was removed on July
6, 2026 in commit dc9edeb0453bd54b7461dae7feb5b35b784914fa.
This patch removes Attribute::isExistingAttribute along with the
ATTRIBUTE_ALL macro, which became unused as a result.
Assisted-by: Antigravity
[mlir] Replace PointerUnion::dyn_cast with llvm::dyn_cast (NFC) (#233681)
PointerUnion::dyn_cast has been soft-deprecated. This patch replaces it
with llvm::dyn_cast:
- PadTilingInterface.cpp: Each OpFoldResult in paddedShape is a
nonnull Attribute or Value, and dynamic dimensions must be valid
Values.
- TestOps.td: CallWithSegmentsOp::setCalleeFromCallable is an unused
CallOpInterface implementation on a test op.
This patch completes the migration away from PointerUnion::dyn_cast
under mlir/.
Assisted-by: Antigravity
[LoongArch] Lower FP_TO_SINT_SAT to [X]VFTINTRZ/FTINTRZ
The ftintrz.{w,l}.{s,d} and [x]vftintrz.{w.s,l.d} instructions saturate
out-of-range inputs to the minimum/maximum integer and convert NaN to 0
(when the invalid operation exception is not trapped), which matches the
semantics of llvm.fptosi.sat when the saturation width is the native
32/64 bits. This was verified on LA664 for all scalar, LSX and LASX variants
and matches QEMU's LoongArch model. Use them instead of the generic
expansion, which emits three compares and selects per element and fully
scalarizes vectors.
Also override shouldConvertFpToSat so DAGCombiner only folds min/max
clamps and NaN-guarded selects into FP_TO_SINT_SAT when that lowers to
a single ftintrz, avoiding the generic expansion for mismatched vector
element widths and for i64 results without the D extension.
This is what Rust emits for `f32 as i32`. rustc-perf svg-render-1 on
LA664: instructions -6.4%, cycles -4.0%.
[AMDGPU] Preserve noclobber after acquire operations from a different addrspace
Follow-up to #219991 to address reviewer comments:
AMDGPUAnnotateUniformValues should annotate uniform global loads as
`!amdgpu.noclobber` if the memory they are reading is never clobbered (since
kernel launch) before the load is executed.
PR #219991 strengthened the logic to consider acquire operations before the
load as clobbering (because they cause stores from other threads to be observed
by the load without a data race).
This PR weakens this criterion to ignore acquire operations that don't affect
the global address space (because they don't affect what the load in question
may or may not see). This addresses the performance-relevant case of s_barriers
with LDS-only fences.
The logic is still overly pessimistic for one-as load acquires, they could also
be ignored when looking for clobbering operations.
Disclaimer: Claude gave review comments and suggestions that I addressed during
[5 lines not shown]
[CIR] Support building vectors from smaller vectors (#231404)
Fix a crash on OpenCL `(float4)(a, b)` and HIP `__half2 data{xx}`
Use shuffles like classic codegen
[CIR] Reject a record redefined with a different body (#230547)
A context holds a single record per kind and name, so parsing a named
record returns any earlier definition of that name. When the earlier
definition was complete, the parser kept it and silently dropped the new
body, so a global declared with a different layout was given the first
one. Report the mismatch instead.
Assisted-by: Claude Opus 5.5
---------
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
[mlir][tblgen] Define enum `FieldParser` specializations out-of-line
The `FieldParser` specializations generated for every enum were defined
inline in the enum declarations header. Since their bodies do not depend
on the `ParserT` template parameter, they are type-checked eagerly in
every translation unit including the header, which instantiates
`FailureOr<E>`, `FailureOr<std::optional<E>>` and their `std::optional`
bases and type traits for every enum. Only a small fraction of these
specializations is ever used, and only from the generated attribute and
op parsers of the dialect itself.
This patch changes `parse` to a non-template taking an `AsmParser &`,
declares it in the header and emits its definition into the enum
definitions file instead.
Compiling all of MLIR (clang 18 with 64 cores) takes 4.5% less CPU time, with the median
translation unit compiling 2.8% faster and some up to 15% faster. The
summed peak memory usage of all translation units drops by 4.5% and the
largest peak memory usage of a single translation unit (`mlir/lib/CAPI/Conversion/Passes.cpp`) drops by 13%
[4 lines not shown]
[CIR][HIP] Match kernel handle linkage, visibility and comdat (#227267)
For HIP, CIRGen emits a host-side kernel handle, i.e. a global named
after the kernel that points to its __device_stub__ function. The handle
was always created with external linkage and default visibility. For
kernels with internal linkage and for linkonce_odr template
instantiations, every TU that launched the kernel therefore defined a
strong external handle, and linking two such TUs failed with
multiple-definition errors.
Matching classic codegen, we give the handle the stub's linkage,
visibility and dso_local, and put it in a trivial comdat under the same
conditions. This is done in emitDeviceStub rather than getKernelHandle
because the handle is created on the kernel's first reference, which can
precede the point where the stub's linkage is set.
Assisted-by: Claude Opus 5.5
---------
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
[LV][RISCV] Add tests for interleaved accesses with gaps. NFC (#233718)
This is to show the RISC-V cost model choosing strided loads when
there's gaps and each segment loaded would take multiple DLEN sized
accesses. See the comment in
https://github.com/llvm/llvm-project/pull/231919#issuecomment-6279062949
[Sema][CUDA][HIP] Reject RTTI operations in device code (#228469)
Currently RTTI operations, i.e. dynamic_cast and typeid, are silently
ignored in device code. This does not match the behavior of NVCC, which
rejects these operations in device code.
Make Sema reject these operations in device code, emitting a diagnostic.
An effect of this is that the compiler may reject code that previously
compiled, such as code that used RTTI in dead device code that the
compiler would remove before reaching the linker. However, cases would
currently fail if optimizations are disabled.
Assisted-by: Claude Opus 5.5
---------
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
[X86] Ignore debug instructions after calls in call frame optimization (#232624)
`collectCallInfo` skips debug instructions while scanning argument
setup, but checks the instruction immediately after a call for
`ADJCALLSTACKUP`. A `DBG_VALUE` at that boundary makes the pass abandon
the sequence, so adding debug information can leave stack argument
stores instead of converting them to pushes.
Use `next_nodbg(..., false)` to ignore only debug instructions when
looking for `ADJCALLSTACKUP`, without skipping pseudo probes. The MIR
test places a `DBG_VALUE` between the call and `ADJCALLSTACKUP` and
covers both pass-manager paths.
Assisted-by: OpenAI Codex 5.6 Sol, co-reviewed by Claude and Yongqiang
Tian.
[SROA] Rename and document `isSafeSelectToSpeculate` (NFC) (#232681)
`isSafeSelectToSpeculate` does much more than the name and documentation
curently suggest:
- It handles stores too, the documentation only mentions loads.
- It checks for predication and returns instructions that need to be
predicated, when we're allowed to modify the CFG.
Rename it as `isSafeSelectToSpeculateOrPredicate` to clarify that the returned
instructions may require predication.
Document what predication is, as used in this pass. Current documentation only
has a speculation example.
Documentation is also updated with a few more details.
