workflows/release-binaries: Move environment declaration to upload job (#212687)
This is the only job that actually needs to use the environment secrets,
so the environment must be declared. We were using secrets in the
prepare job to do a permissions check, but this is unnecessary, because
that job does not do anything that is security sensitive.
Only the upload job needs to have these permission checks and these are
already included in the upload-release-artifact composite action.
[LoopFusion] Allow loop fusion for idempotent output dependency (#206401)
Loop Fusion incorrectly blocks fusion of loops that write the same value
to two arrays. Writing the same value twice produces identical observable results
regardless of execution order, so fusion is safe (whether the two arrays are aliased or not).
Fixes #94676
Co-authored-by: AntonyCJ30 <cj6186609 at gmail@gmail.com>
[AMDGPU] Do not treat bitcast across FP types as canonicality-preserving (#203560)
isCanonicalized recursed through ISD::BITCAST ignoring the type change,
so value canonical as v2bf16 was wrongly treated as canonical when
bitcast to v2f16 (that has different exponent width), dropping a
required fcanonicalize
AMDGPU: Separate image_msaa_load from bvh-ray-tracing-insts (#213270)
Add a separate msaa-load-insts feature so image_msaa_load is available
on gfx13, which has it but not the BVH intersect-ray instructions. These
were assumed to be the same feature previously, but gfx13 does not have the
bvh instructions, but does have image_msaa_load.
Co-authored-by: Claude (Claude-Opus-4.8)
[offload][OpenMP] Add atomic cross-team reductions
Regular cross-team reductions have two phases: the intra-team reduction
and the inter-team reduction. Atomic cross-team reductions replace the
second phase with a atomic instruction which is used by the main thread
of each team to directly fold the result of the intra-team reduction
into the final result. Since this requires a combination of "data type"
and "combine operation" for which an atomic instruction is available,
only some (but very common) reductions can be transformed to atomic
reductions. In cases where multiple reductions are performed on the same
construct, the atomic path is only taken if all reductions can be
transformed. Otherwise, we fall back to the regular cross-team reduction
using a buffer with per-team slots. This is not strictly necessary, but
hybrid reductions would induce more complexity with questionable
benefit.
Selecting an atomic path might not be the best option for every
situation, which is why it is not enabled by default. Instead, it can be
enabled via `-fopenmp-target-fast-reduction`. This flag has already been
[20 lines not shown]
[Offload] add 'olIterateSymbols' runtime function (#213036)
Summary:
This provides users with a way to check all symbols present in a given
program. This is useful for many cases where a power user wants to do
something dynamic with the objects present.
Reland "[HIPSPV] Add in-tree SPIR-V backend support for chipStar" (#213052)
Relands #206910 (reverted in #213088) with the fix for the breakage.
IsIntegratedBackendDefault() was tied to whether the SPIR-V backend is
registered, but it also controls whether clang collapses the compile and
backend
jobs, so builds without the SPIR-V target split every HIPSPV RDC device
compile
and failed hipspv-toolchain-rdc.hip. The default is back to the base
class value
and the translator fallback is decided in
constructLinkAndEmitSpirvCommand.
Third commit adds a test not guarded by spirv-registered-target.
Validated in X86;SPIRV and X86 only builds: driver tests pass in both,
and
non-RDC plus RDC (new driver) flows run for real emit spirv-val clean
modules
via the backend and via llvm-spirv respectively.
[SLP]Fix miscompile from poison base in alias-check versioning
Emit the runtime alias check while the block is still fully connected.
LCSSA-preserving SCEV expansion rewrites out-of-loop uses of
loop-defined bases to poison in predecessor-less blocks, corrupting
both the moved body and the scalar fallback clone.
Reviewers:
Pull Request: https://github.com/llvm/llvm-project/pull/213338
CSKY: Consume "float-abi" module flag (#212975)
Start respecting float-abi, and fall back on the TargetOptions
field if not present.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
[CIR][AMDGPU] Add support for AMDGCN tanh builtins (#197852)
Adds codegen for the following AMDGCN tanh builtins:
- __builtin_amdgcn_tanhf (float)
- __builtin_amdgcn_tanhh (half)
- __builtin_amdgcn_tanh_bf16 (bfloat16)
These are lowered to the corresponding `llvm.amdgcn.tanh` intrinsic.
[CIR][AMDGPU] Add support for AMDGCN trig_preop builtins (#197399)
Adds codegen for the following AMDGCN trigonometric pre-operation
builtins:
- __builtin_amdgcn_trig_preop (double)
- __builtin_amdgcn_trig_preopf (float)
These are lowered to the corresponding `llvm.amdgcn.trig.preop`
intrinsic.
DAG: Skip poison elements in BUILD_VECTOR computeKnownBits
This defends against regressions in future patches. Copies the logic
from the IR version of computeKnownBits's handling of ConstantVector.
I'm not sure why the IR version doesn't directly return a value for poison,
but this follows suit.
Co-authored-by: Claude (Claude-Opus-4.8)
[X86] Don't emit personality info when the personality is not a function (#212803)
Follow-up to #212417, another one in the same corner.
`llc -mtriple=x86_64-pc-windows-msvc` crashes on a function whose
personality isn't a function:
```llvm
define void @a() personality ptr null {
ret void
}
```
`dyn_cast<Function>` gives null, and `beginFunclet` hands that straight
to `getSymbol`, which dereferences it.
`DwarfCFIException` already guards against this, applying the null check
to both branches:
[14 lines not shown]
[flang][NFC] move addPass* template utilities to header (#213288)
Move the addPass* template utilities to the header so external pipelines
can use them.
AMDGPU: Export the TargetParser feature bitset
Previously this bitset was only used to populate the feature
name string map used by clang. Eventually this will replace
the current bitmask integer. AArch64 already has a similar
interface.
Co-authored-by: Claude (Claude-Opus-4.8)
AMDGPU: Tablegenerate TargetParser feature sets
Traditionally we maintained 2 parallel feature mechanisms,
one in clang (later moved to TargetParser), with largely
mirrored subtarget features defined in the backend. Start
directly taking feature information from the backend and putting
it into TargetParser. This is still in a compromise mid-migration
state. We still have both the legacy "ArchAttr" bitfield integer,
plus a new AMDGPUFeatureBitset field stored in the table, which
isn't yet exported.
For the moment, the new bitset is only used to populate the
feature string name map, which is the big maintainability win.
This also lists an explicit subset of exported features to
avoid churn.
Co-authored-by: Claude (Claude-Opus-4.8)
AMDGPU: Separate image_msaa_load from bvh-ray-tracing-insts
Add a separate msaa-load-insts feature so image_msaa_load is available on
gfx13, which has it but not the BVH intersect-ray instructions. These were
assumed to be the same feature previously, but gfx13 does not have the bvh
instructions, but does have image_msaa_load.
Co-authored-by: Claude (Claude-Opus-4.8)
