[CIR][AMDGPU] Add support for AMDGCN class builtins (#213496)
Adds codegen for the following AMDGCN class builtins:
- __builtin_amdgcn_class (double)
- __builtin_amdgcn_classf (float)
- __builtin_amdgcn_classh (half)
These are lowered to the corresponding `llvm.amdgcn.class` intrinsics.
[VPlan] Append recipes created via builder to worklist
The previous PR appended the top most created recipe to the worklist, and this PR extends it to any other nested recipes that were created, similar to InstCombine.
This removes the header mask in a good few more places on RISC-V as measured on SPEC CPU 2017, e.g. for the following loop:
```c
long f(const int *p, const int *q, long n) {
long a = 0, b = 0;
for (long i = 0;; i++) {
if (p[i] && q[i]) { a += i; b += i; }
if (i + 1 == n) break;
}
return a + b;
}
```
Before:
[49 lines not shown]
IR: Add verifier checks and LangRef for llvm.loop.align
Verify the nested !{!"llvm.loop.align", i32 N} tag.
Require exactly two operands, an integer constant of type i32 or smaller,
and a positive power-of-two value
Co-authored-by: Claude (Claude-Opus-4.8)
[AMDGPU][CodeGen] Allow remat with multiple users in multiple regions
This relaxes one of the constraints on rematerialization candidates in
the scheduler's `PreRARematStage`. The current implementation only allows
rematerializing a register if it has users in a single region. This
allows it when a register has multiple users in multiple regions.
In such cases the register is rematerialized as many times as there are
using regions, just before the first user in each using region. The cost
model for assessing rematerialization opportunities now takes into
account that mutliple new instructions may be created for each candidate.
[CodeGen] Correctly classify/mark dead defs when adjusting lane liveness (#215595)
Despite what the documentation of `adjustLaneLiveness` states, the
method never sets dead flags on dead def operands, even when missing
dead flags can later lead to machine verifier errors.
This makes the method identify dead definitions from definitions that
are initially thought to be alive, and makes it add a dead flag on the
last definition of a virtual register, matching the behavior expected by
the machine verifier (ref. "Instruction ending live segment on dead slot
has no dead flag").
`adjustLaneLiveness` and `detectDeadDefs` now also use the same
mechanism to identify dead definitions. It relies on comparing the
defined lanes of a definition with those that stay alive after it.
AMDGPU: Remove llvm.amdgcn.addrspacecast.nonnull
The intrinsic is fully replaced by the nonnull flag on addrspacecast,
so remove it.
Old bitcode/IR is autoupgraded, though this is very conservative. This
intrinsic was only inserted by the backend, and hopefully nobody was
directly emitting it.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
AMDGPU: Set the addrspacecast nonnull flag instead of the intrinsic
AMDGPUCodeGenPrepare proved the source of certain flat<->local/private
addrspacecasts non-null and rewrote them to
llvm.amdgcn.addrspacecast.nonnull. Now that the flag is honored in
codegen, set it in place on the existing instruction instead.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
[X86] Remove x86 PDEP/PEXT clang intrinsics and rely on generic elementwise implementations (#204969)
The clang frontend already converted the x86 builtins to generics, this just moves to using the elementwise builtins directly.
AMDGPU: Use the addrspacecast nonnull flag in codegen (#220926)
Plumb the nonnull flag through to the backend so a flagged addrspacecast
lowers without the runtime null check, matching what
llvm.amdgcn.addrspacecast.nonnull already provides.
Add the NonNull MIFlag with MIR printer/parser support (including the
MIRPrinter path and update_mir_test_checks) so it round-trips on
G_ADDRSPACE_CAST, and preserve it through SelectionDAG vector
scalarization and splitting.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
[ConstantFolding] Implement canConstantFoldCallTo() using TLI (#221903)
This did some odd matching on string names. Use TLI instead, matching
the actual constant folding logic.
I've adjusted callers to pass TLI to canConstantFoldCallTo() if they
also pass TLI to the later constant folding call.
llvm: Remove phantom relocation-model attributes from tests
"relocation-model" was never a real function attribute.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
llvm: Remove phantom ssp-buffers-size attributes from tests
"ssp-buffers-size" was never a real function attribute. There is
"stack-protector-buffer-size". This may have existed in a downstream
fork, but it's also irrelevant for these tests.
Co-authored-by: Claude (Claude-Opus-4.8) <noreply at anthropic.com>
- use ofprint() for debug output so we have a chance to see it
- do not BAT-map anything beyond 0xff000000 - my 43P gets quite grouchy if we do
- copy EDID properties from OF like macppc does, for things like mach64 where
X can program video modes but can't (relably) do DDC
llvm: Remove phantom fp-contract-model attributes from tests
This attribute has never been consumed by upstream llvm,
or emitted by upstream clang. I can only guess this existed in
at least one downstream fork.
[VPlan] Append recipes created via builder to worklist
The previous PR appended the top most created recipe to the worklist, and this PR extends it to any other nested recipes that were created, similar to InstCombine.
This removes the header mask in a good few more places on RISC-V as measured on SPEC CPU 2017, e.g. for the following loop:
```c
long f(const int *p, const int *q, long n) {
long a = 0, b = 0;
for (long i = 0;; i++) {
if (p[i] && q[i]) { a += i; b += i; }
if (i + 1 == n) break;
}
return a + b;
}
```
Before:
[49 lines not shown]
[SelectionDAG] Widen vector math libcalls when no routine is available (#218948)
`tryExpandVecMathCall` currently only checks for a vector math routine
matching the node's exact vector type, unrolling when none is found.
This is suboptimal, and can lead to crashes for scalable types (which
cannot be unrolled).
This PR implements widening to first check if a routine with a wider vec
type exists before falling back to unrolling.
Example:
```
; llc -mtriple=aarch64 -mattr=+sve -vector-library=sleefgnuabi crash.ll
define <vscale x 2 x float> @frem_nxv2f32(<vscale x 2 x float> %a, <vscale x 2 x float> %b) {
%res = frem <vscale x 2 x float> %a, %b
ret <vscale x 2 x float> %res
}
```
[2 lines not shown]
