AMDGPU: Define v_mfma_f32_{16x16x128|32x32x64}_f8f6f4 instructions (#116723)

These use a new VOP3PX encoding for the v_mfma_scale_* instructions,
which bundles the pre-scale v_mfma_ld_scale_b32. Non

AMDGPU: Define v_mfma_f32_{16x16x128|32x32x64}_f8f6f4 instructions (#116723)

These use a new VOP3PX encoding for the v_mfma_scale_* instructions,
which bundles the pre-scale v_mfma_ld_scale_b32. None of the modifiers
are supported yet (op_sel, neg or clamp).

I'm not sure the intrinsic should really expose op_sel (or any of the
others). If I'm reading the documentation correctly, we should be able
to just have the raw scale operands and auto-match op_sel to byte
extract patterns.

The op_sel syntax also seems extra horrible in this usage, especially with the
usual assumed op_sel_hi=-1 behavior.

show more ...

[MC] Replace HasAggressiveSymbolFolding with SetDirectiveSuppressesReloc. NFC

The variable and its comment could lead to confusion (AMDGPU
unnecessarily set it).

[AMDGPU][MC] Use normal ELF syntax for section switching (#77267)

For some reasons `SunStyleELFSectionSwitchSyntax` is set to `true` for
AMDGPU, but according to
https://github.com/llvm/llvm-proje

[AMDGPU][MC] Use normal ELF syntax for section switching (#77267)

For some reasons `SunStyleELFSectionSwitchSyntax` is set to `true` for
AMDGPU, but according to
https://github.com/llvm/llvm-project/issues/64862#issuecomment-1880419239
that syntax is only limited to Sun system.

Fix #64862.

show more ...

[AsmPrinter][AMDGPU] Generate uwtable entries in .eh_frame

Consider only targets where `MCAsmInfo::ExceptionsType == ExceptionHandling::None`
and that support CFI (when `MCAsmInfo::UsesCFIForDebug`

[AsmPrinter][AMDGPU] Generate uwtable entries in .eh_frame

Consider only targets where `MCAsmInfo::ExceptionsType == ExceptionHandling::None`
and that support CFI (when `MCAsmInfo::UsesCFIForDebug` is set to true):
currently, only AMDGPU.

This patch enables the emission of CFI information in the .eh_frame
section when the uwtable attribute is present on a function.

Before, we could generate CFI information for debugging puproses only.

This patch prepares AMDGPU to support collecting GPU stack traces in the future.

I did a first implementation (https://reviews.llvm.org/D139024)
but at the time I had not realized that no other platform used
`UsesCFIForDebug`.

Reviewed By: scott.linder

Differential Revision: https://reviews.llvm.org/D151806

show more ...

Simplify with hasFeature. NFC

[NFC][TargetParser] Remove llvm/ADT/Triple.h

I also ran `git clang-format` to get the headers in the right order for
the new location, which has changed the order of other headers in two
files.

[Target] Remove redundant member initialization (NFC)

Identified with readability-redundant-member-init.

Revert "[llvm] Remove redundant member initialization (NFC)"

This reverts commit fd4808887ee47f3ec8a030e9211169ef4fb094c3.

This patch causes gcc to issue a lot of warnings like:

warning: base cl

Revert "[llvm] Remove redundant member initialization (NFC)"

This reverts commit fd4808887ee47f3ec8a030e9211169ef4fb094c3.

This patch causes gcc to issue a lot of warnings like:

warning: base class ‘class llvm::MCParsedAsmOperand’ should be
explicitly initialized in the copy constructor [-Wextra]

show more ...

[llvm] Remove redundant member initialization (NFC)

Identified with readability-redundant-member-init.

[AMDGPU] Do not generate ELF symbols for the local branch target labels

The compiler was generating symbols in the final code object for local
branch target labels. This bloats the code object, slow

[AMDGPU] Do not generate ELF symbols for the local branch target labels

The compiler was generating symbols in the final code object for local
branch target labels. This bloats the code object, slows down the loader,
and is only used to simplify disassembly.

Use '--symbolize-operands' with llvm-objdump to improve readability of the
branch target operands in disassembly.

Fixes: SWDEV-312223

Reviewed By: scott.linder

Differential Revision: https://reviews.llvm.org/D114273

show more ...

[MCAsmInfo] Support UsesCFIForDebug for targets with no exception handling

This change enables emitting CFI unwind information for debugging purpose
for targets with MCAsmInfo::ExceptionsType == Exc

[MCAsmInfo] Support UsesCFIForDebug for targets with no exception handling

This change enables emitting CFI unwind information for debugging purpose
for targets with MCAsmInfo::ExceptionsType == ExceptionHandling::None.

Currently generating CFI unwind information is entangled with supporting
the exceptions, even when AsmPrinter explicitly recognizes that the unwind
tables are being generated as debug information.

In fact, the unwind information is not generated even if we specify
--force-dwarf-frame-section, unless exceptions are enabled. The LIT test
llvm/test/CodeGen/AMDGPU/debug_frame.ll demonstrates this behavior.

Enable this option for AMDGPU to prepare for future patches which add
complete CFI support.

Reviewed By: dblaikie, MaskRay

Differential Revision: https://reviews.llvm.org/D78778

show more ...

[AMDGPU] Corrected directive to use for ELF weak refs

WeakRefDirective should specify a directive to declare "a global as being a weak undefined symbol".
The directive used by AMDGPU was incorrect -

[AMDGPU] Corrected directive to use for ELF weak refs

WeakRefDirective should specify a directive to declare "a global as being a weak undefined symbol".
The directive used by AMDGPU was incorrect - ".weakref" was intended for other purposes.
The correct directive is ".weak" and it is already defined as default for ELF.
So the redefinition was removed.

Reviewers: arsenm, rampitec

Differential Revision: https://reviews.llvm.org/D87762

show more ...

[MC] Add UseIntegratedAssembler = false. NFC

[AMDGPU] Print DWARF register numbers in AMDGPUInstPrinter

Summary:
Explanation is in a comment in the diff, but essentially printing a
physical register name here is ambiguous. Until we can impleme

[AMDGPU] Print DWARF register numbers in AMDGPUInstPrinter

Summary:
Explanation is in a comment in the diff, but essentially printing a
physical register name here is ambiguous. Until we can implement
printing a DWARF register name here just use the encoding directly.

Tags: #llvm

Differential Revision: https://reviews.llvm.org/D76253

show more ...

[Mips] Use appropriate private label prefix based on Mips ABI

MipsMCAsmInfo was using '$' prefix for Mips32 and '.L' for Mips64
regardless of -target-abi option. By passing MCTargetOptions to MCAsmI

[Mips] Use appropriate private label prefix based on Mips ABI

MipsMCAsmInfo was using '$' prefix for Mips32 and '.L' for Mips64
regardless of -target-abi option. By passing MCTargetOptions to MCAsmInfo
we can find out Mips ABI and pick appropriate prefix.

Tags: #llvm, #clang, #lldb

Differential Revision: https://reviews.llvm.org/D66795

show more ...

MC: Allow getMaxInstLength to depend on the subtarget

Keep it optional in cases this is ever needed in some global
context. Currently it's only used for getting an upper bound inline
asm code size.

MC: Allow getMaxInstLength to depend on the subtarget

Keep it optional in cases this is ever needed in some global
context. Currently it's only used for getting an upper bound inline
asm code size.

For AMDGPU, gfx10 increases the maximum instruction size to
20-bytes. This avoids penalizing older subtargets when estimating code
size, and making some annoying branch relaxation test adjustments.

llvm-svn: 361405

show more ...

Update the file headers across all of the LLVM projects in the monorepo
to reflect the new license.

We understand that people may be surprised that we're moving the header
entirely to discuss the ne

Update the file headers across all of the LLVM projects in the monorepo
to reflect the new license.

We understand that people may be surprised that we're moving the header
entirely to discuss the new license. We checked this carefully with the
Foundation's lawyer and we believe this is the correct approach.

Essentially, all code in the project is now made available by the LLVM
project under our new license, so you will see that the license headers
include that license only. Some of our contributors have contributed
code under our old license, and accordingly, we have retained a copy of
our old license notice in the top-level files in each project and
repository.

llvm-svn: 351636

show more ...

AMDGPU: Set StackGrowsUp in MCAsmInfo

Not sure what this does though.

llvm-svn: 301229

AMDGPU: Set CodePointerSize to 8 for amdgcn

llvm-svn: 300470

Revert "AMDGPU: Set MCAsmInfo::PointerSize"

It breaks line tables because the patch is not complete, working on a complete one at the moment

This reverts commit r294031.

llvm-svn: 297118

AMDGPU: Set MCAsmInfo::PointerSize

llvm-svn: 294031

AMDGPU: Set MinInstAlignment

Not sure this actually changes anything

llvm-svn: 273947

AMDGPU: Fix maximum instruction size for amdgcn

This was causing the conservative estimate of inline asm
size to be twice as big as expected.

llvm-svn: 272956

AMDGPU: Emit functions sizes

Reviewers: arsenm

Subscribers: arsenm, llvm-commits

Differential Revision: http://reviews.llvm.org/D15951

llvm-svn: 257172

AMDGPU/SI: Emit constant arrays in the .hsrodata_readonly_agent section

Summary: This is done only when targeting HSA.

Reviewers: arsenm

Subscribers: arsenm, llvm-commits

Differential Revision: h

AMDGPU/SI: Emit constant arrays in the .hsrodata_readonly_agent section

Summary: This is done only when targeting HSA.

Reviewers: arsenm

Subscribers: arsenm, llvm-commits

Differential Revision: http://reviews.llvm.org/D13807

llvm-svn: 254587

show more ...