From 39c828cb9fdeddd0fa4b757105c85a87643bd018 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Daniel=20Sch=C3=BCrmann?= Date: Wed, 15 Mar 2023 00:14:07 +0100 Subject: [PATCH] aco: remove aco::rt_stack variable Since we initialize scratch in the RT proglog, there is no need for this variable anymore. Part-of: --- src/amd/compiler/aco_instruction_selection.cpp | 1 - src/amd/compiler/aco_ir.h | 1 - src/amd/compiler/aco_live_var_analysis.cpp | 3 +-- src/amd/compiler/aco_lower_to_hw_instr.cpp | 2 +- 4 files changed, 2 insertions(+), 5 deletions(-) diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index 8bab001066f..fb6e4cbc682 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -8934,7 +8934,6 @@ visit_intrinsic(isel_context* ctx, nir_intrinsic_instr* instr) case nir_intrinsic_load_rt_dynamic_callable_stack_base_amd: bld.copy(Definition(get_ssa_temp(ctx, &instr->dest.ssa)), get_arg(ctx, ctx->args->rt_dynamic_callable_stack_base)); - ctx->program->rt_stack = true; break; case nir_intrinsic_overwrite_vs_arguments_amd: { ctx->arg_temps[ctx->args->vertex_id.arg_index] = get_ssa_temp(ctx, instr->src[0].ssa); diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h index 6f42f09483b..847c789f180 100644 --- a/src/amd/compiler/aco_ir.h +++ b/src/amd/compiler/aco_ir.h @@ -2096,7 +2096,6 @@ public: uint16_t min_waves = 0; unsigned workgroup_size; /* if known; otherwise UINT_MAX */ bool wgp_mode; - bool rt_stack = false; bool needs_vcc = false; diff --git a/src/amd/compiler/aco_live_var_analysis.cpp b/src/amd/compiler/aco_live_var_analysis.cpp index 405ad5fdd75..0ecc9575f20 100644 --- a/src/amd/compiler/aco_live_var_analysis.cpp +++ b/src/amd/compiler/aco_live_var_analysis.cpp @@ -310,8 +310,7 @@ uint16_t get_extra_sgprs(Program* program) { /* We don't use this register on GFX6-8 and it's removed on GFX10+. */ - bool needs_flat_scr = - (program->config->scratch_bytes_per_wave || program->rt_stack) && program->gfx_level == GFX9; + bool needs_flat_scr = program->config->scratch_bytes_per_wave && program->gfx_level == GFX9; if (program->gfx_level >= GFX10) { assert(!program->dev.xnack_enabled); diff --git a/src/amd/compiler/aco_lower_to_hw_instr.cpp b/src/amd/compiler/aco_lower_to_hw_instr.cpp index 80db1367dbc..2d92b6321b0 100644 --- a/src/amd/compiler/aco_lower_to_hw_instr.cpp +++ b/src/amd/compiler/aco_lower_to_hw_instr.cpp @@ -2493,7 +2493,7 @@ lower_to_hw_instr(Program* program) } case aco_opcode::p_init_scratch: { 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; Operand scratch_addr = instr->operands[0];