aco: remove aco::rt_stack variable

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>

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);

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;
 

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);

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];