[AMDGPU] Price misaligned flat and LDS accesses with the LDS misaligned bug
In WGP mode on subtargets with the LDS misaligned bug, a flat access is
split into pieces no wider than its alignment and no narrower than 4
bytes, and an LDS access aligned below 4 bytes uses 8-byte pieces as one
aligned to 4 does.
[AMDGPU] Price buffer fat pointer memory operations
Price buffer fat pointer accesses as global accesses, except that an
aligned load is not widened and a vector of odd width elements is not
split element by element. Without relaxed buffer bounds checking, a
vector of elements that are a multiple of 32 bits wide is accessed in
pieces no wider than its alignment and no narrower than 4 bytes, and a
vector of other byte-sized or power of 2 wide elements in pieces of the
largest power of 2 bytes that divides the element size. The width comes
from buffer fat pointer lowering, which now exposes it. When the pointer
is a buffer resource plus a constant offset, the offset and a constant
number of records give the same bounds facts lowering uses to access it
in wider pieces.
[AMDGPU] Price constant address space memory operations
Price constant address space accesses as global accesses, except for
loads selected to scalar loads of up to 64 bytes. These are loads from
the 32-bit constant address space, and loads from the constant address
space whose pointer is an argument passed in SGPRs or a constant, plus
constant offsets, both aligned to 4 bytes or to their size. A scalar
load is widened as its lowering does. An integer whose size is not a
multiple of 4 bytes is loaded as whole dwords, plus a byte load when its
size is 7 modulo 8, since the last byte of such a remainder is left out
of the scalar load. A type with no legal dword vector equivalent is
split into the largest power of 2 pieces that fit, and once the
remainder fits in the alignment it is widened to a power of 2 piece of
up to the alignment, at most 8 bytes for byte vectors. A 3-dword
remainder of a legal dword vector takes a fourth dword when aligned to 8
on a subtarget without scalar dwordx3 loads.
[Offload][Omp] Reimplement interop on top of liboffload (#234165)
Reimplement OpenMP interop support using liboffload interfaces:
* Added a new property liboffload devices to get the native platform of
a device that was missing (could be moved to a sepparate PR).
* Reimplemented the queue creation and logic assuming liboffload inorder
queues.
* Some of the backend decision logic was moved into libomptarget (if
#235844 is merged it should be moved there)
* Renamed some "tgt" interop related stuff to use "omp" prefix.
I also found a bug in which properties were allowed for
`interop(target`) which I fixed here but could do it sepparately.
Assisted by Claude.
[PowerPC] Use 64-bit registers for wide "r" inline asm operands
On PPC64, an integer wider than 64 bits, such as an i128, in an "r" or
"b" operand was given a 32-bit register for each 64 bits of it, so half
of its bits were lost. Give it 64-bit registers instead. AArch64 had the
same problem (#238926).
Assisted-by: Claude Opus 5.5
[AArch64] Use 64-bit registers for wide "r" inline asm operands
An integer wider than 64 bits, such as an i128, in an "r" operand was
given a 32-bit register for each 64 bits of it, so half of its bits were
lost. Give it 64-bit registers instead.
Fixes #238926.
Assisted-by: Claude Opus 5.5
[Clang][RISCV] Record the LMUL and mask ratio in RVV fixed-length vector kinds (#237845)
A riscv_rvv_vector_bits typedef keeps only the element type and count,
and later stages re-derive the LMUL of a data type, or the ratio R of a
vbool<R>_t mask type, from VLEN. That breaks once typedefs may be
shorter than VLEN * LMUL, or VLEN / R for masks.
Split VectorKind::RVVFixedLengthData into RVVFixedLengthData_MF8 to _M8,
and RVVFixedLengthMask{,_1,_2,_4} into RVVFixedLengthMaskRatio<R> and
RVVFixedLengthMaskRatio<R>_{1,2,4}, and read the LMUL or the ratio from
the kind. VecKind grows to 6 bits.
This patch prepares to let riscv_rvv_vector_bits typedefs be shorter. It
does not change any behavior, except that the AST dump prints the LMUL
or the ratio.
Assisted-by: Claude
[libc] Find linux kernel headers in Debian multiarch directories (#224314)
On Debian/Ubuntu systems without gcc-multilib, architecture-specific
kernel headers are installed in /usr/include/<triple>/asm, while generic
headers are in /usr/include/linux. Because the symlink /usr/include/asm
does not exist, full host builds failed to find <asm/errno.h> and other
asm headers.
Instead of requiring root access to create a symlink in /usr, check if
/usr/include/<triple>/asm exists (using LIBC_TARGET_TRIPLE or
CMAKE_LIBRARY_ARCHITECTURE) and prepend it to LIBC_KERNEL_HEADERS.
Also update consumers of LIBC_KERNEL_HEADERS in runtimes and libc++ test
configs so they handle semicolon-separated lists properly.
Assisted-by: Gemini
[Offload][Omp] Add Backend specific OpenMP-only options (#235844)
There's a number growing of cases (as we decouple from liboffload) of
places where libomptarget needs to behave or use a different number
depending on the backend. Instead of start adding a bunuch of ` if
(Backend == XX) YY` this PR introduces
a simple ` BackendConfig` class to hide such behavior. It's a simple
class as I don't think we need a really complicated thing (for now at
least).
Ongoing PR's like #234165 or #234191 could make use of it. For now, I
just used to remove the ` isHostDevice` specific condition.
Assisted by Claude.
[LowerMemIntrinsics] Expand small variable-length memsets inline (#228420)
[LowerMemIntrinsics] Expand small variable-length memsets inline
Previously, a memset whose length is not a constant was always lowered to a call
to memset, even when the length is known to be small. Expand such a memset
inline instead when the upper bound of its length is at most a target-provided
threshold.
Motivating C code snippets:
void bar(void *dst, unsigned n) {
if (!(n <= 12)) __builtin_unreachable(); // __builtin_assume
memset(dst, 0xff, n);
}
void foo(int *dst, int n) {
for (unsigned i = n; i < 4; ++i)
dst[i] = -1;
[20 lines not shown]
Revert "[Mips] Fix MipsJmpLink and MipsTailCall to use compact branch instructions on MIPS32/64 R6" (#239049)
Reverts llvm/llvm-project#232578
I will resubmit again for a kernel compile error.
[NVPTX] Add default value to is_exclusive ImmArg in tcgen05.{de}alloc intrinsics (#236356)
This patch adds default value for the `is_exclusive` `ImmArg` in
`tcgen05.{de}alloc` intrinsics.
This helps to eliminate the AU code for enabling this manually.
---------
Signed-off-by: DharuniRAcharya <dharunira at nvidia.com>
[RISCV] Use a fully covered switch in RISCVABI::getStackAlignment.
This makes it more obvious the alignment has been chosen explicitly
for every ABI.
Not quite NFC because we no longer emit the stack align attribute for
an unknown ABI.
