diff --git a/src/imagination/pco/pco_internal.h b/src/imagination/pco/pco_internal.h
index ab5a0ba4645..b35910c9c38 100644
--- a/src/imagination/pco/pco_internal.h
+++ b/src/imagination/pco/pco_internal.h
@@ -1769,6 +1769,7 @@ bool pco_nir_lower_fs_intrinsics(nir_shader *shader);
 bool pco_nir_lower_images(nir_shader *shader, pco_data *data);
 bool pco_nir_lower_interpolation(nir_shader *shader, pco_fs_data *fs);
 bool pco_nir_lower_io(nir_shader *shader);
+bool pco_nir_lower_subgroups(nir_shader *shader);
 bool pco_nir_lower_tex(nir_shader *shader, pco_data *data, pco_ctx *ctx);
 bool pco_nir_lower_variables(nir_shader *shader, bool inputs, bool outputs);
 bool pco_nir_lower_vk(nir_shader *shader, pco_data *data);
diff --git a/src/imagination/pco/pco_nir.c b/src/imagination/pco/pco_nir.c
index a41312d5e5f..56b007d9f96 100644
--- a/src/imagination/pco/pco_nir.c
+++ b/src/imagination/pco/pco_nir.c
@@ -547,6 +547,8 @@ void pco_preprocess_nir(pco_ctx *ctx, nir_shader *nir)
                });
    }
 
+   NIR_PASS(_, nir, pco_nir_lower_subgroups);
+
    NIR_PASS(_,
             nir,
             nir_lower_io_vars_to_temporaries,
diff --git a/src/imagination/pco/pco_nir_sync.c b/src/imagination/pco/pco_nir_sync.c
index d423f11e840..c9fd5e693e8 100644
--- a/src/imagination/pco/pco_nir_sync.c
+++ b/src/imagination/pco/pco_nir_sync.c
@@ -43,7 +43,8 @@ static nir_def *lower_barrier(nir_builder *b, nir_instr *instr, void *cb_data)
    unsigned wg_size = info->workgroup_size[0] * info->workgroup_size[1] *
                       info->workgroup_size[2];
 
-   if (wg_size <= ROGUE_MAX_INSTANCES_PER_TASK || exec_scope == SCOPE_NONE)
+   if (wg_size <= ROGUE_MAX_INSTANCES_PER_TASK || exec_scope == SCOPE_NONE ||
+       exec_scope == SCOPE_SUBGROUP)
       return NIR_LOWER_INSTR_PROGRESS_REPLACE;
 
    /* TODO: We might be able to re-use barrier counters. */
@@ -171,3 +172,73 @@ bool pco_nir_lower_atomics(nir_shader *shader, pco_data *data)
                                         lower_atomic,
                                         &data->common.uses.usclib);
 }
+
+static nir_def *
+lower_subgroup_intrinsic(nir_builder *b, nir_instr *instr, void *cb_data)
+{
+   nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
+   assert(intr->def.num_components == 1);
+
+   switch (intr->intrinsic) {
+   case nir_intrinsic_load_subgroup_size:
+      return nir_imm_int(b, 1);
+
+   case nir_intrinsic_load_subgroup_invocation:
+      return nir_imm_int(b, 0);
+
+   case nir_intrinsic_load_num_subgroups:
+      return nir_imm_int(b,
+                         b->shader->info.workgroup_size[0] *
+                            b->shader->info.workgroup_size[1] *
+                            b->shader->info.workgroup_size[2]);
+
+   case nir_intrinsic_load_subgroup_id:
+      return nir_load_local_invocation_index(b);
+
+   case nir_intrinsic_first_invocation:
+      return nir_imm_int(b, 0);
+
+   case nir_intrinsic_elect:
+      return nir_imm_true(b);
+
+   default:
+      break;
+   }
+
+   UNREACHABLE("");
+}
+
+static bool is_subgroup_intrinsic(const nir_instr *instr,
+                                  UNUSED const void *cb_data)
+{
+   if (instr->type != nir_instr_type_intrinsic)
+      return false;
+
+   nir_intrinsic_instr *intr = nir_instr_as_intrinsic(instr);
+   switch (intr->intrinsic) {
+   case nir_intrinsic_load_subgroup_size:
+   case nir_intrinsic_load_subgroup_invocation:
+   case nir_intrinsic_load_num_subgroups:
+   case nir_intrinsic_load_subgroup_id:
+   case nir_intrinsic_first_invocation:
+   case nir_intrinsic_elect:
+      return true;
+
+   default:
+      break;
+   }
+
+   return false;
+}
+
+bool pco_nir_lower_subgroups(nir_shader *shader)
+{
+   shader->info.api_subgroup_size = 1;
+   shader->info.min_subgroup_size = 1;
+   shader->info.max_subgroup_size = 1;
+
+   return nir_shader_lower_instructions(shader,
+                                        is_subgroup_intrinsic,
+                                        lower_subgroup_intrinsic,
+                                        NULL);
+}
diff --git a/src/imagination/vulkan/pvr_device.c b/src/imagination/vulkan/pvr_device.c
index 4abd7da5c4c..4fffc03fbb8 100644
--- a/src/imagination/vulkan/pvr_device.c
+++ b/src/imagination/vulkan/pvr_device.c
@@ -585,6 +585,12 @@ static bool pvr_physical_device_get_properties(
       .optimalBufferCopyRowPitchAlignment = PVR_STORAGE_BUFFER_OFFSET_ALIGNMENT,
       .nonCoherentAtomSize = 1U,
 
+      /* Vulkan 1.1 */
+      .subgroupSize = 1,
+      .subgroupSupportedStages = VK_SHADER_STAGE_COMPUTE_BIT,
+      .subgroupSupportedOperations = VK_SUBGROUP_FEATURE_BASIC_BIT,
+      .subgroupQuadOperationsInAllStages = false,
+
       /* Vulkan 1.0 / VK_KHR_maintenance2 */
       .pointClippingBehavior =
          VK_POINT_CLIPPING_BEHAVIOR_USER_CLIP_PLANES_ONLY,