diff options
| author | Timur Kristóf <timur.kristof@gmail.com> | 2020-04-27 19:51:40 +0200 |
|---|---|---|
| committer | Marge Bot <eric+marge@anholt.net> | 2020-04-29 11:51:04 +0000 |
| commit | baa46878d4533f21d12bc93d5eed09436b3cc9fd (patch) | |
| tree | f1acb11338c17d49a9d91d65c5d486f2fcdb77c8 | |
| parent | fdbb2968533be9a1caca731cf11c2ed3b46e6043 (diff) | |
aco: Calculate workgroup size of legacy GS.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4388>
| -rw-r--r-- | src/amd/compiler/aco_instruction_selection_setup.cpp | 6 |
1 files changed, 5 insertions, 1 deletions
diff --git a/src/amd/compiler/aco_instruction_selection_setup.cpp b/src/amd/compiler/aco_instruction_selection_setup.cpp index c09d1459846..bf4b34c6e3e 100644 --- a/src/amd/compiler/aco_instruction_selection_setup.cpp +++ b/src/amd/compiler/aco_instruction_selection_setup.cpp @@ -1293,7 +1293,11 @@ setup_isel_context(Program* program, program->workgroup_size = program->wave_size; } else if (program->stage & hw_gs) { /* If on-chip GS (LDS rings) are enabled on GFX9 or later, merged GS operates in workgroups */ - program->workgroup_size = UINT_MAX; /* TODO: set by VGT_GS_ONCHIP_CNTL, which is not plumbed to ACO */ + assert(program->chip_class >= GFX9); + uint32_t es_verts_per_subgrp = G_028A44_ES_VERTS_PER_SUBGRP(program->info->gs_ring_info.vgt_gs_onchip_cntl); + uint32_t gs_instr_prims_in_subgrp = G_028A44_GS_INST_PRIMS_IN_SUBGRP(program->info->gs_ring_info.vgt_gs_onchip_cntl); + uint32_t workgroup_size = MAX2(es_verts_per_subgrp, gs_instr_prims_in_subgrp); + program->workgroup_size = MAX2(MIN2(workgroup_size, 256), 1); } else if (program->stage == vertex_ls) { /* Unmerged LS operates in workgroups */ program->workgroup_size = UINT_MAX; /* TODO: probably tcs_num_patches * tcs_vertices_in, but those are not plumbed to ACO for LS */ |