[AMDGPU] Set WMMA source-operand reuse bits in SIPreEmitPeephole
gfx1250 WMMA instructions can set matrix_a_reuse / matrix_b_reuse bits
that keep the A or B source operand in a high-temporality state in the
VALU source-operand cache, so a later WMMA reusing the same registers
hits in the cache instead of re-reading the register file.
Add a late, post-RA peephole in the existing pre-emit peephole pass that,
for every WMMA, sets the A or B reuse bit when the next WMMA in the basic
block reads the same physical registers as its A or B operand and those
registers are not redefined in between. Other instructions between the
two WMMAs are allowed.
The cache is only a few entries deep, so at most one reuse bit is set per
WMMA, and none is added if one is already set.
Stale sticky entries in the cache are cleared when a register is used in
an instruction without a reuse bit being set. Therefore, the final WMMA
use of the same source should not set the bit.
[Clang][counted_by] Capture the field l-value instead of re-emitting the load
To share the pointer field's address with the -fsanitize=array-bounds
check, EmitArraySubscriptExpr emitted the field load itself instead of
going through EmitPointerWithAlignment and ScalarExprEmitter. That
duplicated the scalar member-load path by hand, and the copy missed
several things that path does: the 'align_value' assumption (fixed in
the previous commit), the member's debug location, and the
-fdebug-info-for-profiling pseudo variable for the struct base. So
enabling the sanitizer changed debug info and sample-profile
attribution, even for fields where no bounds check was emitted at all.
Instead, emit the base exactly as before, and have
ScalarExprEmitter::EmitLoadOfLValue record the l-value it loads from
when it matches an expression registered in the new
CodeGenFunction::CurLoadedLValueCapture. The bounds check then uses the
recorded field address. Since the base is emitted the same way with and
without the sanitizer, the load keeps everything the scalar path does.
This also covers a cast between the field load and the subscript (e.g.
[7 lines not shown]
[AVR] Support STDSPQRr on reduced tiny cores (#230226)
Currently we always expand `STDSPQRr` to `STDPtrQRr`: which is wrong, because
`STDPtrQRr` doesn't exist on Tiny. This merge request fixes this expansion, same
way we already have `if (STI.hasTinyEncoding()) { ...}` for several other instructions.
Closes https://github.com/llvm/llvm-project/issues/226680.
[IR] Declare command line options in TableGen (#233195)
Move the IR options that code with an LLVMContext reads into IROptions.
-disable-auto-upgrade-debug-info and -disable-i2p-p2i-opt become
-auto-upgrade-debug-info and -i2p-p2i-opt, keeping the old spellings as
aliases (with a TODO for removal).
llvm/unittests/IR/LibraryOptionsContextTest.cpp no longer depends on
InstCombine (a makeshift before LLVMCore is ported)
LLM-aided
[InstCombine] Fold exp2(uitofp x) to ldexp(1.0, x) with ninf
`exp2(uitofp iN x) -> ldexp(1.0, zext x)` currently requires N to be
narrower than `int`, because `ldexp` takes a signed `int` exponent.
When N equals the width of `int`, the fold is still correct under `ninf`.
Every LLVM FP type has $E_{\max} \le 16383 < 2^{15} \le 2^{N-1}$, so:
- $x < 2^{N-1}$: signed and unsigned $x$ agree, so $\operatorname{ldexp}(1.0, x) = 2^{x} = \operatorname{exp2}(x)$.
- $x \ge 2^{N-1}$: $\operatorname{exp2}(x) \ge 2^{2^{N-1}} > 2^{E_{\max}}$ overflows to $+\infty$, which is poison under `ninf`.
This only applies to the intrinsic, since the `exp2` libcall may set
`errno` on overflow.
This also enables `pow(2.0, uitofp x) -> ldexp` via the existing
`pow(2^n, x) -> exp2(n * x)` fold.
[mlir][arith] Fix integer-to-float comparison canonicalization (#233631)
Reject conversions that can produce NaN or saturate, and formats that
cannot represent zero or signed inputs. Preserve safe rewrites when the
integer range fits the floating-point exponent range.
Handle unconditional `cmpf` predicates without reaching an unreachable
case. Add regression coverage for exotic formats and rounding limits.
Assisted-by: Codex
[OpenMP][OMPTarget][test] Fortran tests for host pointers after FROM transfers
From #229966:
-------------------------------------------------+---------------------------
Test | Checks
-------------------------------------------------+---------------------------
derived_type_allocatable_skipped_from_no_restore | skipped FROM: host
| descriptor is not restored
derived_type_pointer_repointed_host_kept | re-pointed host pointer is
| kept after a target region
-------------------------------------------------+---------------------------
Co-authored-by: agozillon <Andrew.Gozillon at amd.com>
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply at anthropic.com>
[OpenMP][OMPTarget] Restore host pointers to their value before the FROM transfer
A FROM transfer of storage that holds an attached pointer/descriptor
overwrites the host copy with the device address, so the runtime writes a
host value back afterwards. That value was the one saved at attach time, and
it was written back for every FROM entry, transferred or not:
-------------------------------+---------------------+---------------------
Case | Before | After
-------------------------------+---------------------+---------------------
FROM skipped (ref count > 1) | host ptr reverted | untouched
FROM of the struct | host ptr reverted | value before FROM
FROM not covering the ptr | host ptr reverted | untouched
target update from(s.x) | host ptr reverted | untouched
-------------------------------+---------------------+---------------------
"Reverted" means set to its value at attach time. Since OpenMP 6.1 (TR15) the
host pointer may be re-pointed while its device pointer stays attached, so
that value can be stale.
[17 lines not shown]
[OpenMP][OMPTarget] Keep only the update-to part; restore FROM behavior
Suggested on top of #229966.
Drop the targetDataEnd FROM-path change. Its per-call copy of the transfer
list misses a FROM that a mapper's nested targetDataEnd skips and the outer
call later performs, leaving a device address in the host pointer:
#pragma omp declare mapper(m : S s) map(tofrom : s, s.p[0:1])
... enter data map(mapper(m), to: s) x2 ; ref 2
#pragma omp target exit data map(delete: s) map(mapper(m), from: s)
; before: s.p == a with the FROM change: s.p == <device address>
Restoring host pointers only when they are actually overwritten is left to a
follow-on change that saves their content right before the transfer.
Keep the update-to change, and add a C test for it, since the Fortran tests
need an AMD GPU. Move the two Fortran tests that check the FROM-path change
(skipped FROM keeps the host descriptor, re-pointed host pointer is kept) to
[3 lines not shown]
[OpenMP][OMPTarget] Avoid Shadow Pointer post-processing for FROM and UPDATE TO unless required
We seem to undergo the shadow pointer post-processng step even in scenarios where the data is
unlikely to have changed this can have significant performance overhead when it comes to
shadow pointers and in certain scenarios over-correction of the original host pointer.
This patch aims to address that by verifying that any transferred ranges (from or to device)
in the selected paths actually cover the shadow pointer before performing any action to touch
up the host shadow pointer or device pointer from shadow pointer data. There is still a number
of other areas that handle shadow pointers that may need this touch up, but I've opted to leave
them be for the time being.
I believe this overhead also can be triggered in C/C++, I just opted to keep it to Fortran
tests to avoid repetition! There however, may be alternative ways to approach this, but from
my limited understanding of the runtime this seems like a reasonable verifiation step.
[CodeGen] Drop dead SlotIndexes before allocation
SlotIndexes keeps the index list entry of an erased instruction and only
clears its instruction pointer. Live range sizes are measured in slot
indexes and greedy ranks ranges by size, so the leftovers inflate some
ranges more than others and reorder allocation, spilling heavily on
register-starved functions.
Add SlotIndexes::compactIndexes() to erase them. Erased entries are
unlinked, so holding an index to one is now a bug; linkedEntry() asserts
on it. The one analysis that did was LiveIntervals, so add
removeUnreferencedIntervals() to retire the intervals responsible.
Off by default behind -greedy-compact-slot-indexes, since it changes
allocation across much of the test suite.
[Modules] Don't treat piped input as file input (#233371)
Close https://github.com/llvm/llvm-project/issues/229431
The root cause of the problem is, clang treat the filename from stdin as
<stdin>. Then when there are multiple input files from stdin, the cache
status is polluted.
This patch fixes this as not treating piped input as file input.
This may show another concern as if we have different file contents with
the same name, it may be problematic.
[scudo] Only lock active TSDs in TSDRegistrySharedT::disable() (#233617)
TSDRegistrySharedT::disable() and enable() run around fork() via
pthread_atfork and lock and unlock all TSDsArraySize TSDs. Only the
first NumberOfTSDs TSDs can ever be handed to a thread, so touching the
rest only dirties their pages.
On Android (TSDsArraySize 8, NumberOfTSDs 2 by default), unlocking the
six inactive TSDs in the child writes to their mutex words and causes
copy-on-write faults on five otherwise untouched .bss pages in every
forked process. On an idle device this is 20 KiB of private dirty
memory per Zygote-forked process, or about 2.3 MiB across the ~120 such
processes.
Bound both loops by NumberOfTSDs, matching drainCaches() and getStats().
setNumberOfTSDs() only grows the count and only while holding Mutex,
which disable() holds until enable(), so both loops see the same value.
Also call disable()/enable() in TSDRegistryTSDsCount before and after
[3 lines not shown]
[clang] Add aarch64-unknown-uefi target
Add the aarch64-unknown-uefi target to clang, building on the backend
support added in the preceding patches. Like x86_64-unknown-uefi it
follows the Microsoft ABI conventions, matching the Windows ARM64 target
which is the only way EFI applications for AArch64 can be built with
Clang/LLVM today:
* add UEFIAArch64TargetInfo: LLP64 types, Microsoft C++ ABI, char*
based va_list and the Win64 calling convention kind,
* use the Windows AArch64 ABI info in CodeGen for UEFI,
* accept aarch64 for *-uefi triples in the driver.
The UEFI toolchain already handles the linker invocation (lld-link with
/subsystem:efi_application), so no further driver changes are needed.
Assisted-by: Claude
[AArch64] Apply the Windows SEH and CRT codegen constraints to UEFI
aarch64-unknown-uefi uses SEH unwind information and has no C runtime
library, so the remaining Windows specific codegen decisions that follow
from this apply to it too:
* do not align loops or basic blocks, since SEH unwind info generation
needs the exact function layout before alignment padding is inserted,
* sign return addresses with the B key by default and reject A key
signing, matching the pac_sign_lr unwind code used by the unwinder;
the branch protection parser defaults to the B key for UEFI as well,
so that -mbranch-protection=pac-ret keeps working,
* do not lower wide shifts to libcalls at -Oz, as those helpers are not
available.
The corresponding Windows tests gain aarch64-unknown-uefi run lines and
the LLVM release notes announce the new target, which completes the
backend side of aarch64-unknown-uefi support.
Assisted-by: Claude
[clang][Driver] Use the Windows ARM64 defaults for aarch64-unknown-uefi
Match the Windows ARM64 driver defaults for aarch64-unknown-uefi, which
follows the same ABI:
* char is signed, so do not pass -fno-signed-char,
* the frame chain must stay valid even when frame pointers are omitted,
so -fomit-frame-pointer yields -mframe-pointer=reserved rather than
none, as the backend reserves the frame pointer for UEFI as well,
* -msign-return-address selects the B key, the only one the backend
accepts for Windows and UEFI.
Assisted-by: Claude
[AArch64] Classify global references on UEFI like Windows COFF
Apply the Windows COFF rules for global references to
aarch64-unknown-uefi, which produces the same PE/COFF images:
* classify dllimport and non-dso-local globals with MO_DLLIMPORT /
MO_COFFSTUB, so extern_weak symbols are referenced through a stub that
can be resolved to zero by the linker,
* keep the existing SelectionDAG, FastISel and GlobalISel handling of
calls and tail calls to extern_weak functions in sync with that,
* default to the PIC relocation model, so jump tables and other
references stay position independent in relocatable PE images,
* lower ifuncs on UEFI as COFF has no support for them.
The extern weak and jump table tests gain aarch64-unknown-uefi run lines,
and a new test checks that the default pipelines lower ifuncs for both
Windows and UEFI.
Assisted-by: Claude
[AArch64] Probe large stack frames with __chkstk on UEFI
Like Windows, aarch64-unknown-uefi has no guard page handling of its
own, so frames larger than the probe size are probed with __chkstk
instead of inline probing, and "no-stack-arg-probe" disables it.
Tests: chkstk.ll, no-stack-arg-probe.ll and win-alloca-no-stack-probe.ll.
Assisted-by: Claude
[AArch64] Use __chkstk for dynamic stack allocation on UEFI
Register the Windows __chkstk stack probe helper for aarch64-unknown-uefi
and lower dynamic stack allocations through it, including the SME ZA
buffer allocation in LowerFormalArguments. GlobalISel falls back to
SelectionDAG for dynamic allocas on UEFI as it does on Windows.
Tests: win-alloca.ll, the vla in win64_vararg.ll (which now runs in full
on UEFI) and sme-lazy-save-windows.ll.
Assisted-by: Claude
[AArch64] Reserve x18 as the platform register on UEFI
Both the Windows ARM64 ABI, which UEFI follows, and the firmware
ecosystem (EDK2 builds AArch64 with -ffixed-x18) treat x18 as a reserved
platform register. Reserve it by default for aarch64-unknown-uefi, as is
already done for Windows, Darwin, Android, Fuchsia and OHOS.
Assisted-by: Claude
[AArch64] Treat the Win64 calling convention as native on UEFI
Functions using the Win64 calling convention on a non-Windows target
back up x18 and are excluded from tail call optimization, because x18
is the platform register there. On UEFI the Win64 convention is the
native one, so apply the Windows behaviour.
Tests: win64cc-backup-x18.ll gains Windows and UEFI run lines showing
that neither backs up x18 and both tail call.
Assisted-by: Claude
[AArch64] Use the Win64 calling convention and va_list on UEFI
UEFI follows the Windows ARM64 procedure call standard, as
x86_64-unknown-uefi already follows the Win64 one. Use it for C, Fast
and Swift calls on aarch64-unknown-uefi: the Win64 variadic convention
(all remaining argument registers are homed and floating point variadic
arguments travel in GPRs), the char* va_list for va_copy, and the
preserve_none variant that follows the Win64 rules.
Tests: win64_vararg_float.ll (SelectionDAG, FastISel and GlobalISel),
win64_vacopy.ll (SelectionDAG and GlobalISel va_copy),
preserve_nonecc_varargs_win64.ll and preserve_nonecc_call.ll.
isArgumentRegister follows for consistency; it only feeds
-fzero-call-used-regs, which is a no-op on Windows ARM64 today.
Assisted-by: Claude
