[AArch64][SME] NFC: Inline getCompatibilityForChangeToStreamingMode.
This removes the need for the `StreamingIncompatibility` enum and
makes the code a little simpler.
This PR also moves hasPossibleIncompatibleOps to AArch64TTIimpl,
so that we can query the costmodel in a follow-up PR.
IR: Speed up AttrBuilder merge (#232204)
AttrBuilder::merge added the attributes of the other builder one at a
time, each with a binary search and a vector insert into the sorted
list. Both lists are sorted, so merge them in one pass, as the TODO
suggested. This speeds up Attributor manifesting attributes.
Instructions spent in amdgpu-attributor
(opt, parse cost subtracted, median of 3 runs):
rocPRIM test_device_scan, gfx942: 1963M -> 1923M (-2.0%)
Blender Cycles HIP kernel, gfx1100: 636M -> 632M (within noise)
rocPRIM 16 tests, full LTO, closed world: 7415M -> 7273M (-1.9%)
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[DebugInfo] Remove more dead debug intrinsic handling (NFC) (#232180)
This patch removes more leftover code for legacy debug intrinsics:
- llvm::skipDebugIntrinsics, which is now an identity function at all
call sites, along with its unit test.
- Debug intrinsic checks in BasicBlock iteration helpers,
User::replaceUsesOfWith, and Verifier::visitDIAssignIDMetadata.
- Unused location-operand and address methods on DbgVariableIntrinsic
and DbgAssignIntrinsic that became dead after dropping the check in
User::replaceUsesOfWith.
Note that debug info is now represented with DbgRecord rather than
debug intrinsics outside the DirectX backend. Bitcode auto-upgrade
converts debug intrinsics to DbgRecords on load, and the DirectX
backend only materializes debug intrinsics late in DXILDebugInfo,
which strips DIAssignID metadata and does not use the removed code
paths.
Assisted-by: Antigravity
[clang] Migrate away from PointerUnion::dyn_cast (NFC) (#232179)
Note that PointerUnion::dyn_cast has been soft deprecated in
PointerUnion.h:
// FIXME: Replace the uses of is(), get() and dyn_cast() with
// isa<T>, cast<T> and the llvm::dyn_cast<T>
Literal migration would result in dyn_cast_if_present (see the
definition of PointerUnion::dyn_cast), but this patch uses dyn_cast.
Note that loadAPINotes() returns false only after storing a nonnull
APINotesReader * in Readers[*Dir].
Assisted-by: Antigravity
[CIR] Fix cir.try successor regions
`TryOp::getSuccessorRegions` sent every exit from the try region back to the
parent, so nothing led from the try region to the handlers. A `cir.resume`
ending the try region, left there once FlattenCFG has flattened an inner
cleanup scope, unwinds into a handler, and `CIRTryOpFlattening` branches it
to the handler dispatch block. Now every exit from the try region can enter
the handlers, and a `cir.resume` leaves the op only when there are none. The
parent keeps its edges to the handlers. An unwind out of a call in the try
region is not a region branch point, so these edges stand in for it, and
without them -remove-dead-values empties a handler that is reached only that
way.
Co-Authored-By: Claude Opus 5.5 <noreply at anthropic.com>
[mlir][nvgpu] Preserve unsigned types in MMA lowering to NVVM (#224837)
Type::isInteger(width) matches signed, unsigned, and signless integers.
As a result, getNvvmMmaType lowered ui8 and ui4 multiplicands to the
signed PTX types s8 and s4. Unsigned values with the high bit set were
therefore interpreted as negative by mma.sync and mma.sp.sync, with no
diagnostic.
Select u8 and u4 for unsigned operands while retaining the existing
signed PTX types for signed and signless operands. The downstream NVVM
and PTX lowering already supports these types. Add coverage for dense
ui8/ui4 and sparse ui8/i4/ui4 MMA lowering.
[CIR][CMake] Add PCH for clangCIR. NFC (#231068)
clangCIR reuses clangAST's precompiled header, so each of its 54 files
reparses the MLIR and CIR dialect headers. Add a private PCH like
clangCodeGen's (#206018).
Rebuilding clangCIR (Release+Asserts, 28 threads):
```
summed per-file user CPU wall
before 312 s 282 s 20.0 s
after 79 s + 6 s PCH 78 s 18.1 s
```
LLM-aided
[Offload][L0] Adjust incorrectly reported limits for Level Zero (#231913)
* Adjust the values of the static SPIRV GridValue with the values of the
actual hw.
* Report correctly to the PluginInterface the MaxBlockSharedMemSize.
* Mark as TODO the value of the timestamp clock used by libdevice.
Assisted by Claude.
virtio_gpu: use an X8R8G8B8 resource on big-endian guests
The virtio-gpu resource formats are defined by byte order in memory,
while vt(4) and the X server (through vt_fb and the fb mmap) write
native 32-bit 0x00RRGGBB pixels into the shadow framebuffer. On a
big-endian guest such as powerpc64 those pixels land in memory as
00 RR GG BB, which the host interprets under B8G8R8X8 with red and
blue swapped and the padding byte taken as blue: white renders yellow
and blue renders black.
Request X8R8G8B8 on big-endian guests instead, which is also what the
Linux driver does (DRM_FORMAT_HOST_XRGB8888). Reproduced with a colour
test pattern on a powerpc64 QEMU pseries guest using the identically
coded out-of-tree fork of this driver (graphics/virtio-gpu-qemu-kmod);
little-endian guests are unchanged.
Differential Revision: https://reviews.freebsd.org/D60085
Approved by: jhibbits
MFC after: 1 week
[3 lines not shown]
[mlir][SPIR-V] Support CPacked decorations in binary round trips (#232192)
The SPIR-V dialect already represents CPacked struct types, but binary
serialization and deserialization reject the decoration. Handle CPacked
as a unit decoration in both directions so packed structs survive binary
round trips.
Add round-trip coverage for literal and identified packed structs.
[SLP][modularisation][NFC] Move loop trip-count helpers to SLPUtils
Move the BoUpSLP-independent helpers findInnermostNonInvariantLoop and
getLoopTripCount out of SLPVectorizer.cpp into the self-contained
SLPVectorizer/SLPUtils.{h,cpp} module. getLoopTripCount reads the file-local
LoopAwareTripCount cl::opt, which stays static in SLPVectorizer.cpp and is
passed to the moved helper as an explicit parameter. NFC.
Part of the SLPVectorizer.cpp modularization effort:
https://discourse.llvm.org/t/modularizing-slpvectorizer-cpp/90922
[CIR] Keep cir.cleanup.scope cleanup regions reachable
RegionBranchOpInterface::getSuccessorRegions must return every region that
may run next, but an unwind out of a call in the body, or a cir.return,
cir.break, cir.continue or cir.goto out of it, runs the cleanup without
passing a region branch point. A cleanup reached only that way had no
predecessor, so generic passes treated it as dead, and -remove-dead-values
erased an EH cleanup that stores to memory. Now the cleanup region is a
successor of the parent and of every body exit, and a body exit that the
cleanup kind does not cover can also leave the op.
Co-Authored-By: Claude Opus 5.5 <noreply at anthropic.com>
[SLP][modularisation][NFC] Move getReductionInstr to SLPReductionUtils, getAggregateSize to SLPUtils
Move the reduction-root helper getReductionInstr out of SLPVectorizer.cpp
into the self-contained SLPVectorizer/SLPReductionUtils.{h,cpp} module.
getAggregateSize is not reduction-specific: its only user is
findBuildAggregate and it only computes the element count of a homogeneous
aggregate from its type. Move it to the generic SLPVectorizer/SLPUtils
module instead.
Part of the SLPVectorizer.cpp modularization effort:
https://discourse.llvm.org/t/modularizing-slpvectorizer-cpp/90922
Assisted by AI
[SLP][modularisation][NFC] Move isFirstInsertElement/getDebugLocFromPHI to SLPUtils
Move the BoUpSLP-independent helpers isFirstInsertElement and
getDebugLocFromPHI out of SLPVectorizer.cpp into the self-contained
SLPVectorizer/SLPUtils.{h,cpp} module.
Part of the SLPVectorizer.cpp modularization effort:
https://discourse.llvm.org/t/modularizing-slpvectorizer-cpp/90922
[AArch64][GlobalISel] Improve MOVI constant splat lower (#229241)
This adds an expandVectorMOVImm, similar to the existing expandMOVImm and
expandMOVAddr methods that can be used to give the optimal breakdown of
constant vectors. This is then used in global isel to perform better lowering
for constant vectors. The generation of fneg has been moved to an earlier
combine in const_build_vector_to_fneg.
This optimal method for lowering vector constants can then be reused for
lowering scalar floating point values where a vector instruction can
materialise the correct value in the low bits of the vector register. The
intent is to make use of the same expandVectorMOVImm for SDAG.
virtio_gpu: use an X8R8G8B8 resource on big-endian guests
The virtio-gpu resource formats are defined by byte order in memory,
while vt(4) and the X server (through vt_fb and the fb mmap) write
native 32-bit 0x00RRGGBB pixels into the shadow framebuffer. On a
big-endian guest such as powerpc64 those pixels land in memory as
00 RR GG BB, which the host interprets under B8G8R8X8 with red and
blue swapped and the padding byte taken as blue: white renders yellow
and blue renders black.
Request X8R8G8B8 on big-endian guests instead, which is also what the
Linux driver does (DRM_FORMAT_HOST_XRGB8888). Reproduced with a colour
test pattern on a powerpc64 QEMU pseries guest using the identically
coded out-of-tree fork of this driver (graphics/virtio-gpu-qemu-kmod);
little-endian guests are unchanged.
Differential Revision: https://reviews.freebsd.org/D60085
Approved by: jhibbits
MFC after: 1 week
[2 lines not shown]
system: add service broker flag for remote handling of Kea; closes #11016
The broker takes over reload for all sub-services (likely id-based),
which is only beneficial for XMLRPC handling.
While in service.inc, add a pinch of style.
[clang][CIR][CUDA][HIP] Unify CUDA decl emission check (NFC) (#231153)
CodeGen and CIRGen both decide which global variables and functions to
emit on the host or device side of a CUDA/HIP compilation with copies of
the same check. Move the check to ASTContext::shouldEmitCUDADecl to
reduce replicated surface.
Assisted-by: Claude Opus 5.5
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
[SPIR-V] Do not emit a declaration for an unused personality (#232200)
The backend counted a personality as a function pointer use and emitted
an invalid declaration for it, without OpFunctionEnd
Skip personality uses, since SPIR-V has no exception handling
The problem was reproduced on CIR -> SPIR-V conversion
