Mesa (main): broadcom/compiler: track if a compute shader uses subgroup functionality

[GitLab Mirror]

Module: Mesa
Branch: main
Commit: 10313b03b5061eba1f8eb46f6576f1108c2554a2
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=10313b03b5061eba1f8eb46f6576f1108c2554a2

Author: Iago Toral Quiroga <itoral at igalia.com>
Date:   Tue Jun 22 12:33:23 2021 +0200

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

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

---

 src/broadcom/compiler/v3d_compiler.h |  5 +++++
 src/broadcom/compiler/vir.c          | 11 +++++++++--
 2 files changed, 14 insertions(+), 2 deletions(-)

diff --git a/src/broadcom/compiler/v3d_compiler.h b/src/broadcom/compiler/v3d_compiler.h
index 564d8b9c8df..4a46e9ee286 100644
--- 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 1b35ed95033..7b77e2172ff 100644
--- 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;
                 }