fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-02-12 15:50:31 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 5

ir3: Split out max_branchstack and branchstack_size

Browse source 
One is the maximum size per wave and the other is the size per uSP (i.e.
"core"). They happened to be the same before, but actually they are
different limits in the HW.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/39468>
This commit is contained in:
Connor Abbott 2026-01-22 18:00:24 -05:00 • committed by Marge Bot
parent c569fb5669
commit f5ea8b9a0b
4 changed files with 7 additions and 3 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

4
src/freedreno/ir3/ir3.c
View file

@ -209,12 +209,12 @@ ir3_should_double_threadsize(struct ir3_shader_variant *v, unsigned regs_count)
if (v->shader_options.real_wavesize == IR3_DOUBLE_ONLY)
return true;
/* We can't support more than compiler->branchstack_size diverging threads
/* We can't support more than compiler->max_branchstack diverging threads
* in a wave. Thus, doubling the threadsize is only possible if we don't
* exceed the branchstack size limit.
*/
if (MIN2(v->branchstack, compiler->info->threadsize_base * 2) >
compiler->branchstack_size) {
compiler->max_branchstack) {
return false;
}

1
src/freedreno/ir3/ir3_compiler.c
View file

@ -167,6 +167,7 @@ ir3_compiler_create(struct fd_device *dev, const struct fd_dev_id *dev_id,
/* TODO see if older GPU's were different here */
compiler->branchstack_size = 64;
compiler->max_branchstack = 64;
compiler->max_variable_workgroup_size = 1024;

3
src/freedreno/ir3/ir3_compiler.h
View file

@ -164,6 +164,9 @@ struct ir3_compiler {
/* The number of total branch stack entries, divided by wave_granularity. */
uint32_t branchstack_size;
/* The maximum number of branch stack entries per wave. */
uint32_t max_branchstack;
/* The byte increment of MEMSIZEPERITEM, the private memory per-fiber allocation. */
uint32_t pvtmem_per_fiber_align;

2
src/freedreno/ir3/ir3_shader.h
View file

@ -1458,7 +1458,7 @@ ir3_shader_branchstack_hw(const struct ir3_shader_variant *v)
if (v->compiler->gen < 5)
return v->branchstack;
return align(MIN2(v->branchstack, v->compiler->branchstack_size), 2);
return align(MIN2(v->branchstack, v->compiler->max_branchstack), 2);
}
ENDC;