[Lanai] Rename -lanai-mem-alu-combiner to avoid a pass name collision (#234513)
opt registers each legacy pass name as an option, and
LanaiMemAluCombinerLegacy, named lanai-mem-alu-combiner, causes a
collision.
[X86] Support -mnop-mcount and -mrecord-mcount on x86 elf targets (#184641)
Support mnop-mcount and mrecord-mcount attributes when used with the
fentry-call=true attribute.
Similar to SystemZ, use the mnop-mcount attribute to place a nop instead
of a call to __fentry__.
Likewise, use the mrecord-mcount attribute to record the location of the
call to fentry (or nop if used with the above) to a section named
__mcount_loc.
This enables the use of -mnop-mcount -mrecord-mcount for x86 in Clang.
[LoongArch] Lower vector sign extension with VSLLWIL and VEXTH
vsllwil.{h.b/w.h/d.w} with a zero shift amount sign-extends the low
half of the source elements, and vexth.{h.b/w.h/d.w} the high half.
Use them for the custom SIGN_EXTEND legalization and for
SIGN_EXTEND_VECTOR_INREG, instead of computing the sign mask with vslti
and interleaving it with vilvl/vilvh, which takes one more instruction
per step. SIGN_EXTEND_VECTOR_INREG now also handles 4x extensions, which
previously fell back to a shuffle followed by a vslli/vsrai pair.
[Clang] Instantiate functions from constant evaluation. (#205557)
This PR is a rework of #173537, which was a rework of #115168; the sole
goal remains to address #73232. I used #173537 as a starting point (for
both code and tests), but attempted to follow the alternative design
requested in the comments therein.
A quick design overview:
- A `SemaProxy` interface is introduced to `AST`, which defines a set of
(possibly mutable) semantic operations that can be triggered by the
evaluation of certain expressions; this class finds its implementation
in `Sema/SemaEval.cpp`. `Sema` creates a unique `SemaProxy` to itself at
construction time, which is made available via
`Sema::getproxyForEval()`.
- Ordinary evaluation (e.g., via `Expr::EvaluateAsConstantExpr`) has no
access to `SemaProxy`, and therefore cannot act on the AST: instead,
variants of these functions (e.g.,
`Expr::EvaluateAsMandatedConstantExpr`) are introduced, which accept a
`SemaProxy &`. This makes clear from the call-side whether an operation
[36 lines not shown]
[IndVarSimplify] Skip exit predication past divergent effects (#183906) (#194202)
This PR makes `IndVarSimplify::predicateLoopExits` skip divergent
effects. As [nikic](https://github.com/nikic) 's suggestion, I reuse the
code about infinite loop detection in `LoopDeletion`
Close #183906
AMDGPU: Include callees in backend_stack_size PAL metadata
The stack size reported in the PAL metadata of a chain function should
include the stack sizes of any callees, since the consumers don't have
access to the call graph (this is especially important for whole wave
functions, because we don't even emit any PAL metadata for them). The
consumers are generally expected to allocate scratch based on the
maximum `.backend_stack_size` emitted for the functions in a pipeline.
They will also have to deal with recursion and indirect calls.
The `.stack_frame_size_in_bytes` metadata remains unchanged, i.e. it
represents just the stack size of the current function (without its
callees).
Assisted-by: AI
AMDGPU: Emit .backend_stack_size for entry points
Currently, the `.backend_stack_size` PAL metadata is only emitted for
functions, but some of the consumers expect it to exist for the entry
points too (without it, if the entry point uses more scratch than any
other functions in the pipeline then we will allocate too little
scratch).
The meaning is the same for entry functions as it is for non-entry
functions: the stack size of the current function plus any of its
non-chain callees.
Unlike `.scratch_memory_size`, this does not get rounded up to 16.
Assisted-by: AI
AMDGPU: Reset SP in entry point before chain call
Chain calls are always tail calls and they are expected to completely
take over the caller's stack frame. Chain functions receive a stack
pointer from their caller, but that is only in order to support
reserving space at the bottom of the stack for CWSR.
Currently, when an entry function has its own stack variables, the SP we
pass on to the chain function points after the stack of the caller. This
is wasteful, so the current patch introduces an epilogue for entry
functions if they contain a chain call. The epilogue will reset SP to
point to the bottom of the stack, either by copying it from FP if we've
been reserving space for CWSR or by setting it directly to zero.
Note that we also modify `requiresStackPointerReference` to return false
for tail calls. This means the prologue will only initialize SP if it's
really needed by the current function, relying on the epilogue to handle
it if we only need to set SP for the sake of our chain callee. This
avoids situations where we might unnecessarily initialize the SP with
[4 lines not shown]
AMDGPU: Don't include stack size of chain callees
Chain callees are always tail called into and overwrite the caller's
stack. Therefore, their stack size should not be added up to that of
their caller when computing the resource usage.
Assisted-by: AI
[clang][OpenMP] Match device and target_device ISA names separately (#232996)
A context selector can name ISAs in both `device={isa(...)}` and
`target_device={isa(...)}`. These are separate selectors, but their ISA
names were checked together, so an ISA that the target does not support
made both selectors fail. The default match_all needs both selectors
anyway, but the `match_any` and `match_none` extensions selected the
wrong variant. For example, on x86-64:
```c
#pragma omp declare variant(variant) \
match(device = {isa(sse2)}, target_device = {isa(bogus_isa)}, \
implementation = {extension(match_any)})
```
did not select variant, although `device={isa(sse2)}` matches. Likewise,
`match_none` accepted a variant whose `target_device` ISA matched.
This patch records which selector each ISA name belongs to and check
[6 lines not shown]
[LSR] Refine target-specific equality loop exits
Query target interest before computing the SCEV exit count.
Also cover samesign comparisons and remove a redundant IV operand check.
Assisted-by: TRAE CLI (GPT-5.6-Sol)
[AtomicExpand][NVPTX] Add single-lane atomic RMW coverage
Record current expansion and code generation before changing single-lane RMW legalization. Check the full pipeline separately from direct instruction selection while their output differs.
[LoongArch] Implement `computeKnownBitsForTargetNode()` cases for `V{ALL,ANY}_{,NON}ZERO` (#230799)
Currently, `emitVecCondBranchPseudo()` loads an constant 0/1 into the
destination GPR for `Pseudo{,X}VB{Z,NZ}{,_B,_H,_W,_D}`:
https://github.com/llvm/llvm-project/blob/762316e293d1db4f495bdbb545c31d588a1a2d62/llvm/lib/Target/LoongArch/LoongArchISelLowering.cpp#L9099-L9112
Yet their corresponding `SDNode`s, `V{ALL,ANY}_{,NON}ZERO`, are not
marked with matching range information in
`computeKnownBitsForTargetNode()`. Doing so elides unnecessary sign
extensions and could also help pattern matching in a later optimization.
[CIR][AMDGPU] Honor __launch_bounds__ on HIP kernels (#233929)
Map max threads to amdgpu-flat-work-group-size and min blocks to
amdgpu-waves-per-eu matching classic codegen
[LoongArch] Lower vector sign extension with VSLLWIL and VEXTH
vsllwil.{h.b/w.h/d.w} with a zero shift amount sign-extends the low
half of the source elements, and vexth.{h.b/w.h/d.w} the high half.
Use them for the custom SIGN_EXTEND legalization and for
SIGN_EXTEND_VECTOR_INREG, instead of computing the sign mask with vslti
and interleaving it with vilvl/vilvh, which takes one more instruction
per step. SIGN_EXTEND_VECTOR_INREG now also handles 4x extensions, which
previously fell back to a shuffle followed by a vslli/vsrai pair.
[mlir][arith] Honor finite min/max reduction identities (#232257)
I hit this while compiling a JAX transformer block through IREE/WebGPU
in JupyterLite. `linalg.softmax` requests a finite reduction
initializer, but `getIdentityValueAttr` returns NaN for `maxnumf`, which
Tint rejects.
Honor `useOnlyFiniteValue` for `maxnumf` and `minnumf`, preserving the
existing NaN identities when the flag is false. Update the existing
softmax check and add finite min/max reduction cases.
