[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>
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]
Switch to SmallVector with space on stack
SmallVector allows for much larger small sizes than SetVector, so use 256 to match InstructionWorklist.
We don't need to worry about duplicate worklist entries until we add users to the worklist.
[VPlan] Process simplifyRecipes in a worklist
This brings simplifyRecipes further in line with InstCombine, and asides from unlocking more simplifications it also helps avoid spurious test churn whenever passes are moved around simplifyRecipes.
For now just push the new recipe onto the worklist, not its users.
This uses a post order traversal so we maintain the same simplification order as before.
I've gone through and checked every simplification we do is a canonicalisation that converges, and I checked on llvm-test-suite + SPEC CPU 2017 in various configurations that we don't hit any cycles.
[SPIRV] Preserve DebugGlobalVariable emission order. (#221759)
Iterate GlobalVariableDebugInfoMap with MapVector so OpString and
DebugGlobalVariable follow insertion order.
The following tests had non-deterministic output, although didn't fail
due to the use of CHECK-DAG:
- llvm/test/CodeGen/SPIRV/debug-info/debug-global-variable-multi-gve.ll
- llvm/test/CodeGen/SPIRV/debug-info/debug-lexical-block-namespace.ll
- llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll
[VPlan] Split simplifyRecipes into simplifyRecipes and combineRecipes
Bringing it in line with InstSimplify and InstCombine, split up simplifyRecipe into a function that modifies and creates recipes (combineRecipe), and one that is analysis-only (simplifyRecipe).
This allows us to avoid adding simplficiation only folds to the worklist in #213899.
The funclet.ll test no longer erases the constant-folded intrinsic call because vputils::isDeadRecipe returns false.
