Remove another build_vector pattern which takes a i16 but placed in a
VGPR_32 from true16 mode. This stop isel from generating illegal
"vgpr_32 = COPY vgpr_16".
ISel will use vgpr16 build vector pattern in true16 mode instead
The arithmetic expansion requires fewer registers, and is often fewer
instructions. The critical path does increase by (up to) one
instruction.
This is a sub-case of the expansion we do without zicond, but restricted
specifically to the simm12 case. In the general case where the other
source is a register using zicond is likely better. (Edit: While
technically true, this is a bit misleading, we do this in
combineSelectToBinOp which is also used in the zicond path, just further
down.)
Add RISCVVLOptimizer supported for unit-stride, strided, and indexed
strided segmented stores. The biggest change was adding the capability
to look through INSERT_SUBREG, which was used for composing segmented
register class values.
Fix#149350
This patch adds a new `TargetLowering` hook `lowerEHPadEntry()` that is
called at the start of lowering EH pads in SelectionDAG. This allows the
insertion of target-specific actions on entry to exception handlers.
This is used on AArch64 to insert SME streaming-mode switches at landing
pads. This is needed as exception handlers are always entered with
PSTATE.SM off, and the function needs to resume the streaming mode of
the function body.
- Add clang built-ins + sema/codegen
- Add IR Intrinsic + verifier
- Add DAG/GlobalISel codegen for the intrinsics
- Add lowering in SIMemoryLegalizer using a MMO flag.
Support tail calls to whole wave functions (trivial) and from whole wave
functions (slightly more involved because we need a new pseudo for the
tail call return, that patches up the EXEC mask).
Move the expansion of whole wave function return pseudos (regular and
tail call returns) to prolog epilog insertion, since that's where we
patch up the EXEC mask.
This hook replaces inline asm with LLVM intrinsics. It was intended to
match inline assembly implementations of bswap in libc headers and
replace them more optimizable implementations.
At this point, it has outlived its usefulness (see
https://github.com/llvm/llvm-project/issues/156571#issuecomment-3247638412),
as libc implementations no longer use inline assembly for this purpose.
Additionally, it breaks the "black box" property of inline assembly,
which some languages like Rust would like to guarantee.
Fixes https://github.com/llvm/llvm-project/issues/156571.
getTypeAllocSize() currently works by taking the type store size and
aligning it to the ABI alignment. However, this ends up doing redundant
work in various cases, for example arrays will unnecessarily repeat the
alignment step, and structs will fetch the StructLayout multiple times.
As this code is rather hot (it is called every time we need to calculate
GEP offsets for example), specialize the implementation. This repeats a
small amount of logic from getAlignment(), but I think that's
worthwhile.
The generated MIR fails machine verifier as stack pointer is being
modified without appropriate attributes in frameInfo. This PR fixes this
issue by adding adjustsStack=true attribute. Fixes the post commit
regression identified in #155673.
Correctly model these without AV_* operands. This is another
step towards removing the special casing in
TargetInstrInfo::getRegClass. Also add some tests for this.
This is similar to the recently added (X & -4096) == 0 -> (X >> 12) ==
0,
but operating only on the lower 32 bits.
This also removes the (X & (1 << 31)) == 0 -> (xor (srliw X, 31), 1)
isel pattern. seqz and xori 1 should have similar cost and encoding
size.
Currently, NVPTX fails to lower some tcgen05 ops for the
sm103 target. It was observed in a downstream project, Triton, in
https://github.com/triton-lang/triton/pull/8045.
So, this patch adds sm103 targets to the hasTcgen05Instructions() method.
All the Lit tests for tcgen05-* are updated as well.
To a slide1up, if the scalar value we're sliding in was extracted from
the first element of a vector, we can use a normal vslideup of 1 instead
with its passthru being that vector. This can eliminate an
extract_element instruction (i.e. vfmv.f.s, vmv.x.s).
---------
Co-authored-by: Craig Topper <craig.topper@sifive.com>
If an instruction's demanded VL is a virtual register defined by a vleff
instruction, it might not dominate and fail to have its VL reduced.
In leiu of the output VL, we can try and use the AVL passed to the vleff
itself since it will be at least greater than or equal the original VL.
I tried to create an LLVM IR test for this in but didn't have any luck
because the scheduler kept on moving the instruction past the vleff, so
it always dominated. So I've just included some mir tests instead.
Make sure the MFMA VGPR to AGPR InstrMapping table is complete.
I think I got everything, except the full cross product of input
types with the mfma scale intrinsics. Also makes sure we have
coverage for smfmac and mfma_scale cases.
This should not require aligned registers. Fixes expensive_checks
test failure. I don't see a better way until the new system
to specify the alignment per register is done.
The instruction definitions for loads and stores do not
accurately model the operand constraints of loads and stores
with AGPRs. They use AV register classes, plus a hack
a hack in getRegClass/getOpRegClass to avoid using AGPRs or
AV classes with the multiple operand cases, but it did not
consider the 3 operand case.
Model this correctly by using separate all-VGPR and all-AGPR
variants for the cases with multiple data operands.
This does regress the assembler errors on gfx908 for the
multi-operand cases. It now reports a generic operand
invalid error for GPU instead of the specific message
that agpr loads and stores aren't supported.
In the future AMDGPURewriteAGPRCopyMFMA should be taught
to replace the VGPR forms with the AGPR ones.
Most of the diff is fighting the DS pseudo structure. The
mnemonic was being used as the key to SIMCInstr, which is a
collision in the AGPR case. We also need to go out of our way
to make sure we are using the gfx9+ variants of the pseudos
without the m0 use. The DS multiclasses could use a lot of
cleanup.
Fixes#155777
The existing test case only exhausted the GPRs so the stack wasn't tested if FPRs were also available for arguments. This new test exhausts the GPRs and FPRs.
fixes#148051
- update EmitHLSLBuiltinExpr in CGHLSLBuiltins.cpp to toggle intrinsics
by target
- Add a GENERATE_HLSL_INTRINSIC_FUNCTION for isinf in CGHLSLRuntime.h
- Update the SPIRVInstructionSelector.cpp to emit the OpSinf instruction
- Updates the isinf.hlsl test to check spirv intrinsic generation
- add OpIsinf.ll tests
According AMX ABI, tile registers (including config) are volatile hence
requiring caller to save/restore config register. This is done in X86's
FastPreTileConfig pass. Currently the PLDTILECFGV instruction is emitted
immediately after the call which can be problematic if call returns a
value in say rax register and AMX tile is configured using the same
register. This PR addresses this issue by ensuring that PLDTILECFGV is
sinked closer to first instruction using a tile after the call.
This patch implements direct N-way splitting for wide scalar shifts
instead
of recursive binary splitting. For example, an i512 G_SHL can now be
split
directly into 8 i64 operations rather than going through i256 -> i128 ->
i64.
The main motivation behind this is to alleviate (although not entirely
fix)
pathological compile time issues with huge types, like i4224. The
problem
we see is that the recursive splitting strategy combined with our messy
artifact combiner ends up with terribly long compiles as tons of
intermediate
artifacts are generated, and then attempted to be combined ad-nauseum.
Going directly from the large shifts to the destination types
short-circuits
a lot of these issues, but it's still an abuse of the backend and
front-ends
should never be doing this sort of thing.
RISCV does not provide a native atomic subtract instruction, so this
patch lowers `G_ATOMICRMW_SUB` by negating the RHS value and performing
an atomic add. The legalization rules in `RISCVLegalizerInfo` are
updated accordingly, with libcall fallbacks when `StdExtA` is not
available, and intrinsic legalization is extended to support
`riscv_masked_atomicrmw_sub`.
For example, lowering
`%1 = atomicrmw sub ptr %a, i32 1 seq_cst`
on riscv32a produces:
```
li a1, -1
amoadd.w.aqrl a0, a1, (a0)
```
On riscv64a, where the RHS type is narrower than XLEN, it currently
produces:
```
li a1, 1
neg a1, a1
amoadd.w.aqrl a0, a1, (a0)
```
There is still a constant-folding or InstConbiner gap. For instance,
lowering
```
%b = sub i32 %x, %y
%1 = atomicrmw sub ptr %a, i32 %b seq_cst
```
generates:
```
subw a1, a1, a2
neg a1, a1
amoadd.w.aqrl a0, a1, (a0)
```
This sequence could be optimized further to eliminate the redundant neg.
Addressing this may require improvements in the Combiner or Peephole
Optimizer in future work.
---------
Co-authored-by: Kane Wang <kanewang95@foxmail.com>
This patch implements a correctly rounded expansion of the frem
instruction in LLVM IR. This is useful for target architectures for
which such an expansion is too involved to be implement in ISel
Lowering. The expansion is based on the code from the AMD device libs
and has been tested successfully against the OpenCL conformance tests on
amdgpu. The expansion is implemented in the preexisting "expand-fp"
pass. It replaces the expansion of "frem" in ISel for the amdgpu target;
it is enabled for targets which do not directly support "frem" and for
which no matching "fmod" LibCall is available.
---------
Co-authored-by: Matt Arsenault <Matthew.Arsenault@amd.com>
By only checking type legality we didn't account for 128/256-bit ops
being run on non-AVX512VL targets, or vXi8/i16 ops being run on
non-AVX512BW targets
This check is cropping up in several places now and I intend to hoist it
out into a common helper, but this initial fix needs to be as clean as
possible to be back ported to 21.X
Fixes#156256
Add the following patterns to performSignExtendInRegCombine:
* SIGN_EXTEND_INREG (CSEL 0, 1, cc), i1 --> CSEL 0, -1, cc
* SIGN_EXTEND_INREG (CSEL 1, 0, cc), i1 --> CSEL -1, 0, cc
The combined forms can be matched to a CSETM.
It specifies the target machine as AMDGPU for
AMDGPUAnnotateUniformValues pass-related test (that uses UA). Before in
its absense, the UA would consider everything Uniform resulting in
setting metadata incorrectly for AMDGPU. Now, after specifying the
AMDGPU, the UA would be rightful sets the right metadata as the test
gets commpiled for AMDGPU.
Reapplying now that buildbot has picked up the new configuration
that does not use -treat-scalable-fixed-error-as-warning.
-----
This removes the `LLVM_ENABLE_STRICT_FIXED_SIZE_VECTORS` cmake option
and the `-treat-scalable-fixed-error-as-warning` opt flag.
We stopped treating these as warnings by default a long time ago
(62f09d788f9fc540db12f3cfa2f98760071fca96), so I don't think it makes
sense to retain these options at this point. Accessing a scalable
TypeSize as fixed should always result in an error.
For IR like this:
%mask = ... @llvm.vector.interleave2(<vscale x 16 x i1> %a, <vscale x 16
x i1> %a)
%vec = ... @llvm.masked.load(..., <vscale x 32 x i1> %mask, ...)
%dvec = ... @llvm.vector.deinterleave2(<vscale x 32 x i8> %vec)
where we're deinterleaving a wide masked load of the supported type
and with an interleaved mask we can lower this directly to a ld2b
instruction. Similarly we can also support other variants of ld2
and ld4.
This PR adds a DAG combine to spot such patterns and lower to ld2X
or ld4X variants accordingly, whilst being careful to ensure the
masked load is only used by the deinterleave intrinsic.
Handle a special case for copies from AGPR VGPR on the MFMA inputs.
If the "input" is really a subregister def, we will not see the
usual copy to VGPR for src2, only the read of the subregister def.
Not sure if this pattern appears in practice.
Before this patch, the selection for `QC_INSB` and `QC_INSBI` entirely
happens in C++, and does not support more than one non-constant input.
This patch seeks to rectify this shortcoming, by moving the C++ into a
target-specific DAGCombine, and adding `RISCV::QC_INSB`. One advantage
is this simplifies the code for handling `QC_INSBI`, as the C++ no
longer needs to choose between the two instructions based on the
inserted value (this is still done, but via ISel Patterns).
Another advantage of the DAGCombine is that this introduction can also
shift the inserted value to the `QC_INSB`, which our patterns need (and
were previously doing to the constant), and this shift can be
CSE'd/optimised with any prior shifts, if they exist. This allows the
inserted value to be variable, rather than a constant.
