[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]
[CIR] Fix record layout for a union with no storage type (#213591)
A union whose CIR type ends up with no members keeps its whole size in
its
padding field, and `UnionType::getTypeSizeInBits` returned early in
exactly that
case, before reaching the padding. A union need not look empty in the
source to
land there: a lone zero-length bitfield is dropped during lowering,
leaving the
same no-storage state.
A record embedding such a union was then laid out wrong. In an unpacked
record
`insertPadding` pads whenever the end of the members placed so far,
rounded up
to the next member's alignment, falls short of that member's offset, so
a union
measuring zero earns a pad the AST layout does not have. In C++,
[22 lines not shown]
[DAGCombine] Fold (select_cc (select cond, x, y), x, a, b, eq) to (select cond, a, b) (#199688)
(select_cc (select cond, x, y), x, a, b, eq) which could be simplified
to (select cond, a, b)
[CodeGen] Fix -fsanitize=array-bounds for __sized_by / _or_null pointers
`EmitCountedByBoundsChecking()` assumed a CountAttributedType is always
a __counted_by pointer. That isn't true, there are four versions of the
attribute:
* `__counted_by`: Already handled correctly.
* `__counted_by_or_null`: Incorectly handled.
* `__sized_by`: Incorreclty handled.
* `__sized_by_or_null`: Incorreclty handled.
In particular:
* __sized_by / __sized_by_or_null: the loaded bound is a byte count, but the
element index was compared against it directly, so an access was only
flagged once the index exceeded the byte count -- missing out-of-bounds
accesses for a pointee larger than one byte. Scale the index to bytes
('index * sizeof(element)') before comparing. counted_by counts elements
and is unchanged; a void (or otherwise zero-sized) pointee uses the GNU
[13 lines not shown]
[CodeGen] Fix __builtin_dynamic_object_size for __sized_by / _or_null pointers
`emitCountedByPointerSize()` assumed a CountAttributedType is always
a __counted_by pointer. That isn't true, there are four versions of the
attribute:
* `__counted_by`: Already handled correctly.
* `__counted_by_or_null`: Incorectly handled.
* `__sized_by`: Incorreclty handled.
* `__sized_by_or_null`: Incorreclty handled.
In particular:
* __sized_by / __sized_by_or_null: the attribute argument is a byte count,
but the object size was computed as count * sizeof(*ptr), over-reporting by
the element size for any pointee larger than one byte. Use the count
directly for the byte-counting variants.
* __counted_by_or_null / __sized_by_or_null: a null pointer describes no
[20 lines not shown]
clang: Use TargetID parsing from AMDGPUTargetParser (#209845)
We had grown 2 parallel parsing implementations for
triple+gpu name+feature flag target ID strings. Mostly
eliminate the redundant clang version.
Co-authored-by: Claude (Opus 4.8)
[CodeGen][NFC] Split __sized_by tests into their own file and rename test cases
In future patches the coverage of the __counted_by family attributes is
going to be increased. To help with this patch refactors the existing
test file.
1. Split `__sized_by` tests into their own file. In later commits files
will be added for each attribute so it makes sense for each attribute
to have its own file.
2. Replace `testN` test case names with human readable descriptions. Not
all test cases that will be added in the future will apply to all
attributes. If we kept on using the `testN` naming convention it
would leave odd gaps in the test numbering because we try to keep
what a test case tests consistent between files (i.e. `testN` would
roughly test the same thing but with a different attribute). Using
named test cases completely avoids this.
