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

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>

diff --git a/src/amd/vulkan/radv_shader_info.c b/src/amd/vulkan/radv_shader_info.c
index 44dec6a..c87f401 100644 (file)
--- a/src/amd/vulkan/radv_shader_info.c
+++ b/src/amd/vulkan/radv_shader_info.c
@@ -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;