[offload][omp] Load and resolve device binaries through liboffload
Migrate DeviceTy::loadBinary and global/kernel symbol resolution off
GenericPluginTy::load_binary/get_global/get_function onto liboffload's
Program/Symbol API, encapsulated in a new ProgramTy abstraction that wraps
an ol_program_handle_t. Kernel symbol resolution still needs the plugin's
opaque GenericKernelTy* handle for the legacy launch path, obtained via a
temporary __ol_tgt_GetKernelFromSymbol helper rather than new public
liboffload API surface. Removes the now-dead __tgt_device_binary type and
the corresponding GenericPluginTy methods and exports entries.
[LoopFlatten] Invalidate SCEV on each widening (#211819)
LoopFlatten can update a loop before deciding on not to flatten it due
to profitability. SCEV has cached loop information during the checks and
the widening, where it can lead to stale data if the pass bails at
flattening. Invalidate it after each successful widening.
Fixes #203176
Co-authored-by: Steve Gustaman <stevegustaman at kaist.ac.kr>
[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.
[OpenMPOpt] Ask the runtime how many of a block's threads can be workers
The custom state machine gates a thread on InitCB < BlockHwSize - WarpSize,
reconstructing the number of worker threads from the block size on the
assumption that the main thread occupies a whole warp above them. The DeviceRTL
already computes that number, in mapping::getMaxTeamThreads(), and its own
generic state machine gates on it in shouldEnterStateMachine(). Export it as
__kmpc_get_max_team_threads() and call that instead, so the compiler's state
machine and the runtime's agree by construction rather than by arithmetic that
has to be kept in step with the launch geometry.
This is NFC here: getMaxTeamThreads() in generic mode is BlockSize - WarpSize,
the same three instructions folded into one call. It is not NFC for a toolchain
whose launch geometry differs. In ROCm, CGOpenMPRuntimeGPU starts a single extra
thread rather than a warp -- "Only one additional thread is started, not an
entire warp" -- so thread_limit(1024) on a 64-lane target launches 961 threads
and the runtime reports 960 workers, while the state machine's own arithmetic
says 961 - 64 = 897. The threads in between are in neither group: the state
machine returns immediately for them, and the parallel region still hands them
[12 lines not shown]
[OpenMPOpt] Look inside the callbacks the loop runtime functions are handed
The __kmpc_{distribute_,for_,distribute_for_}static_loop_* functions receive the
loop body as a callback, so a parallel region written inside that body is
reachable from the kernel through the runtime call. AAKernelInfo could not see
that, and recorded the call as reaching an unknown parallel region. A kernel
using these functions therefore always got a worker state machine whose only
option was to indirectly call whatever work function it was handed.
Describe the callback argument of each of these functions in OMPKinds.def and
attach the corresponding !callback metadata in OpenMPOpt, then fold the
callback's AAKernelInfo state into the caller's. The state machine can now
dispatch directly to the regions the loop body actually reaches. Relax the two
"more than one callee means give up" checks for functions carrying !callback,
since the callback edge is a second edge by construction and is analyzable.
The conservative unknown-region record is kept for the case that motivated it, a
callback we only see a declaration of.
[39 lines not shown]
[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.
