[AMDGPU] Fix out-of-bounds read of a trailing '%' in printf formats (#229088)
locateCStrings checked whether the character after a '%' was another '%'
without checking that there was one. A device printf whose format string
ends in '%' hit the StringRef index assertion in both hostcall and
buffered mode.
Assisted-by: Claude Opus 5.5
---------
Signed-off-by: Steffen Holst Larsen <sholstla at amd.com>
[mlir][x86] Bail out on AMX accumulation without result write (#231144)
Rejects the AMX contraction lowering of zero initialized reduction loops
whose results are not written to memory.
Such cases crashed the lowering instead of being left unchanged.
Assisted-by: Claude
AMDGPU: Legalize SGPR tied src2 for V_FMAC_F16 in moveToVALU (#231146)
When S_FMAC_F16 is moved to the VALU, it becomes V_FMAC_F16_t16_e64 or
V_FMAC_F16_fake16_e64. legalizeOperandsVOP3 only forced the tied src2
into a VGPR for V_FMAC_F32_e64 and V_FMAC_F16_e64, so an SGPR
accumulator fit within the constant bus limit was left in place, producing a
verifier error. Check for any SGPR src2 tied to the result instead.
This was exposed by frem.ll after converting it to use inreg arguments.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[AMDGPU] Allow the AS identifying builtins to take flat pointers (#230677)
Summary:
We have `__builtin_amdgcn_is_shared` and `__builtin_amdgcn_is_private`
which are used to determine the underlying address space of a flat
pointer. There's some magic around AS(0) depending on the language mode,
in OpenCL it's `__generic`, in HIP it's the default AS. In the mode
without a language it's still a distinct AS.
This PR changes the builtin signature to just be a flat pointer. This
should be equivalent on the targets I know of, but makes the expectation
clearer and fixes the type issues w/ plain C/C++. Test file hopefully
verifies that existing OpenCL / HIP usage is unaffected.
[LV] Only fold outer extend into narrow mul if it cannot wrap. (#229230)
For partial reductions, ext(mul(ext(A), ext(B))) is matched and costed
as mul(wide_ext(A), wide_ext(B)) and optimizeExtendsForPartialReduction
rewrites it that way. This is only correct if the narrow mul cannot
wrap.
The current check would incorrectly allow patterns like zext (mul (zext
i8 to i12, zext i8 to i12)) to i32.
The extends of the mul operands must extend wide enough so the mul
cannot wrap or the mul must be marked nuw/nsw.
Alive2 Proofs: https://alive2.llvm.org/ce/z/SY6fCN
PR: https://github.com/llvm/llvm-project/pull/229230
[flang][docs] add design doc for inline runtime checks (#228502)
We currently rely on runtime versions of array assignment and intrinsics
for most runtime checks, and flang does not implement pointer/bound
checking at all.
This document describes a new fir.assert operation to allow inline
implementation to keep conformity and allocation status checking, as
well as to later implement bound and pointer checking.
Assisted-by: AI
Passes: Replace FunctionAttrs with attributor-light when it is enabled
attributor-light is meant to replace FunctionAttrs, but with
-attributor-enable=light both ran. Use AttributorLightCGSCCPass in
place of each PostOrderFunctionAttrsPass, and AttributorLightPass in
place of each ReversePostOrderFunctionAttrsPass, in the default,
ThinLTO and full LTO pipelines.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[clang][bytecode] Do BaseToDerived casts step-by-step (#231131)
Computing the offset of the entire cast path once doesn't work out in
practice if the dynamic type of the expression we're casting doesn't
match its static type. Do it one step at a time. This way we don't have
to compute an offset at all, but we stil need to check the validity of
the cast on every step.
[CIR] Implement Matrix-Insert and do lowering to it for lvalue subscript (#230600)
Matrix element assignment requires an insert operation, which this
provides and implements. Most of this very similar/cribbed from the
previous Extract patch, so hopefully this is simple/right enough.
[CIR] Implement 'matrix-subtraction' lowering. (#230536)
This ends up being basically just a copy/paste from the 'add' patch I
just did, otherthan not having to modify the 'splat' work (as that is
already done!).
SelectionDAG: Mark physreg defs dead on scheduler-cloned nodes
When the list scheduler clones a node to break a physreg dependency,
the copies share one SDNode, so InstrEmitter can't tell that a copy's
implicit physreg def (e.g. EFLAGS) lost its users and didn't mark it
dead. LiveVariables recomputes it later, but dead flags are most useful
to earlier optimizations.
Instead, mark the def dead at emission time if the copy has no
scheduler edge for it. This lets convertToThreeAddress form LEAs again,
and is a step toward removing LiveVariables.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[LV] Add helpers to extend VPWidenMemIntrinsicRecipe to non-VP intrinsics (NFC) (#230486)
This patch adds three new vector utils to query the mask, pointer, and
access type of vector memory intrinsics. It then uses these hooks to
generalize VPWidenMemIntrinsicRecipe to non-VPIntrinsic instructions.
A basic unit test is added to test these hooks work for
compresstore/expandload.
Split off from #214491.
[X86] x87.ll - regenerate checks for all x86/x64 +/-x87 +/-sse1 tests (#231140)
Ensure we check full codegen
Split single large test into int2fp fp32/64/128 tests to keep them
focussed
[VPlan] Model the ExtractValue indices as operands from VPlan0. (#230494)
Add the extractvalue indices as VPlan constant operands when building
VPlan0, instead of only appending a single index when creating
VPWidenRecipes. This gives VPInstruction, VPWidenRecipe and
VPReplicateRecipe the same operand layout, so the indices are printed
for all of them and the result type is derived from the operands.
Includes test coverage for extracts with multiple indices.
As the indices are now proper operands, extractvalues can be CSE'd.
PR: https://github.com/llvm/llvm-project/pull/230494
[BOLT][RISCV] Do not lower a conditional tail call to an unconditional one (#231117)
isCall() is true for a conditional tail call, so FixRISCVCallsPass
replaced it with a PseudoTAIL, dropping the condition and the
fall-through path.
The pass could instead rewrite it as an inverted branch over a
PseudoTAIL, but RISC-V has no conditional tail call instruction to lower
it to. Skipping conditional branches is simpler: the branch stays as it
is, and the assembler still relaxes it when the target is too far.
Found in a BOLT-optimized clang that crashed in
TargetLoweringBase::getSoftFloatCmpLibcallPredicate().
[Mips] Fix symbol names for private functions in relocation hints (#230572)
Fixes errors such as
ld.lld: error: undefined symbol: compiler_rt.rem_pio2.rem_pio2
>>> referenced by compiler_rt
>>>
/home/alexrp/.cache/zig/o/a0062b09f46a198885da4835be133933/compiler_rt_zcu.o:(tan)
in archive
/home/alexrp/.cache/zig/o/a0062b09f46a198885da4835be133933/libcompiler_rt.a
which would happen because the actual symbol is
`$compiler_rt.rem_pio2.rem_pio2` for O32 or
`.Lcompiler_rt.rem_pio2.rem_pio2` for N64/N32.
[ViewOpGraph] Re-use `AsmState` to significantly speed-up printing. (#208240)
I was trying to print some rather large functions (5000+ ops) to a
graphivz/dot file, and noticed that the `ViewOpGraph` pass is rather
slow.
`--view-op-graph` previously always called `Operation::printAsOperand`
with the default `OpPrintingFlags` causing a new `AsmState` to be
constructed every time, which is extremely inefficient.
This patch instead creates a top-level `AsmState` once for each pass
run, and re-uses it for all printing.
This makes printing graphviz/dot files for large graphs *significantly*
faster: A single `func.func` with 5000 operations previously took 33
seconds to print, and can now be printed (essentially) instantly :)
---
[29 lines not shown]
[SCEV] Make operand complexity order total. (#230252)
Update CompareSCEVComplexity to return a total order for SCEVs. The
missing cases added are
* SCEVs with matching pointers are ordered by their use-flags
* SCEVUnknowns get assigned an ID at construction and are ordered by it.
The latter is the main change, ensuring SCEVUnkowns are ordered
determinstically by the ID depending on the construction order. Note
that this makes the order dependent on the order SCEVs are created. We
could try to use other sources to order SCEVUnkowns (comesBefore in
basic block, dominance), but those will be more expensive.
With that, we should only return 0 for equal nodes. Previously we could
return 0 for distinct nodes and IIUC this was the main reason for
MaxSCEVCompareDepth, because the operand compare loop could end up doing
a lot of work for distinct operands.
This came up as part of
[15 lines not shown]
[lldb] Explicitly handle unreachable code (#230372)
This silences a warning with MSVC, which does not treat the switch as
exhaustive unlike gcc and clang.
[LoopFlatten] Avoid asserting on same-width trip counts (#228130)
The backedge-taken count can already have the same type as the widened
compare RHS.
Use getNoopOrZeroExtend() instead of getZeroExtendExpr() (it requires a
strictly wider destination type) to prevent the assert when types are
same.
AMDGPU: Legalize SGPR tied src2 for V_FMAC_F16 in moveToVALU
When S_FMAC_F16 is moved to the VALU, it becomes V_FMAC_F16_t16_e64 or
V_FMAC_F16_fake16_e64. legalizeOperandsVOP3 only forced the tied src2
into a VGPR for V_FMAC_F32_e64 and V_FMAC_F16_e64, so an SGPR accumulator
fit within the constant bus limit was left in place, producing a
verifier error. Check for any SGPR src2 tied to the result instead.
This was exposed by frem.ll after converting it to use inreg arguments.
Co-authored-by: Claude Opus 5.5 <noreply at anthropic.com>
[flang] Enumeration Type: (PR 4/5) Lowering (#193571)
This PR lowers F2023 enumeration types (and their NEXT, PREVIOUS and
HUGE intrinsics). An enumeration type becomes a named record type with a
single i32 component, __ordinal, which holds the enumerator's 1-based
position:
`!fir.type<_QMmodTcolor{__ordinal:i32}>`
The type also gets a runtime type descriptor. As a result:
- SELECT TYPE can tell an enumeration apart from INTEGER.
- ALLOCATE(color :: x) works for a CLASS(*) variable.
- Formatted I/O works through the derived-type path.
Ordinary enumeration scalars are not boxed. Dummy arguments, function
results, PARAMETERs and constants use the standard derived-type
lowering.
[8 lines not shown]
[libc] Implement lchown in unistd (#230851)
Implement the standard POSIX function `lchown(const char *path, uid_t
owner, gid_t group)` in `<unistd.h>`.
Fixes #230837.
[ids-check] Print diff before updating issue (#230315)
This way if the script errors out trying to update the issue, the diff
is still logged to stdout.
[InstCombine] Combine ((ext(A) * ext(B)) >> N) -> ext(mulh(A, B)). (#224311)
The source values must be equally extended to double their width.
The result is sign/zero-extended for an arithmetic/logical shift.
This is a simplified version of DAGCombiner's combineShiftToMULH.
