KonstantinSeurer
/
mesa
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Marek Olšák 2020-09-06 11:38:53 -04:00 committed by Marge Bot
parent 757f790ad8
commit 1dd243d4f5
5 changed files with 13 additions and 16 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
src/gallium/drivers/radeonsi/si_compute.c
View File

@ -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;

15
src/gallium/drivers/radeonsi/si_shader.c
View File

@ -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 =

2
src/gallium/drivers/radeonsi/si_shader.h
View File

@ -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;

5
src/gallium/drivers/radeonsi/si_shader_llvm.c
View File

@ -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);

4
src/gallium/drivers/radeonsi/si_shader_nir.c
View File

@ -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: {