diff --git a/src/freedreno/ir3/ir3_shader.c b/src/freedreno/ir3/ir3_shader.c index d6bd4d68161..5652cd75b37 100644 --- a/src/freedreno/ir3/ir3_shader.c +++ b/src/freedreno/ir3/ir3_shader.c @@ -604,6 +604,7 @@ create_variant(struct ir3_shader *shader, const struct ir3_shader_key *key, if (ir3_shader_compute(v)) { v->cs.force_linear_dispatch = shader->cs.force_linear_dispatch; + v->cs.round_robin_mode = shader->nir->info.occupancy_bounded_workgroup_fairness; v->local_size[0] = shader->nir->info.workgroup_size[0]; v->local_size[1] = shader->nir->info.workgroup_size[1]; diff --git a/src/freedreno/ir3/ir3_shader.h b/src/freedreno/ir3/ir3_shader.h index 45fd7da84bb..ad1390292fc 100644 --- a/src/freedreno/ir3/ir3_shader.h +++ b/src/freedreno/ir3/ir3_shader.h @@ -961,6 +961,7 @@ struct ir3_shader_variant { struct { unsigned req_local_mem; bool force_linear_dispatch; + bool round_robin_mode; uint32_t local_invocation_id; uint32_t work_group_id; } cs; diff --git a/src/freedreno/vulkan/tu_shader.cc b/src/freedreno/vulkan/tu_shader.cc index 0bd6eb97239..4e91274b114 100644 --- a/src/freedreno/vulkan/tu_shader.cc +++ b/src/freedreno/vulkan/tu_shader.cc @@ -1750,6 +1750,7 @@ tu6_emit_xs(struct tu_crb &crb, .fullregfootprint = xs->info.max_reg + 1, .branchstack = ir3_shader_branchstack_hw(xs), .threadsize = thrsz, + .computerrmodeen = xs->cs.round_robin_mode, .earlypreamble = xs->early_preamble, .mergedregs = xs->mergedregs, )); crb.add(A6XX_SP_CS_INSTR_SIZE(xs->instrlen)); diff --git a/src/gallium/drivers/freedreno/a6xx/fd6_program.cc b/src/gallium/drivers/freedreno/a6xx/fd6_program.cc index f958d3df2d8..367eecefea1 100644 --- a/src/gallium/drivers/freedreno/a6xx/fd6_program.cc +++ b/src/gallium/drivers/freedreno/a6xx/fd6_program.cc @@ -230,6 +230,7 @@ emit_shader_regs(struct fd_screen *screen, fd_cs &cs, const struct ir3_shader_va .fullregfootprint = so->info.max_reg + 1, .branchstack = ir3_shader_branchstack_hw(so), .threadsize = thrsz, + .computerrmodeen = so->cs.round_robin_mode, .earlypreamble = so->early_preamble, .mergedregs = so->mergedregs, ));