intel/fs: Clean up variable group size handling in backend

Just use the information from NIR shader_info.

Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Jordan Justen <jordan.l.justen@intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4794>

diff --git a/src/intel/compiler/brw_compiler.h b/src/intel/compiler/brw_compiler.h
index ab39af2..1045ef5 100644 (file)
--- a/src/intel/compiler/brw_compiler.h
+++ b/src/intel/compiler/brw_compiler.h
@@ -917,12 +917,10 @@ struct brw_cs_prog_data {
    struct brw_stage_prog_data base;
 
    unsigned local_size[3];
-   unsigned max_variable_local_size;
    unsigned simd_size;
    unsigned slm_size;
    bool uses_barrier;
    bool uses_num_work_groups;
-   bool uses_variable_group_size;
 
    struct {
       struct brw_push_const_block cross_thread;

diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp
index ccefdb0..d22d2c7 100644 (file)
--- a/src/intel/compiler/brw_fs.cpp
+++ b/src/intel/compiler/brw_fs.cpp
@@ -8981,9 +8981,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
    prog_data->slm_size = src_shader->num_shared;
 
    unsigned local_workgroup_size;
-   if (prog_data->uses_variable_group_size) {
-      prog_data->max_variable_local_size =
-         src_shader->info.cs.max_variable_local_size;
+   if (src_shader->info.cs.local_size_variable) {
       local_workgroup_size = src_shader->info.cs.max_variable_local_size;
    } else {
       prog_data->local_size[0] = src_shader->info.cs.local_size[0];

diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp
index e3149f6..e4fbaa5 100644 (file)
--- a/src/intel/compiler/brw_fs_nir.cpp
+++ b/src/intel/compiler/brw_fs_nir.cpp
@@ -105,7 +105,7 @@ fs_visitor::nir_setup_uniforms()
       assert(uniforms == prog_data->nr_params);
 
       uint32_t *param;
-      if (brw_cs_prog_data(prog_data)->uses_variable_group_size) {
+      if (nir->info.cs.local_size_variable) {
          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);
@@ -3732,7 +3732,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld,
        * invocations are already executed lock-step.  Instead of an actual
        * barrier just emit a scheduling fence, that will generate no code.
        */
-      if (!cs_prog_data->uses_variable_group_size &&
+      if (!nir->info.cs.local_size_variable &&
           workgroup_size() <= dispatch_width) {
          bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE);
          break;
@@ -4297,7 +4297,7 @@ fs_visitor::nir_emit_intrinsic(const fs_builder &bld, nir_intrinsic_instr *instr
        *
        * TODO: Check if applies for many HW threads sharing same Data Port.
        */
-      if (!brw_cs_prog_data(prog_data)->uses_variable_group_size &&
+      if (!nir->info.cs.local_size_variable &&
           slm_fence && workgroup_size() <= dispatch_width)
          slm_fence = false;
 

diff --git a/src/mesa/drivers/dri/i965/brw_cs.c b/src/mesa/drivers/dri/i965/brw_cs.c
index d01fd22..1f2fefc 100644 (file)
--- a/src/mesa/drivers/dri/i965/brw_cs.c
+++ b/src/mesa/drivers/dri/i965/brw_cs.c
@@ -114,11 +114,8 @@ brw_codegen_cs_prog(struct brw_context *brw,
     * the actual size is not known until the dispatch command is issued.
     */
    if (nir->info.cs.local_size_variable) {
-      prog_data.uses_variable_group_size = true;
       nir->info.cs.max_variable_local_size =
          gl_ctx->Const.MaxComputeVariableGroupInvocations;
-   } else {
-      prog_data.uses_variable_group_size = false;
    }
 
    char *error_str;