From 7662a5e9d34515bd44a97b3726490f31490b57c6 Mon Sep 17 00:00:00 2001 From: Emma Anholt Date: Wed, 31 Aug 2022 22:37:29 -0700 Subject: [PATCH] mesa: Remove PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED/lower_cs_derived. We have fine NIR lowering for this (already called from mesa/st), no need for a separate GLSL pass. Reviewed-by: Timothy Arceri Reviewed-by: Alyssa Rosenzweig Part-of: --- docs/gallium/screen.rst | 3 - src/compiler/glsl/ir_optimization.h | 1 - src/compiler/glsl/linker.cpp | 3 - src/compiler/glsl/lower_cs_derived.cpp | 235 ------------------ src/compiler/glsl/meson.build | 1 - src/gallium/auxiliary/nir/nir_to_tgsi.c | 5 + src/gallium/auxiliary/util/u_screen.c | 1 - src/gallium/drivers/asahi/agx_pipe.c | 1 - src/gallium/drivers/crocus/crocus_screen.c | 1 - src/gallium/drivers/iris/iris_screen.c | 1 - .../drivers/nouveau/nv50/nv50_screen.c | 1 - .../drivers/nouveau/nvc0/nvc0_screen.c | 1 - src/gallium/drivers/panfrost/pan_screen.c | 1 - src/gallium/drivers/softpipe/sp_screen.c | 5 + src/gallium/drivers/svga/svga_screen.c | 1 + src/gallium/include/pipe/p_defines.h | 1 - src/mesa/main/consts_exts.h | 3 - src/mesa/state_tracker/st_extensions.c | 2 - 18 files changed, 11 insertions(+), 256 deletions(-) delete mode 100644 src/compiler/glsl/lower_cs_derived.cpp diff --git a/docs/gallium/screen.rst b/docs/gallium/screen.rst index 037f63c0381..2b6bf438d74 100644 --- a/docs/gallium/screen.rst +++ b/docs/gallium/screen.rst @@ -561,9 +561,6 @@ The integer capabilities: OpenMAX should use a compute-based blit instead of pipe_context::blit and compute pipeline for compositing images. * ``PIPE_CAP_FRAGMENT_SHADER_INTERLOCK``: True if fragment shader interlock functionality is supported. -* ``PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED``: True if driver handles - gl_LocalInvocationIndex and gl_GlobalInvocationID. Otherwise, gallium frontends will - lower those system values. * ``PIPE_CAP_ATOMIC_FLOAT_MINMAX``: Atomic float point minimum, maximum, exchange and compare-and-swap support to buffer and shared variables. * ``PIPE_CAP_TGSI_DIV``: Whether opcode DIV is supported diff --git a/src/compiler/glsl/ir_optimization.h b/src/compiler/glsl/ir_optimization.h index 1ac7d8bcaeb..6571ce8c2a1 100644 --- a/src/compiler/glsl/ir_optimization.h +++ b/src/compiler/glsl/ir_optimization.h @@ -122,7 +122,6 @@ void optimize_dead_builtin_variables(exec_list *instructions, enum ir_variable_mode other); bool lower_tess_level(gl_linked_shader *shader); -bool lower_cs_derived(gl_linked_shader *shader); bool lower_blend_equation_advanced(gl_linked_shader *shader, bool coherent); bool lower_builtins(exec_list *instructions); diff --git a/src/compiler/glsl/linker.cpp b/src/compiler/glsl/linker.cpp index 19772918ef3..d61fe0136c2 100644 --- a/src/compiler/glsl/linker.cpp +++ b/src/compiler/glsl/linker.cpp @@ -2544,9 +2544,6 @@ link_intrastage_shaders(void *mem_ctx, } } - if (ctx->Const.LowerCsDerivedVariables) - lower_cs_derived(linked); - /* Set the linked source SHA1. */ if (num_shaders == 1) { memcpy(linked->linked_source_sha1, shader_list[0]->compiled_source_sha1, diff --git a/src/compiler/glsl/lower_cs_derived.cpp b/src/compiler/glsl/lower_cs_derived.cpp deleted file mode 100644 index 9fddd61f033..00000000000 --- a/src/compiler/glsl/lower_cs_derived.cpp +++ /dev/null @@ -1,235 +0,0 @@ -/* - * Copyright © 2017 Ilia Mirkin - * - * Permission is hereby granted, free of charge, to any person obtaining a - * copy of this software and associated documentation files (the "Software"), - * to deal in the Software without restriction, including without limitation - * the rights to use, copy, modify, merge, publish, distribute, sublicense, - * and/or sell copies of the Software, and to permit persons to whom the - * Software is furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice (including the next - * paragraph) shall be included in all copies or substantial portions of the - * Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL - * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING - * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER - * DEALINGS IN THE SOFTWARE. - */ - -/** - * \file lower_cs_derived.cpp - * - * For hardware that does not support the gl_GlobalInvocationID and - * gl_LocalInvocationIndex system values, replace them with fresh - * globals. Note that we can't rely on gl_WorkGroupSize or - * gl_LocalGroupSizeARB being available, since they may only have been defined - * in a non-main shader. - * - * [ This can happen if only a secondary shader has the layout(local_size_*) - * declaration. ] - * - * This is meant to be run post-linking. - */ - -#include "glsl_symbol_table.h" -#include "ir_hierarchical_visitor.h" -#include "ir.h" -#include "ir_builder.h" -#include "linker.h" -#include "program/prog_statevars.h" -#include "builtin_functions.h" -#include "main/shader_types.h" - -using namespace ir_builder; - -namespace { - -class lower_cs_derived_visitor : public ir_hierarchical_visitor { -public: - explicit lower_cs_derived_visitor(gl_linked_shader *shader) - : progress(false), - shader(shader), - local_size_variable(shader->Program->info.workgroup_size_variable), - gl_WorkGroupSize(NULL), - gl_WorkGroupID(NULL), - gl_LocalInvocationID(NULL), - gl_GlobalInvocationID(NULL), - gl_LocalInvocationIndex(NULL) - { - main_sig = _mesa_get_main_function_signature(shader->symbols); - assert(main_sig); - } - - virtual ir_visitor_status visit(ir_dereference_variable *); - - ir_variable *add_system_value( - int slot, const glsl_type *type, const char *name); - void find_sysvals(); - void make_gl_GlobalInvocationID(); - void make_gl_LocalInvocationIndex(); - - bool progress; - -private: - gl_linked_shader *shader; - bool local_size_variable; - ir_function_signature *main_sig; - - ir_rvalue *gl_WorkGroupSize; - ir_variable *gl_WorkGroupID; - ir_variable *gl_LocalInvocationID; - - ir_variable *gl_GlobalInvocationID; - ir_variable *gl_LocalInvocationIndex; -}; - -} /* anonymous namespace */ - -ir_variable * -lower_cs_derived_visitor::add_system_value( - int slot, const glsl_type *type, const char *name) -{ - ir_variable *var = new(shader) ir_variable(type, name, ir_var_system_value); - var->data.how_declared = ir_var_declared_implicitly; - var->data.read_only = true; - var->data.location = slot; - var->data.explicit_location = true; - var->data.explicit_index = 0; - shader->ir->push_head(var); - - return var; -} - -void -lower_cs_derived_visitor::find_sysvals() -{ - if (gl_WorkGroupSize != NULL) - return; - - ir_variable *WorkGroupSize; - if (local_size_variable) - WorkGroupSize = shader->symbols->get_variable("gl_LocalGroupSizeARB"); - else - WorkGroupSize = shader->symbols->get_variable("gl_WorkGroupSize"); - if (WorkGroupSize) - gl_WorkGroupSize = new(shader) ir_dereference_variable(WorkGroupSize); - gl_WorkGroupID = shader->symbols->get_variable("gl_WorkGroupID"); - gl_LocalInvocationID = shader->symbols->get_variable("gl_LocalInvocationID"); - - /* - * These may be missing due to either dead code elimination, or, in the - * case of the group size, due to the layout being declared in a non-main - * shader. Re-create them. - */ - - if (!gl_WorkGroupID) - gl_WorkGroupID = add_system_value( - SYSTEM_VALUE_WORKGROUP_ID, glsl_type::uvec3_type, "gl_WorkGroupID"); - if (!gl_LocalInvocationID) - gl_LocalInvocationID = add_system_value( - SYSTEM_VALUE_LOCAL_INVOCATION_ID, glsl_type::uvec3_type, - "gl_LocalInvocationID"); - if (!WorkGroupSize) { - if (local_size_variable) { - gl_WorkGroupSize = new(shader) ir_dereference_variable( - add_system_value( - SYSTEM_VALUE_WORKGROUP_SIZE, glsl_type::uvec3_type, - "gl_LocalGroupSizeARB")); - } else { - ir_constant_data data; - memset(&data, 0, sizeof(data)); - for (int i = 0; i < 3; i++) - data.u[i] = shader->Program->info.workgroup_size[i]; - gl_WorkGroupSize = new(shader) ir_constant(glsl_type::uvec3_type, &data); - } - } -} - -void -lower_cs_derived_visitor::make_gl_GlobalInvocationID() -{ - if (gl_GlobalInvocationID != NULL) - return; - - find_sysvals(); - - /* gl_GlobalInvocationID = - * gl_WorkGroupID * gl_WorkGroupSize + gl_LocalInvocationID - */ - gl_GlobalInvocationID = new(shader) ir_variable( - glsl_type::uvec3_type, "__GlobalInvocationID", ir_var_temporary); - shader->ir->push_head(gl_GlobalInvocationID); - - ir_instruction *inst = - assign(gl_GlobalInvocationID, - add(mul(gl_WorkGroupID, gl_WorkGroupSize->clone(shader, NULL)), - gl_LocalInvocationID)); - main_sig->body.push_head(inst); -} - -void -lower_cs_derived_visitor::make_gl_LocalInvocationIndex() -{ - if (gl_LocalInvocationIndex != NULL) - return; - - find_sysvals(); - - /* gl_LocalInvocationIndex = - * gl_LocalInvocationID.z * gl_WorkGroupSize.x * gl_WorkGroupSize.y + - * gl_LocalInvocationID.y * gl_WorkGroupSize.x + - * gl_LocalInvocationID.x; - */ - gl_LocalInvocationIndex = new(shader) - ir_variable(glsl_type::uint_type, "__LocalInvocationIndex", ir_var_temporary); - shader->ir->push_head(gl_LocalInvocationIndex); - - ir_expression *index_z = - mul(mul(swizzle_z(gl_LocalInvocationID), swizzle_x(gl_WorkGroupSize->clone(shader, NULL))), - swizzle_y(gl_WorkGroupSize->clone(shader, NULL))); - ir_expression *index_y = - mul(swizzle_y(gl_LocalInvocationID), swizzle_x(gl_WorkGroupSize->clone(shader, NULL))); - ir_expression *index_y_plus_z = add(index_y, index_z); - operand index_x(swizzle_x(gl_LocalInvocationID)); - ir_expression *index_x_plus_y_plus_z = add(index_y_plus_z, index_x); - ir_instruction *inst = - assign(gl_LocalInvocationIndex, index_x_plus_y_plus_z); - main_sig->body.push_head(inst); -} - -ir_visitor_status -lower_cs_derived_visitor::visit(ir_dereference_variable *ir) -{ - if (ir->var->data.mode == ir_var_system_value && - ir->var->data.location == SYSTEM_VALUE_GLOBAL_INVOCATION_ID) { - make_gl_GlobalInvocationID(); - ir->var = gl_GlobalInvocationID; - progress = true; - } - - if (ir->var->data.mode == ir_var_system_value && - ir->var->data.location == SYSTEM_VALUE_LOCAL_INVOCATION_INDEX) { - make_gl_LocalInvocationIndex(); - ir->var = gl_LocalInvocationIndex; - progress = true; - } - - return visit_continue; -} - -bool -lower_cs_derived(gl_linked_shader *shader) -{ - if (shader->Stage != MESA_SHADER_COMPUTE) - return false; - - lower_cs_derived_visitor v(shader); - v.run(shader->ir); - - return v.progress; -} diff --git a/src/compiler/glsl/meson.build b/src/compiler/glsl/meson.build index 55cba69384e..5b0fe725245 100644 --- a/src/compiler/glsl/meson.build +++ b/src/compiler/glsl/meson.build @@ -156,7 +156,6 @@ files_libglsl = files( 'list.h', 'lower_blend_equation_advanced.cpp', 'lower_builtins.cpp', - 'lower_cs_derived.cpp', 'lower_discard.cpp', 'lower_discard_flow.cpp', 'lower_distance.cpp', diff --git a/src/gallium/auxiliary/nir/nir_to_tgsi.c b/src/gallium/auxiliary/nir/nir_to_tgsi.c index dc884bdda03..48a05f0d42f 100644 --- a/src/gallium/auxiliary/nir/nir_to_tgsi.c +++ b/src/gallium/auxiliary/nir/nir_to_tgsi.c @@ -4035,6 +4035,11 @@ static const nir_shader_compiler_options nir_to_tgsi_compiler_options = { .lower_vector_cmp = true, .lower_int64_options = nir_lower_imul_2x32_64, .use_interpolated_input_intrinsics = true, + + /* TGSI doesn't have a semantic for local or global index, just local and + * workgroup id. + */ + .lower_cs_local_index_to_id = true, }; /* Returns a default compiler options for drivers with only nir-to-tgsi-based diff --git a/src/gallium/auxiliary/util/u_screen.c b/src/gallium/auxiliary/util/u_screen.c index 391ec743984..914ffc45833 100644 --- a/src/gallium/auxiliary/util/u_screen.c +++ b/src/gallium/auxiliary/util/u_screen.c @@ -370,7 +370,6 @@ u_pipe_screen_get_param_defaults(struct pipe_screen *pscreen, case PIPE_CAP_IMAGE_STORE_FORMATTED: case PIPE_CAP_PREFER_COMPUTE_FOR_MULTIMEDIA: case PIPE_CAP_FRAGMENT_SHADER_INTERLOCK: - case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED: case PIPE_CAP_ATOMIC_FLOAT_MINMAX: case PIPE_CAP_SHADER_SAMPLES_IDENTICAL: case PIPE_CAP_IMAGE_ATOMIC_INC_WRAP: diff --git a/src/gallium/drivers/asahi/agx_pipe.c b/src/gallium/drivers/asahi/agx_pipe.c index 6cd80506755..968d58dd056 100644 --- a/src/gallium/drivers/asahi/agx_pipe.c +++ b/src/gallium/drivers/asahi/agx_pipe.c @@ -756,7 +756,6 @@ agx_get_param(struct pipe_screen* pscreen, enum pipe_cap param) case PIPE_CAP_TEXTURE_FLOAT_LINEAR: case PIPE_CAP_TEXTURE_HALF_FLOAT_LINEAR: case PIPE_CAP_SHADER_ARRAY_COMPONENTS: - case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED: case PIPE_CAP_PACKED_UNIFORMS: return 1; diff --git a/src/gallium/drivers/crocus/crocus_screen.c b/src/gallium/drivers/crocus/crocus_screen.c index 64e9095b80a..ab804d74dbc 100644 --- a/src/gallium/drivers/crocus/crocus_screen.c +++ b/src/gallium/drivers/crocus/crocus_screen.c @@ -192,7 +192,6 @@ crocus_get_param(struct pipe_screen *pscreen, enum pipe_cap param) case PIPE_CAP_FS_FACE_IS_INTEGER_SYSVAL: case PIPE_CAP_INVALIDATE_BUFFER: case PIPE_CAP_SURFACE_REINTERPRET_BLOCKS: - case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED: case PIPE_CAP_FENCE_SIGNAL: case PIPE_CAP_DEMOTE_TO_HELPER_INVOCATION: case PIPE_CAP_GL_CLAMP: diff --git a/src/gallium/drivers/iris/iris_screen.c b/src/gallium/drivers/iris/iris_screen.c index bc4878e9168..177567b3871 100644 --- a/src/gallium/drivers/iris/iris_screen.c +++ b/src/gallium/drivers/iris/iris_screen.c @@ -277,7 +277,6 @@ iris_get_param(struct pipe_screen *pscreen, enum pipe_cap param) case PIPE_CAP_COMPUTE_SHADER_DERIVATIVES: case PIPE_CAP_INVALIDATE_BUFFER: case PIPE_CAP_SURFACE_REINTERPRET_BLOCKS: - case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED: case PIPE_CAP_TEXTURE_SHADOW_LOD: case PIPE_CAP_SHADER_SAMPLES_IDENTICAL: case PIPE_CAP_GL_SPIRV: diff --git a/src/gallium/drivers/nouveau/nv50/nv50_screen.c b/src/gallium/drivers/nouveau/nv50/nv50_screen.c index db329c40e9a..2ac0c07886b 100644 --- a/src/gallium/drivers/nouveau/nv50/nv50_screen.c +++ b/src/gallium/drivers/nouveau/nv50/nv50_screen.c @@ -381,7 +381,6 @@ nv50_screen_get_param(struct pipe_screen *pscreen, enum pipe_cap param) case PIPE_CAP_ATOMIC_FLOAT_MINMAX: case PIPE_CAP_CONSERVATIVE_RASTER_INNER_COVERAGE: case PIPE_CAP_FRAGMENT_SHADER_INTERLOCK: - case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED: case PIPE_CAP_FBFETCH_COHERENT: case PIPE_CAP_IMAGE_ATOMIC_INC_WRAP: case PIPE_CAP_DEMOTE_TO_HELPER_INVOCATION: diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c b/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c index 33ce857c903..e8805aaa4df 100644 --- a/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c +++ b/src/gallium/drivers/nouveau/nvc0/nvc0_screen.c @@ -409,7 +409,6 @@ nvc0_screen_get_param(struct pipe_screen *pscreen, enum pipe_cap param) case PIPE_CAP_ATOMIC_FLOAT_MINMAX: case PIPE_CAP_CONSERVATIVE_RASTER_INNER_COVERAGE: case PIPE_CAP_FRAGMENT_SHADER_INTERLOCK: - case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED: case PIPE_CAP_FBFETCH_COHERENT: case PIPE_CAP_TGSI_TG4_COMPONENT_IN_SWIZZLE: case PIPE_CAP_OPENCL_INTEGER_FUNCTIONS: /* could be done */ diff --git a/src/gallium/drivers/panfrost/pan_screen.c b/src/gallium/drivers/panfrost/pan_screen.c index 516131640cc..f44d0fdc2ce 100644 --- a/src/gallium/drivers/panfrost/pan_screen.c +++ b/src/gallium/drivers/panfrost/pan_screen.c @@ -166,7 +166,6 @@ panfrost_get_param(struct pipe_screen *screen, enum pipe_cap param) case PIPE_CAP_TEXTURE_FLOAT_LINEAR: case PIPE_CAP_TEXTURE_HALF_FLOAT_LINEAR: case PIPE_CAP_SHADER_ARRAY_COMPONENTS: - case PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED: case PIPE_CAP_TEXTURE_BUFFER_OBJECTS: case PIPE_CAP_TEXTURE_BUFFER_SAMPLER: case PIPE_CAP_PACKED_UNIFORMS: diff --git a/src/gallium/drivers/softpipe/sp_screen.c b/src/gallium/drivers/softpipe/sp_screen.c index cc182867b12..af58c1b47a0 100644 --- a/src/gallium/drivers/softpipe/sp_screen.c +++ b/src/gallium/drivers/softpipe/sp_screen.c @@ -91,6 +91,11 @@ static const nir_shader_compiler_options sp_compiler_options = { .lower_int64_options = nir_lower_imul_2x32_64, .max_unroll_iterations = 32, .use_interpolated_input_intrinsics = true, + + /* TGSI doesn't have a semantic for local or global index, just local and + * workgroup id. + */ + .lower_cs_local_index_to_id = true, }; static const void * diff --git a/src/gallium/drivers/svga/svga_screen.c b/src/gallium/drivers/svga/svga_screen.c index ce41ef410a9..464a448276a 100644 --- a/src/gallium/drivers/svga/svga_screen.c +++ b/src/gallium/drivers/svga/svga_screen.c @@ -761,6 +761,7 @@ vgpu10_get_shader_param(struct pipe_screen *screen, .lower_rotate = true, \ .lower_uniforms_to_ubo = true, \ .lower_vector_cmp = true, \ + .lower_cs_local_index_to_id = true, \ .max_unroll_iterations = 32, \ .use_interpolated_input_intrinsics = true diff --git a/src/gallium/include/pipe/p_defines.h b/src/gallium/include/pipe/p_defines.h index ad79a91c572..f0b60e3ef35 100644 --- a/src/gallium/include/pipe/p_defines.h +++ b/src/gallium/include/pipe/p_defines.h @@ -938,7 +938,6 @@ enum pipe_cap PIPE_CAP_PREFER_COMPUTE_FOR_MULTIMEDIA, PIPE_CAP_FRAGMENT_SHADER_INTERLOCK, PIPE_CAP_FBFETCH_COHERENT, - PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED, PIPE_CAP_ATOMIC_FLOAT_MINMAX, PIPE_CAP_TGSI_DIV, PIPE_CAP_FRAGMENT_SHADER_TEXTURE_LOD, diff --git a/src/mesa/main/consts_exts.h b/src/mesa/main/consts_exts.h index 6fdb89c890e..243020b3237 100644 --- a/src/mesa/main/consts_exts.h +++ b/src/mesa/main/consts_exts.h @@ -907,9 +907,6 @@ struct gl_constants GLuint MaxTessControlTotalOutputComponents; bool LowerTessLevel; /**< Lower gl_TessLevel* from float[n] to vecn? */ bool PrimitiveRestartForPatches; - bool LowerCsDerivedVariables; /**< Lower gl_GlobalInvocationID and - * gl_LocalInvocationIndex based on - * other builtin variables. */ /** GL_OES_primitive_bounding_box */ bool NoPrimitiveBoundingBoxOutput; diff --git a/src/mesa/state_tracker/st_extensions.c b/src/mesa/state_tracker/st_extensions.c index 5a274284d5b..dc1770d1efe 100644 --- a/src/mesa/state_tracker/st_extensions.c +++ b/src/mesa/state_tracker/st_extensions.c @@ -392,8 +392,6 @@ void st_init_limits(struct pipe_screen *screen, screen->get_param(screen, PIPE_CAP_GLSL_TESS_LEVELS_AS_INPUTS); c->LowerTessLevel = !screen->get_param(screen, PIPE_CAP_NIR_COMPACT_ARRAYS); - c->LowerCsDerivedVariables = - !screen->get_param(screen, PIPE_CAP_CS_DERIVED_SYSTEM_VALUES_SUPPORTED); c->PrimitiveRestartForPatches = false; c->MaxCombinedTextureImageUnits =