intel/fs: Clean up variable group size handling in backend
Just use the information from NIR shader_info. Reviewed-by: Kenneth Graunke <kenneth@whitecape.org> Reviewed-by: Jordan Justen <jordan.l.justen@intel.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/4794>
This commit is contained in:
parent
1800e4b58c
commit
0edb58a84e
|
@ -917,12 +917,10 @@ struct brw_cs_prog_data {
|
||||||
struct brw_stage_prog_data base;
|
struct brw_stage_prog_data base;
|
||||||
|
|
||||||
unsigned local_size[3];
|
unsigned local_size[3];
|
||||||
unsigned max_variable_local_size;
|
|
||||||
unsigned simd_size;
|
unsigned simd_size;
|
||||||
unsigned slm_size;
|
unsigned slm_size;
|
||||||
bool uses_barrier;
|
bool uses_barrier;
|
||||||
bool uses_num_work_groups;
|
bool uses_num_work_groups;
|
||||||
bool uses_variable_group_size;
|
|
||||||
|
|
||||||
struct {
|
struct {
|
||||||
struct brw_push_const_block cross_thread;
|
struct brw_push_const_block cross_thread;
|
||||||
|
|
|
@ -8981,9 +8981,7 @@ brw_compile_cs(const struct brw_compiler *compiler, void *log_data,
|
||||||
prog_data->slm_size = src_shader->num_shared;
|
prog_data->slm_size = src_shader->num_shared;
|
||||||
|
|
||||||
unsigned local_workgroup_size;
|
unsigned local_workgroup_size;
|
||||||
if (prog_data->uses_variable_group_size) {
|
if (src_shader->info.cs.local_size_variable) {
|
||||||
prog_data->max_variable_local_size =
|
|
||||||
src_shader->info.cs.max_variable_local_size;
|
|
||||||
local_workgroup_size = src_shader->info.cs.max_variable_local_size;
|
local_workgroup_size = src_shader->info.cs.max_variable_local_size;
|
||||||
} else {
|
} else {
|
||||||
prog_data->local_size[0] = src_shader->info.cs.local_size[0];
|
prog_data->local_size[0] = src_shader->info.cs.local_size[0];
|
||||||
|
|
|
@ -105,7 +105,7 @@ fs_visitor::nir_setup_uniforms()
|
||||||
assert(uniforms == prog_data->nr_params);
|
assert(uniforms == prog_data->nr_params);
|
||||||
|
|
||||||
uint32_t *param;
|
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);
|
param = brw_stage_prog_data_add_params(prog_data, 3);
|
||||||
for (unsigned i = 0; i < 3; i++) {
|
for (unsigned i = 0; i < 3; i++) {
|
||||||
param[i] = (BRW_PARAM_BUILTIN_WORK_GROUP_SIZE_X + 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
|
* invocations are already executed lock-step. Instead of an actual
|
||||||
* barrier just emit a scheduling fence, that will generate no code.
|
* 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) {
|
workgroup_size() <= dispatch_width) {
|
||||||
bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE);
|
bld.exec_all().group(1, 0).emit(FS_OPCODE_SCHEDULING_FENCE);
|
||||||
break;
|
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.
|
* 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 && workgroup_size() <= dispatch_width)
|
||||||
slm_fence = false;
|
slm_fence = false;
|
||||||
|
|
||||||
|
|
|
@ -114,11 +114,8 @@ brw_codegen_cs_prog(struct brw_context *brw,
|
||||||
* the actual size is not known until the dispatch command is issued.
|
* the actual size is not known until the dispatch command is issued.
|
||||||
*/
|
*/
|
||||||
if (nir->info.cs.local_size_variable) {
|
if (nir->info.cs.local_size_variable) {
|
||||||
prog_data.uses_variable_group_size = true;
|
|
||||||
nir->info.cs.max_variable_local_size =
|
nir->info.cs.max_variable_local_size =
|
||||||
gl_ctx->Const.MaxComputeVariableGroupInvocations;
|
gl_ctx->Const.MaxComputeVariableGroupInvocations;
|
||||||
} else {
|
|
||||||
prog_data.uses_variable_group_size = false;
|
|
||||||
}
|
}
|
||||||
|
|
||||||
char *error_str;
|
char *error_str;
|
||||||
|
|
Loading…
Reference in New Issue