RuntimeLibcalls: Add signatures for the SME ABI support routines (#239366)
Use the prototypes from the AArch64 SME ABI. __arm_sme_state returns two
64-bit values.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[clang][Parse][NFC] Share parsing for astype and convertvector builtins (#239107)
Share the parsing logic for __builtin_astype and
__builtin_convertvector, then dispatch to their respective Sema actions.
Preserve diagnostics and error recovery.
[HIP] Use positional ordering for ROCm include paths
Instead of using path-based heuristics to choose between -idirafter and
-internal-isystem for the ROCm include path, split AddHIPIncludeArgs
into two phases:
- Early phase (AddHIPIncludeArgs): emits wrapper headers (cuda_wrappers/)
before C++ stdlib headers, preserving #include_next chains.
- Late phase (AddHIPIncludeArgsLate): emits the ROCm toolkit include
path after C++ stdlib but before C system headers, so the selected
ROCm installation's HIP headers always take precedence over any
distro-installed copies in /usr/include.
This replaces the previous conditional that compared the include path
against {SysRoot}/usr/ to decide the flag. That approach had two bugs
identified in review: a double-slash with trailing-slash sysroots, and
/usr/local/rocm still using -idirafter. The positional approach handles
all install locations uniformly without path comparisons.
Co-Authored-By: Claude Opus 4.6 <noreply at anthropic.com>
[CIR] Add unadjusted align to the CIR record layout (#238913)
The unadjusted align of a record (the alignment calculated by the AST
before adjustments such as attributes are applied) is needed to perform
ABI classification for AArch64 targets. This change adds this field to
the CIR RecordLayoutAttr but does not yet enable the AArch64 calling
convention lowering. That will be done in a future change.
Assisted-by: Cursor / various models
[AMDGPU] Validate prefetch length and cache policy operands (#232143)
Two operands of the AMDGPU prefetch intrinsics, the length and the cache
policy, are not validated. An out-of-range length is encoded differently
from what the assembly says, and an out-of-range cache policy is
silently miscompiled or rejected late by the backend.
Fixes:
- Length: Instruction selection now keeps only the low 5 bits of a
constant length, so the assembly shows the value that is encoded. The
object code is the same as before. As a safeguard for hand-written
assembly, including inline `asm`, the assembler on GFX12 and later now
rejects an immediate length that does not fit in 5 bits instead of
silently truncating it.
- Cache policy: clang and the IR verifier now reject a cache policy that
sets any bit other than th (bits 0-2), scope (bits 3-4) and nv (bit 5),
or that uses the reserved th value 7. The assembler needs no change: it
already rejects th 7.
[clang][CUDA] Search include/cccl of CUDA 13 installations (#235910)
From CUDA 13.0 on, the CUDA C++ Core Libraries (libcu++, CUB and Thrust)
are in `include/cccl` of the CUDA installation rather than in
`include/`. Some of the toolkit's own headers include them, so with CUDA
13 something as simple as `#include <cooperative_groups.h>` fails to
compile with Clang:
```
cooperative_groups/details/info.h: fatal error:
'cuda/std/type_traits' file not found
```
nvcc searches `include/cccl` right after `include/`. This makes Clang do
the same, when the directory exists. `-nogpuinc` still leaves it out,
and the user's own include directories are still searched first.
The HIP toolchain already does the equivalent for
[`include/libhipcxx`](https://github.com/llvm/llvm-project/blob/fe84f6e3b0ae40ef28998cd9cee3220eec2ea261/clang/lib/Driver/ToolChains/AMDGPU.cpp#L617-L624).
Assisted-by: Devin
RuntimeLibcalls: Add signatures for the Darwin sincos stret functions
__sincos_stret and __sincosf_stret return a struct holding the 2
results. Describe the C struct return type in the signature table and
lower for the target ABI. This is a vector on x86-64 for float,
through memory on x86-32 and APCS ARM, and as a struct elsewhere.
This replaces the special case that built the types in C++. Ideally we
would be using the ABI library, but it's incomplete for most targets
and it requires some relayering to use here.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
RuntimeLibcalls: Describe the stack guard and safestack variables
The stack guard and the unsafe stack pointer are global variables, not
functions. Describe them as variables in the signature table, with the
unsafe stack pointer thread local in the alloca address space, and
declare them as globals.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
RuntimeLibcalls: Add a signature for the stack probe functions
The stack probe functions are assembly routines with their own unnamed
calling convention. The allocation size is passed in a fixed register
that is not the standard first argument register (rax/eax on x86, x15
in 16-byte units on AArch64, r4 in 4-byte units on ARM), almost all
registers are preserved, and the i386 _alloca variant also moves the
stack pointer. It's not ideal, but short of defining a custom calling
convention, just declare them with no arguments.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[AMDGPU] Update src2 live range when converting WMMA to three-address form
Since #225174, `LiveIntervals` is computed before the two-address instruction
pass. A two-address WMMA ties `src2` to an early-clobber `vdst`, so if `src2`
dies at the WMMA, its live range ends at the early-clobber slot. When the pass
converts the WMMA to the three-address form, `src2` is no longer tied and is
read at the register slot, but `SIInstrInfo::convertToThreeAddress` only updates
the live range of the def. The machine verifier then reports "Live segment
ending at early clobber slot must be redefined by an EC def in the same
instruction".
This PR moves the end of such a segment to the register slot, for the main range
and the subranges.
Fixes LCOMPILER-3197.
[flang-rt][cuda] Add support to build temporary-stack (#239403)
These functions are used when lowering some forall statements. It would
fails to link with code like:
```fortran
!$acc serial copyin(k) copyout(kk,kkk,kkkk)
forall (n=1:nb)
kk(n,:)=k(n,:)
kk(n,:)=kk(n,:)+1
kkk(n,:)=k(n,:)
kkk(n,:)=kkk(n,:)+1
kkkk(n,:)=k(n,:)
kkkk(n,:)=kkkk(n,:)+1
end forall
!$acc end serial
```
[HLSL] Add WaveMultiPrefixCountBits (#238132)
Add `WaveMultiPrefixCountBits`, with DirectX-only lowering.
This PR does not add SPIR-V support. SPIR-V calls report that the
builtin is not supported on the target.
Fixes #99211
Assisted by: Github Copilot
Co-authored-by: Copilot <223556219+Copilot at users.noreply.github.com>
Copilot-Session: 40997c2f-5aaa-41c6-8d1f-3e2e9ada315f
Revert 216506 Add PreLegalizer pattern matching for degrees GL extension (#239353)
This reverts commit 4e414afe0e3f1566b898c5c5741e2582c5c48514.
The combine rewrites fmul x, 180/pi into spv.degrees for all SPIR-V
targets. OpenCL requires fmul to be correctly rounded (0.5 ulp) while
degrees only has to be accurate to 2 ulp, so this loses precision on the
OpenCL path.
The discussions have been captured in the now closed PR about why we
want to revert it
https://github.com/llvm/llvm-project/pull/237935/changes
[AMDGPU] Infer the global address space for llvm.amdgcn.flat.prefetch (#231968)
InferAddressSpaces rewrites the generic `llvm.prefetch` when it proves
that the prefetched pointer
is global, but it does not handle `llvm.amdgcn.flat.prefetch`, which
`__builtin_amdgcn_flat_prefetch`
emits. A flat prefetch of a pointer that is provably global, such as a
HIP kernel argument,
therefore stays a `flat_prefetch_b8`, while the loads and stores of the
same pointer, and
`__builtin_prefetch` of it, use global instructions.
Add `llvm.amdgcn.flat.prefetch` to the AMDGPU InferAddressSpaces hooks.
When its pointer is
inferred to be global, the call becomes `llvm.amdgcn.global.prefetch`
with the same cache policy
operand. Pointers inferred to be in any other address space are left
alone.
TableGen: Fix advancing past the end of an empty macro fusion list (#239397)
The MacroFusions list could be empty, so drop_begin went off the
end. This trapped on MSVC debug builds for targets without
fusions. Use adjacent_find to locate the duplicate instead.
Fixes regression from #239102.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
RuntimeLibcalls: Add signatures for the MIPS16 return helpers
MIPS16 functions call these to move a floating point return value into
floating point registers. Each takes the value being returned, with
complex values as a pair of floats or doubles.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
RuntimeLibcalls: Add the MIPS16 hard float helpers to "compiler-rt"
The __mips16_* hard float helpers were not in any library, so they
were never available. libgcc's mips16.S provides them for O32, but not
for R6. compiler-rt does not provide them, but we don't get try to
distinguish the 2.
Since these coexist with standard symbol names, this needs explicit
handling to prefer these implementations.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
