[ObjectYAML][NFC] Hoist ContiguousBlobAccumulator into a shared header (#207306)
Move ContiguousBlobAccumulator out of ELFEmitter.cpp's anonymous
namespace so other yaml2obj emitters can reuse it.
[RFC][AMDGPU][lld] Add object linking support
Add AMDGPU ELF object-linking support in lld, including resource propagation,
LDS layout, indirect-call handling, named-barrier updates, target compatibility
checks, and kernel descriptor/metadata patching.
This is a large PR because the linker needs to understand and validate several
AMDGPU object-linking concepts end to end. I tried to keep the changes scoped to
the necessary linker support and related metadata plumbing, but I'm open to
suggestions on how to split or structure the review to make it easier.
[AMDGPU] Introduce ABI occupancy for object linking
This PR introduces ABI occupancy as the contract used to compile functions under
object linking. The default is derived from the occupancy needed for a 1024
workitem workgroup, can be overridden with `-amdgpu-abi-waves-per-eu`, and can
be overridden per function by `amdgpu-flat-work-group-size` or tightened by an
accepted `amdgpu-waves-per-eu` hint.
The backend emits the selected occupancy in `.amdgpu.info` and uses it to
enforce the object linking register budget.
[docs][AMDGPU] move DMA operations to a separate file (#206917)
This includes the various `.async` and `.tensor` operations. Various
scattered bits of information are now in one place. Also added the
missing documentation for tensor intrinsics.
[NewPM] Port StaticDataProfileInfo and wire into x86 AsmPrinter
Ports the analysis to the NewPM and wires it into the NewPM x86
AsmPrinter where we are currently just returning nullptr.
Reviewers: mingmingl-llvm, teresajohnson, arsenm, RKSimon
Pull Request: https://github.com/llvm/llvm-project/pull/207550
[MLIR][NVVM]Import llvm intrinsics for nvvm.barrier (#202862)
This PR adds LLVM IR import support for the `bar.sync` intrinsic family,
lifting the four variants into `nvvm.barrier` with the appropriate
`aligned` attribute and optional thread-count operand.
`LLVM_IntrOpBase` and its auto-generated import is one to one, but four
`bar.sync` variants are many-to-one, where they all lift into
`nvvm.barrier`, differing only in `aligned` and the optional
`numberOfThreads` operand.
An alternative implementation would extend mlir-tblgen with many-to-one
intrinsic-to-op mapping support so the conversion could live in
`NVVMOps.td`. But the change is bigger than manually import `bar.sync`
intrinsic family.
[M68K][GlobalISel] Remove dependency on legal ruleset (#207614)
This fills in always legal rules, to remove the dependency on the legacy
ruleset. This is not guaranteed to be all the rules, just the ones that
appear
in tests.
See #197308
[OpenMPOpt] Null generic-mode wrappers for SPMDized kernels
If a kernel gets SPMDized, it doesn't need the wrapper function that is
passed to __kmpc_parallel_60. Keeping the dead wrapper function can lead
to lots of misleading "local memory global used by non-kernel function"
AMDGPU backend warnings.
Let OpenMPOpt null the wrapper argument such that DCE can then remove
the corresponding dead functions.
Claude assisted with this patch.
[GenericDomTreeConstruction] Store the semidominator value in Label (#207603)
runSemiNCA's eval() stores the vertex with the minimal semidominator in
Label and dereferences NumToInfo[Label]->Semi to compare it. Store the
semidominator (Semi) value directly in Label instead, so eval compares
by number with no NumToInfo lookup.
[ELF] Add target-specific relocation scanning for SPARCV9 (#206284)
Implement SPARCV9::scanSectionImpl, following the pattern established
for x86 and other targets. This merges the getRelExpr and TLS handling
for SHF_ALLOC sections into the target-specific scanner, enabling
devirtualization and eliminating abstraction overhead.
- Inline relocation classification into scanSectionImpl with a switch
on relocation type, replacing the generic rs.scan() path.
- Use processR_PC for PC-relative relocations and processR_PLT_PC for
PLT relocations.
- Handle TLS LE relocations directly with checkTlsLe.
- Simplify getRelExpr to only handle relocations needed by
relocateNonAlloc and preprocessRelocs.
[LifetimeSafety] Support container interior paths and invalidations
This patch completes the implementation of path-sensitive lifetime tracking by supporting container interior paths (`.*`) and deep-nested invalidation.
- Enables `PathElement::getInterior` generation in `FactsGenerator` for GSL Owners and Views (e.g. member functions, function parameters, lambda captures).
- Removes bypass checks in `FactsGenerator::handleInvalidatingCall` to track container invalidation on fields.
- Updates `Checker` to use strict prefix comparison (`isStrictPrefixOf`) for container invalidations, ensuring invalidation of container contents (interior) correctly invalidates iterators but not other sibling fields.
- Reorganizes tests in `invalidations.cpp` by resolving duplicates and distributing them logically.
- Updates unit tests and sema tests with correct expectations for interior paths.
TAG=agy
CONV=2cfd8d00-18d7-4a03-8d78-2aba2f9a8f23
[LifetimeSafety] Support container interior paths and invalidations
This patch completes the implementation of path-sensitive lifetime tracking by supporting container interior paths (`.*`) and deep-nested invalidation.
- Enables `PathElement::getInterior` generation in `FactsGenerator` for GSL Owners and Views (e.g. member functions, function parameters, lambda captures).
- Removes bypass checks in `FactsGenerator::handleInvalidatingCall` to track container invalidation on fields.
- Updates `Checker` to use strict prefix comparison (`isStrictPrefixOf`) for container invalidations, ensuring invalidation of container contents (interior) correctly invalidates iterators but not other sibling fields.
- Reorganizes tests in `invalidations.cpp` by resolving duplicates and distributing them logically.
- Updates unit tests and sema tests with correct expectations for interior paths.
TAG=agy
CONV=2cfd8d00-18d7-4a03-8d78-2aba2f9a8f23
[GlobalISel] Remove LegacyLegalizerInfo (#197308)
After #197238, #197370, #197371, #197374, #197375, #197377, #197378 and
#197379, this removes the legacy rules from global isel.
The above patches attempt to make all _tested_ operations legal, but
have only gone through the operations that have llvm tests. If more
fallbacks are now found to be happening, the other operations from
#197238 might be needed.
[libc++] Add an ABI flag to optimize mersenne_twister_engine (#206423)
Instead of updating the status array every time `operator()` is called,
we can instead update it once every time the entire array has been read.
That allows the compiler to vectorize the updating code, improving
performance.
Apple M4:
```
Benchmark old new Difference % Difference
---------------------------------------- ------- ------- ------------ --------------
std::mt19937::operator() 1.39 0.65 -0.74 -53.26%
std::mt19937::operator()_(into_array) 172.92 80.69 -92.23 -53.34%
std::mt19937_64::operator() 1.46 0.99 -0.47 -31.99%
std::mt19937_64::operator()_(into_array) 187.31 108.10 -79.21 -42.29%
```
Fixes #197221
[libc++] Use std::to_chars to format thread::id and canonicalize the representation across platforms (#181624)
This makes printing `thread::id` faster, since we avoid a bunch of
boilerplate code that isn't optimized away. It also avoids including
`<sstream>` and `<__locale>`, cutting the time to parse `<thread>` in
half.
This also changes the output when printing the id on a platform that
uses a pointer as the underlying type. I don't think that's a problem,
since the thread id isn't in any way stable. It also makes the output
consistent across all platforms we support.
[libcxxabi] Re-enable `thread_local_destruction_order.pass.cpp` on FreeBSD (#186129)
The underlying bug has been fixed in https://reviews.freebsd.org/D55826.
[InstCombine] Fold fpto{u,s}i of int-cast fdiv into {u,s}div (#205853)
Fixes #205305.
Adds an InstCombine fold for the pattern `fpto{u,s}i (fdiv ({u,s}itofp
X), C)` to `{u,s}div X, C`.
Safe when
- Unsigned: C > 0 and the integer width N <= the FP mantissa width p.
- Signed: C != 0 and N - 1 <= p, excluding (X == INT_MIN, C == -1).
See linked issue for detailed reasoning.
