Commit graph

5053 commits

Author SHA1 Message Date
Sviatoslav Peleshko
a6459e0f7b nir/loop_analyze: Don't test non-positive iterations count
Testing negative iterations count makes no sense, and can cause issues
when the unsigned type is used.
Testing 0 iterations is already covered with
will_break_on_first_iteration, so it can be skipped too.

Fixes: 6772a17a ("nir: Add a loop analysis pass")
Closes: https://gitlab.freedesktop.org/mesa/mesa/-/issues/9913
Signed-off-by: Sviatoslav Peleshko <sviatoslav.peleshko@globallogic.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26173>
2023-12-19 12:53:52 +00:00
Job Noorman
6cad2fc230 nir: add helper to create cursor after all @decl_regs
@decl_reg intrinsics must be in the first block so it's convenient to be
able to create an insertion point after all @decl_regs when the first
block needs to be split.

Signed-off-by: Job Noorman <jnoorman@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26737>
2023-12-18 14:52:02 +00:00
Job Noorman
6e7a61df4c nir: add _safe variants of nir_foreach_reg_load/store
Signed-off-by: Job Noorman <jnoorman@igalia.com>
Reviewed-by: Rob Clark <robdclark@chromium.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26175>
2023-12-15 17:19:28 +00:00
Faith Ekstrand
1cf1b9d741 nir: Scalarize bounds checked loads and stores
Fixes: 39da1deb49 ("nir/lower_io: Add a bounds-checked 64-bit global address format")
Reviewed-by: M Henning <drawoc@darkrefraction.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26526>
2023-12-15 03:53:54 +00:00
Caio Oliveira
81e3b28f78 compiler: Remove C++ static member pointers to builtin types
When we moved the bulk of glsl_type to C, these globals were
kept to avoid changes to compiler/glsl code in the MR.  Now that
landed, change the code to use the actual bultins directly.

Acked-by: Timothy Arceri <tarceri@itsqueeze.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26658>
2023-12-15 03:09:19 +00:00
Karol Herbst
8c73b1eb90 nir/algebraic: add support for custom arguments
Those are passed as an optional argument and are declared as a list of
(type, name) tuples.

At the moment this can only be used for conditions.

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26214>
2023-12-12 18:48:11 +00:00
Lionel Landwerlin
f53748c481 nir: fixup nir_printf intrinsic description
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26505>
2023-12-12 11:11:10 +00:00
Lionel Landwerlin
dc3e69af1a nir/serialize: untangle printf serialization from a particular stage
This allows any stage to carry printf instructions.

Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26505>
2023-12-12 11:11:10 +00:00
Lionel Landwerlin
4e4a3820ab nir/divergence: handle printf intrinsic
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Cc: mesa-stable
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26505>
2023-12-12 11:11:10 +00:00
Lionel Landwerlin
f7ae92b868 nir: include printfs from linked shaders
Once lowered low enough, it's not always possible to tell what strings
are used. So include them all when linking another shader.

Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26505>
2023-12-12 11:11:10 +00:00
Lionel Landwerlin
81b3dea993 nir/clone: fix missing printf_info clone
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Cc: mesa-stable
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26505>
2023-12-12 11:11:10 +00:00
Lionel Landwerlin
603f039708 nir: make printf_info (de)serializer available
Signed-off-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26505>
2023-12-12 11:11:10 +00:00
Ian Romanick
7fce0a5598 nir: Handle divergence for decl_reg
Once decl_reg is handled, src[0].ssa->divergent will be properly set, so
load_reg and load_reg_indirect do not need special treatment.

shader-db can run to completion on HSW, IVB, and SNB now. No other
testing was done.

v2: Refactor nir_intrinsic_load_reg and nir_intrinsic_load_reg_indirect
handling. Suggested by Daniel Schürmann.

Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Fixes: 4fd257d20f ("nir: Properly handle divergence for load_reg")
Fixes: 6dbb5f1e07 ("intel/fs: rerun divergence analysis prior to convert_from_ssa")
Closes: #10233
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26436>
2023-12-11 17:10:51 +00:00
Faith Ekstrand
aac1e3f595 nir: Add a new has_fmulz_no_denorms flag
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26569>
2023-12-11 15:29:17 +00:00
Alyssa Rosenzweig
c43c90a5fa asahi: rewrite pointsize handling
In the wise words of Mike Blumenkrantz, "I hate gl_PointSize and so can you".

The mesa/st lowering won't mesh well with vertex shader epilogues, and it falls
over in various circumstances. I am too tired to go against the grain, so let's
just pretend to be a normal gallium driver and trust in the rasterizer CSO,
lowering point size internally. This properly handles transform feedback without
any hacks, both GL and GLES behaviours, etc.

Fixes:

   KHR-GL31.transform_feedback.capture_vertex_separate_test
   gl-2.0-large-point-fs

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26614>
2023-12-09 12:08:39 -04:00
Alyssa Rosenzweig
5987e47a29 asahi: rework GS input assembly
in prep for tessellation (which will share the IA lowering), and for multidraw
indirect (which greatly complicates IA lowering with geom/tess).

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26614>
2023-12-09 12:08:39 -04:00
Karol Herbst
6979a1aa07 nir/opt_preamble: make load_workgroup_size handling optional
not all drivers support it being in the preamble, e.g. asahi.

Signed-off-by: Karol Herbst <git@karolherbst.de>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26614>
2023-12-09 10:56:37 -04:00
Marek Olšák
7d2faa88ab nir,radeonsi: add FLAGS into load_vector_arg_amd to record color input usage
This will be needed for gathering color usage from lowered PS.

Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26307>
2023-12-09 00:05:27 +00:00
Yonggang Luo
e7c614bd20 util,vulkan,mesa,compiler: Generate source files with utf8 encoding from mako template
This is for fixes the following error:
FAILED: src/vulkan/runtime/vk_synchronization_helpers.c src/vulkan/runtime/vk_synchronization_helpers.h
"C:\CI-Tools\msys64\mingw64\bin/python3.EXE" "../../src/vulkan/util/vk_synchronization_helpers_gen.py" "--xml" "../../src/vulkan/registry/vk.xml" "--out-c" "src/vulkan/runtime/vk_synchronization_helpers.c" "--beta" "false"

Traceback (most recent call last):
  File "C:/work/xemu/mesa/src/vulkan/util/vk_synchronization_helpers_gen.py", line 213, in main
    f.write(TEMPLATE_C.render(**environment))
UnicodeEncodeError: 'gbk' codec can't encode character '\xa9' in position 15: illegal multibyte sequence

Signed-off-by: Yonggang Luo <luoyonggang@gmail.com>
Acked-by: Eric Engestrom <eric@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26515>
2023-12-07 12:41:07 +00:00
Timothy Arceri
c468aa6059 nir: add nir_fixup_deref_types()
This will be used to fix up types after arrays have been resized.

Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26534>
2023-12-07 03:05:09 +00:00
Faith Ekstrand
d2ffcb6092 nir: Lower [su]dot_4x8_[ui]add_sat to [su]dot_4x8_[ui]add
Since nir_opt_algebraic runs on its own results, if the driver doesn't
have [su]dot_4x8_[ui]add then the [su]dot_4x8_[ui]add lowering rules
will kick in and lower that to what we had originally.

Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26533>
2023-12-06 23:15:33 +00:00
Faith Ekstrand
09fc5e1c4d nir: Split has_[su]dot_4x8 bits into regular and _sat versions
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Reviewed-by: Ian Romanick <ian.d.romanick@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26533>
2023-12-06 23:15:33 +00:00
Faith Ekstrand
22c60ea384 nir: Respect variable alignments in lower_vars_to_explicit_types
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26522>
2023-12-06 18:37:36 +00:00
Faith Ekstrand
629af540ca spirv: Plumb variable alignments through to NIR
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Jesse Natalie <jenatali@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26522>
2023-12-06 18:37:36 +00:00
Rhys Perry
0cbe0d2968 nir/loop_analyze: support loops with min/max and non-add incrementation
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Timothy Arceri <tarceri@itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26225>
2023-12-06 15:58:36 +00:00
Rhys Perry
fc1ebc67b1 nir/loop_analyze: support umin and {u,i,f}max
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Timothy Arceri <tarceri@itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26225>
2023-12-06 15:58:36 +00:00
Rhys Perry
9591c36666 nir/loop_analyze: check min compatibility with comparison
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Timothy Arceri <tarceri@itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26225>
2023-12-06 15:58:36 +00:00
Rhys Perry
b6c2a5d48d nir/loop_analyze: fix vector basis/limit/comparison
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Timothy Arceri <tarceri@itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26225>
2023-12-06 15:58:36 +00:00
Rhys Perry
2f46773819 nir/loop_analyze: scalarize try_eval_const_alu
This is simpler, and users of this function expected scalar anyway.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Timothy Arceri <tarceri@itsqueeze.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26225>
2023-12-06 15:58:36 +00:00
Rhys Perry
7960588413 nir/loop_analyze: skip if basis/limit/comparison is vector
Replacement in try_eval_const_alu() doesn't work because the replacements
are always scalar. The callers also always give a scalar dest.

This is encountered when compiling a Redout shader under ASan.

Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Acked-by: Timothy Arceri <tarceri@itsqueeze.com>
Fixes: bc170e895f ("nir/loop_analyze: Use try_eval_const_alu and induction variable basis info")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26225>
2023-12-06 15:58:36 +00:00
Faith Ekstrand
06257514c1 nir: Return a mesa_prim from gs_in_prim_for_topology
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24821>
2023-12-05 23:12:32 +00:00
Alyssa Rosenzweig
65e431e61a nir/lower_idiv: Optimize idiv sign calculation
Save a comparison, and move out the comparison to be more backend friendly.
Saves 2 instrs on AGX (as the remaining comparison now fuses with bcsel).

Results on AGX, all affected shaders in asphalt9.

   total instructions in shared programs: 1813003 -> 1812611 (-0.02%)
   instructions in affected programs: 119646 -> 119254 (-0.33%)
   helped: 333
   HURT: 0
   Instructions are helped.

   total bytes in shared programs: 11870344 -> 11867208 (-0.03%)
   bytes in affected programs: 820888 -> 817752 (-0.38%)
   helped: 333
   HURT: 0
   Bytes are helped.

and on Mali-G57:

   total instructions in shared programs: 2677538 -> 2677205 (-0.01%)
   instructions in affected programs: 206923 -> 206590 (-0.16%)
   helped: 333
   HURT: 0
   Instructions are helped.

   total cvt in shared programs: 14667.50 -> 14662.30 (-0.04%)
   cvt in affected programs: 1953.64 -> 1948.44 (-0.27%)
   helped: 333
   HURT: 0
   Cvt are helped.

   total quadwords in shared programs: 1450664 -> 1450544 (<.01%)
   quadwords in affected programs: 5064 -> 4944 (-2.37%)
   helped: 15
   HURT: 0
   Quadwords are helped.

   total threads in shared programs: 53282 -> 53309 (0.05%)
   threads in affected programs: 27 -> 54 (100.00%)
   helped: 27
   HURT: 0
   Threads are helped.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26489>
2023-12-05 21:46:26 +00:00
Faith Ekstrand
eda940c855 nak: Make barriers SSA-friendly
The NIR intrinsics now take and return a barrier whenever one is
modified instead of modifying in-place.  In NAK, we give the internal
instructions the same treatment and convert everything to use barrier
SSA values and RegRefs.  In nak_from_nir, we move all barriers to/from
GPRs.  We'll clean up the massive pile of OpBMov later.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26463>
2023-12-05 18:59:40 +00:00
Alyssa Rosenzweig
d7f0804a47 nir/format_convert: handle clamping smaller bit sizes
asahi will pass in 16bits, works fine if we convert before clamping. note we
don't try to be clever and make a smaller immediate because it would require
extra logic for negatives to make sure we don't have garbage in upper bits
(nir_validate checks that). do the simple, obviously correct thing.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26440>
2023-12-04 15:41:19 +00:00
Qiang Yu
b16017d080 nir: fix load layer id system_values_read info gather
Otherwise radeons+ACO compilation fail.

Fixes: 32ee6376ad ("nir: add lowering from FS LAYER input to LAYER_ID sysval")
Acked-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Signed-off-by: Qiang Yu <yuq825@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26422>
2023-12-01 03:17:57 +00:00
Konstantin Seurer
827b0fa1ef nir/lower_vars_to_scratch: Remove all unused derefs
If the shader passed to nir_lower_vars_to_scratch contains some unused
derefs to a variable that will be lowered, validation will fail because
the variable is not part of the shader after the pass.

cc: mesa-stable

Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26271>
2023-11-29 06:49:48 +00:00
Alyssa Rosenzweig
e14633fa7d nir/lower_tex: Add 1D lowering
From amd/common.

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Faith Ekstrand <faith.ekstrand@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26377>
2023-11-28 20:32:03 +00:00
Marek Olšák
b6e98677c3 nir/print: print PATCH0 and VARn_16BIT names instead of numbers for TCS and TES
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26275>
2023-11-25 09:50:19 -05:00
Marek Olšák
5c8730ebe8 nir: don't declare illegal varyings in nir_create_passthrough_tcs
I called it accidentally with LAYER.

Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26275>
2023-11-25 09:50:19 -05:00
Marek Olšák
7a9b73fcb8 nir: fix gathering TESS_LEVEL_INNER/OUTER usage with lowered IO
Those varyings shouldn't flag patch_inputs_read/patch_outputs_written.

Fixes: 10be706778 - nir: gather indirect info from lowered IO intrinsics

Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26275>
2023-11-25 09:50:19 -05:00
Marek Olšák
27a9ddad28 nir: return progress from nir_remove_sysval_output
Changing IO semantics doesn't affect the SSA structure.

Same as nir_remove_varying.

Reviewed-by: Qiang Yu <yuq825@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26274>
2023-11-24 15:37:24 +00:00
Marek Olšák
32ee6376ad nir: add lowering from FS LAYER input to LAYER_ID sysval
Reviewed-by: Qiang Yu <yuq825@gmail.com>
Reviewed-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26274>
2023-11-24 15:37:24 +00:00
Daniel Schürmann
f1110576d9 nir: add info.fs.require_full_quads
This flag indicates the requirement of helper invocations
in fragment shaders, independent from any present instructions.
This fixes the lowering of OpGroupNonUniformQuad* instructions.

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26026>
2023-11-22 11:31:52 +01:00
Daniel Schürmann
2db0507b5d nir/gather_info: add missing wide subgroup operations
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26026>
2023-11-22 11:31:46 +01:00
Daniel Schürmann
1179d83a89 nir: remove info.fs.needs_all_helper_invocations
Use info.uses_wide_subgroup_intrinsics instead.

Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26026>
2023-11-22 11:31:11 +01:00
Jesse Natalie
c69ca8c5c1 nir_lower_mem_access_bit_sizes: Fix assert (bit -> byte size)
Reviewed-By: Sil Vilerino <sivileri@microsoft.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26293>
2023-11-21 19:11:40 +00:00
Timothy Arceri
4382bd8a34 nir: simplify nir_build_write_masked_store()
Suggested-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25860>
2023-11-21 00:45:21 +00:00
Timothy Arceri
110887de2b glsl/nir: implement a nir based lower distance pass
This is based off the original GLSL IR pass but it is much much
simpler as it doesn't need to do all of the hackery required in
GLSL IR to achieve the lowering.

Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25860>
2023-11-21 00:45:21 +00:00
Timothy Arceri
95da0a9c96 nir: move build_write_masked_stores() to nir builder
We will reused this in the following patch.

Acked-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25860>
2023-11-21 00:45:21 +00:00
Alejandro Piñeiro
c0cfa4f53b nir: add new opcodes to map new v71 packing/conversion instructions
Since v71, broadcom hw include specific packing/conversion
instructions, so this commit adds opcodes to be able to make use of
them, specially for image stores:

   * pack_2x16_to_unorm_2x8 (on backend vftounorm8/vftosnorm8):
     2x16-bit floating point to 2x8-bit unorm/snorm

   * f2unorm_16/f2snorm_16 (on backend ftounorm16/ftosnorm16):
     floating point to 16-bit unorm/snorm

   * pack_2x16_to_unorm_2x10/pack_2x16_to_unorm_10_2 (on backend
     vftounorm10lo/vftounorm10hi): used to convert a floating point to
     a r10g10b10a2 unorm

   * pack_32_to_r11g11b10 (on backend v11fpack): packs 2 2x16 FP into
     R11G11B10.

   * pack_uint_32_to_r10g10b10a2 (on backend v10pack): pack 2 2x16
     integer into R10G10B10A2

   * pack_4x16_to_4x8 (on backend v8pack): packs 2 2x16 bit integer
     into 4x8 bits.

   * pack_2x32_to_2x16 (on backend vpack): 2x32 bit to 2x16 integer
     pack

For the latter, it can be easly confused with the existing
pack_32_2x16_split. But note that this one receives two 16bit integer,
and packs them on a 32bit integer. But broadcom opcode takes two 32bit
integer, takes the lower halfword, and packs them as 2x16 on a 32bit
integer.

Interestingly broadcom also defines a similar one that packs the
higher halfword. Not used yet.

Note that at this point we use agnostic names, even if we add a _v3d
suffix as they are only available for broadcom, in order to follow
current NIR conventions.

Reviewed-by: Iago Toral Quiroga <itoral@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/25726>
2023-11-20 08:20:31 +00:00