Module: Mesa Branch: main Commit: 3bb4c14c759d8fee79a76af25e208cb059177bbd URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=3bb4c14c759d8fee79a76af25e208cb059177bbd
Author: Faith Ekstrand <[email protected]> Date: Thu Dec 7 12:49:28 2023 -0600 nak: Rework barrier handling a bit Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26577> --- src/nouveau/compiler/nak/from_nir.rs | 14 +++++------ src/nouveau/compiler/nak_nir.c | 39 +++++++++++++++++++++++++++-- src/nouveau/compiler/nak_nir_add_barriers.c | 2 +- src/nouveau/compiler/nak_private.h | 13 +--------- 4 files changed, 46 insertions(+), 22 deletions(-) diff --git a/src/nouveau/compiler/nak/from_nir.rs b/src/nouveau/compiler/nak/from_nir.rs index 873968b0e58..cfbe4380327 100644 --- a/src/nouveau/compiler/nak/from_nir.rs +++ b/src/nouveau/compiler/nak/from_nir.rs @@ -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"), } diff --git a/src/nouveau/compiler/nak_nir.c b/src/nouveau/compiler/nak_nir.c index ab3f155143e..c936907c60e 100644 --- a/src/nouveau/compiler/nak_nir.c +++ b/src/nouveau/compiler/nak_nir.c @@ -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); diff --git a/src/nouveau/compiler/nak_nir_add_barriers.c b/src/nouveau/compiler/nak_nir_add_barriers.c index f12b1a65522..328926cc150 100644 --- a/src/nouveau/compiler/nak_nir_add_barriers.c +++ b/src/nouveau/compiler/nak_nir_add_barriers.c @@ -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, diff --git a/src/nouveau/compiler/nak_private.h b/src/nouveau/compiler/nak_private.h index 99133fc335c..eaed9023ec1 100644 --- a/src/nouveau/compiler/nak_private.h +++ b/src/nouveau/compiler/nak_private.h @@ -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);
