KonstantinSeurer
/
mesa
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Iago Toral Quiroga 2021-06-22 12:33:23 +02:00
parent 5081de07f7
commit 10313b03b5
2 changed files with 14 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
src/broadcom/compiler/v3d_compiler.h
View File

@ -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

11
src/broadcom/compiler/vir.c
View File

@ -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;
}