Revert "[HIP] Add __exp2f to the HIP math header (#236328)" (#240703)
This reverts commit dbb0418043e246d31aaab324dbe0c1e7097be391.
This does not exist in CUDA and I do not think it's appropriate to add
extensions in the global namespace with so little discussion
Attributor: Preserve the function analyses in the light module pass
AttributorLightPass preserved FunctionAnalysisManagerCGSCCProxy, which
is meaningless for a module pass, instead of
FunctionAnalysisManagerModuleProxy. When the pass changed anything,
the module proxy was invalidated, which clears all cached function
analyses, so e.g. IPSCCP recomputed the assumption cache, TTI and TLI
of every function. The pass neither adds nor removes functions, and it
already invalidates the analyses of the functions it changes.
CTMark instructions:u with -attributor-enable=light changes by O3 -0.02%,
ThinLTO -0.10% (kimwitu++ -0.50%, sqlite3 -0.38%).
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[IndVars] Rewrite loop-exit comparisons from operands (#223951)
ScalarEvolution cannot model comparison instructions directly, so
rewriteLoopExitValues() may leave a comparison live when it is used both
to
control a loop and after the loop. This can keep an otherwise dead loop
around
until vectorization, as in https://godbolt.org/z/qK7qsK1KG.
When every loop-variant operand has a safe-to-expand exit value, record
those
values in the existing RewritePhi worklist and clone the comparison in
the loop
preheader. Restrict this path to cases where rewriting allows the loop
to be
deleted.
This also handles pointer comparisons and comparisons with two
loop-variant
[2 lines not shown]
RuntimeLibcalls: Move malloc, calloc and free attributes to TableGen
This requires new attribute plumbing for attributes that have an
argument value.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[SelectionDAG] Select the FREEZE target opcode
Selecting freeze as a COPY is wrong if the value may be undefined.
Now we can select ISD::FREEZE as FREEZE instead. Targets that don't
have physical registers can keep selecting COPY since there is no
regalloc to observe the difference.
Fixes #58321, #60429, #144780.
Assisted-By: Claude Opus 5.5 <noreply at anthropic.com>
[SelectionDAG] Don't assume a value from another block is frozen
A virtual register from another block can hold an undefined
value. We can only treat CopyFromReg of a physical register or of a
live-in virtual register as not undef or poison so that a freeze of
a value from another block is kept and selected as FREEZE.
[AMDGPU] Handle FREEZE like COPY in SIFixSGPRCopies
An SGPR FREEZE of a VGPR value cannot be lowered after register
allocation. We should treat FREEZE like the other single-input pseudos
(WQM and friends) when fixing SGPR copies and when moving instructions to
the VALU, so that its result is moved to a VGPR along with its users.
Co-Authored-By: Claude Opus 5.5 <noreply at anthropic.com>
[AMDGPU] Lower FREEZE of i1 in SILowerI1Copies
A FREEZE of a vreg_1 was left as is, which fails the verifier. Freeze
the lane mask, and lower the rest like a copy.
Co-Authored-By: Claude Opus 5.5 <noreply at anthropic.com>
[CodeGen] Hint to register allocator that FREEZE will become a copy
A FREEZE is lowered to a COPY after register allocation, so we should give
its operands the same allocation hint as a COPY. When both get the same
register, the FREEZE can be removed.
[CodeGen] Lower FREEZE of a defined value to COPY
A FREEZE is only needed when its operand may be undefined. FinalizeISel
now turns every other FREEZE into a COPY, right after instruction
selection.
Assisted-By: Claude Opus 5.5 <noreply at anthropic.com>
[SelectionDAG] Don't CSE freeze of undef
Each freeze of undef may have its own arbitrary value.
This also folds a bitcast of a one-use frozen undef into a frozen undef.
Assisted-By: Claude Opus 5.5 <noreply at anthropic.com>
[X86] Don't let BUILD_VECTOR pick a value for multi-use freeze(undef)
LowerBUILD_VECTOR treats freeze(undef) operands as free to choose. It
folds them to zero, blends in a new freeze, or zeroes them in the upper
half. If the freeze has other users, they may see a different value
than the BUILD_VECTOR does. Only treat a freeze(undef) as free when the
BUILD_VECTOR is its only user, and only zero the upper half when each
one has a single use, so that a freeze used by several lanes keeps one
value.
Assisted-By: Claude Opus 5.5 <noreply at anthropic.com>
[NVPTX][MLIR][NVVM] Add 32-bit multicast mask support to mbarrier intrinsics (#239187)
The shared::cluster variants of the mbarrier.expect_tx*/complete_tx* and
mbarrier.arrive* intrinsics now take two additional arguments: a 32-bit
CTA mask and a constant is-multicast flag. The flag selects the
`.multicast::cluster::32b` PTX form, which requires sm_107f or later in
the same family and PTX ISA 9.4; when it is cleared the mask is ignored
and the original instruction is emitted unchanged.
Legacy 2-argument declarations and calls are upgraded automatically by
appending (i32 0, i1 false).
MLIR NVVM ops do not model multicast yet, so the LLVM IR translation
passes 0 for the mask and false for the flag.
[CodeGen] Remove unused TRI in SinkingPreventsImplicitNullCheck (NFC) (#239697)
The last use of TRI in SinkingPreventsImplicitNullCheck was removed on
October 7, 2026 in commit 6afc47d555fb21b780175233f777f1ecff87e603.
Assisted-by: Antigravity
[lldb][test] Remove unused variable always_false (NFC) (#239702)
The last use of always_false in TestWriteOverSoftwareBreakpoint's main.c
was removed on October 14, 2026 in commit
a111af730a345868b8b8d0b180a939e8796ae59f.
Assisted-by: Antigravity
RISCV MC outliner: do not use RegSave in already outlined functions (#236429)
Skip `RegSave` in already outlined functions. `RegSave` inside an
outlined function adds a write to one of `t1..6`. But callers of an
outlined function rely on that outlined function to preserve registers
it doesn't define.
---------
Signed-off-by: Cyrill Leutwiler <bigcyrill at hotmail.com>
Revert #222531 "[clang] Support multiple in-process cc1 jobs in a ToolSession" (#240706)
This reverts commit cbf54d65f620f2efb78ee795ffdec28d269d9b21.
llvm-driver always constructs a ToolSession, so every clang invocation
through the multicall binary (LLVM_TOOL_LLVM_DRIVER_BUILD) takes the
reusable-cc1 path:
* -save-stats files accumulate Statistic counters across inputs
(b.stats includes a.c's counts), and -mllvm -stats output is merged.
* --offload-jobs is ignored: in-process device jobs are excluded from
parallel offload groups.
* After a crash or report_fatal_error in one input, the remaining inputs
are skipped, unlike standalone clang and -fno-integrated-cc1:
`clang -c crash.c ok.c` no longer produces ok.o.
* -disable-free is dropped for every compile, including single-file
compiles (+1.1% instructions at -O0 -g).
Embedding hosts should opt in explicitly rather than keying on
[2 lines not shown]
[RISCV] Combine extractRotateInfo and extractLuiRotateInfo. NFC (#240660)
We can rotate the constant left in the caller and pick ADDI vs LUI based
on the resulting constant.
Assisted-by: Claude
[RISCV] Prevent potentially calling isRegisterReservedByUser with an out of bounds MCRegister. (#240501)
NextReg may be larger than X31. If X31 happens to be the last value in
the register enum, this would index past the valid range for
isRegisterReservedByUser.
I really don't think we should be incrementing past X31 at all, but
that's a more involved change.
[InferAddressSpaces] Use a new value map for each rewritten constant
When the pass rewrites a flat constant expression, it remaps the
instructions that use it through other constant expressions with a
`ValueMapper`. All rewritten constants shared one map, and the mapper
caches every constant that it visits, including the ones that it maps to
themselves. So when one instruction uses two rewritten constants, the
remap for the first one caches the part with the second one as
unchanged, and the remap for the second one reuses that entry. For
example, if a call has `addrspacecast`s of two flat pointers that are
based on the same `addrspace(3)` global, only one argument is folded.
This PR uses a new map for each rewritten constant.
This PR was assisted by AI but I reviewed all the changes.
[InferAddressSpaces] Don't replace uses in other functions (#240628)
`performPointerReplacement` visits every use of a value in the postorder
list. When the value is a constant expression, some of these uses can be
in other functions. It already skipped such uses, but only after
`replaceIfSimplePointerUse`, so when the pass ran on one function, it
could still rewrite a load, store, `atomicrmw`, or `cmpxchg` in another
function, even in an `optnone` function.
This PR skips a use in another function before any replacement, like
#70611 did for the other users of a constant.
This PR was assisted by AI but I reviewed all the changes.
