[AMDGPU] Fix a warning

This patch fixes:

llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp:255:18: error: private
field 'DAG' is not used [-Werror,-Wunused-private-field]

[AMDGPU] Add target hook to isGlobalMemoryObject (#112781)

We want special handing for IGLP instructions in the scheduler but they
should still be treated like they have side effects by other passe

[AMDGPU] Add target hook to isGlobalMemoryObject (#112781)

We want special handing for IGLP instructions in the scheduler but they
should still be treated like they have side effects by other passes. Add
a target hook to the ScheduleDAGInstrs DAG builder so that we have more
control over this.

show more ...

[AMDGPU] Add iglp_opt(3) for simple mfma / exp interleaving (#117269)

Adds a minimal iglp_opt to do simple exp / mfma interleaving.

[AMDGPU] Remove unused includes (NFC) (#116154)

Identified with misc-include-cleaner.

[llvm] Remove redundant calls to std::unique_ptr<T>::get (NFC) (#113415)

[AMDGPU] Qualify auto. NFC. (#110878)

Generated automatically with:
$ clang-tidy -fix -checks=-*,llvm-qualified-auto $(find
lib/Target/AMDGPU/ -type f)

[AMDGPU] Avoid repeated hash lookups (NFC) (#110788)

[AMDGPU] Use llvm::any_of, llvm::all_of, and llvm::none_of (NFC) (#103007)

[AMDGPU] Use llvm::any_of and llvm::none_of (NFC) (#102794)

[AMDGPU] Use range-based for loops. NFC. (#99047)

[AMDGPU] clang-tidy: use emplace_back instead of push_back. NFC.

[AMDGPU] Use using instead of typedef. NFC.

[AMDGPU] Use bool literals for bools. NFC.

Revert "[llvm] Use llvm::sort (NFC) (#96434)"

This reverts commit 05d167fc201b4f2e96108be0d682f6800a70c23d.

Reverting the patch fixes the following under EXPENSIVE_CHECKS:

LLVM :: CodeGen/AMDGPU

Revert "[llvm] Use llvm::sort (NFC) (#96434)"

This reverts commit 05d167fc201b4f2e96108be0d682f6800a70c23d.

Reverting the patch fixes the following under EXPENSIVE_CHECKS:

LLVM :: CodeGen/AMDGPU/sched-group-barrier-pipeline-solver.mir
LLVM :: CodeGen/AMDGPU/sched-group-barrier-pre-RA.mir
LLVM :: CodeGen/PowerPC/aix-xcoff-used-with-stringpool.ll
LLVM :: CodeGen/PowerPC/merge-string-used-by-metadata.mir
LLVM :: CodeGen/PowerPC/mergeable-string-pool-large.ll
LLVM :: CodeGen/PowerPC/mergeable-string-pool-pass-only.mir
LLVM :: CodeGen/PowerPC/mergeable-string-pool.ll

show more ...

[llvm] Use llvm::sort (NFC) (#96434)

[llvm] Use llvm::is_contained (NFC) (#95362)

[AMDGPU] Do not attempt to fallback to default mutations (#83208)

IGLP itself will be in SavedMutations via mutations added during
Scheduler creation, thus falling back results in reapplying IGLP.

[AMDGPU] Do not attempt to fallback to default mutations (#83208)

IGLP itself will be in SavedMutations via mutations added during
Scheduler creation, thus falling back results in reapplying IGLP.

In PostRA scheduling, if we have multiple regions with IGLP
instructions, then we may have infinite loop.

Disable the feature for now.

show more ...

[AMDGPU] Introduce iglp_opt(2): Generalized exp/mfma interleaving for select kernels (#81342)

This implements the basic pipelining structure of exp/mfma interleaving
for better extensibility. While

[AMDGPU] Introduce iglp_opt(2): Generalized exp/mfma interleaving for select kernels (#81342)

This implements the basic pipelining structure of exp/mfma interleaving
for better extensibility. While it does have improved extensibility,
there are controls which only enable it for DAGs with certain
characteristics (matching the DAGs it has been designed against).

show more ...

[AMDGPU][IGLP]: Add SchedGroupMask::TRANS (#75416)

Makes constructing SchedGroups of this type easier, and provides ability
to create them with __builtin_amdgcn_sched_group_barrier

[AMDGPU] [IGLP]: Fix assert (#73710)

We can also re-enter IGLP mutation via later `SchedStage`s in the
`GCNMaxOccupancySchedStrategy` . This is sort of NFC in that there is no
changed behavior for

[AMDGPU] [IGLP]: Fix assert (#73710)

We can also re-enter IGLP mutation via later `SchedStage`s in the
`GCNMaxOccupancySchedStrategy` . This is sort of NFC in that there is no
changed behavior for the only current client of `IsReentry`

show more ...

[AMDGPU] Correct assert that incorrectly chained multiple == operators. (#70291)

I believe this assert was trying to check that 3 variables were equal to
0.

I think it instead got interpreted as

[AMDGPU] Correct assert that incorrectly chained multiple == operators. (#70291)

I believe this assert was trying to check that 3 variables were equal to
0.

I think it instead got interpreted as ((DSWCount == DSWWithPermCount) ==
DSWWithSharedVMEMCount) == 0 I guess (DSWCount == DSWWithPermCount) was
true because both counts were 0. Then true got compared to
DSWWithSharedVMEMCount, and since DSWWithSharedVMEMCount is 0, that
compare was false. And then that false compared equal to the final 0.

show more ...

[llvm] Use llvm::any_of (NFC)

[AMDGPU][IGLP] SingleWaveOpt: Cache DSW Counters from PreRA (#67759)

Save the DSW counters from PreRA scheduling. While this avoids recalculation in the postRA pass, that isn't the main purpose.

[AMDGPU][IGLP] SingleWaveOpt: Cache DSW Counters from PreRA (#67759)

Save the DSW counters from PreRA scheduling. While this avoids recalculation in the postRA pass, that isn't the main purpose.

This is required because of physical register dependencies in PostRA scheduling -- they alter the DAG s.t. our counters may become incorrect -- which alters the layout of the pipeline. By preserving the values from PreRA, we can be sure that we accurately construct the pipeline.

Additionally, remove a bad assert in SharesPredWithPrevNthGroup -- it is possible that we will have an empty cache if OtherGroup has no elements which have a V_PERM pred (possible if the V_PERM SG is empty).

show more ...

[llvm] Use llvm::is_contained (NFC)

[NFC][AMDGPU] assert we've found a value before use

The sync pipeline should always contain the candidate ID. If it doesn't
something's gone awry. assert on that.

Reviewed by: jrbyrnes

Differentia

[NFC][AMDGPU] assert we've found a value before use

The sync pipeline should always contain the candidate ID. If it doesn't
something's gone awry. assert on that.

Reviewed by: jrbyrnes

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

show more ...