Commit graph

7405 commits

Author SHA1 Message Date
Rhys Perry
63659fc15c radv: use byte/word extract/insert instructions
ACO doesn't yet combine extract/insert into instructions, but it seems to
already generate less instructions because NIR optimizes shift+and to
these instructions. Code size is worse in some cases though because we
have to always use a literal when masking.

fossil-db (Sienna Cichlid):
Totals from 14361 (9.58% of 149839) affected shaders:
VGPRs: 850152 -> 850304 (+0.02%); split: -0.02%, +0.04%
SpillSGPRs: 7979 -> 7989 (+0.13%); split: -0.03%, +0.15%
CodeSize: 88031216 -> 88162520 (+0.15%); split: -0.01%, +0.16%
MaxWaves: 269414 -> 269426 (+0.00%)
Instrs: 16695182 -> 16662852 (-0.19%); split: -0.21%, +0.01%
Latency: 375592693 -> 375544364 (-0.01%); split: -0.04%, +0.03%
InvThroughput: 75627700 -> 75607720 (-0.03%); split: -0.07%, +0.04%

fossil-db (Polaris):
Totals from 13816 (9.13% of 151365) affected shaders:
SGPRs: 984896 -> 982512 (-0.24%); split: -0.29%, +0.05%
VGPRs: 809220 -> 809112 (-0.01%); split: -0.02%, +0.01%
SpillSGPRs: 9181 -> 9185 (+0.04%); split: -0.04%, +0.09%
CodeSize: 82017952 -> 82123484 (+0.13%); split: -0.01%, +0.14%
MaxWaves: 65721 -> 65723 (+0.00%)
Instrs: 16008744 -> 15988007 (-0.13%); split: -0.18%, +0.05%
Latency: 439911623 -> 439869622 (-0.01%); split: -0.04%, +0.03%
InvThroughput: 185898770 -> 185841742 (-0.03%); split: -0.08%, +0.05%

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
2021-06-08 08:57:43 +00:00
Rhys Perry
7d76b07d6b ac/llvm: implement byte/word extract/insert instructions
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
2021-06-08 08:57:43 +00:00
Rhys Perry
daa329f664 aco: use byte/word extract pseudo-instructions
fossil-db (Sienna Cichild):
Totals from 1890 (1.26% of 149839) affected shaders:
CodeSize: 5104196 -> 5104300 (+0.00%); split: -0.00%, +0.01%
Latency: 11572943 -> 11572880 (-0.00%); split: -0.00%, +0.00%
InvThroughput: 4876941 -> 4876982 (+0.00%); split: -0.00%, +0.00%
SClause: 26774 -> 26775 (+0.00%)
Copies: 125778 -> 125813 (+0.03%)
PreSGPRs: 56452 -> 56451 (-0.00%)

fossil-db (Polaris):
Totals from 1884 (1.24% of 151365) affected shaders:
CodeSize: 3849340 -> 3849312 (-0.00%); split: -0.00%, +0.00%
Instrs: 741391 -> 741382 (-0.00%)
Latency: 13533815 -> 13533439 (-0.00%)
InvThroughput: 12058777 -> 12058500 (-0.00%)
Copies: 120890 -> 120891 (+0.00%)
PreSGPRs: 48940 -> 48939 (-0.00%)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
2021-06-08 08:57:43 +00:00
Rhys Perry
1f2518ef9f aco: implement nir_op_extract/nir_op_insert
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
2021-06-08 08:57:43 +00:00
Rhys Perry
2f94353735 aco: add p_extract/p_insert
These will let us make the SDWA optimizer much simpler than if we were to
recognize combinations of shift/and/bfe.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
2021-06-08 08:57:42 +00:00
Rhys Perry
e9d1643288 aco: disallow SDWA for instructions with 64-bit definitions/operands
For example, v_cvt_f64_i32. LLVM doesn't seem to allow this either and it
doesn't seem to work correctly.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
2021-06-08 08:57:42 +00:00
Rhys Perry
1cbcfb8b38 nir, nir/algebraic: add byte/word insertion instructions
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3151>
2021-06-08 08:57:42 +00:00
Samuel Pitoiset
736893060f radv: emit PA_SC_CONSERVATIVE_RASTERIZATION_CNTL only on GFX9+
This context register doesn't exist on older generations.

Cc: 21.1 mesa-stable
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11210>
2021-06-08 05:58:01 +00:00
Caio Marcelo de Oliveira Filho
c8a7bd0dc8 nir: Rename WORK_GROUP (and similar) to WORKGROUP
Be consistent with other usages in Vulkan and SPIR-V, and the recently
added workgroup_size field.

Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>
2021-06-07 22:34:42 +00:00
Caio Marcelo de Oliveira Filho
a71a780598 nir: Rename nir_intrinsic_load_local_group_size to nir_intrinsic_load_workgroup_size
Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>
2021-06-07 22:34:42 +00:00
Caio Marcelo de Oliveira Filho
430d2206da compiler: Rename local_size to workgroup_size
Acked-by: Emma Anholt <emma@anholt.net>
Acked-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Acked-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11190>
2021-06-07 22:34:42 +00:00
Timur Kristóf
18d48c01c2 radv: Assert that there is no GS copy shader when the pipeline has NGG.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11092>
2021-06-07 22:05:42 +00:00
Timur Kristóf
7e664a5383 radv: Don't generate GS copy shader when the pipeline has NGG.
Previously the code used radv_pipeline_has_ngg but that always
returned false because the pipeline->shaders was all NULL at the
time when the GS copy shader was created.

Fixes: ca783612e7
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11092>
2021-06-07 22:05:42 +00:00
Timur Kristóf
93b1089d19 radv: Remove duplicate code for getting GS info.
This was my mistake for forgetting to delete this code.

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11092>
2021-06-07 22:05:42 +00:00
Andres Gomez
996a6564cf ci: update some radv trace checksums
After 7d23ea20a0 ("radv: don't allocate DCC predicate if the image doesn't use DCC")
some checksums for the radv driver remained to be updated.

Signed-off-by: Andres Gomez <agomez@igalia.com>
Reviewed-by: Martin Peres <martin.peres@mupuf.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11223>
2021-06-07 20:45:55 +00:00
Tony Wasserka
3c390e2eb6 aco/scheduler: Move cursor handling state to dedicated interfaces
This clarifies the semantics of the index variables compared to the previous
version, which used the same variables in a slightly different way depending
on whether they were used for downwards moves or upwards ones.

Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10885>
2021-06-07 12:09:39 +02:00
Tony Wasserka
81761a311e aco/scheduler: Clean up register demand tracking
Refactoring total_demand and total_demand_clause to cover non-overlapping
instruction intervals makes the code easier to follow and allows the register
demand to be updated more efficiently in some cases.

Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10885>
2021-06-07 12:09:39 +02:00
Daniel Schürmann
d4662e38c4 aco: simplify Phi RegClass selection
Also adds moves validation rules to aco_validate.

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11181>
2021-06-04 16:47:01 +00:00
Daniel Schürmann
dc807dff3e radv,aco: scalarize all phis via nir_lower_phis_to_scalar()
This allows to remove some ACO code which did so previously.

Totals from 93 (0.06% of 149839) affected shaders (Navi2):
CodeSize: 582424 -> 582348 (-0.01%); split: -0.10%, +0.08%
Instrs: 107083 -> 107011 (-0.07%); split: -0.08%, +0.01%
Latency: 483338 -> 484881 (+0.32%); split: -0.09%, +0.40%
InvThroughput: 101129 -> 101532 (+0.40%); split: -0.03%, +0.42%
Copies: 9893 -> 9774 (-1.20%); split: -1.28%, +0.08%
Branches: 2862 -> 2858 (-0.14%)
PreSGPRs: 3342 -> 3339 (-0.09%)
PreVGPRs: 4567 -> 4565 (-0.04%)

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11181>
2021-06-04 16:47:01 +00:00
Samuel Pitoiset
b786c16365 radv/winsys: allow to reserve a VMID
This will be used by SPM and also for configuring the trap handler.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11128>
2021-06-04 14:53:25 +00:00
Rhys Perry
49add985ff nir/unsigned_upper_bound: don't require dominance metadata
Instead, determine if it's a merge or loop exit phi.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9808>
2021-06-04 14:14:00 +00:00
Rhys Perry
aebffc241d aco: don't use nir_block_is_unreachable()
nir_cf_reinsert() can re-create the block, invalidating dominance
metadata.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9808>
2021-06-04 14:14:00 +00:00
Mauro Rossi
e4e4b6bc16 android: aco: add aco_optimizer_postRA.cpp to Makefile.sources
Fixes the following building error:

external/mesa/src/amd/compiler/aco_interface.cpp:155: error: undefined reference to 'aco::optimize_postRA(aco::Program*)'

Fixes: 0e4747d3fb ("aco: Introduce a new, post-RA optimizer.")
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11177>
2021-06-04 09:31:41 +02:00
Mauro Rossi
60e134e83e android: ac: add include src/util path
Fixes the following building error:

external/mesa/src/amd/common/ac_nir_lower_ngg.c:27:10: fatal error: 'u_math.h' file not found
         ^~~~~~~~~~
1 error generated.

Fixes: 3d589b8b46 ("ac: Add new NIR pass to lower NGG VS/TES.")
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11177>
2021-06-04 09:31:36 +02:00
Mauro Rossi
2dea82fc07 android: ac: add ac_nir_lower_ngg.c to Makefile.sources
Fixes the following building errors:

external/mesa/src/amd/vulkan/radv_shader.c:868: error: undefined reference to 'ac_nir_lower_ngg_gs'
external/mesa/src/amd/vulkan/radv_shader.c:851: error: undefined reference to 'ac_nir_lower_ngg_nogs'
external/mesa/src/amd/compiler/aco_interface.cpp:155: error: undefined reference to 'aco::optimize_postRA(aco::Program*)'

Fixes: 3d589b8b46 ("ac: Add new NIR pass to lower NGG VS/TES.")
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11177>
2021-06-04 09:31:28 +02:00
Samuel Pitoiset
aff92f50c6 ac: add ac_thread_trace::data
Instead of passing two different structs to ac_dump_rgp_capture().

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11156>
2021-06-03 15:39:34 +00:00
Samuel Pitoiset
416496a0c4 ac/rgp: fix ac_fill_sqtt_asic_info() name
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11156>
2021-06-03 15:39:34 +00:00
Samuel Pitoiset
ea3f72c9d9 ac: rename ac_dump_thread_trace() to ac_dump_rgp_capture()
RGP captures can contain both SQTT and SPM data. While we are at it,
move it to ac_rgp.h and adjust a message.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11156>
2021-06-03 15:39:34 +00:00
Chia-I Wu
7ebd658e28 radv: use vk_default_allocator
Signed-off-by: Chia-I Wu <olvaffe@gmail.com>
Acked-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11117>
2021-06-03 08:13:26 +00:00
Samuel Pitoiset
380ac28891 ac: import performance counters from RadeonSI
Performance counters will be used by RADV for VK_KHR_performance_query
and also for adding SPM support.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11140>
2021-06-03 07:15:21 +00:00
Rhys Perry
903f814b78 aco: don't create 4 and 5 dword NSA instructions on GFX10
"stability issues", apparently: https://reviews.llvm.org/D103348

fossil-db (Navi10):
Totals from 4512 (3.01% of 149839) affected shaders:
VGPRs: 221516 -> 223308 (+0.81%); split: -0.07%, +0.88%
CodeSize: 23000080 -> 23070672 (+0.31%); split: -0.08%, +0.39%
MaxWaves: 107718 -> 107496 (-0.21%); split: +0.11%, -0.32%
Instrs: 4321890 -> 4362822 (+0.95%); split: -0.00%, +0.95%
Latency: 71495710 -> 71581476 (+0.12%); split: -0.07%, +0.19%
InvThroughput: 11858568 -> 11938960 (+0.68%); split: -0.00%, +0.68%
VClause: 76575 -> 76585 (+0.01%); split: -0.05%, +0.07%
SClause: 168771 -> 168709 (-0.04%); split: -0.06%, +0.02%
Copies: 182305 -> 221948 (+21.75%); split: -0.00%, +21.75%
PreVGPRs: 194657 -> 195635 (+0.50%); split: -0.00%, +0.50%

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Fixes: c353895c92 ("aco: use non-sequential addressing")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10898>
2021-06-03 03:49:07 +00:00
Rhys Perry
bb52484df5 aco/tests: improve reporting of failed code checks
Instead of just reporting the failed statements, print where they
originated. This is useful for tests which have a number of similar
checks.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10898>
2021-06-03 03:49:07 +00:00
Rhys Perry
9bf30c4a5c aco/tests: add tests for form_hard_clauses()
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10898>
2021-06-03 03:49:07 +00:00
Rhys Perry
81162265b1 aco: do not clause NSA instructions
According to LLVM, this has "unpredictable results on GFX10.1".

https://reviews.llvm.org/D102211

fossil-db (Navi10):
Totals from 26690 (17.81% of 149839) affected shaders:
CodeSize: 167935160 -> 167706280 (-0.14%); split: -0.14%, +0.00%
Instrs: 31801427 -> 31744142 (-0.18%); split: -0.18%, +0.00%
Latency: 732672435 -> 732622463 (-0.01%)
InvThroughput: 163361435 -> 163357838 (-0.00%); split: -0.00%, +0.00%
VClause: 546131 -> 546903 (+0.14%); split: -0.00%, +0.14%

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Fixes: c353895c92 ("aco: use non-sequential addressing")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10898>
2021-06-03 03:49:07 +00:00
Andres Gomez
15e41b576b ci: fix the vkd3d-proton runner
Signed-off-by: Andres Gomez <agomez@igalia.com>
Reviewed-by: Martin Peres <martin.peres@mupuf.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11087>
2021-06-02 16:30:57 +00:00
Mike Blumenkrantz
89bac36f09 radv: explicitly load a desc set layout struct member during set allocate
accessing this variable repeatedly like this is a contended hotpath somehow,
so instead just create a const for reference

Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11124>
2021-06-02 13:31:43 +00:00
Mike Blumenkrantz
79742d41c0 radv: declare index_va in a single call for indexed draw packet emit
this is an extreme hotpath, so having a single calculation in a const
variable is slightly better for compiler microoptimizing

Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11124>
2021-06-02 13:31:43 +00:00
Tomeu Vizoso
67af3b6bba ci/lava: Switch LAVA jobs to x86 runners
So we don't need to provision aarch64 servers, which are these days
rarer than x8_64.

In the switch to the new runner tags, switch to one which contains the
device type, so we can dimension the runner jobs taking into account the
number of DUTs available.

Signed-off-by: Tomeu Vizoso <tomeu.vizoso@collabora.com>
Acked-by: Daniel Stone <daniels@collabora.com>
Reviewed-by: Emma Anholt <emma@anholt.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11108>
2021-06-02 08:14:51 +02:00
Samuel Pitoiset
a70c3e5c8a ac/rgp: bump the SQTT file minor version to 5
To match latest RGP spec. Captures generated by RADV still work
with latest RGP (v1.10).

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11063>
2021-06-01 06:09:49 +00:00
Samuel Pitoiset
c3a4ca2908 ac/rgp: mark SQTT_FILE_CHUNK_TYPE_ISA_DATABASE as deprecated
This is now deprecated and reserved for future uses.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11063>
2021-06-01 06:09:49 +00:00
Rhys Perry
0f8fef1261 radv: make attrib_end variable in radv_flush_vertex_descriptors 32-bit
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Fixes: 1e9dc0474e ("radv: make radv_pipeline::attrib_ends 32bit")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11089>
2021-05-31 18:28:12 +00:00
Timur Kristóf
aabe9d2f6e aco: Eliminate SALU comparison when SCC can be used instead.
For example:

s0, scc = s_and_u32 ...
scc = s_cmp_eq_u32 s0, 0
p_cbranch_sccz

is turned into:

s0, scc = s_and_u32 ...
p_cbranch_sccnz

Fossil DB results on Sienna Cichlid:

Totals from 85267 (56.91% of 149839) affected shaders:
CodeSize: 202539256 -> 202237268 (-0.15%)
Instrs: 38964493 -> 38888996 (-0.19%)
Latency: 750062328 -> 749913450 (-0.02%); split: -0.02%, +0.00%
InvThroughput: 167408952 -> 167405157 (-0.00%)

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7779>
2021-05-28 12:14:53 +00:00
Timur Kristóf
a93092d0ed aco: Use s_cbranch_vccz/nz in post-RA optimization.
A simple post-RA optimization which takes advantage of the
s_cbranch_vccz and s_cbranch_vccnz instructions.

It works on the following pattern:

vcc = v_cmp ...
scc = s_and vcc, exec
p_cbranch scc

The result looks like this:

vcc = v_cmp ...
p_cbranch vcc

Fossil DB results on Sienna Cichlid:

Totals from 4814 (3.21% of 149839) affected shaders:
CodeSize: 15371176 -> 15345964 (-0.16%)
Instrs: 3028557 -> 3022254 (-0.21%)
Latency: 21872753 -> 21823476 (-0.23%); split: -0.23%, +0.00%
InvThroughput: 4470282 -> 4468691 (-0.04%); split: -0.04%, +0.00%

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7779>
2021-05-28 12:14:53 +00:00
Timur Kristóf
0e4747d3fb aco: Introduce a new, post-RA optimizer.
This commit adds the skeleton of a new ACO post-RA optimizer,
which is intended to be a simple pass called after RA, and
is meant to do code changes which can only be done
after RA.

It is currently empty, the actual optimizations will be added
in their own commits. It only has a DCE pass, which deletes
some dead code generated by the spiller.

Fossil DB results on Sienna Cichlid:

Totals from 375 (0.25% of 149839) affected shaders:
CodeSize: 2933056 -> 2907192 (-0.88%)
Instrs: 534154 -> 530706 (-0.65%)
Latency: 12088064 -> 12084907 (-0.03%); split: -0.03%, +0.00%
InvThroughput: 4433454 -> 4432421 (-0.02%); split: -0.02%, +0.00%
Copies: 81649 -> 78203 (-4.22%)

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7779>
2021-05-28 12:14:53 +00:00
Timur Kristóf
6f3c472f2e aco: New writeout overloads for the test framework.
These will be used by future tests.

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7779>
2021-05-28 12:14:53 +00:00
Timur Kristóf
8d37aa91d6 aco: Add Operand(Temp, PhysReg) constructor.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7779>
2021-05-28 12:14:53 +00:00
Timur Kristóf
4491b94d58 aco: Don't DCE instructions that write non-temps, eg. exec.
No Fossil DB changes.
This commit makes DCE usable after RA.

Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7779>
2021-05-28 12:14:53 +00:00
Samuel Pitoiset
ea5f1fa279 radv: fix generating hang reports if mutable descriptors are used
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11031>
2021-05-28 10:53:02 +00:00
Samuel Pitoiset
380742b9f3 radv: fix missing default state for DB_DFSM_CONTROL
Fixes: 69ae02151d ("radv: remove DFSM")
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11028>
2021-05-27 11:48:29 +00:00
Samuel Pitoiset
b9ff51f750 radv: move all game workarounds to drirc
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10978>
2021-05-26 18:48:04 +00:00