From 5cb1f46fd19ab9dedae657e13df491cd4935b2be Mon Sep 17 00:00:00 2001 From: Caio Oliveira Date: Fri, 12 Jul 2024 16:36:39 -0700 Subject: [PATCH] intel/brw: Remove workgroup_size() helper from fs_visitor Reviewed-by: Ian Romanick Part-of: --- src/intel/compiler/brw_fs.cpp | 8 -------- src/intel/compiler/brw_fs.h | 2 -- src/intel/compiler/brw_fs_nir.cpp | 13 +++++++++++-- 3 files changed, 11 insertions(+), 12 deletions(-) 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) {