[VPlan] Pass alignment explicitly to widen load/store recipes (NFC). (#237639)
Replace alignment lookup from underlying IR to explicitly storing it for
the recipe.
Preparation for fixing a miscompile with incorrect alignment when
narrowing interleave groups.
NFC except for the printed VPlans.
[VPlan] Store parent via ilist_parent, use iterator-only APIs (NFC) (#236659)
Store the parent VPBasicBlock of recipes in the ilist node via
ilist_parent<VPBasicBlock>, like Instruction/BasicBlock. This also gives
the list sentinel a parent, so any iterator into a recipe list,
including end(), knows its block.
Use this to simplify the APIs that took a block together with an
iterator, mirroring IRBuilder's removal of the separate insertion block.
PR: https://github.com/llvm/llvm-project/pull/236659
[X86] LowerRotate - directly use X86 non-uniform var-shift nodes for generic fallback (#237645)
Skip vector shift legalization and use the X86ISD::VSHLV/VSRLV nodes
directly to handle the RotAmt=0 case, so to correctly handle the out of
bounds shift by EltSizeInBits.
Fixes #236022
Attributor: Don't use callee return attributes for mismatched call types (#231629)
A call site returned position took the state and attributes of the
callee's returned position, even if the call uses a different return
type
than the callee, e.g. calling a void function as returning ptr. This
asserted in AANonNull::isImpliedByIR on the non-pointer callee return
type, and put the callee's nofpclass on calls returning another FP type.
Ignore the callee's returned position, and its returned arguments, if
the
types differ.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[SLP]Postpone small phi groups to build full register vectors
Small phi groups vectorized in the full-width stage took lanes from
wider vectors, producing 2-wide fma accumulators instead of 4-wide.
Postpone them, then try the candidates in prefix-first order at full
width, so phis from different groups with matching leading operands
get vectorized together.
Fixes #47830
Assisted-by: Cursor
Reviewers: RKSimon, bababuck
Pull Request: https://github.com/llvm/llvm-project/pull/234997
[TargetLowering] Don't call willNotOverflowSub in expandABD for ABDU with non-negative inputs. (#237338)
If the inputs are non-negative, then the difference between them
can never overflow so we don't need to call willNotOverflowSub.
[GlobalISel] Don't set disjoint flag in lowerFunnelShiftAsShifts if the shift amount may be undef. (#237061)
lowerFunnelShiftAsShifts uses isNonZeroModBitWidthOrUndef to choose the
X << C | Y >> (BW - C) expansion. Allowing undef there is fine for
picking the expansion, but it does not justify the disjoint flag on the
G_OR. If an undef shift amount is chosen to be 0 modulo BW, the two
operands of the or may have common set bits, and the disjoint flag would
make the result poison.
Only set the disjoint flag on this path if every element of the shift
amount is a constant that is non-zero modulo the bitwidth.
Co-Authored-By: Claude Opus 5.5 <noreply at anthropic.com>
[RISCV][P-ext] Split RV32 v2i32 riscv_pmerge code from v4i16 and v8i8. (#237358)
Lower using extract_vector_elt and build_pair rather than bitcasts and
extract_subvector. This is a more straightforward lowering and may
preserve more optimization opportunities.
Assisted-by: Claude
[compiler-rt][profile] Sum lane-varying GPU counter increments (#235978)
The GPU profiling helper adds each active lane's contribution to a
shared counter. Its `step` argument is the amount that lane contributes.
For an ordinary block counter, every active lane has step 1. For a
select counter, step is 1 when that lane's condition is true and 0
otherwise.
For example, consider a 64-lane wave executing this select:
result = (threadIdx.x % 2 == 0) ? a : b;
All 64 lanes are active, but even lanes have step 1 and odd lanes have
step 0. The counter should increase by 32. The old helper multiplies
the first active lane's step by the number of active lanes, so it adds
64 instead. Being active does not mean a lane's step is 1.
Compiler-generated counter steps are zero or one. Use a ballot to find
the lanes with step 1 and count its set bits. The first active lane adds
that count to the counter. Use the same count for the optional uniform
counter when the whole wave is active.
[libc++][test] Forward LLVM_USE_LINKER to Lit test suites
Serialize LLVM_USE_LINKER as a use_linker Lit parameter in libcxx, libcxxabi, and libunwind so test binaries link with -fuse-ld=<linker> instead of falling back to the compiler's default system linker.
Assisted-by: Gemini
Pull-request: https://github.com/llvm/llvm-project/pull/237635
[Matrix] Align copies of fused load operands to the load's alignment. (#237471)
Make sure the alloca create by getNonAliasingPointer uses the load's
alignemnt, if smaller than the alloca alignemnt. Otherwise the alloca is
not aligned enough and the load is UB.
[llubi] Add support for llvm.matrix.transpose. (#237457)
The input is treated as a column-major Rows * Cols matrix and the result
is its Cols * Rows transpose.
PR: https://github.com/llvm/llvm-project/pull/237457
[MC][NVPTX] Don't emit .addrsig for NVPTX (#200678)
PTX has no `.addrsig` directive, but AsmPrinter emits `.addrsig` and
`.addrsig_sym` whenever `TargetOptions::EmitAddrsig` is set, so `llc
-addrsig` produces PTX that ptxas rejects.
lld always sets `EmitAddrsig` for LTO, so this affects NVPTX bitcode
compiled with `ld.lld --lto-emit-asm` (#200679).
Add `MCAsmInfo::HasAddrsigDirective` and clear it for NVPTX.
NVPTX triples default to the ELF object format, so checking the object
format would not exclude them.
Co-authored-by: Fangrui Song <i at maskray.me>
[ELF] Support NVPTX bitcode with --lto-emit-asm and --lto-emit-llvm (#200679)
rustc links NVPTX bitcode with a custom llvm-link/opt/llc pipeline
(rust-lang/rust#156921), which it would like to replace with lld.
lld rejects nvptx/nvptx64 bitcode because it cannot infer an `e_machine`
from the triple.
Map nvptx/nvptx64 bitcode to `EM_CUDA`, which serves only as an internal
tag for the LTO pipeline, and report an error unless `--lto-emit-asm` or
`--lto-emit-llvm` is specified.
Kernels are referenced only by the CUDA driver at run time, so they must
be exported (e.g. `--export-dynamic-symbol`, or `-shared` with a version
script) or LTO deletes them.
Valid PTX also requires #200678, which stops emitting `.addrsig` for
NVPTX.
Co-authored-by: Fangrui Song <i at maskray.me>
[NFC][AMDGPU] Don't create call site AACallEdges for direct calls
AAAMDGPUMinAGPRAlloc and AAAMDGPUNoLDSDMA asked for a call site AACallEdges for
every call that is not an intrinsic, only to get its callees. For a direct call
whose callee has no callback metadata, that callee is the only one, so use it
directly, like AACallEdgesFunction does since #232616. Calls to callback brokers
still use AACallEdges, because thebroker also calls the callback target.
On a generated module with 2,001 functions and 15,986 calls,
-passes=amdgpu-attributor creates 2,001 instead of 11,987 abstract attributes,
and the median wall time goes from 95 ms to 84 ms.
This PR was assisted by AI but I reviewed all the changes.
[AST] Avoid duplicate isOutOfLine() calls (NFC) (#237621)
This is fairly expensive, so make sure we only call it once.
Improves compile-time for C++ code, by -0.4% on the clang build.
Suggested by LLM. Original motivation was to reduce ASLR noise on
beman-execution.
[NFC][Attributor] Reach fixpoint in call site AACallEdges for direct calls
A direct call whose callee has no callback metadata only calls its callee, so
its call site AACallEdges cannot change after initialization. Add the edge in
initialize and reach a fixpoint there, instead of updating the attribute twice
to find out. AACallEdgesFunction already skips the call site attribute for such
calls since #232616; both now share getOnlyCallee().
Other attributes still create call site AACallEdges for direct calls, e.g.,
AAInterFnReachability and the AMDGPU attributor. On a generated module with
2,001 functions and 15,986 calls, this removes all updates of call site
AACallEdges. The total number of updates goes from 304,263 to 272,291 with
-passes=attributor, from 207,815 to 175,843 with -passes=attributor-light, and
from 35,977 to 16,005 with -passes=amdgpu-attributor. The wall time does not
change measurably on that module, because these updates are cheap.
This PR was assisted by AI but I reviewed all the changes.
[AMDGPU] Allow legal shuffles of four 16-bit lanes
Preserve packed 16-bit pairs so lowering can use op_sel and
v_pk_mov_b32 instead of extracting each half into a separate register.
Keep scalarization for other vector shapes to avoid inhibiting load
narrowing or introducing extra register copies.
