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 3987ecdc66d234dd8c42d4829e1b620470a0c5ef..700e00870d3530c5e44a2618f69c17832fb3b2b1 100644 (file)
--- 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;
                             &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;
 
    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 50b37fb40e184827ca690b3d4bec65cef6903cbb..8a26ac0d06d08e6e9cbbe79169a806bde0e9e4e6 100644 (file)
--- 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;
    }
 
       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];
    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;
 }
 
    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);
       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 =
          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 1cfaab599791ad99fa1bb094f9d30b4641c7d708..805a8b1e87aeb15e194ce3e65d979fc591286c0d 100644 (file)
--- 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_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;
    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 142587b1d538453fe832d1c76dabfaa7a97b028b..0602593ba6e4b435dcc3d33ace4e1817079da731 100644 (file)
--- 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;
    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);
 
       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 46dd6e973fb2ac0b83935ff1124642b36d18b4d0..d88e07c9642cb73eac853b7ba4398c86a1d03033 100644 (file)
--- 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. */
          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: {
          break;
       case nir_intrinsic_load_local_invocation_id:
       case nir_intrinsic_load_work_group_id: {