LLVM/project a0b9a1e — llvm/test/CodeGen/AMDGPU extract_vector_dynelt.ll, llvm/test/CodeGen/NVPTX i32x2-instructions.ll f32x2-instructions.ll

Merge branch 'main' into users/arsenm/aarch64/stack-guard-width-4-no-dead-def
DeltaFile
+725-1,475llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+704-890llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
+522-621llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
+309-438llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
+303-418llvm/test/CodeGen/NVPTX/i32x2-instructions.ll
+336-352llvm/test/CodeGen/AMDGPU/extract_vector_dynelt.ll
+2,899-4,194219 files not shown
+10,443-7,578225 files

LLVM/project 2011a0a — mlir/lib/Dialect/LLVMIR/Transforms InlinerInterfaceImpl.cpp

Update comment
DeltaFile
+2-2mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+2-21 files

LLVM/project fe130f0 — mlir/lib/Dialect/LLVMIR/Transforms InlinerInterfaceImpl.cpp, mlir/test/Dialect/LLVMIR inlining-alias-scopes.mlir

[mlir][LLVM] Use a disjoint scope domain when inlining noalias

This matches recent changes to the LLVM inliner.

AI disclosure: Claude wrote the code, I wrote the commit message and
have done initial review.
DeltaFile
+28-25mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+10-18mlir/test/Dialect/LLVMIR/inlining-alias-scopes.mlir
+38-432 files

LLVM/project 01a19cf — mlir/test/Dialect/LLVMIR inlining-alias-scopes.mlir, mlir/test/Target/LLVMIR attribute-alias-scopes.mlir

[AI test fix] Use property syntax for alias scope attributes

The strict-properties assembly format change means inherent attributes
such as `alignment`, `alias_scopes` and `noalias_scopes` on
`llvm.load`/`llvm.store` can no longer be parsed out of the trailing
attribute dictionary. Switch the newly added tests to the `<...>`
property syntax already used by the rest of these files.
DeltaFile
+2-2mlir/test/Target/LLVMIR/attribute-alias-scopes.mlir
+2-2mlir/test/Dialect/LLVMIR/inlining-alias-scopes.mlir
+4-42 files

LLVM/project f953043 — mlir/test/Dialect/LLVMIR inlining-alias-scopes.mlir

Test that the inliner keeps the flag when it clones a domain

Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
DeltaFile
+35-0mlir/test/Dialect/LLVMIR/inlining-alias-scopes.mlir
+35-01 files

LLVM/project 694246a — mlir/lib/Dialect/LLVMIR/Transforms InlinerInterfaceImpl.cpp

Remove pointless comment
DeltaFile
+0-2mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+0-21 files

LLVM/project a5c919a — mlir/include/mlir/Dialect/LLVMIR LLVMAttrDefs.td, mlir/lib/Dialect/LLVMIR/Transforms InlinerInterfaceImpl.cpp

[mlir][LLVM] Add disjointScopes to AliasScopeDomainAttr

This also updates the MLIR-side inliner to clone disjoint domains
while cloning alias scopes, matching changes to LLVM.

AI disclosure: Claude wrote the code, I wrote the commit message and
looked at the code.
DeltaFile
+25-0mlir/test/Target/LLVMIR/Import/metadata-alias-scopes.ll
+23-0mlir/test/Target/LLVMIR/attribute-alias-scopes.mlir
+17-2mlir/include/mlir/Dialect/LLVMIR/LLVMAttrDefs.td
+4-1mlir/lib/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.cpp
+2-2mlir/lib/Target/LLVMIR/ModuleTranslation.cpp
+3-1mlir/lib/Target/LLVMIR/ModuleImport.cpp
+74-66 files

LLVM/project 542474c — llvm/lib/Target/AMDGPU AMDGPULowerModuleLDSPass.cpp, llvm/test/CodeGen/AMDGPU remove-no-kernel-id-attribute.ll

[AMDGPU] Use a disjoint scope domain for merged LDS structs

When lowering LDS values, all the values are mutually disjoint, so we
can use the newly-added disjoint scopes feature to simplify the IR.

AI disclosure: Claude wrote this and I reviewed it and wrote the
 commit message
DeltaFile
+16-51llvm/lib/Target/AMDGPU/AMDGPULowerModuleLDSPass.cpp
+8-12llvm/test/CodeGen/AMDGPU/remove-no-kernel-id-attribute.ll
+24-632 files

LLVM/project b25db9b — llvm/test/CodeGen/AMDGPU lower-kernel-and-module-lds.ll lower-module-lds-via-hybrid.ll

Test fixups
DeltaFile
+28-19llvm/test/CodeGen/AMDGPU/lower-lds-struct-aa-merge.ll
+36-10llvm/test/CodeGen/AMDGPU/lower-module-lds-precise-allocate-to-module-struct.ll
+17-17llvm/test/CodeGen/AMDGPU/lower-lds-struct-aa-memcpy.ll
+15-18llvm/test/CodeGen/AMDGPU/lower-lds-struct-aa.ll
+11-13llvm/test/CodeGen/AMDGPU/lower-module-lds-via-hybrid.ll
+9-9llvm/test/CodeGen/AMDGPU/lower-kernel-and-module-lds.ll
+116-865 files not shown
+148-12011 files

LLVM/project 5035613 — llvm/lib/Target/AMDGPU AMDGPULowerKernelArguments.cpp, llvm/test/CodeGen/AMDGPU lower-kernargs.ll si-split-load-store-alias-info.ll

[AMDGPU] Use a disjoint scope domain for noalias kernel arguments

All noalias arguments of a kernel are disjoint with each other, so we
can use a disjoint scope to save on metadata construction.

AI disclosure: Claude wrote this, I looked at it and wrote this
message.
DeltaFile
+90-105llvm/test/CodeGen/AMDGPU/lower-noalias-kernargs.ll
+19-13llvm/lib/Target/AMDGPU/AMDGPULowerKernelArguments.cpp
+11-11llvm/test/CodeGen/AMDGPU/lower-kernel-arguments-noalias-call-no-ptr-args.ll
+8-8llvm/test/CodeGen/AMDGPU/si-split-load-store-alias-info.ll
+4-4llvm/test/CodeGen/AMDGPU/lower-kernargs.ll
+132-1415 files

LLVM/project 82354b0 — llvm/lib/Transforms/Utils InlineFunction.cpp

Braces
DeltaFile
+2-1llvm/lib/Transforms/Utils/InlineFunction.cpp
+2-11 files

LLVM/project 8e061e9 — llvm/lib/Transforms/Utils InlineFunction.cpp

Review feedback
DeltaFile
+1-4llvm/lib/Transforms/Utils/InlineFunction.cpp
+1-41 files

LLVM/project 84c08f8 — llvm/test/Transforms/Inline noalias2.ll

Test a callee that has both noalias arguments and its own scopes

Co-Authored-By: Claude Opus 5 (1M context) <noreply at anthropic.com>
DeltaFile
+51-0llvm/test/Transforms/Inline/noalias2.ll
+51-01 files

LLVM/project 3a6ef81 — clang/test/CodeGen arm-v8.2a-neon-intrinsics-generic.c arm_neon_intrinsics.c, llvm/lib/Transforms/Utils InlineFunction.cpp

[Inliner] Use a disjoint scope domain for noalias arguments

InlineFunction creates alias.scope/noalias metadata to represent the
set of `noalias` arguments to a function. We don't need the `!noalias`
now that we have the ability to use disjoint scopes, saving us IR size
and metadata bloat.

TODO move these to a previous commit.
Also changes InstCombine to not drop the experimental.noalias.scope.decl
for disjoint scopes even if they're not mentioned in a `!noalias`, but
do still delete them if they're not used.
DeltaFile
+324-324clang/test/CodeGen/arm_neon_intrinsics.c
+72-72clang/test/CodeGen/arm-v8.2a-neon-intrinsics-generic.c
+35-30llvm/lib/Transforms/Utils/InlineFunction.cpp
+13-14llvm/test/Transforms/Inline/noalias-calls2.ll
+12-12llvm/test/Transforms/Inline/noalias2.ll
+11-11llvm/test/Transforms/PhaseOrdering/pr39282.ll
+467-4639 files not shown
+487-48315 files

LLVM/project 41c8520 — llvm/lib/CodeGen/GlobalISel InlineAsmLowering.cpp, llvm/lib/CodeGen/SelectionDAG SelectionDAGBuilder.cpp

[CodeGen] Report an error for a direct inline asm output in memory

An inline asm output returned by value has no memory to write to, yet a
constraint such as "=rm" picks memory, the most general constraint, as
does "=m". SelectionDAG asserted on that ("Can only indirectify direct
input operands!"), and GlobalISel dereferenced a null pointer. Clang
never emits such an output, since it passes the address of a memory
output, but other IR can. Report "cannot handle direct memory outputs
yet for constraint 'm'" instead, like the other inline asm errors there.

Assisted-by: Claude Opus 5.5
DeltaFile
+31-0llvm/test/CodeGen/X86/inline-asm-direct-mem-output-error.ll
+16-6llvm/lib/CodeGen/SelectionDAG/SelectionDAGBuilder.cpp
+13-0llvm/test/CodeGen/AArch64/inline-asm-direct-mem-output-error.ll
+10-0llvm/lib/CodeGen/GlobalISel/InlineAsmLowering.cpp
+70-64 files

LLVM/project 9976b90 — llvm/lib/CodeGen TargetInstrInfo.cpp, llvm/test/CodeGen/X86 greedy-inline-asm-fold.mir

[TargetInstrInfo] Fix folding inline asm operands next to tied operands

foldInlineAsmMemOperand() swapped a register operand for the target's
memory operands with MachineInstr::removeOperand(), which asserts when a
later operand is tied, because moving it would break the tie. Inline asm
lists every input after every output, so folding any operand that comes
before another tied pair asserted, e.g. an "rm" input followed by the
input of a "+r" operand, or one of two "+rm" operands. Without
assertions, the moved operands kept stale tie indices. Untie the
operands, rebuild the operand list, and re-tie the remaining pairs at
their new positions.

It also gave up when the register appears in more than one operand, e.g.
one value passed to two "rm" operands, which left the greedy allocator
unable to spill that value at all. Fold every such operand into the
stack slot.

Finally, take MayLoad from the folded operands rather than from every
read of the register: a folded def whose tied use is another virtual

    [9 lines not shown]
DeltaFile
+85-32llvm/lib/CodeGen/TargetInstrInfo.cpp
+108-0llvm/test/CodeGen/X86/greedy-inline-asm-fold.mir
+193-322 files

LLVM/project dc90ad7 — llvm/lib/CodeGen/MIRParser MIParser.cpp

Use method predicates rather than testing the flags directly.

Co-authored-by: Nick Desaulniers <ndesaulniers at google.com>
DeltaFile
+1-2llvm/lib/CodeGen/MIRParser/MIParser.cpp
+1-21 files

LLVM/project 3d0eac4 — llvm/lib/CodeGen MIRPrinter.cpp, llvm/lib/CodeGen/MIRParser MIParser.cpp

[MIR] Serialize the "foldable" inline asm register operand flag

The symbolic MIR syntax for inline asm operand flags dropped
InlineAsm::Flag's RegMayBeFolded bit, which marks a register operand
(from an "rm" constraint) that the register allocator may fold to a
stack slot. Printing MIR and parsing it back therefore silently changed
what the allocator was allowed to do, e.g. with -stop-before and
-run-pass, and a test could only set the bit through a raw numeric flag.

Print the bit as a trailing "foldable", as MachineInstr::print() already
does, and parse it back. A tied use stores its matched operand number in
the same bits, so it never prints the marker.

Assisted-by: Claude Opus 5.5
DeltaFile
+20-6llvm/lib/CodeGen/MIRParser/MIParser.cpp
+25-0llvm/test/CodeGen/MIR/X86/inline-asm-foldable.mir
+9-3llvm/lib/CodeGen/MIRPrinter.cpp
+11-0llvm/test/CodeGen/MIR/Generic/inline-asm-foldable-not-reg.mir
+65-94 files

LLVM/project 94e9bb8 — llvm/test/CodeGen/MIR/X86 inline-asm-foldable.mir

Remove (mostly) redundant run line

The second run line showed that the *printed* form of the numeric flag (1076101129 -> reguse:GR32 foldable) parses back. But the first run line already contains the symbolic syntax, so it may be a bit overkill.
DeltaFile
+0-2llvm/test/CodeGen/MIR/X86/inline-asm-foldable.mir
+0-21 files

LLVM/project bd0fda5 — llvm/include/llvm/IR InlineAsm.h, llvm/lib/CodeGen MIRPrinter.cpp

A tied inline asm use is never foldable

A tied use stores its matched operand number in the bits that otherwise
hold RegMayBeFolded. Make InlineAsm::Flag enforce that: setRegMayBeFolded()
asserts the operand isn't tied, and getRegMayBeFolded() returns false for a
tied use. The MIR printer no longer needs to special-case tied uses, so
print 'foldable' independently of 'tiedto' and drop the comments that
explained the shared bits.

Also use F.isMemKind() instead of comparing K in the MIR parser.

Assisted-by: Claude Opus 5.5
DeltaFile
+7-0llvm/include/llvm/IR/InlineAsm.h
+1-4llvm/lib/CodeGen/MIRPrinter.cpp
+2-3llvm/lib/CodeGen/MIRParser/MIParser.cpp
+10-73 files

LLVM/project 292e3d5 — llvm/lib/CodeGen/MIRParser MIParser.cpp, llvm/test/CodeGen/MIR/X86 inline-asm-foldable.mir

Parse 'foldable' in one place and test the class-less form

Replace the ParseOptionalFoldable lambda with a single check at the end of
parseSymbolicInlineAsmOperand by making the ':' constraint parsing
conditional. Add class-less "regdef foldable" and "reguse foldable" operands
to inline-asm-foldable.mir, which take the path without a ':' constraint
that wasn't tested before.

Assisted-by: Claude Opus 5.5
DeltaFile
+64-68llvm/lib/CodeGen/MIRParser/MIParser.cpp
+5-5llvm/test/CodeGen/MIR/X86/inline-asm-foldable.mir
+69-732 files

LLVM/project 81205e0 — llvm/lib/CodeGen/MIRParser MIParser.cpp

Reformat
DeltaFile
+2-1llvm/lib/CodeGen/MIRParser/MIParser.cpp
+2-11 files

LLVM/project adc7cba — llvm/lib/CodeGen/MIRParser MIParser.cpp

Use methods on Flag to query the register kind.
DeltaFile
+2-2llvm/lib/CodeGen/MIRParser/MIParser.cpp
+2-21 files

LLVM/project 21200de — llvm/lib/Target/NVPTX NVPTXISelLowering.cpp, llvm/test/CodeGen/NVPTX i16x2-instructions.ll

Revert "[NVPTX] Mark SMUL_LOHI and UMUL_LOHI as Expand for v2i16 (#225808)" (#232505)

This reverts commit 0de93f30409d245efa221d41d161f01d476cc15a.

It breaks CodeGen/NVPTX/i16x2-instructions.ll on main. The new test_smulh and
test_umulh CHECK lines were generated before #232309 (3db057a82c92), which
stopped emitting unused argument registers at -O0, so they still expect '.reg .b32 %r<3>;'. See
https://github.com/llvm/llvm-project/pull/225808#issuecomment-6202608974

The change can reland with those two functions' CHECK lines regenerated.

🤖 Generated with [Claude Code](https://claude.com/claude-code)
DeltaFile
+0-38llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
+3-4llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+3-422 files

LLVM/project 01a1ab1 — llvm/lib/Target/AMDGPU AMDGPUAttributor.cpp

AMDGPU: Create the implicit argument attributes once

AAAMDAttributes created all 21 string attributes again for every
function it manifested, which hashes the attribute name to find the
uniqued attribute. Create them once sorted in the information cache.

Instructions spent in amdgpu-attributor:
  rocPRIM test_device_scan, gfx942:          1376M -> 1313M (-4.5%)
  Blender Cycles HIP kernel, gfx1100:        503M -> 502M (-0.1%)
  rocPRIM 16 tests, full LTO, closed world:  5148M -> 4982M (-3.2%)

Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
DeltaFile
+22-6llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp
+22-61 files

LLVM/project 1c3cb88 — clang/include/clang/Basic CodeGenOptions.def, clang/include/clang/Options Options.td

[SPIRV] Add size optimisation for AMDGCN flavoured SPIR-V (#228188)

For AMDGCN flavoured SPIR-V we disable optimisations on SPIR-V IR, to
prevent introducing changes that are detrimental to the subsequent
lowering to AMDGPU. However, this can lead to fairly voluminous
binaries, in extreme cases. This patch adds a minimal set of size
optimising passes that *should* be safe to run. They execute for
`spirv64-amd-amdhsa` codegen when passing the dedicated
`-fenable-spirv-size-optimization` flag, for optimisation levels greater
than `O0`.
DeltaFile
+53-0clang/test/CodeGen/amdgcnspirv-size-optimization.c
+50-0clang/lib/CodeGen/BackendUtil.cpp
+7-0clang/include/clang/Options/Options.td
+5-1clang/lib/Driver/ToolChains/Clang.cpp
+5-0clang/test/Driver/hip-options.hip
+1-0clang/include/clang/Basic/CodeGenOptions.def
+121-16 files

FreeNAS/freenas 1470f6c — tests/api2/zfs_tier test_shares.py

Merge branch 'sharing-nfs-mypy' of github.com:truenas/middleware into sharing-nfs-mypy
DeltaFile
+0-11tests/api2/zfs_tier/test_shares.py
+0-111 files

FreeNAS/freenas adc34f7 — src/middlewared/middlewared/alert/source nfs_host.py, src/middlewared/middlewared/migration 0021_nfs_share_alerts.py

Fix NFS share alerts
DeltaFile
+70-0tests/api2/test_300_nfs.py
+37-0src/middlewared/middlewared/plugins/nfs_/share_alerts.py
+27-7src/middlewared/middlewared/alert/source/nfs_host.py
+4-26src/middlewared/middlewared/plugins/nfs_/share_validate.py
+16-0src/middlewared/middlewared/migration/0021_nfs_share_alerts.py
+4-8src/middlewared/middlewared/plugins/nfs_/share_crud.py
+158-411 files not shown
+164-417 files

LLVM/project c4b3930 — llvm/include/llvm/CodeGen SlotIndexes.h, llvm/lib/CodeGen RegAllocGreedy.cpp LiveIntervals.cpp

[CodeGen] Drop dead SlotIndexes before allocation

SlotIndexes keeps the index list entry of an erased instruction and only
clears its instruction pointer. Live range sizes are measured in slot
indexes and greedy ranks ranges by size, so the leftovers inflate some
ranges more than others and reorder allocation, spilling heavily on
register-starved functions.

Add SlotIndexes::compactIndexes() to erase them. Erased entries are
unlinked, so LiveIntervals first reports the indexes it holds via
appendReferencedIndexes().

Off by default behind -greedy-compact-slot-indexes, since it changes
allocation across much of the test suite.
DeltaFile
+221-0llvm/unittests/CodeGen/SlotIndexesTest.cpp
+46-0llvm/lib/CodeGen/SlotIndexes.cpp
+32-0llvm/test/CodeGen/AMDGPU/greedy-compact-slot-indexes.ll
+23-6llvm/include/llvm/CodeGen/SlotIndexes.h
+29-0llvm/lib/CodeGen/LiveIntervals.cpp
+12-1llvm/lib/CodeGen/RegAllocGreedy.cpp
+363-71 files not shown
+367-77 files

LLVM/project d760136 — clang/docs ReleaseNotes.md

[clang][CodeGen] Add release note for EmitDeferred re-emission fix

Follow-up to #223877, which fixed a use-after-free when an already-emitted deferred declaration was emitted a second time. The release note was missed before that PR merged.
DeltaFile
+2-0clang/docs/ReleaseNotes.md
+2-01 files