diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp
index f4b8c285b2b..9248176cb1a 100644
--- a/src/intel/compiler/brw_fs.cpp
+++ b/src/intel/compiler/brw_fs.cpp
@@ -1977,14 +1977,6 @@ brw_cs_get_dispatch_info(const struct intel_device_info *devinfo,
    return info;
 }
 
-unsigned
-fs_visitor::workgroup_size() const
-{
-   assert(gl_shader_stage_uses_workgroup(stage));
-   const struct brw_cs_prog_data *cs = brw_cs_prog_data(prog_data);
-   return cs->local_size[0] * cs->local_size[1] * cs->local_size[2];
-}
-
 bool brw_should_print_shader(const nir_shader *shader, uint64_t debug_flag)
 {
    return INTEL_DEBUG(debug_flag) && (!shader->info.internal || NIR_DEBUG(PRINT_INTERNAL));
diff --git a/src/intel/compiler/brw_fs.h b/src/intel/compiler/brw_fs.h
index 627a2f67a2d..ffb4059f379 100644
--- a/src/intel/compiler/brw_fs.h
+++ b/src/intel/compiler/brw_fs.h
@@ -466,8 +466,6 @@ public:
 
    struct shader_stats shader_stats;
 
-   unsigned workgroup_size() const;
-
    void debug_optimizer(const nir_shader *nir,
                         const char *pass_name,
                         int iteration, int pass_num) const;
diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp
index 8967f444ca5..c13a2ff0a15 100644
--- a/src/intel/compiler/brw_fs_nir.cpp
+++ b/src/intel/compiler/brw_fs_nir.cpp
@@ -4410,6 +4410,15 @@ fs_nir_emit_fs_intrinsic(nir_to_brw_state &ntb,
    }
 }
 
+static unsigned
+brw_workgroup_size(fs_visitor &s)
+{
+   assert(gl_shader_stage_uses_workgroup(s.stage));
+   assert(!s.nir->info.workgroup_size_variable);
+   const struct brw_cs_prog_data *cs = brw_cs_prog_data(s.prog_data);
+   return cs->local_size[0] * cs->local_size[1] * cs->local_size[2];
+}
+
 static void
 fs_nir_emit_cs_intrinsic(nir_to_brw_state &ntb,
                          nir_intrinsic_instr *instr)
@@ -4435,7 +4444,7 @@ fs_nir_emit_cs_intrinsic(nir_to_brw_state &ntb,
           * barrier just emit a scheduling fence, that will generate no code.
           */
          if (!s.nir->info.workgroup_size_variable &&
-             s.workgroup_size() <= s.dispatch_width) {
+             brw_workgroup_size(s) <= s.dispatch_width) {
             bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE);
             break;
          }
@@ -6264,7 +6273,7 @@ fs_nir_emit_intrinsic(nir_to_brw_state &ntb,
        * TODO: Check if applies for many HW threads sharing same Data Port.
        */
       if (!s.nir->info.workgroup_size_variable &&
-          slm_fence && s.workgroup_size() <= s.dispatch_width)
+          slm_fence && brw_workgroup_size(s) <= s.dispatch_width)
          slm_fence = false;
 
       switch (s.stage) {