broadcom/compiler: track if a compute shader uses subgroup functionality

Reviewed-by: Alejandro Piñeiro <apinheiro@igalia.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11620>

diff --git a/src/broadcom/compiler/v3d_compiler.h b/src/broadcom/compiler/v3d_compiler.h
index 564d8b9..4a46e9e 100644 (file)
--- a/src/broadcom/compiler/v3d_compiler.h
+++ b/src/broadcom/compiler/v3d_compiler.h
@@ -728,6 +728,9 @@ struct v3d_compile {
         struct qreg cs_shared_offset;
         int local_invocation_index_bits;
 
+        /* If the shader uses subgroup functionality */
+        bool has_subgroups;
+
         uint8_t vattr_sizes[V3D_MAX_VS_INPUTS / 4];
         uint32_t vpm_output_size;
 
@@ -947,6 +950,8 @@ struct v3d_compute_prog_data {
         /* Size in bytes of the workgroup's shared space. */
         uint32_t shared_size;
         uint16_t local_size[3];
+        /* If the shader uses subgroup functionality */
+        bool has_subgroups;
 };
 
 static inline bool

diff --git a/src/broadcom/compiler/vir.c b/src/broadcom/compiler/vir.c
index 1b35ed9..7b77e21 100644 (file)
--- a/src/broadcom/compiler/vir.c
+++ b/src/broadcom/compiler/vir.c
@@ -808,6 +808,8 @@ v3d_cs_set_prog_data(struct v3d_compile *c,
         prog_data->local_size[0] = c->s->info.workgroup_size[0];
         prog_data->local_size[1] = c->s->info.workgroup_size[1];
         prog_data->local_size[2] = c->s->info.workgroup_size[2];
+
+        prog_data->has_subgroups = c->has_subgroups;
 }
 
 static void
@@ -1384,11 +1386,16 @@ lower_subgroup_intrinsics(struct v3d_compile *c,
                         continue;
 
                 switch (intr->intrinsic) {
-                case nir_intrinsic_load_num_subgroups: {
+                case nir_intrinsic_load_num_subgroups:
                         lower_load_num_subgroups(c, b, intr);
                         progress = true;
+                        FALLTHROUGH;
+                case nir_intrinsic_load_subgroup_id:
+                case nir_intrinsic_load_subgroup_size:
+                case nir_intrinsic_load_subgroup_invocation:
+                case nir_intrinsic_elect:
+                        c->has_subgroups = true;
                         break;
-                }
                 default:
                         break;
                 }