fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2025-12-20 20:20:18 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

aco: remove aco::rt_stack variable

Browse source 
Since we initialize scratch in the RT proglog,
there is no need for this variable anymore.

Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/21780>
This commit is contained in:
Daniel Schürmann 2023-03-15 00:14:07 +01:00 • committed by Marge Bot
parent f123d65e9f
commit 39c828cb9f
4 changed files with 2 additions and 5 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
src/amd/compiler/aco_instruction_selection.cpp
View file

@ -8934,7 +8934,6 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr)
case nir_intrinsic_load_rt_dynamic_callable_stack_base_amd: case nir_intrinsic_load_rt_dynamic_callable_stack_base_amd:
bld.copy(Definition(get_ssa_temp(ctx, &instr->dest.ssa)), bld.copy(Definition(get_ssa_temp(ctx, &instr->dest.ssa)),
get_arg(ctx, ctx->args->rt_dynamic_callable_stack_base)); get_arg(ctx, ctx->args->rt_dynamic_callable_stack_base));
ctx->program->rt_stack = true;
break; break;
case nir_intrinsic_overwrite_vs_arguments_amd: { case nir_intrinsic_overwrite_vs_arguments_amd: {
ctx->arg_temps[ctx->args->vertex_id.arg_index] = get_ssa_temp(ctx, instr->src[0].ssa); ctx->arg_temps[ctx->args->vertex_id.arg_index] = get_ssa_temp(ctx, instr->src[0].ssa);

1
src/amd/compiler/aco_ir.h
View file

@ -2096,7 +2096,6 @@ public:
uint16_t min_waves = 0; uint16_t min_waves = 0;
unsigned workgroup_size; /* if known; otherwise UINT_MAX */ unsigned workgroup_size; /* if known; otherwise UINT_MAX */
bool wgp_mode; bool wgp_mode;
bool rt_stack = false;
bool needs_vcc = false; bool needs_vcc = false;

3
src/amd/compiler/aco_live_var_analysis.cpp
View file

@ -310,8 +310,7 @@ uint16_t
get_extra_sgprs(Program* program) get_extra_sgprs(Program* program)
{ {
/* We don't use this register on GFX6-8 and it's removed on GFX10+. */ /* We don't use this register on GFX6-8 and it's removed on GFX10+. */
bool needs_flat_scr = bool needs_flat_scr = program->config->scratch_bytes_per_wave && program->gfx_level == GFX9;
(program->config->scratch_bytes_per_wave || program->rt_stack) && program->gfx_level == GFX9;
if (program->gfx_level >= GFX10) { if (program->gfx_level >= GFX10) {
assert(!program->dev.xnack_enabled); assert(!program->dev.xnack_enabled);

2
src/amd/compiler/aco_lower_to_hw_instr.cpp
View file

@ -2493,7 +2493,7 @@ lower_to_hw_instr(Program* program)
} }
case aco_opcode::p_init_scratch: { case aco_opcode::p_init_scratch: {
assert(program->gfx_level >= GFX8 && program->gfx_level <= GFX10_3); assert(program->gfx_level >= GFX8 && program->gfx_level <= GFX10_3);
if (!program->config->scratch_bytes_per_wave && !program->rt_stack) if (!program->config->scratch_bytes_per_wave)
break; break;
Operand scratch_addr = instr->operands[0]; Operand scratch_addr = instr->operands[0];