[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>
RuntimeLibcalls: Add signatures for the MIPS16 return helpers
MIPS16 functions call these to move a floating point return value into
floating point registers. Each takes the value being returned, with
complex values as a pair of floats or doubles.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
RuntimeLibcalls: Add the MIPS16 hard float helpers to "compiler-rt"
The __mips16_* hard float helpers were not in any library, so they
were never available. libgcc's mips16.S provides them for O32, but not
for R6. compiler-rt does not provide them, but we don't get try to
distinguish the 2.
Since these coexist with standard symbol names, this needs explicit
handling to prefer these implementations.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[lldb] Find the end of a gdb-remote packet in linear time (#238530)
GDBRemoteCommunication::CheckForPacket appends each read, at most 8 KiB,
to m_bytes and then searches all of m_bytes for the '#' that ends the
first packet. A packet of S bytes is searched about S / 8 KiB times, so
receiving it takes time quadratic in S. With "process plugin packet
speed-test", a 32 MiB reply took 1.7 s, a 64 MiB reply 8.5 s and a 128
MiB reply 38.6 s. jThreadsInfo replies reach these sizes when a process
has tens of thousands of threads, as a GPU process can.
Resume the search where the previous call stopped. A packet cannot
contain an unescaped '#', so bytes already searched never need searching
again. Only the two paths that wait for more bytes record the position,
and neither changes m_bytes; every other path, including all of those
that remove bytes, leaves it at zero. DecompressPacket can rewrite
m_bytes before the search, so compressed connections still search from
the start. The 128 MiB reply now takes 0.73 s.
Searching only the bytes that the current call appends would not be
[8 lines not shown]
Merge branch 'users/void/ppc-asm-wide-gpr' into users/void/asm-goto-_Complext-types
Take main's InlineAsmLowering.cpp and arm64-fallback.ll. #238927 landed
the same GlobalISel fallback for tied multi-register inputs, so this PR
no longer carries its own copy.
WasmEHPrepare: Use RuntimeLibcallsInfo
Take the function type, attributes and calling convention from
RuntimeLibcallsInfo instead of constructing them by hand. Also
validate the required functions are available.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[MemProf] Keep the ambiguous attribute on allocations cloning could not disambiguate (#239087)
MemProfUse adds a "memprof"="ambiguous" attribute to allocations whose
matched contexts have both cold and not cold types, to be replaced once
a single type is identified, possibly after cloning. When cloning cannot
separate the contexts (unmatched or indirect callers, recursion, tail
calls, ...), the two LTO modes currently disagree on what is left on the
allocation:
* Regular LTO always rewrites it to "notcold".
* ThinLTO keeps "ambiguous", but only if the containing function was not
cloned for some other allocation; otherwise it is rewritten to
"notcold".
Make the result consistent: an allocation that still has both types
after cloning keeps the ambiguous attribute in both modes, so the
ambiguous hint reliably means "mixed contexts the compiler could not
separate" rather than "resolved not cold".
