f6b2db298f
ac/nir: Refactor and optimize the repacking sequence.
...
According to feedback, the terminology with "exclusive scan"
and "reduction" is difficult. Change it to use "repack" instead,
which better fits what this sequence is actually used for.
The new sequence stores only 1 byte / wave to LDS, and uses packed
instructions to produce the results. This has lower latency and
fewer instructions than what we previously had.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Tony Wasserka <tony.wasserka@gmx.de >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11072 >
2021-06-09 16:48:51 +00:00
b4e22eb482
aco: Keep VGPR destinations for uniform shared loads when beneficial.
...
When the result of these loads is only used by cross-lane instructions,
it is beneficial to use a VGPR destination. This is because this allows
to put the s_waitcnt further down, which decreases latency.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Tony Wasserka <tony.wasserka@gmx.de >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11072 >
2021-06-09 16:48:51 +00:00
ce141e4c5f
aco: Implement byte and lane permute intrinsics.
...
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Tony Wasserka <tony.wasserka@gmx.de >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11072 >
2021-06-09 16:48:51 +00:00
5713e059ea
aco: Add validation for v_permlane instructions.
...
Previously there hasn't been any validation for these instructions,
but after shooting myself in the leg with it a few times, I decided
to add the validation now.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Tony Wasserka <tony.wasserka@gmx.de >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11072 >
2021-06-09 16:48:51 +00:00
fd6605367d
aco: Implement nir_op_sad_u8x4.
...
Fix up the operand size for v_sad instructions, and implement
the new NIR horizontal add. There is no viable way to do this
in SALU, so let's always use a VGPR destination.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Tony Wasserka <tony.wasserka@gmx.de >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11072 >
2021-06-09 16:48:51 +00:00
228169c87c
aco: Add note about v_alignbyte in the ISA README.
...
We tried to use this instruction for a more optimal sequence,
but it turned out that it doesn't exactly work as it was
supposed to. This note is to help others who want to use it.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Tony Wasserka <tony.wasserka@gmx.de >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11072 >
2021-06-09 16:48:51 +00:00
c129ede523
aco: use ds_read_{u8,u16}_d16
...
This allows partial writes and writes to the upper half of the destination.
fossil-db (Sienna Cichlid):
Totals from 135 (0.09% of 149839) affected shaders:
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/11113 >
2021-06-09 12:06:50 +00:00
6334d73fc9
aco: don't ever widen 8/16-bit sgpr load_shared
...
Doesn't seem to create incorrect code, but it is suboptimal.
No fossil-db changes.
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/11113 >
2021-06-09 12:06:50 +00:00
d2b9c7e982
radv: improve LDS alignment check for load/store vectorization
...
Previously, this could vectorize two scalar 16-bit loads into a u8vec4
load.
No fossil-db changes.
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/11113 >
2021-06-09 12:06:50 +00:00
4870d7d829
aco: use v1b/v2b for ds_read_u8/ds_read_u16
...
The p_extract_vector isn't necessary.
For ds_read_u8 and ds_read_u16, we used a 32-bit regclass, but did't load
32 bits, and used dst_hint for vector loads when we shouldn't have.
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com >
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/4863
Cc: mesa-stable
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11113 >
2021-06-09 12:06:50 +00:00
2fb436e92a
ci: update list of expected failures for Pitcairn/Oland (RADV)
...
The robustness2 failures were a mistake because they are actually
not supported (no VK_EXT_scalar_block_layout on GFX6).
The sparse related failures are no longer supported since sparse
is only enabled for Polaris10+.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11243 >
2021-06-09 11:27:44 +00:00
d169dad393
aco: fix emitting literal offsets with SMEM on GFX7
...
When the offset is negative, reg() isn't 255. Fix this by splitting
SGPR and literal emission. While we are at it, adjust a comment
saying that literals are also accepted on GFX6 which is wrong.
Fixes another batch of robustness tests.
Cc: 21.1 mesa-stable
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11247 >
2021-06-09 11:10:38 +00:00
13efad3086
radv: dump SPIR-V instead of using spirv-dis when generating a hang report
...
Useful when spirv-dis isn't found.
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/11034 >
2021-06-09 10:07:17 +00:00
3149eccc1c
radv: Implement VK_EXT_global_priority_query.
...
Signed-off-by: Georg Lehmann <dadschoorse@gmail.com >
Reviewed-by: Bas Nieuwenhuizen <bas@basnieuwenhuizen.nl >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11215 >
2021-06-09 08:25:25 +00:00
3761d994f6
aco: fix range checking for SSBO loads/stores with SGPR offset on GFX6-7
...
GFX6-7 are affected by a hw bug that prevents address clamping to work
correctly when the SGPR offset is used. Use the VGPR offset to fix it.
Fixes various hangs with dEQP-VK.robustness.robustness2.* on Bonaire.
Cc: 21.1 mesa-stable
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11238 >
2021-06-09 06:40:16 +00:00
8af6766062
nir: Move workgroup_size and workgroup_variable_size into common shader_info
...
Move it out the "cs" sub-struct, since these will be used for other
shader stages in the future.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11225 >
2021-06-08 09:23:55 -07:00
b5f6fc442c
nir: Move zero_initialize_shared_memory into common shader_info
...
Move it out the "cs" sub-struct, since the bit will be used for other
shader stages in the future.
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org >
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11225 >
2021-06-08 09:23:55 -07:00
3b81f53e34
aco/ra: Split print_regs by lines of 64 registers
...
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10517 >
2021-06-08 17:03:08 +02:00
69584478c9
aco/ra: Clean up print_regs output and support byte-allocated variables
...
Example output:
00 03 06 09 12 15 18 21 24 27 30 33 36 39 42
sgprs: ·▉█▉███▉▉█··████···········▉████············
00 03 06 09 12 15 18 21 24 27 30 33 36 39 42
vgprs: ▉▉··▉▉▉▉▘▀▉▉▉···▉▘▘▉▉▉▉···▉▉▉▀▀▉············
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10517 >
2021-06-08 17:03:08 +02:00
5bfef2de66
aco/ra: Fix off-by-one-error in print_regs
...
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Fixes: 3675aefa84https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/10517 >
2021-06-08 17:03:08 +02:00
43c04ab2b4
radeonsi: separate video hw info based on HW engine individually
...
This removes previous "has_hw_decode" and "uvd_enc_supported" and
makes information more accuate for cases where HW decode, HW encode,
and HW JPEG decode might partially available.
Signed-off-by: Leo Liu <leo.liu@amd.com >
Reviewed-by: James Zhu <James.Zhu@amd.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11201 >
2021-06-08 09:32:48 -04:00
9f7e63e12a
ac/debug: fix color printing PKT3 when count in header is too low
...
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/11211 >
2021-06-08 11:19:00 +02:00
c768d7d8f2
aco/tests: add SDWA tests
...
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
24418304b0
aco/tests: add tests for p_extract/p_insert lowering
...
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
8e0c6e196e
aco: disallow literals with some instruction formats
...
Because isVOPn() is true for many VOP3, SDWA and DPP instructions, this
would often not complain.
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
cf22eabc68
aco: make validate_ir() output usable in tests
...
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
54292e99c7
aco: optimize 32-bit extracts and inserts using SDWA
...
Still need to use dst_u=preserve field to optimize packs
fossil-db (Sienna Cichlid):
Totals from 15974 (10.66% of 149839) affected shaders:
VGPRs: 1009064 -> 1008968 (-0.01%); split: -0.03%, +0.02%
SpillSGPRs: 7959 -> 7964 (+0.06%)
CodeSize: 101716436 -> 101159568 (-0.55%); split: -0.55%, +0.01%
MaxWaves: 284464 -> 284490 (+0.01%); split: +0.02%, -0.01%
Instrs: 19334216 -> 19224241 (-0.57%); split: -0.57%, +0.00%
Latency: 375465295 -> 375230478 (-0.06%); split: -0.14%, +0.08%
InvThroughput: 79006105 -> 78860705 (-0.18%); split: -0.25%, +0.07%
fossil-db (Polaris):
Totals from 11369 (7.51% of 151365) affected shaders:
SGPRs: 787920 -> 787680 (-0.03%); split: -0.04%, +0.01%
VGPRs: 681056 -> 681040 (-0.00%); split: -0.01%, +0.00%
CodeSize: 68127288 -> 67664120 (-0.68%); split: -0.69%, +0.01%
MaxWaves: 54370 -> 54371 (+0.00%)
Instrs: 13294638 -> 13214109 (-0.61%); split: -0.62%, +0.01%
Latency: 373515759 -> 373214571 (-0.08%); split: -0.11%, +0.03%
InvThroughput: 166529524 -> 166275291 (-0.15%); split: -0.20%, +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
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
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
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
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
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
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
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
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
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
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
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
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
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: ca783612e7timur.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
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
996a6564cf
ci: update some radv trace checksums
...
After 7d23ea20a0agomez@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
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
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
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
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
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
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
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
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: 0e4747d3fbtimur.kristof@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11177 >
2021-06-04 09:31:41 +02:00
Timur Kristóf