[AMDGPU] Lower uniform uaddsat.i16 to SALU instructions
Promote uniform i16 uadd.sat to i32 in promoteUniformOpToI32 so it
lowers to SALU (s_add_i32 + s_min_u32) instead of VALU + readfirstlane.
The saturating add on zero-extended operands reduces to
umin(add(lhs, rhs), 0xffff).
Co-authored-by: Cursor <cursoragent at cursor.com>
[AMDGPU] Add tests for uniform uaddsat.i16 (NFC)
Precommit tests covering uniform (SALU) uadd.sat for i16 and vector
i16 (v2i16, v3i16, v4i16) using the amdgpu_ps calling convention so
results are returned in SGPRs. These currently lower to VALU
(v_add_u16 clamp / v_pk_add_u16 clamp + v_readfirstlane).
Co-authored-by: Cursor <cursoragent at cursor.com>
[Offload] Save jit image (#212384)
This PR adds the envar LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME
to indicate that the device image generated after JIT compilation
should be saved to a file.
workflows: Remove unnecessary checkouts before uses of upload-release-artifact
The workflow is on self-contained and checks out its own scripts, so we
don't need to do this in the calling workflow. The '$' prefix in the
uses tag tells github actions to load the action from the repository
directly rather than searching for it on the local file system.
https://github.blog/changelog/2026-07-30-reference-same-repository-actions-with-self-repository-syntax/
[SimplifyCFG] Do not thread branches into uncontrolled convergent regions
SimplifyCFG's foldCondBranchOnValueKnownInPredecessor can thread an edge past
a block that acts as a reconvergence point. If the threaded destination reaches
an uncontrolled convergent operation before returning to the threaded-through
block, the transform can change which dynamic instance of the convergent
operation is executed.
Add a conservative destination scan for this fold and skip the threading
candidate when it can reach an uncontrolled convergent call before returning
to the original block. Controlled convergent operations using convergence
control tokens are left alone.
Fixes ROCM-26496.
[clang][bytecode] Add DeclOrExpr (#213686)
And replace the previous usages of llvm::PointerUnion with this new
struct. This is currently just a refactoring but we'll later need to
have a proper type for the previos typedef so we can pass it to emit*
functions.
[AMDGPU] Support partial and empty WWM pools for SGPR spills (#213491)
SGPR lane spilling currently treats the WWM VGPR pool as all-or-nothing.
This can fail compilation when the requested pool cannot be formed, even
though scratch spilling or a smaller spillable pool could make progress.
This PR lets ordinary SGPR spills fall back to scratch when the pool is
empty and lets WWM register allocation use a nonempty partial pool. It
keeps the full-pool requirement for strict WWM/WQM and explicit
spill-carrier preallocation.
The no-pool fallback is recorded in SIMachineFunctionInfo so frame
lowering can provide enough emergency scavenging slots. The state is
also serialized to preserve the behavior across MIR round trips.
Fixes LCOMPILER-2542.
[AMDGPU] Lower uniform uaddsat.i16 to SALU instructions
Promote uniform i16 uadd.sat to i32 in promoteUniformOpToI32 so it
lowers to SALU (s_add_i32 + s_min_u32) instead of VALU + readfirstlane.
The saturating add on zero-extended operands reduces to
umin(add(lhs, rhs), 0xffff).
Co-authored-by: Cursor <cursoragent at cursor.com>
[AMDGPU] Add tests for uniform uaddsat.i16 (NFC)
Precommit tests covering uniform (SALU) uadd.sat for i16 and vector
i16 (v2i16, v3i16, v4i16) using the amdgpu_ps calling convention so
results are returned in SGPRs. These currently lower to VALU
(v_add_u16 clamp / v_pk_add_u16 clamp + v_readfirstlane).
Co-authored-by: Cursor <cursoragent at cursor.com>
[RISCV] Pick correct lround libcall for fp128 (#213220)
**LROUND** with an i32 result on rv64 chose the libcall as `f64 ?
LROUND_F64 : LROUND_F32`, so **fp128** fell through to the f32 libcall
(lroundf) and lost precision. Use `RTLIB::getLROUND`, which maps
**fp128** to lroundl.
[lit] Run builtin cat / diff in-process instead of spawning (#208024)
cat and diff are the only two builtins that still spawn a subprocess:
every cat/diff on a RUN line spawns a fresh Python interpreter, which
dominates wall time given how small lit's typical inputs are. Run them
in-process instead.
The spawned-script path stays as a fallback for now. Removing it
entirely is a follow-up PR. Output is byte-identical to the spawn path
either way.
[OpenMP][Clang] Enable `ATTACH`-style maps for mappers. (#210213)
This is a follow-up to #153683 to support OpenMP compliant
pointer-attachment
in `declare_mappers` via `ATTACH`-style maps.
In addition to enabling attach-style maps, we also need to propagate
information about
which map entries are for "pointee" data, i.e. have an "attach-ptr", and
thus occupy a different storage block than the base variable for which
the mapper is being generated. e.g.
```c
S sa[10];
#pragma omp declare_mapper (default: S s) map (s.x, s.p[0:10])
#pragma omp target_enter_data map(sa)
```
[18 lines not shown]
[OpenMP][NFC] Drop `omptarget` prefix from CHECK lines in mapper present tests
Upstream 58f386207ac8 ("[offload] Remove `omptarget` references from
tests") made offload test CHECK lines generic so libomptarget components
can be moved/renamed -- some debug prints will come from `ompaccsupport`
rather than `omptarget`.
The two tests updated here re-add CHECK lines to files whose other
`omptarget`-prefixed lines that commit had already rewritten, so they
merged cleanly while reintroducing the old prefix. Match the convention
used by every other offload test. The address/size captures are
unchanged; only the component prefix is dropped.
[flang][cuda] Only register module variables defined in the current translation unit (#213720)
CUFAddConstructor registered every CUF module variable mirrored in the
GPU module, including variables brought in with USE from another
translation unit. Without relocatable device code each translation unit
has its own device module, so registering a declaration binds the host
address to a module that does not contain the symbol, and a later
cudaGetSymbolAddress fails with cudaErrorInvalidSymbol or
cudaErrorSymbolNotFound. Under unified memory the same declaration is
registered as host memory, overriding the defining unit's device
registration.
Skip globals that have no initializer and no body. Non-allocatable
managed globals still register, since they go through a companion
pointer local to the translation unit.
[Offload] Add `--kernel <name>` command to `llvm-gpu-loader` (#213738)
Summary:
This makes it easier to test a single function without opting-in to the
whole `crt1.o` infra for `libc` that was originally intended to test
existing CPU tests. Good for possible future unit tests, cheap tool that
can launch a kernel. Only provides `foo(void)` kernels for now, can be
improved.
[ADT] Fix - gcc warning: buf may be used uninitialized [-Werror=maybe-uninitialized] (#213578)
When compiled with `gcc`, the uninitialized `buf` triggers a warning.
Since some external projects (e.g., `TPP-MLIR`) treat warnings as
errors, the build fails. This patch initializes `buf` to eliminate the
warning and prevent those build failures
[SimplifyCFG] Do not thread branches into uncontrolled convergent regions
SimplifyCFG's foldCondBranchOnValueKnownInPredecessor can thread an edge past
a block that acts as a reconvergence point. If the threaded destination reaches
an uncontrolled convergent operation before returning to the threaded-through
block, the transform can change which dynamic instance of the convergent
operation is executed.
Add a conservative destination scan for this fold and skip the threading
candidate when it can reach an uncontrolled convergent call before returning
to the original block. Controlled convergent operations using convergence
control tokens are left alone.
Fixes ROCM-26496.
[MC][X86] Reintroduce aligned instruction bundling (#175830)
Aligned bundling partitions instructions into fixed-size, naturally
aligned groups called bundles and guarantees that no instruction
crosses a bundle boundary, giving the instruction stream a single
canonical decoding. It is a building block for software-based fault
isolation: control flow cannot jump into the middle of an instruction
to manufacture a different, unchecked sequence, and when combined with
masking of indirect branch targets it constrains control flow to a
statically verifiable set of locations.
The previous target-independent implementation was removed in #148781,
which simplified MC by eliminating per-fragment BundlePadding, the
virtual emitInstToData, and BundleGroupBeforeFirstInst. This change
reimplements the feature in the X86 backend on top of the existing
MCBoundaryAlignFragment infrastructure added for branch alignment,
keeping the generic MC surface smaller:
* AsmParser parses .bundle_align_mode, .bundle_lock and .bundle_unlock
[36 lines not shown]
[Flang][OpenMP][OpenMPIRBuilder] Implement module scope declare target use rewrite mechanism (#212920)
During lowering of declare target'd variables we generate new global
variables for device that replace the use of the pre-existing global
variable. In Flang we currently rewrite this for each target region, but
that's not enough to cover indirect use cases inside of declare target
functions which can be imported into the module and utilised inside of a
target region. This PR tries to extend the scope of the rewriting to the
module than a per target region rewrite.
It does so by creating a mechanism where we can register globals for
replacement which will trigger on finalization of the OMPIRBuilder. This
is required as due to the ordering of lowering for MLIR, where we
generate the replacement global at the beginning of the module before
any uses have been generated, effectively meaning we cannot replace the
uses at that point. So, we defer the replacement to the OMPIRBuilder as
there is no deferral mechanism directly in the OpenMP MLIR lowering.
The alternative might be to rebind the global maps in ModuleTranslation
[7 lines not shown]
