Module: Mesa Branch: main Commit: aba16211a8a7f62a6022545a6240f0f2e217ec04 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=aba16211a8a7f62a6022545a6240f0f2e217ec04
Author: Samuel Pitoiset <[email protected]> Date: Tue Aug 15 10:52:53 2023 +0200 aco: disable shared VGPRs for non-monolithic shaders on GFX9+ For unmerged shaders on GFX9+, we would need to jump to the second shader part which means shared VGPRs can't be enabled. Signed-off-by: Samuel Pitoiset <[email protected]> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/24697> --- src/amd/compiler/aco_instruction_selection.cpp | 9 +++++---- src/amd/compiler/aco_shader_info.h | 1 + src/amd/vulkan/radv_aco_shader_info.h | 1 + 3 files changed, 7 insertions(+), 4 deletions(-) diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index b69cb238b15..17e738cb1ec 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -196,10 +196,11 @@ emit_bpermute(isel_context* ctx, Builder& bld, Temp index, Temp data) * of multiple binaries, because the VGPR use is not known when choosing * which registers to use for the shared VGPRs. */ - const bool avoid_shared_vgprs = ctx->options->gfx_level >= GFX10 && - ctx->options->gfx_level < GFX11 && - ctx->program->wave_size == 64 && - (ctx->program->info.has_epilog || ctx->stage == raytracing_cs); + const bool avoid_shared_vgprs = + ctx->options->gfx_level >= GFX10 && ctx->options->gfx_level < GFX11 && + ctx->program->wave_size == 64 && + (ctx->program->info.has_epilog || !ctx->program->info.is_monolithic || + ctx->stage == raytracing_cs); if (ctx->options->gfx_level <= GFX7 || avoid_shared_vgprs) { /* GFX6-7: there is no bpermute instruction */ diff --git a/src/amd/compiler/aco_shader_info.h b/src/amd/compiler/aco_shader_info.h index 241af7544f2..b713f1c5017 100644 --- a/src/amd/compiler/aco_shader_info.h +++ b/src/amd/compiler/aco_shader_info.h @@ -102,6 +102,7 @@ struct aco_shader_info { bool image_2d_view_of_3d; unsigned workgroup_size; bool has_epilog; /* Only for TCS or PS. */ + bool is_monolithic; struct { bool tcs_in_out_eq; uint64_t tcs_temp_only_input_mask; diff --git a/src/amd/vulkan/radv_aco_shader_info.h b/src/amd/vulkan/radv_aco_shader_info.h index 468e8d3b884..a8ee61ae749 100644 --- a/src/amd/vulkan/radv_aco_shader_info.h +++ b/src/amd/vulkan/radv_aco_shader_info.h @@ -49,6 +49,7 @@ radv_aco_convert_shader_info(struct aco_shader_info *aco_info, const struct radv ASSIGN_FIELD(has_ngg_early_prim_export); ASSIGN_FIELD(workgroup_size); ASSIGN_FIELD(has_epilog); + ASSIGN_FIELD(is_monolithic); ASSIGN_FIELD(vs.tcs_in_out_eq); ASSIGN_FIELD(vs.tcs_temp_only_input_mask); ASSIGN_FIELD(vs.has_prolog);
