[libc][complex] Enable cabs and cabsf in libc (#206677)
This PR adds basic implementation for cabs and cabsf in LLVM libc and
also extend current MPC infrastructure to support cabs operation.
---------
Signed-off-by: jinge90 <ge.jin at intel.com>
[SPARC] Make SPARC_REGISTER Sparc specific in llvm-readobj and LLDB (#206316)
The commit 4418a8e5 failed to limit SPARC_REGISTER to Sparc targets
only, which causes conflict when adding new dynamic tags for other
targets.
[NewPM] Port StaticDataAnnotator
Standard NewPM pass porting. The PassBuilder wireup/test coverage will
come in the next PR.
Reviewers: arsenm, teresajohnson, mingmingl-llvm
Pull Request: https://github.com/llvm/llvm-project/pull/207629
[NewPM] Port StaticDataSplitter
Standard pass porting. Remove the MBPI member given it was unused.
Reviewers: arsenm, teresajohnson, mingmingl-llvm
Pull Request: https://github.com/llvm/llvm-project/pull/207627
[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.