intel/compiler: Drop variable group size lowering

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 <marcin.slusarz@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/18055>

diff --git a/src/intel/compiler/brw_compiler.h b/src/intel/compiler/brw_compiler.h
index 95ac367..42385df 100644 (file)
--- a/src/intel/compiler/brw_compiler.h
+++ b/src/intel/compiler/brw_compiler.h
@@ -106,12 +106,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
     * up with VK_FORMAT_R32G32B32A32_FLOAT whereas if it's going through the

diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp
index 0caa741..9372fd1 100644 (file)
--- 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 ac5aba3..a234a80 100644 (file)
--- 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 d248a83..67cd5eb 100644 (file)
--- 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;
    }