fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-04-05 06:00:36 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

nak: Rework barrier handling a bit

Browse source 
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26577>
This commit is contained in:
Faith Ekstrand 2023-12-07 12:49:28 -06:00 • committed by Marge Bot
parent eafc8f58c6
commit 3bb4c14c75
4 changed files with 46 additions and 22 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

14
src/nouveau/compiler/nak/from_nir.rs
View file

@ -2228,13 +2228,13 @@ impl<'a> ShaderFromNir<'a> {
match intrin.execution_scope() {
SCOPE_NONE => (),
SCOPE_WORKGROUP => {
if self.nir.info.stage() == MESA_SHADER_COMPUTE {
// OpBar needs num_barriers > 0 but, as far as we
// know, it doesn't actually use a barrier.
self.info.num_barriers = 1;
b.push_op(OpBar {});
b.push_op(OpNop { label: None });
}
assert!(
self.nir.info.stage() == MESA_SHADER_COMPUTE
|| self.nir.info.stage() == MESA_SHADER_KERNEL
);
self.info.num_barriers = 1;
b.push_op(OpBar {});
b.push_op(OpNop { label: None });
}
_ => panic!("Unhandled execution scope"),
}

39
src/nouveau/compiler/nak_nir.c
View file

@ -19,6 +19,41 @@
#define OPT_V(nir, pass, ...) NIR_PASS_V(nir, pass, ##__VA_ARGS__)
bool
nak_nir_workgroup_has_one_subgroup(const nir_shader *nir)
{
switch (nir->info.stage) {
case MESA_SHADER_VERTEX:
case MESA_SHADER_TESS_EVAL:
case MESA_SHADER_GEOMETRY:
case MESA_SHADER_FRAGMENT:
unreachable("Shader stage does not have workgroups");
break;
case MESA_SHADER_TESS_CTRL:
/* Tessellation only ever has one subgroup per workgroup. The Vulkan
* limit on the number of tessellation invocations is 32 to allow for
* this.
*/
return true;
case MESA_SHADER_COMPUTE:
case MESA_SHADER_KERNEL: {
if (nir->info.workgroup_size_variable)
return false;
uint16_t wg_sz = nir->info.workgroup_size[0] *
nir->info.workgroup_size[1] *
nir->info.workgroup_size[2];
return wg_sz <= 32;
}
default:
unreachable("Unknown shader stage");
}
}
static void
optimize_nir(nir_shader *nir, const struct nak_compiler *nak, bool allow_copies)
{
@ -204,7 +239,7 @@ nak_nir_lower_subgroup_id_intrin(nir_builder *b, nir_intrinsic_instr *intrin,
b->cursor = nir_instr_remove(&intrin->instr);
nir_def *num_subgroups;
if (nak_nir_has_one_subgroup(b->shader)) {
if (nak_nir_workgroup_has_one_subgroup(b->shader)) {
num_subgroups = nir_imm_int(b, 1);
} else {
assert(b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE);
@ -225,7 +260,7 @@ nak_nir_lower_subgroup_id_intrin(nir_builder *b, nir_intrinsic_instr *intrin,
b->cursor = nir_instr_remove(&intrin->instr);
nir_def *subgroup_id;
if (nak_nir_has_one_subgroup(b->shader)) {
if (nak_nir_workgroup_has_one_subgroup(b->shader)) {
subgroup_id = nir_imm_int(b, 0);
} else {
assert(b->shader->info.cs.derivative_group == DERIVATIVE_GROUP_NONE);

2
src/nouveau/compiler/nak_nir_add_barriers.c
View file

@ -113,7 +113,7 @@ lower_control_barriers_block(nir_block *block,
"Control barrier with scope > WORKGROUP");
if (exec_scope == SCOPE_WORKGROUP &&
nak_nir_has_one_subgroup(state->builder.shader))
nak_nir_workgroup_has_one_subgroup(state->builder.shader))
exec_scope = SCOPE_SUBGROUP;
/* Because we're guaranteeing maximal convergence with this pass,

13
src/nouveau/compiler/nak_private.h
View file

@ -108,18 +108,7 @@ enum PACKED nak_sv {
NAK_SV_CLOCK = 0x50,
};
static bool
nak_nir_has_one_subgroup(const nir_shader *nir)
{
if (nir->info.workgroup_size_variable)
return false;
uint16_t wg_sz = nir->info.workgroup_size[0] *
nir->info.workgroup_size[1] *
nir->info.workgroup_size[2];
return wg_sz <= 32;
}
bool nak_nir_workgroup_has_one_subgroup(const nir_shader *nir);
struct nak_xfb_info
nak_xfb_from_nir(const struct nir_xfb_info *nir_xfb);