LLVM/project 02bde07.github/workflows release-binaries.yml

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.
DeltaFile
+3-11.github/workflows/release-binaries.yml
+3-111 files

LLVM/project 8631e82llvm/lib/Transforms/Scalar LoopFuse.cpp, llvm/test/Transforms/LoopFusion loop_invariant.ll

[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>
DeltaFile
+82-1llvm/test/Transforms/LoopFusion/loop_invariant.ll
+8-0llvm/lib/Transforms/Scalar/LoopFuse.cpp
+90-12 files

LLVM/project e0162eallvm/lib/Target/AMDGPU SIISelLowering.h SIISelLowering.cpp, llvm/test/CodeGen/AMDGPU fcanonicalize-elimination.bf16.ll

[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
DeltaFile
+41-15llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+21-0llvm/test/CodeGen/AMDGPU/fcanonicalize-elimination.bf16.ll
+9-0llvm/lib/Target/AMDGPU/SIISelLowering.h
+71-153 files

LLVM/project b53cccfllvm/lib/Target/AMDGPU AMDGPUImageIntrinsicOptimizer.cpp, llvm/test/CodeGen/AMDGPU llvm.amdgcn.image.load.2dmsaa.ll

[AMDGPU] Fix ImageIntrinsicOptimizer DMask check starting at wrong arg index (#213264)
DeltaFile
+27-0llvm/test/CodeGen/AMDGPU/llvm.amdgcn.image.load.2dmsaa.ll
+1-1llvm/lib/Target/AMDGPU/AMDGPUImageIntrinsicOptimizer.cpp
+28-12 files

LLVM/project 55f3ddfllvm/lib/Target/AMDGPU MIMGInstructions.td AMDGPU.td, llvm/test/CodeGen/AMDGPU llvm.amdgcn.intersect_ray.ll

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)
DeltaFile
+14-3llvm/lib/Target/AMDGPU/AMDGPU.td
+13-0llvm/test/MC/AMDGPU/gfx13_asm_vimage_err.s
+5-3llvm/lib/Target/AMDGPU/MIMGInstructions.td
+2-0llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll
+1-0llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll
+35-65 files

LLVM/project 2292716clang/lib/CodeGen CGOpenMPRuntimeGPU.cpp, clang/lib/Driver/ToolChains Clang.cpp

implement feedback
DeltaFile
+0-90clang/test/OpenMP/target_teams_fast_reduction_codegen.cpp
+90-0clang/test/OpenMP/target_teams_atomic_reduction_codegen.cpp
+0-18clang/test/Driver/openmp-target-fast-reduction-flag.c
+18-0clang/test/Driver/openmp-target-atomic-reduction-flag.c
+7-7clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+4-4clang/lib/Driver/ToolChains/Clang.cpp
+119-1194 files not shown
+127-12710 files

LLVM/project fa3d882clang/include/clang/Options Options.td, clang/lib/CodeGen CGOpenMPRuntimeGPU.cpp

[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]
DeltaFile
+107-2clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
+90-0clang/test/OpenMP/target_teams_fast_reduction_codegen.cpp
+64-0offload/test/offloading/xteam_atomic_reduction_usm.cpp
+32-0llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp
+18-0clang/test/Driver/openmp-target-fast-reduction-flag.c
+9-0clang/include/clang/Options/Options.td
+320-26 files not shown
+340-212 files

LLVM/project 14ec3c4offload/liboffload/API Symbol.td, offload/liboffload/src OffloadImpl.cpp

[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.
DeltaFile
+118-0offload/unittests/OffloadAPI/symbol/olIterateSymbols.cpp
+59-16offload/liboffload/src/OffloadImpl.cpp
+49-0offload/plugins-nextgen/common/src/GlobalHandler.cpp
+31-2offload/plugins-nextgen/amdgpu/src/rtl.cpp
+32-0offload/liboffload/API/Symbol.td
+21-1offload/plugins-nextgen/common/include/GlobalHandler.h
+310-1910 files not shown
+350-2916 files

LLVM/project 7ef0ca2clang/lib/Driver/ToolChains HIPSPV.h HIPSPV.cpp, clang/test/Driver hipspv-pass-plugin.hip hipspv-no-spirv-backend.hip

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.
DeltaFile
+126-35clang/lib/Driver/ToolChains/HIPSPV.cpp
+39-27clang/test/Driver/hipspv-toolchain.hip
+41-0clang/test/Driver/hipspv-no-spirv-backend.hip
+13-9clang/test/Driver/hipspv-pass-plugin.hip
+10-2clang/lib/Driver/ToolChains/HIPSPV.h
+2-1llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp
+231-741 files not shown
+233-757 files

LLVM/project 54fedf4llvm/lib/Transforms/Vectorize SLPVectorizer.cpp, llvm/test/Transforms/SLPVectorizer/X86 runtime-alias-checks.ll

[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
DeltaFile
+17-17llvm/test/Transforms/SLPVectorizer/X86/runtime-alias-checks.ll
+12-15llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp
+29-322 files

LLVM/project c3fe695llvm/lib/Target/CSKY CSKYAsmPrinter.cpp CSKYTargetMachine.cpp, llvm/lib/Target/CSKY/MCTargetDesc CSKYELFStreamer.cpp

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>
DeltaFile
+45-0llvm/test/CodeGen/CSKY/float-abi-module-flag.ll
+41-0llvm/test/CodeGen/CSKY/fpu-abi-attribute.ll
+4-11llvm/lib/Target/CSKY/CSKYSubtarget.cpp
+7-3llvm/lib/Target/CSKY/MCTargetDesc/CSKYELFStreamer.cpp
+8-1llvm/lib/Target/CSKY/CSKYTargetMachine.cpp
+5-3llvm/lib/Target/CSKY/CSKYAsmPrinter.cpp
+110-186 files not shown
+121-2712 files

LLVM/project 93a1495llvm/test/Transforms/SLPVectorizer/X86 runtime-alias-checks.ll

[SLP][NFC]Add a test with the runtime checks miscompilation, NFC



Reviewers: 

Pull Request: https://github.com/llvm/llvm-project/pull/213336
DeltaFile
+259-0llvm/test/Transforms/SLPVectorizer/X86/runtime-alias-checks.ll
+259-01 files

LLVM/project 1e089f1llvm/utils/gn/secondary/lldb/source/Core BUILD.gn

[gn build] Port d9278ad4e640 (#213335)
DeltaFile
+1-0llvm/utils/gn/secondary/lldb/source/Core/BUILD.gn
+1-01 files

LLVM/project ae14df3llvm/utils/gn/secondary/clang/unittests/Analysis BUILD.gn

[gn build] Port 0a1f0aafed6d (#213334)
DeltaFile
+1-0llvm/utils/gn/secondary/clang/unittests/Analysis/BUILD.gn
+1-01 files

LLVM/project 00e7c35clang/lib/CIR/CodeGen CIRGenBuiltinAMDGPU.cpp, clang/test/CIR/CodeGenHIP builtins-amdgcn-gfx1250.hip

[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.
DeltaFile
+24-0clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
+1-4clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+25-42 files

LLVM/project 3ece94dclang/lib/CIR/CodeGen CIRGenBuiltinAMDGPU.cpp, clang/test/CIR/CodeGenHIP builtins-amdgcn.hip

[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.
DeltaFile
+16-0clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+3-6clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+19-62 files

LLVM/project 460c1d3llvm/test/CodeGen/PowerPC nofpclass.ll fneg.ll

[NFC][PowerPC] add `ppcf128` tests for fneg and nofpclass (#213307)
DeltaFile
+189-30llvm/test/CodeGen/PowerPC/fneg.ll
+53-11llvm/test/CodeGen/PowerPC/nofpclass.ll
+242-412 files

LLVM/project 22308c4.github/workflows release-sources.yml

workflows/release-sources: Pass release-version to upload-release-artifacts (#212660)

This was omitted from a8ccd42ab23af6848929a638cd6b099953c7e491.
DeltaFile
+1-0.github/workflows/release-sources.yml
+1-01 files

LLVM/project 55b693cllvm/test/CodeGen/AMDGPU amdgcn.bitcast.96bit.ll, llvm/test/CodeGen/X86 build-vector-known-bits-poison.ll kmov.ll

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)
DeltaFile
+67-64llvm/test/CodeGen/X86/ifma-combine-vpmadd52.ll
+40-41llvm/test/CodeGen/X86/srem-vector-lkk.ll
+29-21llvm/test/CodeGen/AMDGPU/amdgcn.bitcast.96bit.ll
+8-18llvm/test/CodeGen/X86/pr120906.ll
+9-9llvm/test/CodeGen/X86/kmov.ll
+3-11llvm/test/CodeGen/X86/build-vector-known-bits-poison.ll
+156-1647 files not shown
+183-18613 files

LLVM/project 4a31762llvm/test/CodeGen/X86 build-vector-known-bits-poison.ll

X86: Add baseline test for computeKnownBits poison+vector handling

Currently computeKnownBits is too conservative with poison inputs.
DeltaFile
+44-0llvm/test/CodeGen/X86/build-vector-known-bits-poison.ll
+44-01 files

LLVM/project 6d7c951llvm/test/Transforms/SLPVectorizer/RISCV folded-broadcast-cost.ll

[SLP] Add tests for overcosted scalar splats for RISCV (#213142)

Tests for #213104.
DeltaFile
+95-0llvm/test/Transforms/SLPVectorizer/RISCV/folded-broadcast-cost.ll
+95-01 files

LLVM/project eea8ef2libc/src/__support/math ceilf128.h

nit
DeltaFile
+2-0libc/src/__support/math/ceilf128.h
+2-01 files

LLVM/project ae22616llvm/lib/CodeGen/AsmPrinter WinException.cpp, llvm/test/CodeGen/X86 win-null-personality.ll

[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]
DeltaFile
+37-0llvm/test/CodeGen/X86/win-null-personality.ll
+6-2llvm/lib/CodeGen/AsmPrinter/WinException.cpp
+43-22 files

LLVM/project 32e5b7dllvm/include/llvm-c/CAS PluginAPI_types.h PluginAPI_functions.h, llvm/lib/CAS PluginAPI.h PluginCAS.cpp

[𝘀𝗽𝗿] initial version

Created using spr 1.3.7
DeltaFile
+779-0llvm/tools/libCASPluginTest/libCASPluginTest.cpp
+514-0llvm/lib/CAS/PluginCAS.cpp
+416-0llvm/include/llvm-c/CAS/PluginAPI_functions.h
+129-0llvm/lib/CAS/PluginAPI.h
+119-0llvm/include/llvm-c/CAS/PluginAPI_types.h
+91-0llvm/unittests/CAS/PluginCASTest.cpp
+2,048-010 files not shown
+2,233-2016 files

LLVM/project f3827c2libc/src/__support/math ceilf128.h, libc/src/math/generic ceilf128.cpp

structured and applied suggestions
DeltaFile
+1-12libc/src/__support/math/ceilf128.h
+3-9libc/src/math/generic/ceilf128.cpp
+4-2libc/test/shared/shared_math_constexpr_test.cpp
+0-3libc/test/src/math/smoke/ceilf128_test.cpp
+0-3libc/test/src/math/ceilf128_test.cpp
+1-1libc/utils/MPFRWrapper/CMakeLists.txt
+9-305 files not shown
+10-3611 files

LLVM/project f8d6aafflang/include/flang/Optimizer/Passes Pipelines.h, flang/lib/Optimizer/Passes Pipelines.cpp

[flang][NFC] move addPass* template utilities to header (#213288)

Move the addPass* template utilities to the header so external pipelines
can use them.
DeltaFile
+0-21flang/lib/Optimizer/Passes/Pipelines.cpp
+18-2flang/include/flang/Optimizer/Passes/Pipelines.h
+18-232 files

LLVM/project eef072dllvm/include/llvm/TargetParser AMDGPUTargetParser.h, llvm/lib/TargetParser AMDGPUTargetParser.cpp

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)
DeltaFile
+17-14llvm/lib/TargetParser/AMDGPUTargetParser.cpp
+29-0llvm/unittests/TargetParser/TargetParserTest.cpp
+16-0llvm/include/llvm/TargetParser/AMDGPUTargetParser.h
+62-143 files

LLVM/project 9dbcd54llvm/lib/Target/AMDGPU AMDGPU.td, llvm/lib/TargetParser AMDGPUTargetParser.cpp

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)
DeltaFile
+51-518llvm/lib/TargetParser/AMDGPUTargetParser.cpp
+115-3llvm/utils/TableGen/Basic/AMDGPUTargetDefEmitter.cpp
+43-0llvm/unittests/TargetParser/TargetParserTest.cpp
+39-0llvm/lib/Target/AMDGPU/AMDGPU.td
+248-5214 files

LLVM/project 68351e3llvm/lib/Target/AMDGPU MIMGInstructions.td AMDGPU.td, llvm/test/CodeGen/AMDGPU llvm.amdgcn.intersect_ray.ll

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)
DeltaFile
+14-3llvm/lib/Target/AMDGPU/AMDGPU.td
+13-0llvm/test/MC/AMDGPU/gfx13_asm_vimage_err.s
+5-3llvm/lib/Target/AMDGPU/MIMGInstructions.td
+2-0llvm/test/CodeGen/AMDGPU/llvm.amdgcn.intersect_ray.ll
+1-0llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.intersect_ray.ll
+35-65 files

LLVM/project 8c88b6fllvm/lib/Target/CSKY CSKYAsmPrinter.cpp, llvm/lib/Target/CSKY/MCTargetDesc CSKYTargetStreamer.h CSKYTargetStreamer.cpp

Remove untested hard-float-abi feature
DeltaFile
+41-0llvm/test/CodeGen/CSKY/fpu-abi-attribute.ll
+7-3llvm/lib/Target/CSKY/MCTargetDesc/CSKYELFStreamer.cpp
+5-3llvm/lib/Target/CSKY/CSKYAsmPrinter.cpp
+2-1llvm/lib/Target/CSKY/MCTargetDesc/CSKYTargetStreamer.h
+2-1llvm/lib/Target/CSKY/MCTargetDesc/CSKYTargetStreamer.cpp
+2-1llvm/lib/Target/CSKY/MCTargetDesc/CSKYELFStreamer.h
+59-92 files not shown
+60-138 files