diff options
| author | Caio Marcelo de Oliveira Filho <[email protected]> | 2020-04-28 21:04:04 -0700 |
|---|---|---|
| committer | Caio Marcelo de Oliveira Filho <[email protected]> | 2020-05-01 12:50:28 -0700 |
| commit | 0edb58a84eb4a2b74b1ce55fea9dc06386c56bf6 (patch) | |
| tree | 57c7a2805b31d2ba550cb5fb8bf07003b58a837a /src/intel/compiler | |
| parent | 1800e4b58caaa89acfe45c95d0d22e533b50ee03 (diff) | |
intel/fs: Clean up variable group size handling in backend
Just use the information from NIR shader_info.
Reviewed-by: Kenneth Graunke <[email protected]>
Reviewed-by: Jordan Justen <[email protected]>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4794>
Diffstat (limited to 'src/intel/compiler')
| -rw-r--r-- | src/intel/compiler/brw_compiler.h | 2 | ||||
| -rw-r--r-- | src/intel/compiler/brw_fs.cpp | 4 | ||||
| -rw-r--r-- | src/intel/compiler/brw_fs_nir.cpp | 6 |
3 files changed, 4 insertions, 8 deletions
diff --git a/src/intel/compiler/brw_compiler.h b/src/intel/compiler/brw_compiler.h index ab39af22684..1045ef5076f 100644 --- a/src/intel/compiler/brw_compiler.h +++ b/src/intel/compiler/brw_compiler.h @@ -917,12 +917,10 @@ struct brw_cs_prog_data { struct brw_stage_prog_data base; unsigned local_size[3]; - unsigned max_variable_local_size; unsigned simd_size; unsigned slm_size; bool uses_barrier; bool uses_num_work_groups; - bool uses_variable_group_size; struct { struct brw_push_const_block cross_thread; diff --git a/src/intel/compiler/brw_fs.cpp b/src/intel/compiler/brw_fs.cpp index ccefdb081fd..d22d2c7a905 100644 --- a/src/intel/compiler/brw_fs.cpp +++ b/src/intel/compiler/brw_fs.cpp @@ -8981,9 +8981,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data, prog_data->slm_size = src_shader->num_shared; unsigned local_workgroup_size; - if (prog_data->uses_variable_group_size) { - prog_data->max_variable_local_size = - src_shader->info.cs.max_variable_local_size; + if (src_shader->info.cs.local_size_variable) { local_workgroup_size = src_shader->info.cs.max_variable_local_size; } else { prog_data->local_size[0] = src_shader->info.cs.local_size[0]; diff --git a/src/intel/compiler/brw_fs_nir.cpp b/src/intel/compiler/brw_fs_nir.cpp index e3149f6254c..e4fbaa51050 100644 --- a/src/intel/compiler/brw_fs_nir.cpp +++ b/src/intel/compiler/brw_fs_nir.cpp @@ -105,7 +105,7 @@ fs_visitor::nir_setup_uniforms() assert(uniforms == prog_data->nr_params); uint32_t *param; - if (brw_cs_prog_data(prog_data)->uses_variable_group_size) { + if (nir->info.cs.local_size_variable) { param = brw_stage_prog_data_add_params(prog_data, 3); for (unsigned i = 0; i < 3; i++) { param[i] = (BRW_PARAM_BUILTIN_WORK_GROUP_SIZE_X + i); @@ -3732,7 +3732,7 @@ fs_visitor::nir_emit_cs_intrinsic(const fs_builder &bld, * invocations are already executed lock-step. Instead of an actual * barrier just emit a scheduling fence, that will generate no code. */ - if (!cs_prog_data->uses_variable_group_size && + if (!nir->info.cs.local_size_variable && workgroup_size() <= dispatch_width) { bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE); break; @@ -4297,7 +4297,7 @@ fs_visitor::nir_emit_intrinsic(const fs_builder &bld, nir_intrinsic_instr *instr * * TODO: Check if applies for many HW threads sharing same Data Port. */ - if (!brw_cs_prog_data(prog_data)->uses_variable_group_size && + if (!nir->info.cs.local_size_variable && slm_fence && workgroup_size() <= dispatch_width) slm_fence = false; |