[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.
[SelectionDAG] Reject promoted extracts in build-vector shuffle folds
A promoted EXTRACT_VECTOR_ELT result can be wider than the source vector element. Require matching types before constructing the shuffle mask to avoid an incorrect lane count and assertions.
[scudo] Use memmove to move leftover elements in drain. (#237444)
Rather than loop over the array and move one by one, use the memmove
function. On a 64 bit Android device, in the system_server process, the
average time went from 68ns to 55ns. The worst case also went from
~250,000ns to ~45,000ns.
[clang] Remove unused local variables (NFC) (#237612)
This patch removes unused local variables of container, string, and
value types across clang/.
Note that these variables did not trigger -Wunused-variable because
their types have non-trivial constructors or destructors, and their
initializers have no side effects.
Assisted-by: Antigravity
[SLP] Propagate lost IR flags to vectorized casts (#237240)
Vectorized casts lost the fast-math flags, nneg and trunc nuw/nsw of the scalar casts. Propagate them as for other nodes. nneg is dropped when the zext or uitofp operand is narrowed, or when a copyable zext lane is a constant that is not known non-negative in the source type. Trunc wrap flags are dropped when the trunc result is demoted and kept when only the source is narrowed, which still holds the low bits of the original value.
[LV][RISCV] Fix crash in emitInvalidCostRemarks with EVL load/store recipes (#237063)
emitInvalidCostRemarks' TypeSwitch handled VPWidenLoadRecipe/
VPWidenStoreRecipe but not their EVL counterparts, causing
llvm_unreachable when remarks are enabled on RISC-V +v with minsize.
Add the two missing cases for
VPWidenLoadEVLRecipe/VPWidenStoreEVLRecipe.
Fixes #232281
[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.
[SelectionDAG] Reject promoted extracts in build-vector shuffle folds
A promoted EXTRACT_VECTOR_ELT result can have a wider type than the
source vector element. Using that result type to construct a shuffle
mask produces the wrong lane count and can assert. Require matching
types before forming the mask.
Construct the promoted extract directly in a DAG unit test because IR
lowering can replace its zero extension with a mask first.
[Analysis] Port exported cl::opt variables to AnalysisOptions (#236844)
CodeGen, Passes, Instrumentation, MemProfContextDisambiguation and
llvm-ir2vec read Analysis options through extern declarations. Move
these options to AnalysisOptions, with getters that take an LLVMContext.
Remove IR2Vec Embedder's unused weight members, and the unittest's
writes to the weights, which only affect vocabularies loaded from a
file.
LLM-aided
