15.2: create releng/15.2 branch
* Update from PRERELEASE to BETA1
* Ship quarterly packages on DVDs
* Point pkg at stable/15 kmod packages since releng/15.2 packages
don't exist yet
* Bump __FreeBSD_version to 1502000
Approved by: re (implicit)
xdr_mbuf.c: Add a check for the return value for xdrmbuf_getbytes()
The xdrmbuf_getlong() function failed to check for the return
value from xdrmbuf_getbytes(). This could have resulted in
a failure during NFS RPC message processing.
Reviewed by: rmacklem
MFC after: 2 weeks
Sponsored by: VersatusHPC
[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: 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
libc/locale: Check EILSEQ return
_citrus_iconv_std_iconv_convert can return EILSEQ without setting
invlen. As a result, we should check EILSEQ in the return value or it
will incorrectly return -2
Reviewed by: kevans
Fixes: e2030ca246a7 ("libc: fix c*rtomb/mbrtoc*")
MFC after: 2 days
Differential Revision: https://reviews.freebsd.org/D61151
(cherry picked from commit b138d017c70421e9d074096fcc3ac676f7e1c846)
[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>
