From: Marek Olšák <marek.olsak@amd.com>
Date: Sun, 6 Sep 2020 15:38:53 +0000 (-0400)
Subject: radeonsi: use shader_info::cs::local_size_variable to clean up some code
X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=76f9eae7bb2b503e46ebe2847619bba1cec393b6;p=mesa.git

radeonsi: use shader_info::cs::local_size_variable to clean up some code

Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6624>
---

diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c
index 3987ecdc66d..700e00870d3 100644
--- a/src/gallium/drivers/radeonsi/si_compute.c
+++ b/src/gallium/drivers/radeonsi/si_compute.c
@@ -132,8 +132,7 @@ static void si_create_compute_state_async(void *job, int thread_index)
                             &sel->active_samplers_and_images);
 
    program->shader.is_monolithic = true;
-   program->reads_variable_block_size =
-      sel->info.uses_block_size && sel->info.base.cs.local_size[0] == 0;
+   program->reads_variable_block_size = sel->info.uses_variable_block_size;
    program->num_cs_user_data_dwords =
       sel->info.base.cs.user_data_components_amd;
 
diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 50b37fb40e1..8a26ac0d06d 100644
--- a/src/gallium/drivers/radeonsi/si_shader.c
+++ b/src/gallium/drivers/radeonsi/si_shader.c
@@ -211,17 +211,15 @@ unsigned si_get_max_workgroup_size(const struct si_shader *shader)
       return 0;
    }
 
+   /* Compile a variable block size using the maximum variable size. */
+   if (shader->selector->info.base.cs.local_size_variable)
+      return SI_MAX_VARIABLE_THREADS_PER_BLOCK;
+
    uint16_t *local_size = shader->selector->info.base.cs.local_size;
    unsigned max_work_group_size = (uint32_t)local_size[0] *
                                   (uint32_t)local_size[1] *
                                   (uint32_t)local_size[2];
-
-   if (!max_work_group_size) {
-      /* This is a variable group size compute shader,
-       * compile it for the maximum possible group size.
-       */
-      max_work_group_size = SI_MAX_VARIABLE_THREADS_PER_BLOCK;
-   }
+   assert(max_work_group_size);
    return max_work_group_size;
 }
 
@@ -695,8 +693,7 @@ void si_create_function(struct si_shader_context *ctx, bool ngg_cull_shader)
       declare_per_stage_desc_pointers(ctx, true);
       if (shader->selector->info.uses_grid_size)
          ac_add_arg(&ctx->args, AC_ARG_SGPR, 3, AC_ARG_INT, &ctx->args.num_work_groups);
-      if (shader->selector->info.uses_block_size &&
-          shader->selector->info.base.cs.local_size[0] == 0)
+      if (shader->selector->info.uses_variable_block_size)
          ac_add_arg(&ctx->args, AC_ARG_SGPR, 3, AC_ARG_INT, &ctx->block_size);
 
       unsigned cs_user_data_dwords =
diff --git a/src/gallium/drivers/radeonsi/si_shader.h b/src/gallium/drivers/radeonsi/si_shader.h
index 1cfaab59979..805a8b1e87a 100644
--- a/src/gallium/drivers/radeonsi/si_shader.h
+++ b/src/gallium/drivers/radeonsi/si_shader.h
@@ -362,7 +362,7 @@ struct si_shader_info {
    bool uses_invocationid;
    bool uses_thread_id[3];
    bool uses_block_id[3];
-   bool uses_block_size;
+   bool uses_variable_block_size;
    bool uses_grid_size;
    bool uses_subgroup_info;
    bool writes_position;
diff --git a/src/gallium/drivers/radeonsi/si_shader_llvm.c b/src/gallium/drivers/radeonsi/si_shader_llvm.c
index 142587b1d53..0602593ba6e 100644
--- a/src/gallium/drivers/radeonsi/si_shader_llvm.c
+++ b/src/gallium/drivers/radeonsi/si_shader_llvm.c
@@ -370,9 +370,10 @@ LLVMValueRef si_llvm_get_block_size(struct ac_shader_abi *abi)
    LLVMValueRef values[3];
    LLVMValueRef result;
    unsigned i;
-   uint16_t *local_size = ctx->shader->selector->info.base.cs.local_size;
 
-   if (local_size[0] != 0) {
+   if (!ctx->shader->selector->info.base.cs.local_size_variable) {
+      uint16_t *local_size = ctx->shader->selector->info.base.cs.local_size;
+
       for (i = 0; i < 3; ++i)
          values[i] = LLVMConstInt(ctx->ac.i32, local_size[i], 0);
 
diff --git a/src/gallium/drivers/radeonsi/si_shader_nir.c b/src/gallium/drivers/radeonsi/si_shader_nir.c
index 46dd6e973fb..d88e07c9642 100644
--- a/src/gallium/drivers/radeonsi/si_shader_nir.c
+++ b/src/gallium/drivers/radeonsi/si_shader_nir.c
@@ -262,8 +262,8 @@ static void scan_instruction(const struct nir_shader *nir, struct si_shader_info
          break;
       case nir_intrinsic_load_local_group_size:
          /* The block size is translated to IMM with a fixed block size. */
-         if (info->base.cs.local_size[0] == 0)
-            info->uses_block_size = true;
+         if (info->base.cs.local_size_variable)
+            info->uses_variable_block_size = true;
          break;
       case nir_intrinsic_load_local_invocation_id:
       case nir_intrinsic_load_work_group_id: {