[CodeGen] Drop dead SlotIndexes before allocation
SlotIndexes keeps the index list entry of an erased instruction and only
clears its instruction pointer. Live range sizes are measured in slot
indexes and greedy ranks ranges by size, so the leftovers inflate some
ranges more than others and reorder allocation, spilling heavily on
register-starved functions.
Add SlotIndexes::compactIndexes() to erase them. Erased entries are
unlinked, so holding an index to one is now a bug; linkedEntry() asserts
on it. The one analysis that did was LiveIntervals, so add
removeUnreferencedIntervals() to retire the intervals responsible.
Off by default behind -greedy-compact-slot-indexes, since it changes
allocation across much of the test suite.
AMDGPU: Reject 16-bit physical register inline asm constraints without true16 (#235063)
Constraints naming 16-bit physical registers such as {v1.l}, {v1.h},
{s1.l} or {a1.l} were accepted on subtargets without true16, which
then hit an assertion in copyPhysReg. Report the constraint as
unallocatable instead.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
WebAssembly: Use reportFatalUsageError for eh option errors (#233031)
Avoid using the deprecated report_fatal_error, and add missing
error test coverage.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[SimplifyCFG] Fix profiles in hoistCommonCodeFromSuccessors (#234833)
We cannot keep branch weights here given the select condition/control
flow condition we're hoisting across might be dependent.
Found in a bootstrapping build of LLVM with -mllvm
-enable-profcheck-in-pipeline.
[SimplifyCFG] Propagate unknown profiles in performValueComparisonIntoPredecessorFolding (#234827)
Otherwise we get profcheck verification failures as the verifier assumes
we just did not handle this case.
This came up in a bootstrapping build of LLVM with -mllvm
-enable-profcheck-in-pipeline.
[IndVars] Synthesize profile information for urem->select folding (#231990)
In these cases we can derive the profile information from the loop trip
count assuming a couple of assumptions are met.
[MemProf] Fix remapping of matching calls when assigning function clones
When a function has several calls with identical callsite metadata, they
are recorded on one context node with the extra calls in
Node->MatchingCalls. If that function is cloned by context
disambiguation, assignFunctions must remap each version's matching calls
through the function clone's CallMap, the same way it remaps the primary
Call.
Two problems in the IR path:
* AssignCallsiteCloneToFuncClone rewrote the matching calls of the
original node (Node->MatchingCalls) instead of those of the callsite
clone being assigned. The clone node (whose MatchingCalls were copied
from the original at creation) kept the original instructions, and the
original node's list was overwritten with the clone's. UpdateCalls then
updated each function version's matching calls with the other version's
callee clone, i.e. swapped.
[16 lines not shown]
PowerPC: Stop adding kill flags when promoting instructions to 64-bit
Continue purging kill flags from virtual register cases
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
LoongArch: Stop adding kill flags in LoongArchPreRAExpandPseudo
Continue purging kill flags on virtual registers.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[GitHub] Use client-id in require-team-membership (#235060)
app-id is deprecated and presumably going away soon, and client-id is a
drop-in replacement for our needs.
[CIR] Remove 'dead' test that no longer has a purpose. (#235069)
The patch #225814 stops us from rewriting unreachable blocks, and the
purpose of the original patch that wrote this test (#229191) intended to
make sure we handled unreachable blocks correctly.
So this test no longer matters. I left the original code in 229191 as it
is reasonable defensive code.
[libc] Suppress GCC's <mm_malloc.h> in freestanding builds (#234863)
Unlike Clang, GCC's <xmmintrin.h> (included via <immintrin.h>)
unconditionally includes <mm_malloc.h> even with -ffreestanding, which
in turn includes <stdlib.h>.
In full-build mode, targets that use x86 intrinsics (such as op_x86.h)
don't declare a dependency on libc.include.stdlib (and they don't really
need it). If these targets are built before our stdlib.h header is
generated, this will end up including the system stdlib.h -- and fail
miserably because of mixing of standard libraries. This can happen
already, but the problem gets more acute if we remove (#234150)
transitive full-header dependencies from the proxy headers (as it lets
us build these targets much sooner).
Since we don't use _mm_malloc or _mm_free, I'm suppressing <mm_malloc.h>
by defining its include guard (_MM_MALLOC_H_INCLUDED) on x86/x86_64 GCC
full builds.
Assisted-by: Gemini
RuntimeLibcalls: Split __atomic_* into a separate libatomic library
The __atomic_* functions are not part of libgcc or the default
compiler-rt builtins. GCC ships them in libatomic, and compiler-rt in the
optional clang_rt.atomic. Every system that uses compiler-rt also lists
libatomic, so the available functions are unchanged.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
WebAssembly: Stop adding kill flags on virtual registers (#234968)
Continue purging new kill flags on virtual registers.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[OpenMP][offload] Trade threads for blocks in Generic-SPMD kernels
The threads of a Generic-SPMD block run the inner parallel loops, whose
trip counts are unknown. Halve the threads down to one warp while the
distribute loop has an iteration for every block, and default to 6
blocks per CU on AMDGPU (LIBOMPTARGET_AMDGPU_GENERIC_SPMD_TEAMS_PER_CU).
libomptarget reads the number of compute units and the platform backend
of the device through liboffload, as part of its geometry properties.
This is an upstream port of the corresponding downstream.
Claude assisted with this patch.
[OpenMP][offload] Size the grid of cross-team reductions by their size
Every block adds to the cost of combining the partial results of a
cross-team reduction. Launch one block per compute unit for reductions
that work with a reduction data size of more than 32 bytes, and enough
blocks for a per-device number of waves per compute unit (16 on AMDGPU)
for the smaller reduction data sizes, capped by the trip count. Such
kernels keep their number of threads: the low-trip-count reduction of
threads does not apply to them. libomptarget reads the number of compute
units and the platform backend of the device through liboffload, as part
of its geometry properties.
This is an upstream port of the corresponding downstream.
Claude assisted with this patch.
[OpenMP][offload] Size only reductions that fill the device by their size
A cross-team reduction whose loop is too short for one block per compute
unit now gets the grid of a kernel without a reduction. In SPMD mode,
that is the low-trip-count rule, which lowers the number of threads per
block to spread the iterations over more compute units. Keeping all
threads made such reductions, e.g. one of 1792 iterations, up to 2.5x
slower on MI210.
Claude assisted with this patch.
RuntimeLibcalls: Specify system libraries by name (#233569)
Each system is now described by a list of default library names.
Per-target exclusions become standalone LibraryExclusion defs.
Co-Authored-By: Claude Opus 5.5 <noreply at anthropic.com>
[CIR] Set dead_on_return for non-byval indirect arguments (#234773)
This change updates the CIR calling convention rewrite context to set
the dead_on_return attribute when lowering non-byval indirect arguments
to match the behavior of classic codegen.
Assisted-by: Cursor / various models
