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

ir3: Initialize local size earlier 

We need the local size in RA for occupancy calculations. Not
initializing these had the unfortunate consequence of
ir3_get_reg_independent_max_waves() returning 0 for compute shaders with
shared variables, disabling the register limiting logic.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/13143>
This commit is contained in:
Connor Abbott 2021-09-30 11:53:14 +02:00 committed by Marge Bot
parent ce3dd1375f
commit decbced48f
1 changed files with 7 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
src/freedreno/ir3/ir3_compiler_nir.c
View File

@ -3992,6 +3992,13 @@ ir3_compile_shader_nir(struct ir3_compiler *compiler,
ir = so->ir = ctx->ir;
if (so->type == MESA_SHADER_COMPUTE) {
so->local_size[0] = ctx->s->info.workgroup_size[0];
so->local_size[1] = ctx->s->info.workgroup_size[1];
so->local_size[2] = ctx->s->info.workgroup_size[2];
so->local_size_variable = ctx->s->info.workgroup_size_variable;
}
/* Vertex shaders in a tessellation or geometry pipeline treat END as a
* NOP and has an epilogue that writes the VS outputs to local storage, to
* be read by the HS. Then it resets execution mask (chmask) and chains
@ -4348,13 +4355,6 @@ ir3_compile_shader_nir(struct ir3_compiler *compiler,
ctx->s->info.fs.needs_quad_helper_invocations)
so->need_pixlod = true;
if (so->type == MESA_SHADER_COMPUTE) {
so->local_size[0] = ctx->s->info.workgroup_size[0];
so->local_size[1] = ctx->s->info.workgroup_size[1];
so->local_size[2] = ctx->s->info.workgroup_size[2];
so->local_size_variable = ctx->s->info.workgroup_size_variable;
}
out:
if (ret) {
if (so->ir)