[flang-rt] - Reduce ShallowCopy template instantiations to improve LTO time (#209915)
Limit `ShallowCopy` rank specializations from 1-15 to 1-4, using a
generic fallback for higher ranks. This significantly reduces code size
and LTO compile time when the Fortran runtime is linked into GPU
offloading code.
Before: 15 ranks × 7 types × 3 scenarios = 315 template instantiations
After: 4 ranks × 7 types × 3 scenarios = 84 template instantiations +
fallbacks
Trade-off: Arrays with rank > 4 use a generic runtime loop instead of
compile-time specialized iteration. This is acceptable because:
- Most real-world Fortran arrays are rank 1-3
- Rank 5-15 arrays are rare in practice
- The generic loop is still efficient (just not unrollable)
## Results
### Compilation Time (from `time` command)
[35 lines not shown]
[flang][cuda] Use cuf.register_variable_static instead of CUFRegisterExternalVariable (#211886)
This gives more flexibility to the backend to select a different
registration mechanism
[Pipeliner] Use VRMapPhi to generate phi in epilog (#211723)
When generating a phi in the epilog block corresponding to an existing
phi in the loop, the loop value should come from the new phi generated
in the kernel block (for a non-phi instruction in the loop) rather than
directly from the value map of the previous stage.
The new phis generated by generatePhis for non-phi instructions are
stored in VRMapPhi. Thread VRMapPhi through generateExistingPhis and add
a helper getMapPhiReg that returns the phi-generated register when
available, falling back to VRMap otherwise.
Without this fix, the epilog can pick up the initial prolog value of a
loop-carried register instead of the last kernel-iteration value,
producing incorrect results for pipelined loops with two or more stages
that carry values across iterations. This has been observed on Hexagon
with the modulo scheduled epilog of a right-shift-with-carry loop after
loop unrolling.
Fixes #208943
[AMDGPU] Fix instruction size of LDS-DMA buffer loads (#211302)
Mark i1imm_0 as OPERAND_IMMEDIATE type, so its packed 1-bit value is not
mistaken for a trailing literal (same for InstFlag).
Applies to the swz/IsAsync fields of buffer (MUBUF) and FLAT LDS-DMA
instructions.
Changes estimated size for the corresponding instructions; codegen stays
untouched.
Add unit test, checking encoded size is 8 bytes (was: 12).
AI-assisted.
[SCEV] Use SCEVPtrToAddr instead of SCEVPtrToInt in SCEV. (#180244)
Use SCEVPtrToAddr instead of SCEVPtrToInt for pointer -> int
conversion in SCEV, like when computing trip count expressions.
getPtrToIntExpr now also creates SCEVPtrToAddrExpr.
Removal of SCEVPtrToIntExpr out of this PR, to limit the scope and
will be done as NFC follow-up.
PR: https://github.com/llvm/llvm-project/pull/180244
DAG: Stop legalizing poison to undef
Previously all targets would expand poison to undef, which was a
shortcut taken when poison was introduced.
Introduces a new dedicated poison matcher in tablegen; this avoids
multiple opcode changes, and avoids some problematic contexts that
couldn't deal with PatFrags.
This probably should be legal by default for all legal types.
Co-authored-by: Claude (Claude-Opus-4.8)
AMDGPU: Avoid dropping poison when extracting a high element (#211872)
Haven't found a case where this matters, but I want to stop finding
this instance of the is-undef-get-undef pattern.
AMDGPU: Don't pack a build_vector with an undef/poison lane (#211871)
Once ISD::POISON is legal it reaches isel and is treated as canonical, so
is_canonicalized_2<build_vector> matched build_vector<x, poison> and
selected a v_pack for a don't-care high lane. Avoids regressions in future
patches.
Co-authored-by: Claude (Claude-Opus-4.8)
[clang][Fuchsia] Add __Fuchsia_Compiler_ABI__ predefine (#211744)
This establishes a new canonical predefine for *-fuchsia targets:
`__Fuchsia_Compiler_ABI__` is set to an integer for the version
number of the Fuchsia Compiler ABI the compiler supports. The
value will change as appropriate with changes to *-fuchsia target
ABI expectations. It is not anticipated that the compiler will
support selecting target ABI versions, only the latest as the
canonical target-specific behavior. Fuchsia Compiler ABI
versions are unrelated to the Fuchsia API levels supported by
Fuchsia systems; newly-built programs may target older API levels
with `-ffuchsia-api-level`, but everything relying on the Fuchsia
Compiler ABI is built and deployed together as an isolated unit.
[mlir][xegpu] Allow create_mem_desc from ND memref (#211836)
Relax the create_mem_desc source operand constraint to accept a
statically shaped shared-memory memref of any rank, replacing the
1D/2D-only StaticShared{1,2}DMemRefOf classes with a rank-agnostic
StaticSharedMemRefOf.
Add a verifier requiring the source memref to be contiguous row-major,
update the op documentation, and add valid/invalid lit tests.
assisted-by-claude
Co-authored-by: Claude Opus 4.8 <noreply at anthropic.com>
[flang][OpenMP] Fix alias analysis for omp.private copy region block arguments to help inline hlfir.assign to improve time taken in LTO. (#209539)
## Summary
This PR fixes https://github.com/llvm/llvm-project/issues/200922 - a
compile-time regression affecting firstprivate clauses on simple arrays
in OpenMP target regions.
When `InlineHLFIRAssign` checks whether to inline `hlfir.assign`
operations in `omp.private` copy regions, the alias analysis
conservatively returns `MayAlias` for the copy region block arguments
(`%arg0` mold vs `%arg1` private). This prevents inlining and forces a
fallback to `fir.call @_FortranAAssign`, which:
- Pulls in 89 runtime functions via LTO
- Creates 550K+ abstract attributes in OpenMPOpt
- Causes severe compile-time degradation (60%+ increase in some cases)
**Example:**
[32 lines not shown]
[libc++] Rename pre-commit CI workflows (#211844)
Since we're expanding libc++'s pre-commit CI to add other types of
tests, (e.g. tools tests, performance tests, etc), it makes sense to be
a bit more precise about what each workflow does, and to use a
consistent pattern across workflow names.
Also, update stale references to old names in the documentation and in
some infrastructure scripts.
[libc++][CI] Add a Github workflow to test libc++ tools (#211067)
We have a few scripts under libcxx/utils, and they wouldn't be tested on
a regular basis. As a result, changes to e.g. the build system could
break these scripts, which in turn could break some Github workflows
like the libc++ PR benchmarking workflow.
This patch adds a new `test-tools` configuration to `run-buildbot` which
allows testing these tools locally, and a new dedicated Github workflow
to run it.
clang/AMDGPU: Migrate cc1 tests to subarch triples (5) (#211840)
Rewrite OpenMP cc1 test RUN lines to the new subarch triple form,
dropping the redundant -target-cpu, and regenerate autogenerated CHECK
lines. Tests that intentionally assert the "target-cpu" attribute are
left unchanged.
Co-authored-by: Claude (Opus 4.8) <noreply at anthropic.com>
AMDGPU: Reject unrecognized amdgpu subarch in TargetID parsing (#211866)
A bare "amdgpu" or an unrecognized "amdgpu<x>" both parse to the
amdgpu arch with no subarch; TargetID::parse previously accepted them,
taking a named processor and ignoring the missing subarch.
Reject triples with no subarch, unless it uses the legacy spelling.
This is another workaround for Triple not having a proper unknown value
for subarches.
Co-authored-by: Claude (Opus 4.8) <noreply at anthropic.com>
[NVPTXAsmPrinter] Allow self-referential device global initializers (#197838)
This change adds support for cyclic device global initializers in NVPTX,
including self-references and references between distinct globals. The
backend’s global-emission ordering walk previously treated every cycle
as an error and aborted with “Circular dependency found in global
variable set.” As a result, valid patterns such as __device__ void
*arr[] = {(void *)"foo", &arr[0]}; were rejected. The NVPTX asm printer
now models initializer dependencies as a graph and emits its strongly
connected components in dependency-first order. Acyclic components are
emitted normally. For cyclic components, it first emits compatible PTX
.extern declarations and then emits the definitions in a valid order.
Cycles that cannot be represented this way, such as cycles containing
only static globals, are still rejected. The forward declarations reuse
the normal definition emitter so their address spaces, alignments,
types, and aggregate layouts match the definitions. Added new tests to
cover self-references, cycles between distinct globals, mixed linkage,
packed aggregates, dependency ordering, and unrepresentable static-only
cycles.
[MC][NFC] Remove some virtual function from MCAsmInfo (#211858)
The MCAsmInfo vtable is currently 144B (no anchor)/152B (anchor) large
-- and there are 49 of these in an all-target libLLVM.so. Some of these
are not needed or can be easily removed. This removes 5 vtable entries.
Some notes on the remaining functions -- I think this vtable can be
removed entirely without too much effort:
- getStackSection can also be removed, this is ELF-only; unfortunately
this will require code duplication between AsmPrinter and
MCELFStreamer.
- getExprForPersonalitySymbol is single-use and two-overrides (Darwin on
x86-64, AArch64), likely removable.
- shouldOmitSectionDirective has three override modes -- COFF (disallow
for COMDAT/unique section), AMDGPU (4 .hsa sections), and always-omit
(NVPTX, SPIRV). This is only called from printSwitchToSection, so COFF
[20 lines not shown]
[lldb] Recover a WebAssembly function's mangled name from its declaration (#211707)
The Wasm "name" section stores only demangled names, so SymbolFileWasm
recovers each function's mangled name from the DWARF and attaches it to
the code symbol.
Currently, it only does so from a defining DIE, matched by address. When
we only have a declaration DIE (e.g. TU compiled without debug info), it
still carries the linkage name, but wouldn't get associated without the
address.
Match a declaration-only DIE to its code symbol by the demangled name
that the name section already carries, and attach the linkage name that
way.
Assisted-by: Claude
[OpenMP] Propagate PRESENT to pointee entries in mapper codegen
Extend map-type-modifier propagation in emitUserDefinedMapper to the PRESENT
modifier, but only for entries that have an attach ptr (the pointee data, whose
storage differs from the struct's own). A present modifier on the outer clause
must require that pointee to be present on the device.
This is gated on a new PropagatePresentToPointee argument, which Clang sets from
CGM.getLangOpts().OpenMP >= 60. Before 6.0 the present modifier is treated as
not applying to the pointee: the spec committee confirmed the divergence
between the present motion modifier (to/from) and the present map-type modifier
(map) was unintentional, to be fixed as an OpenMP 6.0 erratum. Only propagation
is gated; present written directly in a mapper's own clause applies at all
versions.
A TODO notes PRESENT should also propagate to the struct's own members, which
is blocked while pointer members use PTR_AND_OBJ.
Update the present-check tests to their final 6.0-gated behavior.
[2 lines not shown]
[ConstraintElim] use condition from latch in addInfoForInductions. (#210984)
Also use condition in latch to inject %iv <= %B and co into the header,
if the latch condition compares a post-inc IV like %iv + %step != %B.
This allows reasoning about loops that have been rotated before
ConstraintElimination, e.g. during the LTO phase or Swift where the
frontend already rotates loops before generating LLVM IR.
Alive2 Proof for injecting %iv <= B from %iv + %step != %B in latch:
https://alive2.llvm.org/ce/z/pH8Czb
This enables a number of additional simplifications in practice
(https://github.com/dtcxzyw/llvm-opt-benchmark-nightly/pull/736).
It comes with a compile-time impact when doing LTO, as all loops will be
rotated
and will hit the new code path.
* stage1-ReleaseThinLTO: +0.08%
[9 lines not shown]
[OpenMP][Clang] Enable ATTACH-style maps for mappers.
Track per-entry attach-ptr info (HasAttachPtr) through mapper codegen so that
emitUserDefinedMapper does not add a new outer MEMBER_OF to pointee/combined
entries (which occupy different storage than the struct) or to ATTACH entries.
Clang and the MLIR translator populate the per-entry array in parallel with the
other MapInfosTy arrays.
Address review:
- Rename MapSkipMemberOfArrayTy to MapHasAttachPtrArrayTy to match the
HasAttachPtr field it backs.
- Restructure the emitUserDefinedMapper comment into a bulleted (*)/(**)/(***)
list keyed to the example entries.
- Reword the Clang comments: HasAttachPtr marks pointee entries that have a
base attach-ptr; a combined entry has a base attach-ptr if its constituents
do; cross-reference emitUserDefinedMapper for the MEMBER_OF rationale.
- Update the moved present-check tests to their now-correct behavior (the
attach-style maps make the inbounds present checks pass and remove the
"explicit extension" errors).
[2 lines not shown]
[flang][OpenMP] Lower DO and SIMD variants in metadirectives
Lower DO, SIMD, and DO SIMD replacement directives selected by a
metadirective. Support standalone and begin/end forms with static or runtime
selection.
A standalone metadirective and its associated loop are represented as sibling
PFT evaluations. For example:
```fortran
!$omp metadirective &
!$omp& when(user={condition(flag)}: do) &
!$omp& otherwise(nothing)
do i = 1, n
a(i) = i
end do
```
has the following evaluation shape:
[54 lines not shown]
[lldb] Remove ConstString from FileSpec (#211394)
This commit completely removes ConstString from FileSpec, replacing it
with llvm::SmallString instead.
I considered combining the directory and the filename together into one
field, but then it became impossible to distinguish between a
partially-constructed FileSpec's last directory and a fully-constructed
FileSpec's filename.
The sizes of the SmallStrings are somewhat arbitrary. I tested out a few
other configurations on my machine locally and this yielded the best
memory/runtime tradeoffs.
I measured the impact of this change in two ways:
(1) Runtime performance
I measured the runtime impact by using LLDB's statistics with `stat
enable` and `stat dump -f` (forcing the parsing of symbols) for a debug
[15 lines not shown]
[Clang] Avoid querying tag definitions for invalid DeclSpecs (#210085)
Guard hasTagDefinition() against invalid type-specifier state so
recovery from invalid auto/tag combinations does not assert.
Fixes #210014
---------
Signed-off-by: Osama Abdelkader <osama.abdelkader at gmail.com>
[OpenMP] Propagate ALWAYS/DELETE/CLOSE map-type modifiers to mapper entries
When a map clause uses a user-defined mapper, the map-type-modifying bits
(ALWAYS, DELETE, CLOSE) on the outer clause must apply to each map the mapper
inserts (OpenMP 6.0:281:34). Propagate them in emitUserDefinedMapper by OR-ing
the imported modifier bits into each pushed component, except ATTACH entries
(ATTACH|ALWAYS is reserved for attach(always) and the other bits are
meaningless for ATTACH).
PRESENT is intentionally not propagated here yet (a TODO notes it is handled in
a follow-on, since it requires distinguishing pointee entries from the struct's
own storage).
Add tests:
- offload mapper_map_always_from.c: ALWAYS forces a member copy-back that the
ref count would otherwise suppress.
- offload Fortran target-custom-mapper-delete.f90: DELETE is propagated to the
mapper's member entry so it is removed by target exit data.
Co-Authored-By: Claude Opus 4.8 <noreply at anthropic.com>
