[BFI] Simplify irreducible header discovery. NFC (#213213)
IrreducibleGraph is the explicit graph of a region -- a loop, or the
whole function -- that BFI searches for irreducible SCCs. It maintains
predecessor lists only so findIrreducibleHeaders can ask of each node
whether a predecessor lies outside its SCC (an entry), and whether a
non-entry in the same SCC reaches it via a backedge (an extra header).
Answer both from the successor lists instead; IrrNode then needs only a
successor vector. findIrreducibleHeaders is left partitioning the SCC
its sole caller is packaging, so fold it into createIrreducibleLoop.
Headers and members are sorted, so the changed iteration order does not
affect output.
Once BFI uses CycleInfo, a region's irreducible SCCs are its
non-reducible child cycles, which should let IrreducibleGraph go away
entirely. That removal has no predecessor lists to walk and
getResolvedNode has no inverse, so it must sweep successors regardless;
settling the formulation here leaves it changing only the graph.
Aided by Claude Opus 5
[lldb] Launch the Wasm runtime with the file the target has (#213384)
PlatformWasm hands the runtime the module to run as a path on the host
it launches the runtime on. It takes that path from the launch info,
whose executable is the name the module goes by on the platform. That
name is whatever a stub reported the module under, which need not be a
path that resolves on this host, so a relaunch runs a file that does not
exist:
```
(lldb) run
error: WebAssembly runtime exited with exit code 255
```
Run the file the target has instead. Only a runtime launched on this
host is affected, since a connection to a remote Wasm platform delegates
the launch to that platform.
[mlir-c] Add RewriterBase insertion point save/restore (#206531)
Exposes `OpBuilder::saveInsertionPoint` / `restoreInsertionPoint` through the MLIR C API for `MlirRewriterBase`, continuing the buildout of the rewrite/conversion C bindings.
Assisted by: Claude
Disable this test on Darwin to give the author a chance to fix it. (#213388)
The test is passing elsewhere.
I filed:
https://github.com/llvm/llvm-project/issues/213386
to cover fixing the issue and re-enabling the tests.
[AMDGPU] Fix canonicalization of undef in v2f16 (#213373)
When processing canonicalization with an undef in v2f16,
1. If the other element is constant, create a splat vector
2. Otherwise, convert undef to 0.0.
Update comment to make this behavior clear. Do not imply that a v2f16
with two undef's should be unconverted. Do not state that conversion of
undef to qNan is normal. `ConstantFoldScalarCall1` in
`llvm/lib/Analysis/ConstantFolding.cpp` already canonicalizes an undef
to 0.0.
Previous code only converted element 0 if element 1 was a constant.
Signed-off-by: John Lu <John.Lu at amd.com>
Disable the tests in TestOSIndSYM.py (#213382)
I filed: github.com/llvm/llvm-project/issues/213380
to cover this issue. I'm disabling this one test till that is fixed.
AMDGPU/GlobalISel: Switch more FP opcodes to extended LLTs (part 4) (#213353)
Migrate G_FFREXP, G_FMODF, rounding operations, and the LOG/EXP families
to extended LLTs.
Remove f64 for G_FLOG and G_FLOG10 since there is no custom lowering for
them. Adjust G_AND masking for LOG/EXP using extended LLTs.
Update the relevant MIR tests and *.ll tests.
[libc] Move generated elf header to include (#211428)
The fact that `hdr/elf_proxy.h` was generated with headergen was causing
dependency issues. This shouldn't be necessary, and this PR moves it out
into `include` to avoid needing headergen for an overlay build.
Assisted-by: Automated tooling, human reviewed.
[CIR][AMDGPU] Add support for AMDGCN ldexp builtins (#198087)
Adds support for the following AMDGCN ldexp builtins:
- __builtin_amdgcn_ldexp (double)
- __builtin_amdgcn_ldexpf (float)
- __builtin_amdgcn_ldexph (half)
These are lowered to the `llvm.ldexp` intrinsic.
[DropAssumes] Print drop-deref in pipeline where appropriate (#211905)
drop-deref was added as an option in #166947 and is parsed correctly,
but is not serialized. This can make it difficult to reduce test cases,
especially automatically with llvm/utils/reduce_pipeline.py.
This patch implements printPipeline() to serialize drop-deref where
appropriate.
[OpenACC][NFC] Minor clean up in ACCRoutineLowering. (#213333)
Minor NFC clean up after recent changes to remove nohost handling from
ACCRoutineLowering.
Assisted-by: Codex
[IPO] Remove IR Outliner (#211971)
The IR Outliner has major bugs and no active maintainer, and is disabled
by default. The new LLVM Policy states the pass should be removed.
This commit removes:
- The IROutliner pass
- The IRSimilarity analysis
- The `llvm-sim` executable, used for understanding the latter
- All tests of the above
Related discussion:
https://discourse.llvm.org/t/ir-outliner-status-interest/89672
[flang-rt][cuda] Skip scope-exit cleanup when the context has a sticky error (#213184)
A sticky CUDA error (e.g. an illegal memory access in a kernel) leaves
the primary context active but unusable, so CUFDeviceIsActive() reports
it as fine and the compiler-generated scope-exit frees abort a program
that ran to completion: 'cudaFree(p)' failed with
'cudaErrorIllegalAddress'.
Detect this by freeing a null pointer, a no-op that still reports the
sticky error. It runs only once the primary context is known active, so
it cannot lazily create one.
[SLP]Keep scheduled order when moving body in alias-check versioning
The scheduler physically reorders the block before versioning, and
emission insertion points depend on that order. Moving the body by the
pre-scheduling snapshot scrambled it and produced use-before-def IR.
Move in the block's current order instead, filtered to the original
body so the emitted check instructions stay in the header block.
Reviewers:
Pull Request: https://github.com/llvm/llvm-project/pull/213358
[NVPTX] Pass `-Xcuda-ptxas` to the nvlink wrapper for LTO (#213351)
Summary:
This allows `-foffload-lto -fgpu-rdc -Xcuda-ptxas` to work properly by
forwarding it.
[lldb] Avoid returning a stale AddressOf (#212915)
The 'ValueObject::AddressOf()' method assumes that the address of a
value object cannot change, so when it is calculated once, it does not
need to be updated afterwards. However, this is not the case if the
'ValueObject' is a dependent object obtained by calling 'Dereference()'
of another 'ValueObject'. If the latter object is changed, the dependent
value object should return a new address from the 'AddressOf()' method
to reflect the change.
clang: Replace Is*OffloadArch free functions with OffloadArch methods
Drop the IsNVIDIAOffloadArch/IsAMDOffloadArch/IsIntel*OffloadArch free
functions in favor of the OffloadArch member predicate functions.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
clang: Store vendor GPU kinds in OffloadArch instead of re-listing GPUs
OffloadArch was a flat enum that hand-duplicated every AMDGPU and NVPTX
targets, plus a few edge cases. This was yet another place that needed
updating every time a new target is added, which should now be avoided.
Replace with a tagged union-like scheme.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
