llvm-project.git/llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp, branch users/nico/python-2

AMDGPU: Avoid report_fatal_error on ds ordered intrinsics (#145202)

2025-06-23T04:09:09+00:00

AMDGPU/GFX12: Fix s_barrier_signal_isfirst for single-wave workgroups (#143634)

2025-06-19T18:22:49+00:00

Barrier instructions are no-ops in single-wave workgroups. This includes
s_barrier_signal_isfirst, which will leave SCC unmodified.

Model this correctly (via an implicit use of SCC) and ensure SCC==1
before the barrier instruction (if the wave is the only one of the
workgroup, then it is the first).

---------

Co-authored-by: Matt Arsenault

Warn on misuse of DiagnosticInfo classes that hold Twines (#137397)

2025-05-28T19:26:39+00:00

This annotates the `Twine` passed to the constructors of the various
DiagnosticInfo subclasses with `[[clang::lifetimebound]]`, which causes
us to warn when we would try to print the twine after it had already
been destructed.

We also update `DiagnosticInfoUnsupported` to hold a `const Twine &`
like all of the other DiagnosticInfo classes, since this warning allows
us to clean up all of the places where it was being used incorrectly.

[AMDGPU] Add a new amdgcn.load.to.lds intrinsic (#137425)

2025-05-19T14:15:04+00:00

This PR adds a amdgns_load_to_lds intrinsic that abstracts over loads to
LDS from global (address space 1) pointers and buffer fat pointers
(address space 7), since they use the same API and "gather from a
pointer to LDS" is something of an abstract operation.

This commit adds the intrinsic and its lowerings for addrspaces 1 and 7,
and updates the MLIR wrappers to use it (loosening up the restrictions
on loads to LDS along the way to match the ground truth from target
features).

It also plumbs the intrinsic through to clang.

AMDGPU: Fix -Wextra (#138539)

2025-05-05T18:02:39+00:00

Another stupid gcc warning. Ideally we would directly use the enum type,
but subregister indexes are emitted as an anonymous enum.

Fixes #125548

[AMDGPU] Support arbitrary types in amdgcn.dead (#134841)

2025-05-05T12:08:00+00:00

Legalize the amdgcn.dead intrinsic to work with types other than i32. It
still generates IMPLICIT_DEFs.

Remove some of the previous code for selecting/reg bank mapping it for
32-bit types, since everything is done in the legalizer now.

[AMDGPU] Use variadic isa<>. NFC. (#137016)

2025-04-24T07:19:09+00:00

Revert "[AMDGPU] Implement vop3p complex pattern optmization for gisel" (#136249)

2025-04-18T03:45:30+00:00

Reverts llvm/llvm-project#130234

[AMDGPU] Implement vop3p complex pattern optmization for gisel (#130234)

2025-04-18T02:56:20+00:00

Seeking opportunities to optimize VOP3P instructions by altering opsel,
opsel_hi, neg, neg_hi bits

Tests differences:
1. fix op_sel_hi bit for inline constant:
   1. `CodeGen/AMDGPU/packed-fp32.ll`
2. use neg bit to remove xor with 0x80008000
   1. `CodeGen/AMDGPU/strict_fsub.f16.ll`
   2. `CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.fdot2.ll`
   3. `CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.sdot4.ll`
   4. `CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.sdot8.ll`
   5. `CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.udot2.ll`
   6. `CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.udot4.ll`
   7. `CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.udot8.ll`
3. Remove xor 0x80008000, and use opsel, opsel_hi to remove alignbit
   1. `CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.sdot2.ll`

[AMDGPU] Use a target feature to enable __builtin_amdgcn_global_load_lds on gfx9/10 (#133055)

2025-04-02T18:00:09+00:00

This patch introduces the `vmem-to-lds-load-insts` target feature, which
can be used to enable builtins `__builtin_amdgcn_global_load_lds` and
`__builtin_amdgcn_raw_ptr_buffer_load_lds` on platforms which have this
feature.

This feature is only available on gfx9/10.

A limitation of using a common target feature for both builtins is that
we could have made `__builtin_amdgcn_raw_ptr_buffer_load_lds` available
on gfx6,7,8.