From 2cea0d6ef6039c407a08fdd558de057cc8a7afaf Mon Sep 17 00:00:00 2001 From: Kenneth Graunke Date: Mon, 15 Aug 2022 00:18:59 -0700 Subject: [PATCH] intel/compiler: Drop variable group size lowering MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit This backend lowering code has been dead since the removal of i965 - nothing in the current source tree ever sets the flag. This is handled by iris_setup_uniforms() and crocus_setup_uniforms(). Variable group size does not appear to be a feature in anv. Reviewed-by: Marcin Ślusarz Part-of: --- src/intel/compiler/brw_compiler.h | 6 ------ src/intel/compiler/brw_fs.cpp | 2 -- src/intel/compiler/brw_fs.h | 1 - src/intel/compiler/brw_fs_nir.cpp | 26 +++++--------------------- 4 files changed, 5 insertions(+), 30 deletions(-) diff --git a/src/intel/compiler/brw_compiler.h b/src/intel/compiler/brw_compiler.h index 95ac367431a..42385df0d9e 100644 --- a/src/intel/compiler/brw_compiler.h +++ b/src/intel/compiler/brw_compiler.h @@ -105,12 +105,6 @@ struct brw_compiler { */ bool supports_shader_constants; - /** - * Whether or not the driver wants variable group size to be lowered by the - * back-end compiler. - */ - bool lower_variable_group_size; - /** * Whether indirect UBO loads should use the sampler or go through the * data/constant cache. For the sampler, UBO surface states have to be set diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp index 0caa7413375..9372fd1bf52 100644 --- a/src/intel/compiler/brw_fs.cpp +++ b/src/intel/compiler/brw_fs.cpp @@ -1148,8 +1148,6 @@ fs_visitor::import_uniforms(fs_visitor *v) this->push_constant_loc = v->push_constant_loc; this->uniforms = v->uniforms; this->subgroup_id = v->subgroup_id; - for (unsigned i = 0; i < ARRAY_SIZE(this->group_size); i++) - this->group_size[i] = v->group_size[i]; } void diff --git a/src/intel/compiler/brw_fs.h b/src/intel/compiler/brw_fs.h index ac5aba38ea7..a234a80c619 100644 --- a/src/intel/compiler/brw_fs.h +++ b/src/intel/compiler/brw_fs.h @@ -372,7 +372,6 @@ public: int *push_constant_loc; fs_reg subgroup_id; - fs_reg group_size[3]; fs_reg scratch_base; fs_reg frag_depth; fs_reg frag_stencil; diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp index d248a83bd50..67cd5ebef1b 100644 --- a/src/intel/compiler/brw_fs_nir.cpp +++ b/src/intel/compiler/brw_fs_nir.cpp @@ -112,21 +112,11 @@ fs_visitor::nir_setup_uniforms() /* Add uniforms for builtins after regular NIR uniforms. */ assert(uniforms == prog_data->nr_params); - uint32_t *param; - if (nir->info.workgroup_size_variable && - compiler->lower_variable_group_size) { - param = brw_stage_prog_data_add_params(prog_data, 3); - for (unsigned i = 0; i < 3; i++) { - param[i] = (BRW_PARAM_BUILTIN_WORK_GROUP_SIZE_X + i); - group_size[i] = fs_reg(UNIFORM, uniforms++, BRW_REGISTER_TYPE_UD); - } - } - /* Subgroup ID must be the last uniform on the list. This will make * easier later to split between cross thread and per thread * uniforms. */ - param = brw_stage_prog_data_add_params(prog_data, 1); + uint32_t *param = brw_stage_prog_data_add_params(prog_data, 1); *param = BRW_PARAM_BUILTIN_SUBGROUP_ID; subgroup_id = fs_reg(UNIFORM, uniforms++, BRW_REGISTER_TYPE_UD); } @@ -3976,16 +3966,10 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld, } case nir_intrinsic_load_workgroup_size: { - /* For non-variable case, this should've been lowered already. */ - assert(nir->info.workgroup_size_variable); - - assert(compiler->lower_variable_group_size); - assert(gl_shader_stage_is_compute(stage)); - - for (unsigned i = 0; i < 3; i++) { - bld.MOV(retype(offset(dest, bld, i), BRW_REGISTER_TYPE_UD), - group_size[i]); - } + /* Should have been lowered by brw_nir_lower_cs_intrinsics() or + * crocus/iris_setup_uniforms() for the variable group size case. + */ + unreachable("Should have been lowered"); break; }