e10e74a7af
aco: Implement new buffer load/store intrinsics.
...
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9201 >
2021-03-17 12:42:23 +00:00
8f71be0a7b
aco: simplify loop_nest_depth tracking in isel
...
Keep track of the current loop depth in Program and set the depth inside
Program::insert_block() instead of repeating it every time we insert one.
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/8994 >
2021-03-11 15:35:24 +00:00
a6e1178f91
aco: Disallow LSHS temp-only I/O when VS output is written indirectly.
...
Cc: mesa-stable
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9100 >
2021-02-18 12:10:56 +00:00
48f349971f
aco: Fix LDS statistics of tess control shaders.
...
The calculate_tess_lds_size function already returns the size in blocks
of the encoding granule, but we forgot to adjust config->lds_size.
This variable is not used to actually set the LDS size used for TCS,
but by ACO to make scheduling decisions.
Fossil DB stats on Sienna Cichlid:
Please note that the +3729.43% is NOT a regression.
The real LDS size used didn't change, it was just reported incorrectly.
Totals from 1342 (0.96% of 139391) affected shaders:
VGPRs: 60880 -> 80240 (+31.80%); split: -0.05%, +31.85%
CodeSize: 3378456 -> 3381224 (+0.08%); split: -0.23%, +0.31%
LDS: 687104 -> 26312192 (+3729.43%)
MaxWaves: 29794 -> 23962 (-19.57%)
Instrs: 644194 -> 644610 (+0.06%); split: -0.32%, +0.39%
Cycles: 2675068 -> 2676804 (+0.06%); split: -0.31%, +0.38%
VMEM: 428840 -> 517418 (+20.66%); split: +22.53%, -1.88%
SMEM: 91831 -> 88587 (-3.53%); split: +5.70%, -9.23%
VClause: 22740 -> 19384 (-14.76%); split: -16.18%, +1.42%
SClause: 19116 -> 18373 (-3.89%); split: -4.34%, +0.46%
Copies: 66662 -> 63448 (-4.82%); split: -5.55%, +0.73%
Fixes: cf89bdb9batimur.kristof@gmail.com >
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9098 >
2021-02-18 11:57:22 +00:00
3d4c13f3b8
aco: add DeviceInfo
...
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/8761 >
2021-02-15 13:44:22 +00:00
b759557cac
aco: consider that GFX10.3 allocates LDS in 1024 byte blocks
...
fossil-db (GFX10.3):
Totals from 3 (0.00% of 139391) affected shaders:
VMEM: 513 -> 511 (-0.39%)
SMEM: 94 -> 92 (-2.13%)
VClause: 31 -> 30 (-3.23%)
fossil-db (GFX10.3, wave32):
Totals from 4 (0.00% of 139391) affected shaders:
VClause: 82 -> 81 (-1.22%)
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/8761 >
2021-02-15 13:35:38 +00:00
7ff805a19d
radv,aco: add radv_nir_compiler_options::wgp_mode
...
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com >
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8761 >
2021-02-15 13:35:36 +00:00
b98a4d4dd7
aco: refactor GPR limit calculation
...
This patch delays the calculation of GPR limits in order to
precisely incorporate extra registers (VCC etc.) and shared VGPRs.
Additionally, the allocation granularity is used to set the config.
This has some effect on the reported SGPR stats.
Totals (Navi10):
SGPRs: 6971787 -> 17753642 (+154.65%)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8921 >
2021-02-12 19:00:18 +00:00
b2dbe2b87b
aco: implement non-uniform get_ssbo_size
...
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com >
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/3711
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7969 >
2021-01-27 13:00:33 +00:00
454bbf8f23
aco: emit packed 16bit instructions
...
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6680 >
2021-01-13 17:46:56 +00:00
38da379b3e
aco: Note if rasterization can start early.
...
When there are no param exports in an NGG (or legacy VS) shader,
the NO_PC_EXPORT=1 is set by RADV, which means PS waves can launch
before the current stage finishes.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7868 >
2021-01-12 16:43:27 +00:00
0bd14be962
aco: implement sparse image loads
...
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/7775 >
2021-01-08 14:27:07 +00:00
77343576eb
aco: implement a workaround for gl_FragCoord.z with VRS on GFX10.3
...
Without it, FragCoord.z will have the value of one of the fine pixels
instead of the center of the coarse pixel.
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/7837 >
2020-12-14 16:22:39 +00:00
c587eaadf6
aco: implement fragment shading rate
...
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/7837 >
2020-12-14 16:22:38 +00:00
2d12991e01
aco: use FALLTHROUGH macro
...
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com >
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7844 >
2020-12-01 10:28:36 +00:00
4eac442217
aco/ngg: fix division-by-zero in assertion
...
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com >
Reviewed-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/7576 >
2020-11-25 13:41:04 +00:00
867323379e
aco: don't use SMEM for SSBO stores
...
fossil-db (Navi):
Totals from 70 (0.05% of 138791) affected shaders:
SGPRs: 2324 -> 2097 (-9.77%)
VGPRs: 1344 -> 1480 (+10.12%)
CodeSize: 157872 -> 154836 (-1.92%); split: -1.93%, +0.01%
MaxWaves: 1288 -> 1260 (-2.17%)
Instrs: 29730 -> 29108 (-2.09%); split: -2.13%, +0.04%
Cycles: 394944 -> 391280 (-0.93%); split: -0.94%, +0.01%
VMEM: 5288 -> 5695 (+7.70%); split: +11.97%, -4.27%
SMEM: 2680 -> 2444 (-8.81%); split: +1.34%, -10.15%
VClause: 291 -> 502 (+72.51%)
SClause: 1176 -> 918 (-21.94%)
Copies: 3549 -> 3517 (-0.90%); split: -1.80%, +0.90%
Branches: 1230 -> 1228 (-0.16%)
PreSGPRs: 1675 -> 1491 (-10.99%)
PreVGPRs: 1101 -> 1223 (+11.08%)
Totals from 70 (0.05% of 139517) affected shaders (RAVEN):
SGPRs: 2368 -> 2121 (-10.43%)
VGPRs: 1344 -> 1480 (+10.12%)
CodeSize: 156664 -> 153252 (-2.18%)
MaxWaves: 636 -> 622 (-2.20%)
Instrs: 29968 -> 29226 (-2.48%)
Cycles: 398284 -> 393492 (-1.20%)
VMEM: 5544 -> 5930 (+6.96%); split: +11.72%, -4.76%
SMEM: 2752 -> 2502 (-9.08%); split: +1.20%, -10.28%
VClause: 292 -> 504 (+72.60%)
SClause: 1236 -> 940 (-23.95%)
Copies: 3907 -> 3852 (-1.41%); split: -2.20%, +0.79%
Branches: 1230 -> 1228 (-0.16%)
PreSGPRs: 1671 -> 1487 (-11.01%)
PreVGPRs: 1102 -> 1225 (+11.16%)
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com >
Reviewed-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/6143 >
2020-11-16 15:52:22 +00:00
3a72021d7c
aco: store NIR range analysis data to the isel context
...
It will be used to optimize some ALU instructions.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7405 >
2020-11-03 13:47:40 +00:00
f74ef15879
aco/ngg: Incorporate GS invocations into workgroup size calculation.
...
If the workgroup_size variable is lower than the actual workgroup size,
that means it's possible that ACO won't emit some s_barrier instructions
when in fact it should. This can possibly cause a GPU hang.
This is just for the sake of general correctness, currently this
can't cause a real problem because the maximum vertex count is always
greater than (or equal to) the primitive count in GS, and already
takes into account the number of GS invocations.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7232 >
2020-10-28 21:55:54 +01:00
73449f9a62
aco: Add a few assertions about LDS usage.
...
This is to make sure we don't compile a shader which doesn't
fit the available LDS space.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7232 >
2020-10-28 21:47:22 +01:00
27ce5d921e
aco: remove isel_context::allocated
...
Now that we have Program::temp_rc, we can replace it with the first
temporary id allocated for NIR's ssa defs.
No fossil-db changes on Navi.
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/7067 >
2020-10-26 15:14:32 +00:00
eb6877d3af
radv,aco: fix use of texop_samples_identical in the resolve meta path
...
The return value of this texture intrinsic should be a NIR 1-bit bool.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7236 >
2020-10-21 13:06:53 +02:00
34bc9477de
aco: Clean up symbol names and comments related to NGG
...
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com >
Acked-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7094 >
2020-10-21 09:49:38 +00:00
86c227c10c
aco: Use strong typing to model SW<->HW stage mappings
...
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com >
Acked-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7094 >
2020-10-21 09:49:38 +00:00
4ca1030774
radv: move all NIR pass outside of ACO
...
This has several advantages:
- it generates roughly the same NIR for both compiler backends
(this might help for debugging purposes)
- it might allow to move around some NIR pass to improve compile time
- it might help for RadeonSI support
- it improves fossils-db stats for RADV/LLVM (this shouldn't matter
much but it's a win for free)
fossil-db (Navi/LLVM):
Totals from 80732 (59.18% of 136420) affected shaders:
SGPRs: 5390036 -> 5382843 (-0.13%); split: -3.38%, +3.24%
VGPRs: 3910932 -> 3890320 (-0.53%); split: -2.38%, +1.85%
SpillSGPRs: 319212 -> 283149 (-11.30%); split: -17.69%, +6.39%
SpillVGPRs: 14668 -> 14324 (-2.35%); split: -7.53%, +5.18%
CodeSize: 265360860 -> 267572132 (+0.83%); split: -0.47%, +1.30%
Scratch: 5338112 -> 6134784 (+14.92%); split: -2.65%, +17.57%
MaxWaves: 1077230 -> 1086902 (+0.90%); split: +2.79%, -1.90%
No fossils-db changes on RADV/ACO.
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/7077 >
2020-10-20 10:21:39 +00:00
2f125908b3
radv,aco: lower_pack_half_2x16
...
This patch also optimizes pack_half_2x16(a, 0.0).
Totals from 1949 (1.43% of 136546) affected shaders (RAVEN):
SGPRs: 83376 -> 83336 (-0.05%)
CodeSize: 3532144 -> 3512352 (-0.56%)
Instrs: 660746 -> 660682 (-0.01%); split: -0.01%, +0.00%
Cycles: 6780716 -> 6780472 (-0.00%); split: -0.00%, +0.00%
VMEM: 990886 -> 990883 (-0.00%); split: +0.00%, -0.00%
SMEM: 150506 -> 150538 (+0.02%); split: +0.05%, -0.03%
SClause: 30595 -> 30594 (-0.00%); split: -0.01%, +0.00%
Copies: 40801 -> 40729 (-0.18%)
PreSGPRs: 52335 -> 52341 (+0.01%); split: -0.03%, +0.04%
PreVGPRs: 45104 -> 45097 (-0.02%)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6777 >
2020-10-14 15:31:38 +00:00
112e66fa09
aco: compute the CS workgroup size from the shader NIR info
...
cs.block_size is copied from cs.local_size during the shader info pass.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7061 >
2020-10-14 15:09:34 +00:00
bcf7a70008
aco: use nir_opt_uniform_atomics
...
Significantly improves performance of a Control compute shader. Also seems
to increase FPS at the very start of the game by ~9% (RX 580, 1080p,
medium settings, no MSAA).
fossil-db (Navi):
Totals from 315 (0.23% of 135946) affected shaders:
SGPRs: 18296 -> 18336 (+0.22%); split: -0.26%, +0.48%
VGPRs: 11856 -> 11844 (-0.10%); split: -0.81%, +0.71%
CodeSize: 2233800 -> 2457508 (+10.01%)
MaxWaves: 4506 -> 4497 (-0.20%); split: +0.04%, -0.24%
Instrs: 438766 -> 486215 (+10.81%); split: -0.00%, +10.81%
Cycles: 7880180 -> 8963340 (+13.75%); split: -0.00%, +13.75%
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/6558 >
2020-10-13 12:47:21 +00:00
e1120f274f
nir: move divergence analysis options to nir_shader_compiler_options
...
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/6558 >
2020-10-13 12:47:21 +00:00
5ae3656890
aco/ngg: Calculate workgroup size of NGG shaders.
...
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6964 >
2020-10-09 15:26:15 +02:00
61280bb4b6
aco/ngg: Allocate NGG GS space early for const vertex/primitive counts.
...
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6964 >
2020-10-09 15:26:15 +02:00
c29e288fb5
aco/ngg: Create LDS layout for NGG GS.
...
For NGG GS, we need to store the following in LDS:
1. The ESGS ring, similarly to legacy ESGS.
2. Emitted vertices from the GS threads.
3. Temporary space used by the workgroup scan.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6964 >
2020-10-09 15:26:15 +02:00
2680329fb7
aco/ngg: Setup NGG GS.
...
Make it possible for ACO to recognize when to use HW NGG GS.
Also add a few notes about the various GS stages in the comments.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6964 >
2020-10-09 15:26:15 +02:00
b57b1a06e4
aco/ngg: Clean up and reorganize NGG VS/TES code.
...
Make the NGG VS/TES code easier to follow, give better names to
some functions and make ngg_nogs_early_prim_export a variable.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6964 >
2020-10-09 15:26:14 +02:00
df63491594
radv/aco: lower IO for all stages outside of ACO
...
Lowering IO for VS, TCS, TES and GS still have to be done for LLVM.
No fossils-db change on NAVI10.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6897 >
2020-10-01 14:58:25 +00:00
618556a8cb
nir: Drop the high_offset argument to the load_store_vectorizer filter.
...
Nothing uses it, and it's not clear to me what it provides over
alignment/num_components/bit_size.
Reviewed-by: Rob Clark <robdclark@chromium.org >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6612 >
2020-09-30 19:53:43 +00:00
5f757bb95c
nir: Make the load_store_vectorizer provide align_mul + align_offset.
...
It was passing an encoding of the two that wasn't good for ensuring "Don't
combine loads that would make us straddle a vec4 boundary" for
nir_lower_ubo_vec4.
Reviewed-by: Rob Clark <robdclark@chromium.org >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6612 >
2020-09-30 19:53:43 +00:00
291cfb1e41
radv: move lowering of FS outputs outside of ACO
...
This enables lowering of FS outputs for RADV/LLVM.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com >
Acked-by: Daniel Schürmann <daniel@schuermann.dev >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6865 >
2020-09-29 14:44:05 +00:00
85074ec5f6
radv/aco: Set I/O variable locations outside ACO.
...
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Acked-by: Daniel Schürmann <daniel@schuermann.dev >
Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6865 >
2020-09-29 14:44:05 +00:00
1588644543
radv: lower deref operations for global memory for both backends
...
To match ACO.
Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5316 >
2020-09-29 07:24:35 +00:00
9750164c09
nir: Rename get_buffer_size to get_ssbo_size
...
This makes it explicit that this intrinsic is only for SSBOs. For the
v3dv driver, we'll be adding a get_ubo_size intrinsic and we want to be
able to distinguish between the two.
Reviewed-by: Marek Olšák <marek.olsak@amd.com >
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com >
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6812 >
2020-09-22 13:34:12 +00:00
f100cf0d30
aco: stop multiplying driver_location by 4
...
This didn't really serve any purpose, doesn't match how FS inputs are
currently done, and prevented us from using
nir_io_add_const_offset_to_base in the future.
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/6689 >
2020-09-22 12:38:43 +00:00
9bba79088d
aco: use io semantics to get an intrinsic's slot
...
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/6689 >
2020-09-22 12:38:43 +00:00
d58a1a87cc
aco: Use NIR IO semantics for tess factor IO locations.
...
Previously we relied on looping over the NIR output variables
to remember the driver location of the tess factors, now use
the new NIR IO semantics instead.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com >
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6689 >
2020-09-22 12:38:43 +00:00
ec2185c598
aco: keep track of temporaries' regclasses in the Program
...
A future change will switch the liveness sets to bit vectors, which don't
contain regclass information.
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/6733 >
2020-09-21 13:47:28 +00:00
47de553283
aco/isel: Move context initialization code to a dedicated file
...
aco_instruction_selection_setup.cpp (previously used as a header) has
been split into a header and an implementation file. The latter "only"
implements init_context and setup_isel_context, but since these files
carry a long trail of helper functions, this cleans up the isel header
a lot.
Reduces library size by 3.1% due to more functions being compiled with
static linkage. Makes aco_instruction_selection.cpp compile 3% faster.
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6504 >
2020-09-08 20:13:51 +00:00
1eac0b52e3
aco/isel: Remove unused definitions
...
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev >
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6504 >
2020-09-08 20:13:51 +00:00
70cbddc4a7
nir: use enum operator helper for nir_variable_mode and nir_metadata
...
those are used quite a bit
Signed-off-by: Karol Herbst <kherbst@redhat.com >
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6520 >
2020-09-01 17:45:08 +00:00
a79dad950b
nir,amd: remove trinary_minmax opcodes
...
These consist of the variations nir_op_{i|u|f}{min|max|med}3 which are either
lowered in the backend (LLVM) anyway or can be recombined by the backend (ACO).
Reviewed-by: Marek Olšák <marek.olsak@amd.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6421 >
2020-08-24 20:56:11 +00:00
2133e64203
aco: use nir_intrinsic_has_access
...
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com >
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com >
Reviewed-by: Rob Clark <robdclark@gmail.com >
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6402 >
2020-08-21 16:47:00 +00:00
Timur Kristóf