fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-01-01 18:20:10 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 3

radeonsi: improve the heuristic when to use Wave32 for compute shaders

Browse source 
Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24759>
This commit is contained in:
Marek Olšák 2023-08-16 16:04:10 -04:00 • committed by Marge Bot
parent 7f25f30443
commit 1c82067b60
1 changed files with 5 additions and 6 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

11
src/gallium/drivers/radeonsi/si_state_shaders.cpp
View file

@ -38,12 +38,11 @@ unsigned si_determine_wave_size(struct si_screen *sscreen, struct si_shader *sha
(stage == MESA_SHADER_GEOMETRY && !shader->key.ge.as_ngg))
return 64;
/* Small workgroups use Wave32 unconditionally. */
if (stage == MESA_SHADER_COMPUTE && info &&
!info->base.workgroup_size_variable &&
info->base.workgroup_size[0] *
info->base.workgroup_size[1] *
info->base.workgroup_size[2] <= 32)
/* Workgroup sizes that are not divisible by 64 use Wave32. */
if (stage == MESA_SHADER_COMPUTE && info && !info->base.workgroup_size_variable &&
(info->base.workgroup_size[0] *
info->base.workgroup_size[1] *
info->base.workgroup_size[2]) % 64 != 0)
return 32;
/* Debug flags. */