Commit graph

4476 commits

Author SHA1 Message Date
Caio Marcelo de Oliveira Filho
8548fe19f0 nir: Make nir_deref_path_init skip trivial casts
In a NIR generated using SPIR-V initializers to variables, copy
propagation can end up transforming

    vec1 32 ssa_33 = deref_var &@1 (shared mat2x4)
    vec1 32 ssa_35 = mov ssa_33
    vec1 32 ssa_7 = deref_cast (mat2x4 *)ssa_35 (shared mat2x4)  /* ptr_stride=0 */

into

    vec1 32 ssa_33 = deref_var &@1 (shared mat2x4)
    vec1 32 ssa_7 = deref_cast (mat2x4 *)ssa_33 (shared mat2x4)  /* ptr_stride=0 */

Before the optimization, the "head" of a path of deref that uses ssa_7
will be the cast.  After, it will be the variable in ssa_33.  Since
the types are the same, this is a trivial cast that would be picked up
by nir_opt_deref.

If we need to compare such deref-chain after optimization with another
deref-chain for the same variable, the compare function will get
confused by the cast in the middle.

One alternative would be to add nir_opt_deref to places that compare
derefs, but that might not scale well, so skip the trivial casts when
generating the paths instead.

Motivated by the discussion in
https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3047#note_383660.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3420>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3420>
2020-01-29 18:25:36 +00:00
Rhys Perry
1f72857739 nir/algebraic: add some half packing optimizations
pipeline-db (ACO):
Totals from affected shaders:
SGPRS: 29200 -> 29200 (0.00 %)
VGPRS: 17372 -> 17372 (0.00 %)
Spilled SGPRs: 105 -> 105 (0.00 %)
Spilled VGPRs: 0 -> 0 (0.00 %)
Private memory VGPRs: 0 -> 0 (0.00 %)
Scratch size: 0 -> 0 (0.00 %) dwords per thread
Code Size: 1406576 -> 1389256 (-1.23 %) bytes
LDS: 83 -> 83 (0.00 %) blocks
Max Waves: 3976 -> 3976 (0.00 %)

pipeline-db (LLVM):
Totals from affected shaders:
SGPRS: 21320 -> 21320 (0.00 %)
VGPRS: 17056 -> 17036 (-0.12 %)
Spilled SGPRs: 22 -> 22 (0.00 %)
Spilled VGPRs: 503 -> 487 (-3.18 %)
Private memory VGPRs: 0 -> 0 (0.00 %)
Scratch size: 396 -> 396 (0.00 %) dwords per thread
Code Size: 1441244 -> 1423292 (-1.25 %) bytes
LDS: 463 -> 463 (0.00 %) blocks
Max Waves: 3609 -> 3611 (0.06 %)

v2: add pattern for ishr

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/2271>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/2271>
2020-01-29 14:30:33 +00:00
Rhys Perry
5476d18183 nir/algebraic: add patterns for a >> #b << #b
Fixes compilation of a Battlefront 2 shader with ACO by removing VGPR
spilling. The reassociation makes it worse on LLVM though.

pipeline-db (ACO):
Totals from affected shaders:
SGPRS: 10704 -> 10688 (-0.15 %)
VGPRS: 18736 -> 18528 (-1.11 %)
Spilled SGPRs: 70 -> 70 (0.00 %)
Spilled VGPRs: 0 -> 0 (0.00 %)
Private memory VGPRs: 0 -> 0 (0.00 %)
Scratch size: 0 -> 0 (0.00 %) dwords per thread
Code Size: 909696 -> 885796 (-2.63 %) bytes
LDS: 225 -> 225 (0.00 %) blocks
Max Waves: 1115 -> 1129 (1.26 %)

pipeline-db (LLVM):
Totals from affected shaders:
SGPRS: 8472 -> 8424 (-0.57 %)
VGPRS: 14284 -> 14368 (0.59 %)
Spilled SGPRs: 0 -> 0 (0.00 %)
Spilled VGPRs: 442 -> 503 (13.80 %)
Private memory VGPRs: 0 -> 0 (0.00 %)
Scratch size: 268 -> 396 (47.76 %) dwords per thread
Code Size: 862568 -> 853028 (-1.11 %) bytes
LDS: 0 -> 0 (0.00 %) blocks
Max Waves: 971 -> 964 (-0.72 %)

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Connor Abbott <cwabbott0@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/2271>
2020-01-29 14:30:33 +00:00
Samuel Pitoiset
531a26d5aa spirv: implement SPV_AMD_shader_explicit_vertex_parameter
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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
cf6cae832c nir: lower interp_deref_at_vertex to load_input_vertex
This introduces a new NIR intrinsic for loading inputs at a specific
vertex index.

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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
d29f10a7ca nir: add nir_intrinsic_interp_deref_at_vertex
From the SPV_AMD_shader_explicit_vertex_parameter extension:
   "Returns the value of the input <interpolant> without any
    interpolation, i.e. the raw output value of previous shader
    stage."

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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
687f170311 nir: lower SYSTEM_VALUE_BARYCENTRIC_* to nir_load_barycentric()
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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
9021b45b35 nir: add nir_intrinsic_load_barycentric_model
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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
df8dd12e5b spirv: add support for SpvBuiltInBaryCoord*
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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
61d24080bb compiler: add new SYSTEM_VALUE_BARYCENTRIC_*
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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
15d53d8294 compiler: add PERSP to the existing barycentric system values
We need the LINEAR versions for AMD_shader_explicit_vertex_parameter.

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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
5c053cc6ec spirv: add support for SpvDecorationExplicitInterpAMD
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/3578>
2020-01-29 09:49:50 +00:00
Samuel Pitoiset
746e9e5d66 compiler: add a new explicit interpolation mode
This introduces one more interpolation mode INTERP_MODE_EXPLICIT,
which is needed for AMD_shader_explicit_vertex_parameter.

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/3578>
2020-01-29 09:49:50 +00:00
Vasily Khoruzhick
3c754900b5 nir: don't emit ishl in _nir_mul_imm() if backend doesn't support bitops
Otherwise we'll have to lower it later.

Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Reviewed-by: Erico Nunes <nunes.erico@gmail.com>
Signed-off-by: Vasily Khoruzhick <anarsoul@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3529>
2020-01-23 21:16:22 +00:00
Anthony Pesch
1496cc92f6 util/hash_table: added hash functions for integer types
A few hash_table users roll their own integer hash functions which
call _mesa_hash_data to perform the hashing which ultimately calls
into XXH32 with a dynamic key length. When using small keys with a
constant size the hash rate can be greatly improved by inlining
XXH32 and providing it a constant key length, see:
https://fastcompression.blogspot.com/2018/03/xxhash-for-small-keys-impressive-power.html

Additionally, this patch removes calls to _mesa_key_hash_string and
makes them instead call _mesa_has_string directly, matching the new
integer hash functions.

Reviewed-by: Eric Anholt <eric@anholt.net>
Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3475>
2020-01-23 17:06:57 +00:00
Samuel Pitoiset
84b08971fb nir/lower_input_attachments: lower nir_texop_fragment_{mask}_fetch
These instructions are allowed to fetch from multisampled
subpass input attachments.

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/3304>
2020-01-23 10:48:02 +00:00
Samuel Pitoiset
76a34f5d3f spirv: add support for SpvOpFragment{Mask}FetchAMD operations
nir_tex_src_ms_index is re-used for the fragment index with
nir_texop_fragment_fetch to avoid introducing a new texture source type.

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/3304>
2020-01-23 10:48:02 +00:00
Samuel Pitoiset
603e6ba972 nir: add two new texture ops for multisample fragment color/mask fetches
This introduces:
   - nir_texop_fragment_mask_fetch (fetch a fragment mask from a
     compressed multisampled color surface)
   - nir_texop_fragment_fetch (fetch a color fragment for a
     particular sample at corresponding fragment mask index).

These two texture operations are necessary for implementing
SPV_AMD_shader_fragment_mask.

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/3304>
2020-01-23 10:48:02 +00:00
Samuel Pitoiset
dea29b3818 spirv: add SpvCapabilityFragmentMaskAMD
This new capability is for SPV_AMD_shader_fragment_mask.

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/3304>
2020-01-23 10:48:02 +00:00
Matt Turner
4413537c80 util: Remove tmp argument from BITSET_FOREACH_SET macro
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3499>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3499>
2020-01-23 01:52:43 +00:00
Timothy Arceri
bf830250a7 glsl_to_nir: update interface type properly
Since 76ba225184 the member variable types were being redefined
but we assigned the old interface type to the variable.

In a following patch series we will use the types to check if we
are dealing with an interface instance when apply GLSL linking
rules.

Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3468>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3468>
2020-01-23 01:02:25 +00:00
Timothy Arceri
d3a4d1775e glsl: count uniform components and storage better in nir linking
This helps avoid incorrect validation error when linking glsl
shaders and avoids assigning uniform storage slots that will
never be used.

Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3468>
2020-01-23 01:02:25 +00:00
Timothy Arceri
e5b3cf433e glsl: fix check for matrices in blocks when using nir uniform linker
We need to stripe any arrays before checking the type. Here we
just use the uniform type which has already be stripped.

Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3468>
2020-01-23 01:02:25 +00:00
Timothy Arceri
55e4410b34 glsl: remove bogus assert in nir uniform linking
I'm not sure why this was first added but it causes an assert
on any uniform matrix.

Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3468>
2020-01-23 01:02:25 +00:00
Ian Romanick
b065d8fb8c nir/algebraic: Optimize some 64-bit integer comparisons involving zero
I noticed that we can do better for these kinds of comparisons while
working on the lowering for iadd_sat@64 and isub_sat@64.  This
eliminated 11 instruction from the fs-addSaturate-int64.shader_test.

My hope is that this will improve the run-time of int64 tests on Ice
Lake.  I have no data to support or refute this.

Unsurprisingly, no changes on shader-db.

v2: Condition the min and max patterns with nir_lower_minmax64.
Suggested by Caio.  Very long discussion in the MR. :)

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
4fcddb55f2 spirv: Add support for IntegerFunctions2INTEL capability
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
aa56934e2a spirv: Silence a bunch of unused parameter warnings
The change to get_uniform_nir_atomic_op make it look like the other
get_*_nir_atomic_op functions.  The rest just add UNUSED or ASSERTED
to parameters required for some of the interfaces.

src/compiler/spirv/spirv_to_nir.c: In function ‘struct_member_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:673:47: warning: unused parameter ‘val’ [-Wunused-parameter]
                             struct vtn_value *val, int member,
                                               ^~~
src/compiler/spirv/spirv_to_nir.c: In function ‘struct_member_matrix_stride_cb’:
src/compiler/spirv/spirv_to_nir.c:778:50: warning: unused parameter ‘val’ [-Wunused-parameter]
                                struct vtn_value *val, int member,
                                                  ^~~
src/compiler/spirv/spirv_to_nir.c: In function ‘type_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:805:61: warning: unused parameter ‘ctx’ [-Wunused-parameter]
                     const struct vtn_decoration *dec, void *ctx)
                                                             ^~~
src/compiler/spirv/spirv_to_nir.c: In function ‘spec_constant_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:1359:70: warning: unused parameter ‘v’ [-Wunused-parameter]
 spec_constant_decoration_cb(struct vtn_builder *b, struct vtn_value *v,
                                                                      ^
src/compiler/spirv/spirv_to_nir.c: In function ‘handle_workgroup_size_decoration_cb’:
src/compiler/spirv/spirv_to_nir.c:1407:43: warning: unused parameter ‘data’ [-Wunused-parameter]
                                     void *data)
                                           ^~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_function_call’:
src/compiler/spirv/spirv_to_nir.c:1806:55: warning: unused parameter ‘opcode’ [-Wunused-parameter]
 vtn_handle_function_call(struct vtn_builder *b, SpvOp opcode,
                                                       ^~~~~~
src/compiler/spirv/spirv_to_nir.c:1807:54: warning: unused parameter ‘count’ [-Wunused-parameter]
                          const uint32_t *w, unsigned count)
                                                      ^~~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘get_uniform_nir_atomic_op’:
src/compiler/spirv/spirv_to_nir.c:2548:47: warning: unused parameter ‘b’ [-Wunused-parameter]
 get_uniform_nir_atomic_op(struct vtn_builder *b, SpvOp opcode)
                                               ^
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_atomics’:
src/compiler/spirv/spirv_to_nir.c:2633:48: warning: unused parameter ‘count’ [-Wunused-parameter]
                    const uint32_t *w, unsigned count)
                                                ^~~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_barrier’:
src/compiler/spirv/spirv_to_nir.c:3197:48: warning: unused parameter ‘count’ [-Wunused-parameter]
                    const uint32_t *w, unsigned count)
                                                ^~~~~
src/compiler/spirv/spirv_to_nir.c: In function ‘vtn_handle_execution_mode’:
src/compiler/spirv/spirv_to_nir.c:3618:68: warning: unused parameter ‘data’ [-Wunused-parameter]
                           const struct vtn_decoration *mode, void *data)
                                                                    ^~~~

Acked-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
44471a76e9 nir/spirv: Translate SPIR-V to NIR for new INTEL_shader_integer_functions2 opcodes
v2: Rebase on 272e927d0e ("nir/spirv: initial handling of OpenCL.std
extension opcodes")

v3: Add missing SpvOpUCountTrailingZerosINTEL case to switch in
vtn_handle_body_instruction. Remove stray semicolon in
vtn_nir_alu_op_for_spirv_opcode. Use umin instead of umax for
SpvOpUCountTrailingZerosINTEL "lowering" in vtn_handle_alu.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
db649fd582 compiler: Translate GLSL IR to NIR for new INTEL_shader_integer_functions2 expressions
v2: Rebase on 272e927d0e ("nir/spirv: initial handling of OpenCL.std
extension opcodes")

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
d3d970166c nir/algebraic: Add lowering for 64-bit iadd_sat and isub_sat
v2: Rearranged and expand the comment about the optimizations applied to
the lowering.  Suggested by Caio.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
dcadbd2dd2 nir/algebraic: Add lowering for 64-bit uadd_sat
Fixes piglit fs-addsaturate-uint64 and vs-addsaturate-uint64 on Ice
Lake.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
1bdfc6d7cb nir/algebraic: Add lowering for 64-bit usub_sat
v2: Rebase on 272e927d0e ("nir/spirv: initial handling of OpenCL.std
extension opcodes")

v3: Add a new lower_usub_sat64 flag that only applies to the 64-bit
version of the nir_op_usub_sat instruction.

v4: Also enable the lowering when nir_lower_iadd64 is set.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> [v3]
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
a483771045 nir/algebraic: Add lowering for 64-bit hadd and rhadd
v2: Rebase on 272e927d0e ("nir/spirv: initial handling of OpenCL.std
extension opcodes")

v3: Add a new lower_hadd64 flag that only applies to the 64-bit versions
of the instructions.

v4: Also enable the lowering when nir_lower_iadd64 is set.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> [v3]
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
ea435560ee nir/algebraic: Add lowering for uabs_usub and uabs_isub
v2: Remove some rebase failures noticed by Caio.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
21f0d020fe nir: Add new instructions for INTEL_shader_integer_functions2
uctz isn't added because it will implemented in the GLSL path and the
SPIR-V path using other pre-existing instructions.

v2: Avoid signed integer overflow for uabs_isub(0, INT_MIN).  Noticed by
Caio.

v3: Alternate fix for signed integer overflow for abs_sub(0, INT_MIN).
I tried the previous methon in a small test program with -ftrapv, and it
failed.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com> [v1]
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
cb518df775 glsl: Add built-in functions for INTEL_shader_integer_functions2
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
5eda9f5832 glsl_types: Add function to get an unsigned base type from a signed type
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
1d165b0548 glsl: Add new expressions for INTEL_shader_integer_functions2
v2: Re-write iadd64_saturate and isub64_saturate to avoid undefined
overflow behavior.  Also fix copy-and-paste bug in isub64_saturate.
Suggested by Caio.

v3: Avoid signed integer overflow for abs_sub(0, INT_MIN).  Noticed by
Caio.

v4: Alternate fix for signed integer overflow for abs_sub(0, INT_MIN).
I tried the previous methon in a small test program with -ftrapv, and it
failed.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Ian Romanick
20d34c4ebf mesa: Extension boilerplate for INTEL_shader_integer_functions2
Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767>
2020-01-23 00:18:57 +00:00
Eric Anholt
7558b5da13 compiler: Add a note about how num_ssbos works in the program info.
These numbers are always confusing, and it's particularly so for this
field where it has a different meaning in different info structs.

Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3240>
2020-01-21 10:06:23 -08:00
Eric Anholt
d0975bfc4a nir: Drop the ssbo_offset to atomic lowering.
The arguments passed in were:
- prog->info.num_ssbos
- prog->nir->info.num_ssbos
- arbitrary values for standalone compilers

The num_ssbos should match between the prog's info and prog->nir's info
until this lowering happens.

Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3240>
2020-01-21 10:06:23 -08:00
Eric Anholt
10dc4ac4c5 mesa: Make atomic lowering put atomics above SSBOs.
Gallium arbitrarily (it seems) put atomics below SSBOs, resulting in a
bunch of extra index management, and surprising shader code when you would
see your SSBOs up at index 16.  It makes a lot more sense to see atomics
converted to SSBOs appear as magic high numbers.

Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3240>
2020-01-21 10:06:23 -08:00
Eric Anholt
d55573aac6 nir: Fix printing of ~0 .locations.
I kept wondering what "429" meant in variable declarations, when it was
just a truncated ~0 snprintf.

Reviewed-by: Rob Clark <robdclark@chromium.org>
Reviewed-by: Alyssa Rosenzweig <alyssa.rosenzweig@collabora.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3423>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3423>
2020-01-16 23:29:10 +00:00
Jason Ekstrand
721666e52a anv,nir: Lower quad_broadcast with dynamic index in NIR
This is required for the subgroupBroadcastDynamicId feature that was
added in Vulkan 1.2.

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira@intel.com>
2020-01-15 08:34:57 -06:00
Elie Tournier
22c5c54a4f nir/algebraic: sqrt(x)*sqrt(x) -> fabs(x)
total instructions in shared programs: 12840840 -> 12839341 (-0.01%)
instructions in affected programs: 122581 -> 121082 (-1.22%)
helped: 559
HURT: 0

total cycles in shared programs: 302505756 -> 302490031 (<.01%)
cycles in affected programs: 2022900 -> 2007175 (-0.78%)
helped: 1090
HURT: 130

Signed-off-by: Elie Tournier <elie.tournier@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/948>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/948>
2020-01-15 00:30:52 +00:00
Elie Tournier
6f394343b1 nir/algebraic: i2f(f2i()) -> trunc()
total instructions in shared programs: 12840968 -> 12840784 (<.01%)
instructions in affected programs: 17886 -> 17702 (-1.03%)
helped: 77
HURT: 0

total cycles in shared programs: 302508917 -> 302505592 (<.01%)
cycles in affected programs: 249964 -> 246639 (-1.33%)
helped: 70
HURT: 7

Signed-off-by: Elie Tournier <elie.tournier@collabora.com>
Reviewed-by: Matt Turner <mattst88@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/948>
2020-01-15 00:30:52 +00:00
Eric Anholt
3d9a3d0be0 i965: Reuse the new core glsl_count_dword_slots().
The only difference I could see was treating interfaces like structs.
Maintain that case.

Reviewed-by: Kristian H. Kristensen <hoegsberg@google.com>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3297>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3297>
2020-01-14 23:55:00 +00:00
Eric Anholt
bc4f089d01 mesa/st: Move the dword slot counting function to glsl_types as well.
To implement NIR-to-TGSI, we need to be able to get the size of the
uniform variable for the TGSI declaration, not just the
.driver_location.  With its location in mesa/st, drivers couldn't link
to it from nir-to-tgsi.

This feels like a common enough function to want, so let's share it in
the core compiler.

Reviewed-by: Kristian H. Kristensen <hoegsberg@google.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3297>
2020-01-14 23:55:00 +00:00
Eric Anholt
74ee3f76de mesa/st: Move the vec4 type size function into core GLSL types.
The only bit that gallium varied on was handling of bindless.  We can
retain previous behavior for count_attribute_slots() by passing in
"true" (though I suspect this is just giving a silly answer to a silly
question), and delete our recursive function from mesa/st.

Reviewed-by: Kristian H. Kristensen <hoegsberg@google.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3297>
2020-01-14 23:55:00 +00:00
Caio Marcelo de Oliveira Filho
3cc501be69 nir: Add missing nir_var_mem_global to various passes
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Tested-by: Marge Bot <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3322>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/3322>
2020-01-14 14:42:12 -08:00