[yaml2obj] Migrate to OptTable (#224964)
yaml2obj is one of three tools that parse with
cl::ParseCommandLineOptions(LongOptionsUseDoubleDash=true). Parse its
options with an OptTable instead, so that the mode can be removed from
cl::. Spellings follow the other migrated tools: -D<macro>=<v> and
-D <macro>=<v>, -o<file> and -o <file>, --docnum=<n> and --docnum <n>,
--max-size likewise, -E, -h/--help and --version. `-o=<file>` is no
longer accepted; the tests that used it are updated. `--help` lists
only yaml2obj's options.
LLM-aided
[lldb] Add plugin for Clang offload bundle object container (#222362)
## Summary
Clang supports embedding binaries using the clang offload bundler:
https://clang.llvm.org/docs/ClangOffloadBundler.html
This patch creates an ObjectContainer plug-in that allows us to get the
contained binaries within any executable.
### why we need this?
AMD/HIP host binaries may contain a .hip_fatbin section holding a Clang
offload bundle with one or more AMDGPU ELF code objects, typically for
different GPU architectures. Each ELF code object can contain multiple
ml kernels.
### Test
[11 lines not shown]
[FMV][AIX] Implement target_clones part 2 (target-features) (#206786)
This is part 2 (of 2) of FMV on AIX that would allow specifying
target-feature strings on the ``target_clones`` attribute. Part 1 was in
#177428.
The list of attributes supported is a subset of the attribute accepted
on the `target` attribute on AIX, specifically those that are runtime
detectable via the `__builtin_cpu_supports` builtin.
---------
Co-authored-by: Wael Yehia <wyehia at ca.ibm.com>
[DirectX] Support `cmpxchg` on buffer resources in DXIL lowering (#222162)
This PR is a backend change alone. It teaches DXIL lowering to handle a
`cmpxchg` on a buffer resource pointer, which the
`InterlockedCompareStore`
and `InterlockedCompareExchange` PRs above it in this stack need.
The PR adds:
- the `llvm.dx.resource.atomic.compare.exchange` intrinsic,
- its DXIL operation and op lowering,
- the `DXILResourceAccess` handling that turns a `cmpxchg` on a value
returned by `llvm.dx.resource.getpointer` into that intrinsic.
DXIL has no atomic compare exchange for texture resources, so the pass
reports an error for them. A test covers that path.
There is no Clang change here, and no HLSL function becomes available.
Assisted by: Github Copilot
[HLSL] Add `InterlockedExchange` function and resource methods (#222161)
This PR adds the `InterlockedExchange` standalone function and resource
methods. The operation lowers to `atomicrmw xchg`.
`InterlockedExchange` is the first interlocked operation that reports
the
previous value. It takes that value through an `original_value`
reference
parameter, so this PR adds a `RequiresOriginalValue` flag to
`HLSLBuiltinTypeDeclBuilder` and a matching overload shape to
`HLSLExternalSemaSource`. The later compare operations reuse both.
The float overload of `InterlockedExchange` follows in a separate PR in
this
stack. This PR covers the integer overloads alone, so it links
https://github.com/llvm/llvm-project/issues/99129 without closing it.
Assisted by: Github Copilot
AArch64: Mark the LR def of GlobalISel calls dead
In SelectionDAG InstrEmitter would have marked this dead.
Co-Authored-By: Claude Opus 5 <noreply at anthropic.com>
[docs] Require MyST for Sphinx documentation builds (#223829)
Fedora plans to use the separate llvm_man_pages package that we build as part of our release process now:
https://discourse.llvm.org/t/rfc-make-myst-markdown-the-llvm-docs-format-rip-rest/90840/37
I'm assuming that other distributions are either going to be able to
install the necessary Python modules that we depend on (myst_parser
etc), or will use this man page package if they need it. Therefore, I think we
can remove this enum and associated complexity.
Assisted-by: codex
[mlir][bufferization] Relax happens-before analysis (#223777)
On unstructured control-flow, the one-shot analysis is rather
conservative in regards to computing RaW conflicts, requiring strict
dominance which may not always be satisfied.
Relax this requirement to allow more cases to not be marked as
conflicting. Notably, readers/writers within a region that contains a
"diamond" or multiple loops are now handled.
This is a simpler version of
https://github.com/llvm/llvm-project/pull/223006 without the CFGLoopInfo
analysis, meaning it is a bit more conservative when analyzing regions
with nested loops.
[InstCombine] Fold variable extract of a constant-strided pointer vector (#217078)
A vector built from pointers that all point into the same object at a
constant stride is a table of addresses, so a variable extract from it
is address arithmetic:
%v0 = insertelement <4 x ptr> poison, ptr %b, i64 0
%v1 = insertelement <4 x ptr> %v0, ptr %b + 8, i64 1
%v2 = insertelement <4 x ptr> %v1, ptr %b + 16, i64 2
%v3 = insertelement <4 x ptr> %v2, ptr %b + 24, i64 3
%r = extractelement <4 x ptr> %v3, i32 %i
-->
%r = getelementptr i8, ptr %b, i64 %i * 8
Fold it to a single GEP. Every element must be reachable from one base
pointer at a constant byte offset, every lane of the vector must be
defined, and the offsets must form an arithmetic sequence; otherwise the
pattern is left alone.
[CodeGen] Use validated wave counts for AMDGPU spill costs
Lane-based block frequencies can understate the cost of a spill in
divergent GPU code: a wave still executes a block with only some lanes
active. Use measured block-wave counts to weight SpillPlacement's costs
relative to the original entry-wave count.
Opt in on AMDGPU only. Accept a measured count only when its IR block
maps uniquely to a machine block with matching predecessors and
successors. Keep the existing MBFI cost for unmeasured or rejected
blocks, and leave branch probabilities and general BFI unchanged.
[lld][MachO] Remove `__thread_ptrs` section so TLV can relocate against dynamic-lookup symbols (#221364)
lld/MachO rejects a relocation whenever `hasAttr(TLV) != sym->isTlv()`, but
`isTlv()` cannot separate "not thread-local" from "unknown", so that one
inequality rejects three unrelated cases and only one of them is ill-formed.
C++20 `constinit` made another of them common: it lets the compiler elide the
thread-local wrapper and access a cross-TU object directly ([folly 2c666105](https://github.com/facebook/folly/commit/2c666105e9a694aef7290cce3da14ed1ebf14462)),
producing a TLV relocation against a dynamic-lookup symbol. The check was also
masking a defect — `Symbol::gotIndex` is one field indexing whichever of `__got`
or `__thread_ptrs` holds the symbol — and never enforced it anyway, since
`UnwindInfoSection` and the stub paths call `in.got->addEntry` directly, so a
thread-local personality already collided on stock lld.
Both pools hold the same value, the descriptor's address, and ld-prime never
synthesizes `__thread_ptrs` at all. Make `__got` the sole non-lazy pointer
section, deleting `TlvPointerSection`, `NonLazyPointerSectionBase`,
`in.tlvPointers`, `Symbol::getTlvVA` and the indirect-symbol-table
special-casing, and collapsing `resolveGotVA`/`resolveTlvVA` into
`resolveNonLazyPtrVA`. One section, one index, so the collision becomes
[9 lines not shown]
[VPlan] Compute execution probability via APFloat for printing. (#226135)
Similar to #201824 (8e517e1d17b1), use APFloat to compute the execution
probability. This ensures the output is independent of the underlying
libc/OS/architecture.
Suggested in https://github.com/llvm/llvm-project/pull/203982.
[SLP][NFC] Use BoUpSLP::getNumberOfParts() where appropriate (#225995)
Caches the results for better performance rather than directly calling
slpvectorizer::getNumberOrParts().
Small improvement in compile time.
[SPIRV] Support structurizing nested break chains (#222459)
This PR updates the SPIRVStructurizer to handle nested break chains.
This pattern occurs when switch statements are lowered instead as branch
instructions and break statements need to exit out to an outer
containing block.
This is a required change to support #112056.
Assisted-by: GitHub Copilot
[Flang][OpenMP][MLIR] Attach source location to target/target-data runtime c… (#222389)
…alls
This commit adds ident info to the OpenMP target call so the runtime now
reports the kernels source location in -g and non -g builds
Of note this does diverge from clang where the file info is only
provided with `-g` builds, however similar work is done in
`createMappingInformation` in
`mlir/lib/Target/LLVMIR/Dialect/OpenMPCommon.cpp`.
AMDGPU: Mark the SCC def dead when expanding CF pseudos
The control flow pseudos start with dead flags on the implicit scc def,
but did not transfer to the replacement instruction
Co-Authored-By: Claude Opus 5 <noreply at anthropic.com>
[obj2yaml] Migrate to OptTable (#224965)
obj2yaml is one of three tools that parse with
cl::ParseCommandLineOptions(LongOptionsUseDoubleDash=true). Parse its
options with an OptTable instead, so that the mode can be removed from
cl::. Spellings follow the other migrated tools: -o<file> and -o <file>,
--raw-segment=<segment> and --raw-segment <segment>, -h/--help and
--version. `--help` lists only obj2yaml's options.
LLM-aided
[split-file] Migrate to OptTable (#224966)
split-file is one of three tools that parse with
cl::ParseCommandLineOptions(LongOptionsUseDoubleDash=true). Parse its
options with an OptTable instead, so that the mode can be removed from
cl::. --no-leading-lines now overrides an earlier --leading-lines
rather than being ignored, and `--help` lists only split-file's
options. `--version` still exits 0, which libc++'s test configuration
relies on to detect the tool.
LLM-aided
[test] Use -o <file> for yaml2obj (#226245)
`-o <file>` is the preferred form. Prepare for the OptTable migration of
yaml2obj, where -o=<file> is no longer accepted.
[libc] Add the get and set sys/xattr.h entrypoints. (#224729)
This is a follow-up to #224357, implementing more of sys/xattr.h. As
before, these are non-POSIX APIs.
The list of entrypoints is: `fgetxattr`, `getxattr`, `lgetxattr`,
`fsetxattr`, `setxattr`, and `lsetxattr`.
Assisted-by: Automated tooling, human reviewed.
[lldb][NativePDB] Don't duplicate public symbols that restate a PE export (#226200)
`SymbolFileNativePDB::AddSymbols` adds every
[`S_PUB32`](https://llvm.org/docs/PDB/CodeViewSymbols.html#s-pub32-0x110e)
to the symtab, but `ObjectFilePECOFF` has usually already added the same
name and address from the PE export table, so the symbol resolves to two
entries.
On the other hand, `ObjectFilePECOFF` already solves this for its own
`export`/`symtab` overlap by marking the duplicate entry as
`eSymbolTypeAdditional`:
https://github.com/llvm/llvm-project/blob/d1106deb71cc81016406a1917d0508d2df296266/lldb/source/Plugins/ObjectFile/PECOFF/ObjectFilePECOFF.cpp#L839-L844
This patch implements the same for `SymbolFileNativePDB::AddSymbols`.
In swiftlang, this fixes
`lang/swift/expression/exclusivity_suppression/TestExclusivitySuppression.py`
on Windows.
[llvm-cov] Migrate gcov to OptTable (#224955)
llvm-cov gcov is the last user of cl::Grouping. Parse its options with
an OptTable that enables grouped short options, as the other
binutils-style tools do, so that cl::Grouping can be removed.
Spellings follow the other migrated tools: long options take `--` only,
so `-gcno` and `-gcda=` become `--gcno` and `--gcda=`. `--help` now
lists only gcov's options rather than every cl:: option linked into
llvm-cov, and an unknown option or a missing source file is reported
as `llvm-cov gcov: error: ...`.
LLM-aided
llvm-reduce: Error on -j with MIR inputs (#226225)
Parallel chunk processing distributes work by serializing the program to
bitcode and reparsing it in each worker thread. Bitcode cannot represent
MachineFunctions, so readBitcode leaves ReducerWorkItem::MMI null and
the reparsed item reports isMIR() == false. The delta pass list is still the
MIR one, chosen from the original program, so the worker runs a MIR pass
over an item with no MachineModuleInfo and dereferences null.
This broke every MIR reduction using -j > 1, though not immediately: the
parallel path is only taken once there is more than one chunk left to
process, so the crash appeared after the first granularity increase.
Reject the combination up front instead of crashing partway through a
reduction. Supporting it would mean exchanging MIR text rather than
bitcode, and giving each worker its own TargetMachine.
Co-Authored-By: Claude Opus 5 <noreply at anthropic.com>
[flang][cuda][openacc] Resolve DEVICE dummies against !$acc data mapped objects (#225899)
CUDA Fortran generic resolution rejects a call when an actual argument
with
no CUDA data attribute is associated with an ATTRIBUTES(DEVICE) dummy.
That
makes cuBLAS-style generic interfaces unusable inside a structured
OpenACC
data region, even though the mapped objects do have a device copy there:
```
!$acc data copyin(a, b) copyout(c)
call cublaszgemm(transa, transb, m, n, k, alpha, a, lda, b, ldb, beta, c, ldc)
!$acc end data
error: No specific subroutine of generic 'cublaszgemm' matches the actual
arguments
dummy argument 'a=' has ATTRIBUTES(DEVICE) but its associated actual
argument has no CUDA data attribute
[32 lines not shown]
[SystemZ][z/OS] Fix product version format in PPA2 on z/OS. (#225749)
The product version in PPA2 should be right-aligned (e.g. 090807)
instead of left-aligned (e.g. 908070) on z/OS.
