From: Marek Olšák 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 Part-of: --- 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: {