59987 Commits

Author SHA1 Message Date
Shilei Tian
e801a10b44
[AMDGPU] Add the code generation support for llvm.[sin/cos].bf16 (#149631)
This is a partial support because some other instructions have not been upstreamed yet.
2025-07-21 11:01:59 -04:00
Shilei Tian
ba81903196
[gfx1250][SDAG] Lower unsafe bf16 divisions (#149628)
Co-authored-by: Kosarev, Ivan <Ivan.Kosarev@amd.com>
2025-07-21 10:58:08 -04:00
Shilei Tian
ce345cc793
[NFC][AMDGPU] Add an IR test for v_cvt_f16_bf8 (#149627)
This was left during the upstream.

Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
2025-07-21 10:54:55 -04:00
Shilei Tian
5d3b0578d8
[NFC][AMDGPU] Add an IR test for v_sqrt_bf16 (#149629)
This was left during the upstream.

Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
2025-07-21 09:40:14 -04:00
Nikita Popov
92c55a315e
[IR] Only allow lifetime.start/end on allocas (#149310)
lifetime.start and lifetime.end are primarily intended for use on
allocas, to enable stack coloring and other liveness optimizations. This
is necessary because all (static) allocas are hoisted into the entry
block, so lifetime markers are the only way to convey the actual
lifetimes.

However, lifetime.start and lifetime.end are currently *allowed* to be
used on non-alloca pointers. We don't actually do this in practice, but
just the mere fact that this is possible breaks the core purpose of the
lifetime markers, which is stack coloring of allocas. Stack coloring can
only work correctly if all lifetime markers for an alloca are
analyzable.

* If a lifetime marker may operate on multiple allocas via a select/phi,
we don't know which lifetime actually starts/ends and handle it
incorrectly (https://github.com/llvm/llvm-project/issues/104776).
* Stack coloring operates on the assumption that all lifetime markers
are visible, and not, for example, hidden behind a function call or
escaped pointer. It's not possible to change this, as part of the
purpose of lifetime markers is that they work even in the presence of
escaped pointers, where simple use analysis is insufficient.

I don't think there is any way to have coherent semantics for lifetime
markers on allocas, while also permitting them on arbitrary pointer
values.

This PR restricts lifetimes to operate on allocas only. As a followup, I
will also drop the size argument, which is superfluous if we always
operate on an alloca. (This change also renders various code handling
lifetime markers on non-alloca dead. I plan to clean up that kind of
code after dropping the size argument as well.)

In practice, I've only found a few places that currently produce
lifetimes on non-allocas:

* CoroEarly replaces the promise alloca with the result of an intrinsic,
which will later be replaced back with an alloca. I think this is the
only place where there is some legitimate loss of functionality, but I
don't think this is particularly important (I don't think we'd expect
the promise in a coroutine to admit useful lifetime optimization.)
* SafeStack moves unsafe allocas onto a separate frame. We can safely
drop lifetimes here, as SafeStack performs its own stack coloring.
* Similar for AddressSanitizer, it also moves allocas into separate
memory.
* LSR sometimes replaces the lifetime argument with a GEP chain of the
alloca (where the offsets ultimately cancel out). This is just
unnecessary. (Fixed separately in
https://github.com/llvm/llvm-project/pull/149492.)
* InferAddrSpaces sometimes makes lifetimes operate on an addrspacecast
of an alloca. I don't think this is necessary.
2025-07-21 15:04:50 +02:00
Simon Pilgrim
2b826dff01 [RISCV][rvv] setcc-fp-vp.ll - regenerate with missing riscv32/64 checks 2025-07-21 13:57:23 +01:00
Marcos Maronas
1d02de28dd
[SPIRV] Implement translation for llvm.modf.* intrinsics (#147556)
Based on
https://github.com/KhronosGroup/SPIRV-LLVM-Translator/pull/3100, I'm
adding translation for `llvm.modf.*` intrinsics.
2025-07-21 14:51:34 +02:00
Haohai Wen
43a829a7e8
[X86] Add assertion required for pseudo probe desc check test (#149800)
Fix failed test introduced in #148564.
2025-07-21 19:28:14 +08:00
Alex Bradbury
fc69f25a8f
[RISCV] Convert LWU to LW if possible in RISCVOptWInstrs (#144703)
After the refactoring in #149710 the logic change is trivial.

Motivation for preferring sign-extended 32-bit loads (LW) vs
zero-extended (LWU):
* LW is compressible while LWU is not.
* Helps to minimise the diff vs RV32 (e.g. LWU vs LW)
* Helps to minimise distracting diffs vs GCC. I see this come up
frequently when comparing GCC code and in these cases it's a red
herring.

Similar normalisation could be done for LHU and LH, but this is less
well motivated as there is a compressed LHU (and if performing the
change in RISCVOptWInstrs it wouldn't be done for RV32). There is a
compressed LBU but not LB, meaning doing a similar normalisation for
byte-sized loads would actually be a regression in terms of code size.
Load narrowing when allowed by hasAllNBitUsers isn't explored in this
patch.

This changes ~20500 instructions in an RVA22 build of the
llvm-test-suite including SPEC 2017. As part of the review, the option
of doing the change at ISel time was explored but was found to be less
effective.
2025-07-21 11:48:33 +01:00
David Green
3b8adcfd92
[GlobalISel] Add computeNumSignBits for ASHR (#139503) 2025-07-21 10:23:26 +01:00
Jay Foad
5a9c201bc9
[CodeGen] Fix INLINEASM regclass numbers to match names in tests (#142359)
INLINEASM operands like "2097162 /* regdef:SReg_32 */" tend to get
broken over time as the register class definitions get updated, so the
numbers change. Fix the numbers to match the names in the comments.
2025-07-21 10:23:08 +01:00
Jeremy Morse
040bffc633
[DebugInfo][AMDGPU] Convert a debug-intrinsic method to debug records (#149505)
It appears this wasn't handled in the initial migration a year ago,
seemingly because it didn't lead to any test failures. Find and interpret
debug records in the same way the original code handled intrinsics. Note
that we drop a call to copyMetadata: debug records can't carry additional
metadata like instructions, nothing relies on this in AMDGPU AFAIUI.
2025-07-21 10:07:14 +01:00
Diana Picus
20d8398825
[AMDGPU] ISel & PEI for whole wave functions (#145858)
Whole wave functions are functions that will run with a full EXEC mask.
They will not be invoked directly, but instead will be launched by way
of a new intrinsic, `llvm.amdgcn.call.whole.wave` (to be added in
a future patch). These functions are meant as an alternative to the
`llvm.amdgcn.init.whole.wave` or `llvm.amdgcn.strict.wwm` intrinsics.

Whole wave functions will set EXEC to -1 in the prologue and restore the
original value of EXEC in the epilogue. They must have a special first
argument, `i1 %active`, that is going to be mapped to EXEC. They may
have either the default calling convention or amdgpu_gfx. The inactive
lanes need to be preserved for all registers used, active lanes only for
the CSRs.

At the IR level, arguments to a whole wave function (other than
`%active`) contain poison in their inactive lanes. Likewise, the return
value for the inactive lanes is poison.

This patch contains the following work:
* 2 new pseudos, SI_SETUP_WHOLE_WAVE_FUNC and SI_WHOLE_WAVE_FUNC_RETURN
  used for managing the EXEC mask. SI_SETUP_WHOLE_WAVE_FUNC will return
  a SReg_1 representing `%active`, which needs to be passed into
  SI_WHOLE_WAVE_FUNC_RETURN.
* SelectionDAG support for generating these 2 new pseudos and the
  special handling of %active. Since the return may be in a different
  basic block, it's difficult to add the virtual reg for %active to
  SI_WHOLE_WAVE_FUNC_RETURN, so we initially generate an IMPLICIT_DEF
  which is later replaced via a custom inserter.
* Expansion of the 2 pseudos during prolog/epilog insertion. PEI also
  marks any used VGPRs as WWM registers, which are then spilled and
  restored with the usual logic.

Future patches will include the `llvm.amdgcn.call.whole.wave` intrinsic
and a lot of optimization work (especially in order to reduce spills
around function calls).

---------

Co-authored-by: Matt Arsenault <Matthew.Arsenault@amd.com>
Co-authored-by: Shilei Tian <i@tianshilei.me>
2025-07-21 10:39:09 +02:00
hev
8a307ae619
[LoongArch] Fix failure to widen operand for [X]VMSK{LT,GE,NE}Z (#149442)
Reported-by: tangyan <tangyan01@loongson.cn>
2025-07-21 16:36:49 +08:00
Sudharsan Veeravalli
84e689b1db
[RISCV] Swap source register operands in QC_SHLADD ISEL patterns (#149697)
The instruction does `rd = (rs1 << shamt) + rs2` but the ISEL patterns
had `rs1` and `rs2` the other way around which is incorrect.
2025-07-21 12:03:55 +05:30
Haohai Wen
6193dd5553
[PseudoProbe] Warn on illegal guid (#148564)
Check whether guid exists in pseudo probe desc when emitting pseudo
probe.
2025-07-21 13:56:09 +08:00
Pecco
8ae4dee4d0
[NVPTX] Lower stmatrix intrinsics to PTX (#148561)
Lower stmatrix intrinsics defined in #148377 to PTX. See [PTX
Doc](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-store-instruction-stmatrix).

---------

Co-authored-by: peterbell10 <peterbell10@openai.com>
2025-07-20 21:06:45 -07:00
Jasmine Tang
343f7475be
[WebAssembly] Add support for memcmp expansion (#148298)
Fixes https://github.com/llvm/llvm-project/issues/61400

Added test case in llvm/test/CodeGen/WebAssembly/memcmp-expand.ll
2025-07-20 10:27:42 -07:00
Tom Vijlbrief
9f29007999
[AVR] Force to emit relocation slots for relative branch instructions (#145291)
fixes https://github.com/llvm/llvm-project/issues/133579
2025-07-20 22:10:31 +08:00
Simon Pilgrim
17c7c2ebe8
[DAG] Add missing Depth argument to isGuaranteedNotToBeUndefOrPoison calls inside SimplifyDemanded methods (#149550)
Ensure we don't exceed the maximum recursion depth
2025-07-20 13:06:55 +01:00
Alex Bradbury
c58225f757
[RISCV] Add RISCV::SUBW to RISCVOptWInstrs::stripWSuffixes (#149071)
This is purely a benefit for reducing unnecessary diffs between RV32 and
RV64, as RVC does have a compressed form of SUBW (so SUB isn't more
compressible). This affects ~57.2k instructions in an rva22u64 build of
llvm-test-suite with SPEC CPU 2017 included.
2025-07-20 09:59:41 +01:00
Jasmine Tang
e7ac49977a
[InstSimplify] Add poison propagation for trivially vectorizable intrinsics (#149243)
Fixes https://github.com/llvm/llvm-project/issues/146769

Test cases added to
`llvm/test/Transforms/InstSimplify/fold-intrinsics.ll`
2025-07-19 19:37:21 -07:00
David Green
07100c6658 [AArch64] Clean up abd lowering. NFC
The neon ABD intrinsics are lowering to ISD::ABDS/ISD::ABDU nodes in SDAG. This
adds the same for GISel so that we can remove the AArch64sabd/AArch64uabd
PatFrags, reducing the number of patterns needed. Some very basic legalization
info is added for G_ABD nodes.
2025-07-19 22:13:13 +01:00
David Green
269e3091a6 [AArch64][GlobalISel] Add test coverage for some abd cases. NFC 2025-07-19 21:56:28 +01:00
Durgadoss R
3866e4e7f8
[NVPTX] Add im2colw/w128 modes support to TMA intrinsics (#148863)
This patch adds support for the im2col-w/w128 and scatter/gather modes
for TMA Copy and Prefetch intrinsics, completing support for all the
available modes. These are lowered through tablegen, building
on top of earlier patches.

* lit tests are added for all the combinations and verified with a
  12.8 ptxas executable.
* Documentation is updated in the NVPTXUsage.rst file.

Signed-off-by: Durgadoss R <durgadossr@nvidia.com>
2025-07-19 17:16:21 +05:30
Alex MacLean
b3c9ed151f
[NVPTX][test] regenerate some tests broken by successive changes (#149611)
#149393 and #149571 landed in quick succession requiring
some tests to be regenerated to account for their interactions.
2025-07-18 18:49:49 -07:00
Alex MacLean
965b68e8f2
[NVPTX] Prevent fptrunc of v2f32 from being folded into store (#149571) 2025-07-18 14:20:13 -07:00
Princeton Ferro
d63ab5467d
[NVPTX] don't erase CopyToRegs when folding movs into loads (#149393)
We may still need to keep CopyToReg even after folding uses into vector
loads, since the original register may be used in other blocks.

Partially reverts 1fdbe6984976d9e85ab3b1a93e8de434a85c5646
2025-07-18 14:11:31 -07:00
Shilei Tian
d46de86ca4
[NFC][AMDGPU] Re-enable two tests previously disabled due to missing upstream features (#149568)
This PR re-enables two tests that were previously disabled because they
depended on features not yet upstreamed.
2025-07-18 17:04:34 -04:00
Jeffrey Byrnes
695660cdfd
[AMDGPU] Provide control to force VGPR MFMA form (#148079)
This gives an override to the user to force select VGPR form of MFMA.
Eventually we will drop this in favor of compiler making better
decisions, but this provides a mechanism for users to address the cases
where MayNeedAGPRs favors the AGPR form and performance is degraded due
to poor RA.
2025-07-18 13:53:17 -07:00
Peter Collingbourne
9878ef3abd
CodeGen: Respect function align attribute if less than preferred alignment.
Reviewers: arsenm, efriedma-quic

Reviewed By: arsenm

Pull Request: https://github.com/llvm/llvm-project/pull/149444
2025-07-18 13:33:46 -07:00
Philip Reames
c5f0c4ad37 [RISCV][IA] Add test coverage for vp.store of interleaveN with one active 2025-07-18 13:33:23 -07:00
Alex MacLean
9d9662e4bd
[NVPTX][test] fixup version for ptxas on trunc-tofp.ll (#149558) 2025-07-18 13:27:31 -07:00
Stanislav Mekhanoshin
cfa918bec1
[AMDGPU] Select flat GVS atomics on gfx1250 (#149554) 2025-07-18 12:31:29 -07:00
Tobias Decking
10b0dee97d
[X86] Ensure that bit reversals of byte vectors are properly lowered on pure GFNI targets (#148304)
Fixes #148238.

When GFNI is present, custom bit reversal lowerings for scalar integers
become active. They work by swapping the bytes in the scalar value and
then reversing bits in a vector of bytes. However, the custom bit
reversal lowering for a vector of bytes is disabled if GFNI is present
in isolation, resulting messed up code.

---------

Co-authored-by: Simon Pilgrim <llvm-dev@redking.me.uk>
2025-07-18 19:14:34 +01:00
Philip Reames
87c2adbb58 [RISCV][IA] Precommit tests for deinterleaveN of masked.load 2025-07-18 10:39:11 -07:00
Shilei Tian
2c50e4cac2
[AMDGPU] Add support for v_sat_pk4_i4_[i8,u8] on gfx1250 (#149528)
Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
Co-authored-by: Foad, Jay <Jay.Foad@amd.com>
2025-07-18 13:08:50 -04:00
Shilei Tian
e11d28faee
[AMDGPU] Add support for v_permlane16_swap_b32 on gfx1250 (#149518)
Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
2025-07-18 13:05:08 -04:00
Brox Chen
5138b61a25
[AMDGPU][True16][Codegen] remove packed build_vector pattern from true16 (#148715)
Some of the packed build_vector use vgpr_32 for i16/f16/bf16. 

In gfx11, bf16 arithmetic get promoted to f32 and this is done via v2i16
pack. In true16 mode this v2i16 pack is selected to a
build_vector/v_lshlrev pattern which only accepts VGPR32. This causes
isel to insert an illegal copy "vgpr32 = copy vgpr16" between def and
use. In the end this illegal copy confuses cse pass and trigger wrong
code elimination.

Remove the packed build_vector pattern from true16. After removal, ISel
will use vgpr16 build_vector patterns instead.
2025-07-18 12:55:11 -04:00
Annu Singh
148fd6ed0a
[DAG] Adding abdu/abds to canCreateUndefOrPoison (#149017)
Fixes #147695 
- [Alive2 test - freeze abdu](https://alive2.llvm.org/ce/z/aafeJs)
- [Alive 2 test - freeze abds](https://alive2.llvm.org/ce/z/XrSmP4)

---------

Co-authored-by: Simon Pilgrim <llvm-dev@redking.me.uk>
2025-07-18 17:00:44 +01:00
Alex MacLean
f73e163278
[DAGCombiner] Fold [us]itofp of truncate (#149391) 2025-07-18 08:10:20 -07:00
Shilei Tian
95b69e0e70
[AMDGPU] Add support for v_prng_b32 on gfx1250 (#149450)
Co-authored-by: Mekhanoshin, Stanislav <Stanislav.Mekhanoshin@amd.com>
2025-07-18 10:59:47 -04:00
Paul Walker
44cd5027f8
[LLVM][CodeGen][SVE] List MVTs that are desirable for extending loads. (#149153)
Extend AArch64TargetLowering::isVectorLoadExtDesirable to specify the
set of MVT for which load extension is desirable.

Fixes https://github.com/llvm/llvm-project/issues/148939
2025-07-18 15:34:48 +01:00
Sergei Barannikov
6112ebde0c
[RISCV] Guard CFI emission code with MF.needsFrameMoves() (#136060)
Currently, AsmPrinter skips CFI instructions created by a backend if
they are not needed. I'd like to change that so that it always
prints/encodes CFI instructions if a backend created them.

This change should slightly (perhaps negligibly) improve compile time as
post-PEI passes no longer need to skip over these instructions in
no-exceptions no-debug builds, and will allow to simplify convoluted
logic in AsmPrinter once other targets stop emitting CFI instructions
when they are not needed (that's my final goal).

The changes in a test seem to be caused by slightly different post-RA
scheduling in the absence of CFI instructions.
2025-07-18 16:49:30 +03:00
Matt Arsenault
d883d5fecf AMDGPU: Add testcase with bad regalloc behavior
This demonstrates poor allocation due to not ordering
AV classes relative to the A and V classes
2025-07-18 17:46:54 +09:00
Matt Arsenault
8f3e78f971
AMDGPU: Add pass to replace constant materialize with AV pseudos (#149292)
If we have a v_mov_b32 or v_accvgpr_write_b32 with an inline immediate,
replace it with a pseudo which writes to the combined AV_* class. This
relaxes the operand constraints, which will allow the allocator to
inflate the register class to AV_* to potentially avoid spilling.

The allocator does not know how to replace an instruction to enable
the change of register class. I originally tried to do this by changing
all of the places we introduce v_mov_b32 with immediate, but it's along
tail of niche cases that require manual updating. Plus we can restrict
this to only run on functions where we know we will be allocating AGPRs.
2025-07-18 17:15:38 +09:00
tangaac
64a0478e08
[LoongArch] Strengthen stack size estimation for LSX/LASX extension (#146455)
This patch adds an emergency spill slot when ran out of registers.
PR #139201 introduces `vstelm` instructions with only 8-bit imm offset, 
it causes no spill slot to store the spill registers.
2025-07-18 16:12:11 +08:00
Fabian Ritter
efedd49a22
[AMDGPU][SDAG] Test ISD::PTRADD handling in VOP3 patterns (#143880)
Pre-committing tests to show improvements in a follow-up PR.
2025-07-18 10:04:51 +02:00
Fabian Ritter
daa6de37ba
[AMDGPU][SDAG] Add target-specific ISD::PTRADD combines (#143673)
This patch adds several (AMDGPU-)target-specific DAG combines for
ISD::PTRADD nodes that reproduce existing similar transforms for
ISD::ADD nodes. There is no functional change intended for the existing
target-specific PTRADD combine.

For SWDEV-516125.
2025-07-18 10:00:54 +02:00
Matt Arsenault
176ae32de0
AMDGPU: Fix introducing use of killed vgpr in gfx908 agpr copy (#149291)
When searching for an existing VGPR source for an AGPR to AGPR
copy on gfx908, this wasn't verifying the vgpr wasn't killed by
other prior uses.
2025-07-18 15:34:47 +09:00