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);