URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=b065d8fb8cf55373bfdd80994417f1ac60976158 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Nov 27 16:26:03 2019 -0800
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.olive...@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> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=c57338b924710b93193f921cd1e95d6de6b398ef Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Jan 2 16:07:59 2019 -0800 anv: Enable SPV_INTEL_shader_integer_functions2 and VK_INTEL_shader_integer_functions2 Currently only implemented in the scalar backend, so only enable for Gen8+. If support for the other opcodes is added to the vec4 backend, Gen7 could be supported. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=76970940a6b22d25fcff42ab6c779dc646b2d9d4 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Thu Nov 14 14:20:48 2019 -0800 iris: Enable INTEL_shader_integer_functions2 Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Reviewed-by: Kenneth Graunke <kenn...@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=b14e718e68019019ea241d7e7a7f1bbdb8fcf21e Author: Ian Romanick <ian.d.roman...@intel.com> Date: Thu Nov 14 14:16:26 2019 -0800 gallium: Add a cap bit for integer multiplication between 32-bit and 16-bit Driver supports integer multiplication between a 32-bit integer and a 16-bit integer. If the second operand is 32-bits, the upper 16-bits are ignored, and the low 16-bits are possibly sign extended as necessary. Iris will eventually enable this. Not sure about other drivers. v2: Add default value to u_screen.c. Suggested by Caio. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Reviewed-by: Kenneth Graunke <kenn...@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=9db20748fd1af930920424a95321ee11b6eae16e Author: Ian Romanick <ian.d.roman...@intel.com> Date: Thu Nov 14 14:12:30 2019 -0800 gallium: Add a cap bit for OpenCL-style extended integer functions Iris will eventually enable this. Looking at the header files, it looks like Midgard could also enable it. Basically, any GPU that fully supports OpenCL can. v2: Add default value to u_screen.c. Suggested by Caio. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Reviewed-by: Kenneth Graunke <kenn...@whitecape.org> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=4e9079d0c71e42e152a00678bbe2665882849a43 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Tue Sep 11 16:50:06 2018 -0700 i965: Enable INTEL_shader_integer_functions2 on Gen8+ v2: Use new lower_hadd64 and lower_usub_sat64 flags. v3: Enable SPIR-V capability. v4: Move lowering options to COMMON_SCALAR_OPTIONS. Suggested by Caio. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=4fcddb55f27e29d78c6937c20d91e7f9962ce875 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Mon Sep 24 06:46:48 2018 -0700 spirv: Add support for IntegerFunctions2INTEL capability Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=aa56934e2ae75b31fbc22a5e03f95628c38f8d84 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Mon Sep 24 06:44:38 2018 -0700 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.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=44471a76e9b32410d02c202c67ce48a3b69770a8 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Fri Sep 21 00:35:18 2018 -0700 nir/spirv: Translate SPIR-V to NIR for new INTEL_shader_integer_functions2 opcodes v2: Rebase on 272e927d0e9 ("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.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=de6c0f848797d26d28223bcad25da03654461874 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Tue Sep 11 16:49:51 2018 -0700 intel/fs: Implement support for NIR opcodes for INTEL_shader_integer_functions2 v2: Remove smashing type to D for nir_op_irhadd. Caio noticed it was odd, and removing it fixes an assertion failure in the crucible func.shader.averageRounded.int64_t test (because the source should be W). v3: Emit BRW_OPCODE_MUL directly for nir_op_umul_32x16 and nir_op_imul_32x16. Suggested by Curro. v4: Smash types of MUL instruction generated for nir_op_umul_32x16 and nir_op_imul_32x16. With this change, I get the same assembly now as I did with v2. v5: Remove support for pre-Gen7. The integer multiply path was incorrect, and, since the extension isn't enabled pre-Gen7, there's no way to test it. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=58907568ec526df87fa87177441743fa0d1d0a66 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Sep 19 01:28:06 2018 -0700 intel/fs: Add SHADER_OPCODE_[IU]SUB_SAT pseudo-ops v2: Add a big comment explaining the [IU]SUB_SAT lowering. Suggested by Caio. v3: Use get_fpu_lowered_simd_width in get_lowered_simd_width. Suggested by Ken on IRC. v4: Fix a typo in a comment. Noticed by Caio. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=74cd0964d684d7b91207b442eb3237fee1e4b4e0 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Thu Dec 5 08:25:34 2019 -0800 intel/fs: Don't lower integer multiplies that don't need lowering v2: Move the check to fs_visitor::lower_integer_multiplication. Previously the cases where lowering was skipped, the original instruction was removed by fs_visitor::lower_integer_multiplication. Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=db649fd5822569e1f33b058c0f4d38c27a47b083 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Sep 12 17:05:14 2018 -0700 compiler: Translate GLSL IR to NIR for new INTEL_shader_integer_functions2 expressions v2: Rebase on 272e927d0e9 ("nir/spirv: initial handling of OpenCL.std extension opcodes") Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=d3d970166cad7d834811fd1f8bcd6c53db706bce Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Nov 27 13:22:38 2019 -0800 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.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=dcadbd2dd222ecaa4f14bdadf90587bda83deba7 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Nov 27 12:55:27 2019 -0800 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.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=1bdfc6d7cb4c897ae6fe826d7f778574c8ca7551 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Sep 19 01:17:31 2018 -0700 nir/algebraic: Add lowering for 64-bit usub_sat v2: Rebase on 272e927d0e9 ("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.olive...@intel.com> [v3] Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=a483771045f49cb549e098394ceea309b5312c65 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Wed Sep 19 01:17:09 2018 -0700 nir/algebraic: Add lowering for 64-bit hadd and rhadd v2: Rebase on 272e927d0e9 ("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.olive...@intel.com> [v3] Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=ea435560ee26b2669d923c8af6077cd0c7ac0ff4 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Tue Sep 11 15:38:36 2018 -0700 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.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=21f0d020fe2518da008c86cb111579e7f0e636b5 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Tue Sep 11 00:13:36 2018 -0700 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.olive...@intel.com> [v1] Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=cb518df775f222e34a4e724be2fa825df9388b0a Author: Ian Romanick <ian.d.roman...@intel.com> Date: Mon Sep 10 23:17:49 2018 -0700 glsl: Add built-in functions for INTEL_shader_integer_functions2 Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=5eda9f5832105bea0b3f75bcc866c760dd801718 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Mon Sep 17 08:53:24 2018 -0700 glsl_types: Add function to get an unsigned base type from a signed type Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=1d165b0548639df48c406fd6514298309e68aba9 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Mon Sep 10 22:38:29 2018 -0700 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.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=20d34c4ebf07c98a40ea43b0cccc95537c176fa9 Author: Ian Romanick <ian.d.roman...@intel.com> Date: Mon Sep 10 17:54:56 2018 -0700 mesa: Extension boilerplate for INTEL_shader_integer_functions2 Reviewed-by: Caio Marcelo de Oliveira Filho <caio.olive...@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/merge_requests/767> _______________________________________________ mesa-commit mailing list mesa-commit@lists.freedesktop.org https://lists.freedesktop.org/mailman/listinfo/mesa-commit