diff --git a/src/amd/vulkan/radv_pipeline.c b/src/amd/vulkan/radv_pipeline.c
index b9c9304d5b5..0ff2ab986fd 100644
--- a/src/amd/vulkan/radv_pipeline.c
+++ b/src/amd/vulkan/radv_pipeline.c
@@ -3420,34 +3420,6 @@ radv_fill_shader_info(struct radv_pipeline *pipeline,
 
    radv_nir_shader_info_link(device, pipeline_key, stages, last_vgt_api_stage);
 
-   if (stages[MESA_SHADER_COMPUTE].nir) {
-      unsigned subgroup_size = pipeline_key->cs.compute_subgroup_size;
-      unsigned req_subgroup_size = subgroup_size;
-      bool require_full_subgroups = pipeline_key->cs.require_full_subgroups;
-
-      if (!subgroup_size)
-         subgroup_size = device->physical_device->cs_wave_size;
-
-      unsigned local_size = stages[MESA_SHADER_COMPUTE].nir->info.workgroup_size[0] *
-                            stages[MESA_SHADER_COMPUTE].nir->info.workgroup_size[1] *
-                            stages[MESA_SHADER_COMPUTE].nir->info.workgroup_size[2];
-
-      /* Games don't always request full subgroups when they should,
-       * which can cause bugs if cswave32 is enabled.
-       */
-      if (device->physical_device->cs_wave_size == 32 &&
-          stages[MESA_SHADER_COMPUTE].nir->info.cs.uses_wide_subgroup_intrinsics && !req_subgroup_size &&
-          local_size % RADV_SUBGROUP_SIZE == 0)
-         require_full_subgroups = true;
-
-      if (require_full_subgroups && !req_subgroup_size) {
-         /* don't use wave32 pretending to be wave64 */
-         subgroup_size = RADV_SUBGROUP_SIZE;
-      }
-
-      stages[MESA_SHADER_COMPUTE].info.cs.subgroup_size = subgroup_size;
-   }
-
    for (int i = 0; i < MESA_VULKAN_SHADER_STAGES; i++) {
       if (stages[i].nir) {
          stages[i].info.wave_size = radv_get_wave_size(device, i, &stages[i].info);
diff --git a/src/amd/vulkan/radv_shader_info.c b/src/amd/vulkan/radv_shader_info.c
index ce5dc9d1e99..29a8c1f7cec 100644
--- a/src/amd/vulkan/radv_shader_info.c
+++ b/src/amd/vulkan/radv_shader_info.c
@@ -506,9 +506,34 @@ gather_shader_info_fs(const nir_shader *nir, const struct radv_pipeline_key *pip
 }
 
 static void
-gather_shader_info_cs(const nir_shader *nir, struct radv_shader_info *info)
+gather_shader_info_cs(struct radv_device *device, const nir_shader *nir,
+                      const struct radv_pipeline_key *pipeline_key, struct radv_shader_info *info)
 {
    info->cs.uses_ray_launch_size = BITSET_TEST(nir->info.system_values_read, SYSTEM_VALUE_RAY_LAUNCH_SIZE_ADDR_AMD);
+
+   unsigned subgroup_size = pipeline_key->cs.compute_subgroup_size;
+   unsigned req_subgroup_size = subgroup_size;
+   bool require_full_subgroups = pipeline_key->cs.require_full_subgroups;
+
+   if (!subgroup_size)
+      subgroup_size = device->physical_device->cs_wave_size;
+
+   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.
+    */
+   if (device->physical_device->cs_wave_size == 32 && nir->info.cs.uses_wide_subgroup_intrinsics &&
+       !req_subgroup_size && local_size % RADV_SUBGROUP_SIZE == 0)
+      require_full_subgroups = true;
+
+   if (require_full_subgroups && !req_subgroup_size) {
+      /* don't use wave32 pretending to be wave64 */
+      subgroup_size = RADV_SUBGROUP_SIZE;
+   }
+
+   info->cs.subgroup_size = subgroup_size;
 }
 
 static void
@@ -648,7 +673,7 @@ radv_nir_shader_info_pass(struct radv_device *device, const struct nir_shader *n
 
    switch (nir->info.stage) {
    case MESA_SHADER_COMPUTE:
-      gather_shader_info_cs(nir, info);
+      gather_shader_info_cs(device, nir, pipeline_key, info);
       break;
    case MESA_SHADER_TASK:
       gather_shader_info_task(nir, info);