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

radv: Don't transparently use wave32 with cooperative matrices.

Browse source 
The instruction has different regsizes for wave32 vs. wave64.

To ensure cases with cooperative matrix load/store without any
actual wmma instructions get handled correctly, also require
full subgroups if subgroup invocation/id are used. Not sure
those could be transparently changed anyway.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24683>
This commit is contained in:
Bas Nieuwenhuizen 2023-07-23 22:14:04 +02:00 • committed by Marge Bot
parent d8458c0559
commit 31863aa48c
1 changed files with 3 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
src/amd/vulkan/radv_shader_info.c
View file

@ -911,10 +911,11 @@ gather_shader_info_cs(struct radv_device *device, const nir_shader *nir, const s
unsigned local_size = nir->info.workgroup_size[0] * nir->info.workgroup_size[1] * nir->info.workgroup_size[2];
/* Games don't always request full subgroups when they should, which can cause bugs if cswave32
* is enabled.
* is enabled. Furthermore, if cooperative matrices or subgroup info are used, we can't transparently change
* the subgroup size.
*/
const bool require_full_subgroups =
pipeline_key->stage_info[MESA_SHADER_COMPUTE].subgroup_require_full ||
pipeline_key->stage_info[MESA_SHADER_COMPUTE].subgroup_require_full || nir->info.cs.has_cooperative_matrix ||
(default_wave_size == 32 && nir->info.uses_wide_subgroup_intrinsics && local_size % RADV_SUBGROUP_SIZE == 0);
const unsigned required_subgroup_size = pipeline_key->stage_info[MESA_SHADER_COMPUTE].subgroup_required_size * 32;