AlexIndustrial/mesa
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
186,192 Commits 1 Branch 0 Tags
31920cb60c3cf487bc29ebd1d8ad8b1825e09fab
Commit Graph

186192 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
José Roberto de Souza
31920cb60c intel: Enable Xe KMD support by default 
Xe KMD landed on drm-next, uAPI is now stable and we can remove
the build time parameter to enable support to it but platforms
older than Lunar lake will have experimental support with Xe KMD.

Signed-off-by: José Roberto de Souza <jose.souza@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/20418>
 2024-03-12 22:22:50 +00:00
Kenneth Graunke
97aec40111 intel/brw: Emit better code for read_invocation(x, constant) 
For something as basic as read_invocation(x, 0), we were emitting:

   mov(8) vgrf67:D, 0d
   find_live_channel(8) vgrf236:UD, NoMask
   broadcast(8) vgrf237:D, vgrf67:D, vgrf236+0.0<0>:UD NoMask
   broadcast(8) vgrf235+0.0:W, vgrf197+0.0:W, vgrf237+0.0<0>:D NoMask
   mov(8) vgrf234+0.0:W, vgrf235+0.0<0>:W

This is way overcomplicated - if the invocation is a constant, we can
simply emit a single MOV which reads the desired channel index.  Not
only that, but it's difficult to clean up:

1. If this expression appears multiple times, CSE will find all the
   redundant emit_uniformize(invocation) and get rid of the duplicate
   (find_live_channel+broadcast) on future instructions.
2. Copy propagation will put the 0d directly in the first broadcast.
3. Dead code elimination will get rid of the vgrf67 temp holding 0.
4. Algebraic will replace the first broadcast(x, 0) with a MOV.
5. Copy propagation will put the 0d directly in the second broadcast.
6. Dead code elimination will get rid of the vgrf237 temp.
7. Algebraic will replace the second broadcast(x, 0) with a MOV.
8. Copy propagation will finally combine the two MOVs

That's at least 7-8 optimization passes and several loops through the
same passes just to clean up something we can do trivially.

Cuts 25% of the of the optimizer steps in pipeline 22200210259a2c9c
of fossil-db/google-meet-clvk/BgBlur.1f58fdf742c27594.1 (31 to 23).

Shortens compilation time of the google-meet-clvk/Relight pipeline by
-2.87717% +/- 0.509162% (n=150).

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28097>
 2024-03-12 21:58:27 +00:00
Ian Romanick
e87881f616 intel/brw: Avoid a silly add with zero in assign_curb_setup 
No shader-db changes.

fossil-db:

DG2
Totals:
Instrs: 161008251 -> 161004452 (-0.00%)
Cycles: 13894249509 -> 13893050101 (-0.01%); split: -0.01%, +0.00%

Totals from 3804 (0.58% of 652145) affected shaders:
Instrs: 2232984 -> 2229185 (-0.17%)
Cycles: 7124966553 -> 7123767145 (-0.02%); split: -0.02%, +0.00%

No fossil-db changes on any other platform.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27552>
 2024-03-12 21:31:30 +00:00
Ian Romanick
d9674cbe7d intel/brw: Combine constants for src0 of POW instructions too 
I tried this when I was working on MR !7698, and it didn't have much
affect back then. Maybe I've added more stuff to my fossil-db?

Gfx12 platforms (Tiger Lake and DG2) are unaffected because the POW
instruction was removed.

shader-db:

Ice Lake and Skylake had similar results. (Ice Lake shown)
total instructions in shared programs: 20301933 -> 20301900 (<.01%)
instructions in affected programs: 9077 -> 9044 (-0.36%)
helped: 33 / HURT: 0

total cycles in shared programs: 842797624 -> 842799471 (<.01%)
cycles in affected programs: 1361911 -> 1363758 (0.14%)
helped: 35 / HURT: 111

LOST:   0
GAINED: 9

fossil-db:

Ice Lake and Skylake had similar results. (Ice Lake shown)
Totals:
Instrs: 165510222 -> 165510163 (-0.00%)
Cycles: 15125195835 -> 15125194484 (-0.00%); split: -0.00%, +0.00%
Spill count: 45204 -> 45196 (-0.02%)
Fill count: 74157 -> 74149 (-0.01%)

Totals from 65 (0.01% of 656118) affected shaders:
Instrs: 57426 -> 57367 (-0.10%)
Cycles: 1667918 -> 1666567 (-0.08%); split: -0.11%, +0.03%
Spill count: 137 -> 129 (-5.84%)
Fill count: 515 -> 507 (-1.55%)

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27552>
 2024-03-12 21:31:30 +00:00
Ian Romanick
e7480f94c1 intel/brw: Combine constants for src0 of integer multiply too 
The majority of cases that would have been affected by this actually
had both sources as integer constants. The earlier commit "intel/rt:
Don't directly generate umul_32x16" allowed those to be constant
folded.

v2: Move the a*-1 block to be near the existing a*-1 block.

No shader-db changes on any Intel platform.

fossil-db results:

All Intel platforms had similar results. (Ice Lake shown)
Totals:
Instrs: 165510246 -> 165510222 (-0.00%)
Cycles: 15125198238 -> 15125195835 (-0.00%); split: -0.00%, +0.00%

Totals from 46 (0.01% of 656118) affected shaders:
Instrs: 36010 -> 35986 (-0.07%)
Cycles: 2613658 -> 2611255 (-0.09%); split: -0.17%, +0.07%

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27552>
 2024-03-12 21:31:30 +00:00
Ian Romanick
dd3bed1d92 intel/brw: Integer multiply w/ DW and W sources is not commutative 
The DW source must be first on all platforms since Gfx7. On previous
platforms it's the other way around.

Unsurprisingly, no shader-db or fossil-db changes. This change is
necessary for the next commit.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27552>
 2024-03-12 21:31:30 +00:00
Ian Romanick
93478c095e intel/compiler: Enforce 64-bit RepCtrl restriction in eu_validate 
For some reason, this wasn't always caught in fs_visitor::validate.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27552>
 2024-03-12 21:31:30 +00:00
Ian Romanick
31f640bc5f intel/brw: Correctly dump subnr for FIXED_GRF in INTEL_DEBUG=optimizer 
v2: Also update printing FIXED_GRF as destionation. Suggested by Lionel.

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27552>
 2024-03-12 21:31:30 +00:00
Ian Romanick
f89d9cc53d intel/brw: Silence "statement may fall through" warning 
src/intel/compiler/brw_lower_logical_sends.cpp: In member function ‘bool fs_visitor::lower_logical_sends()’:
src/intel/compiler/brw_lower_logical_sends.cpp:3170:10: warning: this statement may fall through [-Wimplicit-fallthrough=]
 3170 |          if (devinfo->has_lsc) {
      |          ^~
src/intel/compiler/brw_lower_logical_sends.cpp:3174:7: note: here
 3174 |       case SHADER_OPCODE_DWORD_SCATTERED_READ_LOGICAL:
      |       ^~~~

Reviewed-by: Lionel Landwerlin <lionel.g.landwerlin@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27552>
 2024-03-12 21:31:30 +00:00
Mike Blumenkrantz
c85768e380 zink: delete faked_e5sparse 
this never actually worked, and I'm not sure why I added it from testing

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28115>
 2024-03-12 20:39:56 +00:00
Mike Blumenkrantz
8fa413fef0 zink: set the sparse format usage flags directly based on queried props 
this should yield more consistent results and avoid weird cases where
various formats are queried for things they don't support and won't use

Fixes: 9a412c10b7 ("zink: set all usage flags when querying sparse features")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28115>
 2024-03-12 20:39:56 +00:00
Mike Blumenkrantz
ba2731941f zink: nvk baseline updates 
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28115>
 2024-03-12 20:39:56 +00:00
Mike Blumenkrantz
c248411907 zink: only add STORAGE bit for sparse images based on multisample usage 
if the driver doesn't support multisample storage images, don't create them

Fixes: d320e8328d ("zink: set up image create bits for sparse textures")
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28115>
 2024-03-12 20:39:56 +00:00
Mike Blumenkrantz
63f5a6cbe3 u/inlines: constify util_res_sample_count() 
no functional changes

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28115>
 2024-03-12 20:39:56 +00:00
Mike Blumenkrantz
61e5b6ad9d zink: try getting sparse page size again without storage bit on fail 
only certain formats are required to have the storage bit, so be more
tolerant of failure in the case where drivers actually check flags
and reject storage usage when it's actually unsupported

cc: mesa-stable

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28115>
 2024-03-12 20:39:56 +00:00
Mary Guillemard
80eac1337d nvk: Always copy conditional rendering value before compare 
The spec requires a compare on 32-bit but the hardware actually compare 64-bit.

As such, we are required to copy the value to a temporary buffer before
the compare.

Signed-off-by: Mary Guillemard <mary.guillemard@collabora.com>
Fixes: 8c25cd307a ("nvk: EXT_conditional_rendering")
Reviewed-by: M Henning <drawoc@darkrefraction.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28106>
 2024-03-12 20:22:27 +00:00
Alyssa Rosenzweig
8a3e85776a glsl_types: derive ht 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28056>
 2024-03-12 19:41:50 +00:00
Alyssa Rosenzweig
c29110ef99 v3d: derive ht 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28056>
 2024-03-12 19:41:49 +00:00
Alyssa Rosenzweig
2d32239b7a radeonsi: derive ht 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28056>
 2024-03-12 19:41:49 +00:00
Alyssa Rosenzweig
3f287d100c nvk: derive ht 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28056>
 2024-03-12 19:41:49 +00:00
Alyssa Rosenzweig
da59672e2b asahi: derive ht 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28056>
 2024-03-12 19:41:49 +00:00
Alyssa Rosenzweig
be5c288dbf panfrost: derive ht 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28056>
 2024-03-12 19:41:49 +00:00
Alyssa Rosenzweig
93879b1920 util/hash_table: add DERIVE macro 
we typically use a hash table with a fixed struct key, but this requires tedious
boilerplate. add a macro that generates all the boilerplate for you so you can
just create a table and go.

naming inspired by Rust #![derive].

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Reviewed-by: Boris Brezillon <boris.brezillon@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28056>
 2024-03-12 19:41:49 +00:00
Alyssa Rosenzweig
49a89911c4 nir/print: do not print empty lists on intrinsics 
before:

  32     %0 = @load_vertex_id () ()

after:

   32     %0 = @load_vertex_id

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Konstantin Seurer <konstantin.seurer@gmail.com>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Caio Oliveira <caio.oliveira@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27925>
 2024-03-12 19:00:26 +00:00
Alyssa Rosenzweig
ec9b6b5cfa util: add _mesa_hash_table_u64_num_entries 
Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Marek Olšák <marek.olsak@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27925>
 2024-03-12 19:00:26 +00:00
Alyssa Rosenzweig
a6123a80da nir/opt_shrink_vectors: shrink some intrinsics from start 
If the backend supports it, intrinsics with a component() are straightforward to
shrink from the start. Notably helps vectorized I/O.

v2: add an option for this and enable only on grown up backends, because some
backends ignore the component() parameter.

RADV GFX11:
Totals from 921 (1.16% of 79439) affected shaders:
Instrs: 616558 -> 615529 (-0.17%); split: -0.30%, +0.14%
CodeSize: 3099864 -> 3095632 (-0.14%); split: -0.25%, +0.11%
Latency: 2177075 -> 2160966 (-0.74%); split: -0.79%, +0.05%
InvThroughput: 299997 -> 298664 (-0.44%); split: -0.47%, +0.02%
VClause: 16343 -> 16395 (+0.32%); split: -0.01%, +0.32%
SClause: 10715 -> 10714 (-0.01%)
Copies: 24736 -> 24701 (-0.14%); split: -0.37%, +0.23%
PreVGPRs: 30179 -> 30173 (-0.02%)
VALU: 353472 -> 353439 (-0.01%); split: -0.03%, +0.02%
SALU: 40323 -> 40322 (-0.00%)
VMEM: 25353 -> 25352 (-0.00%)

AGX:

total instructions in shared programs: 2038217 -> 2038049 (<.01%)
instructions in affected programs: 10249 -> 10081 (-1.64%)

total alu in shared programs: 1593094 -> 1592939 (<.01%)
alu in affected programs: 7145 -> 6990 (-2.17%)

total fscib in shared programs: 1589254 -> 1589102 (<.01%)
fscib in affected programs: 7217 -> 7065 (-2.11%)

total bytes in shared programs: 13975666 -> 13974722 (<.01%)
bytes in affected programs: 65942 -> 64998 (-1.43%)

total regs in shared programs: 592758 -> 591187 (-0.27%)
regs in affected programs: 6936 -> 5365 (-22.65%)

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev> (v1)
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28004>
 2024-03-12 18:17:17 +00:00
Alyssa Rosenzweig
aa99753a28 nir/opt_shrink_vectors: hoist alu helpers 
to be used earlier in the file in the next commit

Signed-off-by: Alyssa Rosenzweig <alyssa@rosenzweig.io>
Reviewed-by: Timur Kristóf <timur.kristof@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28004>
 2024-03-12 18:17:17 +00:00
José Roberto de Souza
d1916432ab intel/dev: Nuke display_ver 
It is not used.

Signed-off-by: José Roberto de Souza <jose.souza@intel.com>
Reviewed-by: Jordan Justen <jordan.l.justen@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28128>
 2024-03-12 17:44:46 +00:00
José Roberto de Souza
b09ffe48f2 intel/dev: Nuke 'ver == 10' check 
There is no intel_device_info with ver 10 anymore.

Signed-off-by: José Roberto de Souza <jose.souza@intel.com>
Reviewed-by: Jordan Justen <jordan.l.justen@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/28128>
 2024-03-12 17:44:46 +00:00
Rhys Perry
a977a51a21 radv: stop using 5/8 component SSBO stores 
These apparently work, but I'm not sure they were supposed to.

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/28108>
 2024-03-12 17:23:29 +00:00
Rhys Perry
cc7e3efc7c radv: don't advertise DGC with LLVM 
The meta shaders for this feature don't compile with LLVM because of 5/8
component SSBO stores. I'm not sure this was ever expected to work.

This seemed to break vkd3d-proton.

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/28108>
 2024-03-12 17:23:29 +00:00
Christian Gmeiner
3409c60099 etnaviv: isa: Support multiple encodings for texldb 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
20678b9204 etnaviv: isa: Fix #instruction-tex-src0-src1-src2 bitset 
src1 got wrong values.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
ccc99bd42d etnaviv: isa: Support multiple encodings for texldl 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
8d117b46ea etnaviv: isa: Support unary texkill instruction 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
86de104d07 etnaviv: isa: Support unary branch instruction 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
0aa737018e etnaviv: isa: Combine branch and branch_if 
As we want to use the new asm as a drop-in replacement we
need to combine branch and branch_if back to one bitset.

This is caused by the fact that we need to replicate the defines
in isa.xml.h.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
48e1589b44 etnaviv: isa: Correct #instruction-alu-no-dst-has-src0-src1 expr name 
This expression only checks if src0 and src1 are in use.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
10a7cf3121 etnaviv: isa: Correct #instruction-alu-no-dst-maybe-src1-src2 name 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
dfb2fcf652 etnaviv: isa: Correct #instruction-cf-src1-src2 bitset name 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
c2ffc7a09b etnaviv: isa: Correct SRC0_AMODE 
It is 3 bit long and not one.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
b1cbd35bb5 etnaviv: isa: Move {TEX_SWIZ} 
Should have never been there.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
d8f6de7314 etnaviv: isa: Add movar opcode 
I was unable to grab this opcode from blob, so lets just
document it as the Gallium driver makes use of it.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
e77fbe2bcc etnaviv: isa: Add internal register group 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
f416bb3f8c etnaviv: isa: Rename reg_group u2 to u 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
cf3fa2fd8c etnaviv: isa: Reorder instructions 
Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
9c6378abec etnaviv: isa: Add div opcode 
Encoded instruction is taken from blob running the following CL kernel:

 __kernel void simple(__global float *out, __global float *in)
{
    int iGID = get_global_id(0);
    out[iGID] = 4.5f / in[iGID];
}

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
e2a9bc73f5 etnaviv: isa: Remove note about GC3000 
All the encoded instructions in the Opcodes test are comming
from blob running on different GPU models.

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
fa3d2bc486 etnaviv: isa: Add texldd opcode 
Encoded instruction is taken from blob running:
- dEQP-GLES3.functional.shaders.texture_functions.texturegrad.sampler2d_float_vertex

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00
Christian Gmeiner
0701f3ef9b etnaviv: isa: Add texldl opcode 
Encoded instruction is taken from blob running:
- dEQP-GLES3.functional.shaders.texture_functions.texturegrad.isampler2d_vertex

Signed-off-by: Christian Gmeiner <cgmeiner@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/27871>
 2024-03-12 17:02:44 +00:00