pvr, pco: add minimal support required for Vulkan 1.2 subgroups

Signed-off-by: Simon Perretta <simon.perretta@imgtec.com>
Acked-by: Erik Faye-Lund <erik.faye-lund@collabora.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37512>

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);

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,

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);
+}

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,