| Age | Commit message (Collapse) | Author |
|---|
|
|
|
|
|
Also adapt hazard and wait handling.
|
|
|
|
|
|
This assert was failing in a fuzzing test. I consulted with @jrbyrnes
who said:
The MFMASmallGemmSingleWaveOpt::apply() method is invoked if and only if
the user has inserted an intrinsic llvm.amdgcn.iglp.opt(i32 1) into
their source code. This intrinsic applies a highly specialized DAG
mutation to result in specific scheduling for a specific set of kernels.
These assertions are really just confirming that the characteristics of
the kernel match what is expected (i.e. The kernels are similar to the
ones this DAG mutation strategy were designed against).
However, if we apply this DAG mutation to kernels for which is was not
designed, then we may not find the types of instructions we are looking
for, and may end up with empty caches.
I think it should be fine to just return false if the cache is empty
instead of the assert.
|
|
This patch partially reverts:
commit 5e1b0f97735083b6762834b83fdbb35e76002e03
Author: Kazu Hirata <kazu@google.com>
Date: Fri Apr 18 10:05:55 2025 -0700
to fix:
LLVM :: CodeGen/AMDGPU/sched-group-barrier-pipeline-solver.mir
LLVM :: CodeGen/AMDGPU/sched-group-barrier-pre-RA.mir
under LLVM_ENABLE_EXPENSIVE_CHECKS.
|
|
|
|
|
|
- Remove calls to pass initialization from pass constructors.
- https://github.com/llvm/llvm-project/issues/111767
|
|
This patch fixes:
llvm/lib/Target/AMDGPU/AMDGPUIGroupLP.cpp:255:18: error: private
field 'DAG' is not used [-Werror,-Wunused-private-field]
|
|
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.
|
|
Adds a minimal iglp_opt to do simple exp / mfma interleaving.
|
|
Identified with misc-include-cleaner.
|
|
|
|
Generated automatically with:
$ clang-tidy -fix -checks=-*,llvm-qualified-auto $(find
lib/Target/AMDGPU/ -type f)
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
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
|
|
|
|
|
|
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.
|
|
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).
|
|
Makes constructing SchedGroups of this type easier, and provides ability
to create them with __builtin_amdgcn_sched_group_barrier
|
|
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`
|
|
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.
|
|
|
|
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).
|
|
|
|
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
|
|
This adds the IGLP strategy for single-wave gemms. The SchedGroup pipeline is laid out in multiple phases, with each phase corresponding to a distinct pattern present in gemm kernels. The resilience of the optimization is dependent upon IR (as seen by pre-RA scheduling) continuing to have these patterns (as defined by instruction class and dependencies) in their current relative ordering.
The kernels of interest have these specific phases:
NT: 1, 2a, 2c
NN: 1, 2a, 2b
TT: 1, 2b, 2c
TN: 1, 2b
The general approach taken was to have a long SchedGroup pipeline. In this way the scheduler will have less capability of doing the wrong thing. In order to resolve the challenge of correctly fitting these long pipelines, we leverage the rules infrastructure to help the solver.
Differential Revision: https://reviews.llvm.org/D149773
Change-Id: I1a35962a95b4bdf740602b8f110d3297c6fb9d96
|
|
Differential Revision: https://reviews.llvm.org/D146774
Change-Id: Icd7aaaa0b257a25713c22ead0813777cef7d5859
|
|
Currently the PipelineSolver processes SchedGroups in bottom up manner. However, there is no compelling reason to require this. Providing the option to toggle this affords greater experimentation capability, and make usage a bit more intuitive. Importantly, it makes designing rules much easier.
Differential Revision: https://reviews.llvm.org/D149393
Change-Id: Ic4abd3408f9faa105c0eef72eab7873d46083ee4
|
|
This reverts commit 3c42a58c4f20ae3b621733bf5ee6d57c912994a9.
Breaks tests on mac, see https://reviews.llvm.org/rG3c42a58c4f20ae3b621733bf5ee6d57c912994a9#1191724
|
|
|
|
MFMA and WMMA essentially the same thing, but apear on different ASICs.
Differential Revision: https://reviews.llvm.org/D142062
|
|
C++17 allows us to call constructors pair and tuple instead of helper
functions make_pair and make_tuple.
Differential Revision: https://reviews.llvm.org/D139828
|
|
|
|
Based on experiments this does better with target small GEMM kernels.
Reviewed By: jrbyrnes
Differential Revision: https://reviews.llvm.org/D139227
|
|
This patch mechanically replaces None with std::nullopt where the
compiler would warn if None were deprecated. The intent is to reduce
the amount of manual work required in migrating from Optional to
std::optional.
This is part of an effort to migrate from llvm::Optional to
std::optional:
https://discourse.llvm.org/t/deprecating-llvm-optional-x-hasvalue-getvalue-getvalueor/63716
|
|
Identified with modernize-loop-convert.
|
|
|
|
|
|
Adds a builtin that serves as an optimization hint to apply specific optimized
DAG mutations during scheduling. This also disables any other mutations or
clustering that may interfere with the desired pipeline. The first optimization
strategy that is added here is designed to improve the performance of small gemm
kernels on gfx90a.
Reviewed By: jrbyrnes
Differential Revision: https://reviews.llvm.org/D132079
|
|
Requested SchedGroup pipelines may be non-trivial to satisify. A minimimal example is if the requested pipeline is {2 VMEM, 2 VALU, 2 VMEM} and the original order of SUnits is {VMEM, VALU, VMEM, VALU, VMEM}. Because of existing dependencies, the choice of which SchedGroup the middle VMEM goes into impacts how closely we are able to match the requested pipeline. It seems minimizing the degree of misfit (as measured by the number of edges we can't add) w.r.t the choice we make when mapping an instruction -> SchedGroup is an NP problem. This patch implements the PipelineSolver class which produces a solution for the defined problem for the sched_group_barrier mutation. The solver has both an exponential time exact algorithm and a greedy algorithm. The patch includes some controls which allows the user to select the greedy/exact algorithm.
Differential Revision: https://reviews.llvm.org/D130797
|
