From 10313b03b5061eba1f8eb46f6576f1108c2554a2 Mon Sep 17 00:00:00 2001 From: Iago Toral Quiroga Date: Tue, 22 Jun 2021 12:33:23 +0200 Subject: [PATCH] broadcom/compiler: track if a compute shader uses subgroup functionality MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reviewed-by: Alejandro Piñeiro Part-of: --- 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; }