mirror of
https://gitlab.freedesktop.org/mesa/mesa.git
synced 2026-05-05 18:18:06 +02:00
aco: move live var information into struct Program
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29713>
This commit is contained in:
Daniel Schürmann
Marge Bot
parent
2322ab427e
commit
a497d105e3
10 changed files with 84 additions and 95 deletions
|
|
@ -100,7 +100,6 @@ aco_postprocess_shader(const struct aco_compiler_options* options,
|
|||
ASSERTED bool is_valid = validate_cfg(program.get());
|
||||
assert(is_valid);
|
||||
|
||||
live live_vars;
|
||||
if (!info->is_trap_handler_shader) {
|
||||
dominator_tree(program.get());
|
||||
lower_phis(program.get());
|
||||
|
|
@ -124,10 +123,10 @@ aco_postprocess_shader(const struct aco_compiler_options* options,
|
|||
validate(program.get());
|
||||
|
||||
/* spilling and scheduling */
|
||||
live_vars = live_var_analysis(program.get());
|
||||
live_var_analysis(program.get());
|
||||
if (program->collect_statistics)
|
||||
collect_presched_stats(program.get());
|
||||
spill(program.get(), live_vars);
|
||||
spill(program.get());
|
||||
}
|
||||
|
||||
if (options->record_ir) {
|
||||
|
|
@ -146,15 +145,15 @@ aco_postprocess_shader(const struct aco_compiler_options* options,
|
|||
}
|
||||
|
||||
if ((debug_flags & DEBUG_LIVE_INFO) && options->dump_shader)
|
||||
aco_print_program(program.get(), stderr, live_vars, print_live_vars | print_kill);
|
||||
aco_print_program(program.get(), stderr, print_live_vars | print_kill);
|
||||
|
||||
if (!info->is_trap_handler_shader) {
|
||||
if (!options->optimisations_disabled && !(debug_flags & DEBUG_NO_SCHED))
|
||||
schedule_program(program.get(), live_vars);
|
||||
schedule_program(program.get());
|
||||
validate(program.get());
|
||||
|
||||
/* Register Allocation */
|
||||
register_allocation(program.get(), live_vars);
|
||||
register_allocation(program.get());
|
||||
|
||||
if (validate_ra(program.get())) {
|
||||
aco_print_program(program.get(), stderr);
|
||||
|
|
|
|||
|
|
@ -2063,6 +2063,13 @@ public:
|
|||
/* For shader part with previous shader part that has lds access. */
|
||||
bool pending_lds_access = false;
|
||||
|
||||
struct {
|
||||
/* live temps out per block */
|
||||
std::vector<IDSet> live_out;
|
||||
/* register demand (sgpr/vgpr) per instruction per block */
|
||||
std::vector<std::vector<RegisterDemand>> register_demand;
|
||||
} live;
|
||||
|
||||
struct {
|
||||
FILE* output = stderr;
|
||||
bool shorten_messages = false;
|
||||
|
|
@ -2112,13 +2119,6 @@ private:
|
|||
uint32_t allocationID = 1;
|
||||
};
|
||||
|
||||
struct live {
|
||||
/* live temps out per block */
|
||||
std::vector<IDSet> live_out;
|
||||
/* register demand (sgpr/vgpr) per instruction per block */
|
||||
std::vector<std::vector<RegisterDemand>> register_demand;
|
||||
};
|
||||
|
||||
struct ra_test_policy {
|
||||
/* Force RA to always use its pessimistic fallback algorithm */
|
||||
bool skip_optimistic_path = false;
|
||||
|
|
@ -2158,7 +2158,7 @@ void lower_phis(Program* program);
|
|||
void lower_subdword(Program* program);
|
||||
void calc_min_waves(Program* program);
|
||||
void update_vgpr_sgpr_demand(Program* program, const RegisterDemand new_demand);
|
||||
live live_var_analysis(Program* program);
|
||||
void live_var_analysis(Program* program);
|
||||
std::vector<uint16_t> dead_code_analysis(Program* program);
|
||||
void dominator_tree(Program* program);
|
||||
void insert_exec_mask(Program* program);
|
||||
|
|
@ -2166,14 +2166,14 @@ void value_numbering(Program* program);
|
|||
void optimize(Program* program);
|
||||
void optimize_postRA(Program* program);
|
||||
void setup_reduce_temp(Program* program);
|
||||
void lower_to_cssa(Program* program, live& live_vars);
|
||||
void register_allocation(Program* program, live& live_vars, ra_test_policy = {});
|
||||
void lower_to_cssa(Program* program);
|
||||
void register_allocation(Program* program, ra_test_policy = {});
|
||||
void ssa_elimination(Program* program);
|
||||
void lower_to_hw_instr(Program* program);
|
||||
void schedule_program(Program* program, live& live_vars);
|
||||
void schedule_program(Program* program);
|
||||
void schedule_ilp(Program* program);
|
||||
void schedule_vopd(Program* program);
|
||||
void spill(Program* program, live& live_vars);
|
||||
void spill(Program* program);
|
||||
void insert_wait_states(Program* program);
|
||||
bool dealloc_vgprs(Program* program);
|
||||
void insert_NOPs(Program* program);
|
||||
|
|
@ -2216,8 +2216,6 @@ void aco_print_operand(const Operand* operand, FILE* output, unsigned flags = 0)
|
|||
void aco_print_instr(enum amd_gfx_level gfx_level, const Instruction* instr, FILE* output,
|
||||
unsigned flags = 0);
|
||||
void aco_print_program(const Program* program, FILE* output, unsigned flags = 0);
|
||||
void aco_print_program(const Program* program, FILE* output, const live& live_vars,
|
||||
unsigned flags = 0);
|
||||
|
||||
void _aco_err(Program* program, const char* file, unsigned line, const char* fmt, ...);
|
||||
|
||||
|
|
|
|||
|
|
@ -107,14 +107,14 @@ instr_needs_vcc(Instruction* instr)
|
|||
}
|
||||
|
||||
void
|
||||
process_live_temps_per_block(Program* program, live& lives, Block* block, unsigned& worklist,
|
||||
process_live_temps_per_block(Program* program, Block* block, unsigned& worklist,
|
||||
std::vector<PhiInfo>& phi_info)
|
||||
{
|
||||
std::vector<RegisterDemand>& register_demand = lives.register_demand[block->index];
|
||||
std::vector<RegisterDemand>& register_demand = program->live.register_demand[block->index];
|
||||
RegisterDemand new_demand;
|
||||
|
||||
register_demand.resize(block->instructions.size());
|
||||
IDSet live = lives.live_out[block->index];
|
||||
IDSet live = program->live.live_out[block->index];
|
||||
|
||||
/* initialize register demand */
|
||||
for (unsigned t : live)
|
||||
|
|
@ -239,7 +239,7 @@ process_live_temps_per_block(Program* program, live& lives, Block* block, unsign
|
|||
|
||||
if (fast_merge) {
|
||||
for (unsigned pred_idx : block->linear_preds) {
|
||||
if (lives.live_out[pred_idx].insert(live))
|
||||
if (program->live.live_out[pred_idx].insert(live))
|
||||
worklist = std::max(worklist, pred_idx + 1);
|
||||
}
|
||||
} else {
|
||||
|
|
@ -254,7 +254,7 @@ process_live_temps_per_block(Program* program, live& lives, Block* block, unsign
|
|||
#endif
|
||||
|
||||
for (unsigned pred_idx : preds) {
|
||||
auto it = lives.live_out[pred_idx].insert(t);
|
||||
auto it = program->live.live_out[pred_idx].insert(t);
|
||||
if (it.second)
|
||||
worklist = std::max(worklist, pred_idx + 1);
|
||||
}
|
||||
|
|
@ -276,7 +276,7 @@ process_live_temps_per_block(Program* program, live& lives, Block* block, unsign
|
|||
if (operand.isFixed() && operand.physReg() == vcc)
|
||||
program->needs_vcc = true;
|
||||
/* check if we changed an already processed block */
|
||||
const bool inserted = lives.live_out[preds[i]].insert(operand.tempId()).second;
|
||||
const bool inserted = program->live.live_out[preds[i]].insert(operand.tempId()).second;
|
||||
if (inserted) {
|
||||
worklist = std::max(worklist, preds[i] + 1);
|
||||
if (insn->opcode == aco_opcode::p_phi && operand.getTemp().type() == RegType::sgpr) {
|
||||
|
|
@ -455,12 +455,12 @@ update_vgpr_sgpr_demand(Program* program, const RegisterDemand new_demand)
|
|||
}
|
||||
}
|
||||
|
||||
live
|
||||
void
|
||||
live_var_analysis(Program* program)
|
||||
{
|
||||
live result;
|
||||
result.live_out.resize(program->blocks.size());
|
||||
result.register_demand.resize(program->blocks.size());
|
||||
program->live.live_out.clear();
|
||||
program->live.live_out.resize(program->blocks.size());
|
||||
program->live.register_demand.resize(program->blocks.size());
|
||||
unsigned worklist = program->blocks.size();
|
||||
std::vector<PhiInfo> phi_info(program->blocks.size());
|
||||
RegisterDemand new_demand;
|
||||
|
|
@ -471,19 +471,20 @@ live_var_analysis(Program* program)
|
|||
* program->blocks vector */
|
||||
while (worklist) {
|
||||
unsigned block_idx = --worklist;
|
||||
process_live_temps_per_block(program, result, &program->blocks[block_idx], worklist,
|
||||
phi_info);
|
||||
process_live_temps_per_block(program, &program->blocks[block_idx], worklist, phi_info);
|
||||
}
|
||||
|
||||
/* Handle branches: we will insert copies created for linear phis just before the branch. */
|
||||
for (Block& block : program->blocks) {
|
||||
result.register_demand[block.index].back().sgpr += phi_info[block.index].linear_phi_defs;
|
||||
result.register_demand[block.index].back().sgpr -= phi_info[block.index].linear_phi_ops;
|
||||
program->live.register_demand[block.index].back().sgpr +=
|
||||
phi_info[block.index].linear_phi_defs;
|
||||
program->live.register_demand[block.index].back().sgpr -=
|
||||
phi_info[block.index].linear_phi_ops;
|
||||
|
||||
/* update block's register demand */
|
||||
if (program->progress < CompilationProgress::after_ra) {
|
||||
block.register_demand = RegisterDemand();
|
||||
for (RegisterDemand& demand : result.register_demand[block.index])
|
||||
for (RegisterDemand& demand : program->live.register_demand[block.index])
|
||||
block.register_demand.update(demand);
|
||||
}
|
||||
|
||||
|
|
@ -493,8 +494,6 @@ live_var_analysis(Program* program)
|
|||
/* calculate the program's register demand and number of waves */
|
||||
if (program->progress < CompilationProgress::after_ra)
|
||||
update_vgpr_sgpr_demand(program, new_demand);
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
} // namespace aco
|
||||
|
|
|
|||
|
|
@ -514,14 +514,14 @@ emit_parallelcopies(cssa_ctx& ctx)
|
|||
} /* end namespace */
|
||||
|
||||
void
|
||||
lower_to_cssa(Program* program, live& live_vars)
|
||||
lower_to_cssa(Program* program)
|
||||
{
|
||||
reindex_ssa(program, live_vars.live_out);
|
||||
cssa_ctx ctx = {program, live_vars.live_out};
|
||||
reindex_ssa(program, program->live.live_out);
|
||||
cssa_ctx ctx = {program, program->live.live_out};
|
||||
collect_parallelcopies(ctx);
|
||||
emit_parallelcopies(ctx);
|
||||
|
||||
/* update live variable information */
|
||||
live_vars = live_var_analysis(program);
|
||||
live_var_analysis(program);
|
||||
}
|
||||
} // namespace aco
|
||||
|
|
|
|||
|
|
@ -1007,7 +1007,7 @@ print_stage(Stage stage, FILE* output)
|
|||
|
||||
void
|
||||
aco_print_block(enum amd_gfx_level gfx_level, const Block* block, FILE* output, unsigned flags,
|
||||
const live& live_vars)
|
||||
const Program* program)
|
||||
{
|
||||
fprintf(output, "BB%d\n", block->index);
|
||||
fprintf(output, "/* logical preds: ");
|
||||
|
|
@ -1022,7 +1022,7 @@ aco_print_block(enum amd_gfx_level gfx_level, const Block* block, FILE* output,
|
|||
|
||||
if (flags & print_live_vars) {
|
||||
fprintf(output, "\tlive out:");
|
||||
for (unsigned id : live_vars.live_out[block->index])
|
||||
for (unsigned id : program->live.live_out[block->index])
|
||||
fprintf(output, " %%%d", id);
|
||||
fprintf(output, "\n");
|
||||
|
||||
|
|
@ -1034,7 +1034,7 @@ aco_print_block(enum amd_gfx_level gfx_level, const Block* block, FILE* output,
|
|||
for (auto const& instr : block->instructions) {
|
||||
fprintf(output, "\t");
|
||||
if (flags & print_live_vars) {
|
||||
RegisterDemand demand = live_vars.register_demand[block->index][index];
|
||||
RegisterDemand demand = program->live.register_demand[block->index][index];
|
||||
fprintf(output, "(%3u vgpr, %3u sgpr) ", demand.vgpr, demand.sgpr);
|
||||
}
|
||||
if (flags & print_perf_info)
|
||||
|
|
@ -1047,7 +1047,7 @@ aco_print_block(enum amd_gfx_level gfx_level, const Block* block, FILE* output,
|
|||
}
|
||||
|
||||
void
|
||||
aco_print_program(const Program* program, FILE* output, const live& live_vars, unsigned flags)
|
||||
aco_print_program(const Program* program, FILE* output, unsigned flags)
|
||||
{
|
||||
switch (program->progress) {
|
||||
case CompilationProgress::after_isel: fprintf(output, "After Instruction Selection:\n"); break;
|
||||
|
|
@ -1061,7 +1061,7 @@ aco_print_program(const Program* program, FILE* output, const live& live_vars, u
|
|||
print_stage(program->stage, output);
|
||||
|
||||
for (Block const& block : program->blocks)
|
||||
aco_print_block(program->gfx_level, &block, output, flags, live_vars);
|
||||
aco_print_block(program->gfx_level, &block, output, flags, program);
|
||||
|
||||
if (program->constant_data.size()) {
|
||||
fprintf(output, "\n/* constant data */\n");
|
||||
|
|
@ -1081,10 +1081,4 @@ aco_print_program(const Program* program, FILE* output, const live& live_vars, u
|
|||
fprintf(output, "\n");
|
||||
}
|
||||
|
||||
void
|
||||
aco_print_program(const Program* program, FILE* output, unsigned flags)
|
||||
{
|
||||
aco_print_program(program, output, live(), flags);
|
||||
}
|
||||
|
||||
} // namespace aco
|
||||
|
|
|
|||
|
|
@ -1691,7 +1691,7 @@ alloc_linear_vgpr(ra_ctx& ctx, const RegisterFile& reg_file, aco_ptr<Instruction
|
|||
}
|
||||
|
||||
bool
|
||||
should_compact_linear_vgprs(ra_ctx& ctx, live& live_vars, const RegisterFile& reg_file)
|
||||
should_compact_linear_vgprs(ra_ctx& ctx, const RegisterFile& reg_file)
|
||||
{
|
||||
if (!(ctx.block->kind & block_kind_top_level) || ctx.block->linear_succs.empty())
|
||||
return false;
|
||||
|
|
@ -1709,7 +1709,7 @@ should_compact_linear_vgprs(ra_ctx& ctx, live& live_vars, const RegisterFile& re
|
|||
ctx.program->blocks[next_toplevel].instructions;
|
||||
if (!instructions.empty() && is_phi(instructions[0])) {
|
||||
max_vgpr_usage =
|
||||
MAX2(max_vgpr_usage, (unsigned)live_vars.register_demand[next_toplevel][0].vgpr);
|
||||
MAX2(max_vgpr_usage, (unsigned)ctx.program->live.register_demand[next_toplevel][0].vgpr);
|
||||
}
|
||||
|
||||
for (unsigned tmp : find_vars(ctx, reg_file, get_reg_bounds(ctx, RegType::vgpr, true)))
|
||||
|
|
@ -2971,9 +2971,9 @@ emit_parallel_copy(ra_ctx& ctx, std::vector<std::pair<Operand, Definition>>& par
|
|||
} /* end namespace */
|
||||
|
||||
void
|
||||
register_allocation(Program* program, live& live_vars, ra_test_policy policy)
|
||||
register_allocation(Program* program, ra_test_policy policy)
|
||||
{
|
||||
std::vector<IDSet>& live_out_per_block = live_vars.live_out;
|
||||
std::vector<IDSet>& live_out_per_block = program->live.live_out;
|
||||
ra_ctx ctx(program, policy);
|
||||
get_affinities(ctx, live_out_per_block);
|
||||
|
||||
|
|
@ -3342,7 +3342,7 @@ register_allocation(Program* program, live& live_vars, ra_test_policy policy)
|
|||
ASSERTED PhysRegInterval sgpr_bounds = get_reg_bounds(ctx, RegType::sgpr, false);
|
||||
assert(register_file.count_zero(vgpr_bounds) == ctx.vgpr_bounds);
|
||||
assert(register_file.count_zero(sgpr_bounds) == ctx.sgpr_bounds);
|
||||
} else if (should_compact_linear_vgprs(ctx, live_vars, register_file)) {
|
||||
} else if (should_compact_linear_vgprs(ctx, register_file)) {
|
||||
aco_ptr<Instruction> br = std::move(instructions.back());
|
||||
instructions.pop_back();
|
||||
|
||||
|
|
|
|||
|
|
@ -1165,12 +1165,12 @@ schedule_VMEM_store(sched_ctx& ctx, Block* block, Instruction* current, int idx)
|
|||
}
|
||||
|
||||
void
|
||||
schedule_block(sched_ctx& ctx, Program* program, Block* block, live& live_vars)
|
||||
schedule_block(sched_ctx& ctx, Program* program, Block* block)
|
||||
{
|
||||
ctx.last_SMEM_dep_idx = 0;
|
||||
ctx.last_SMEM_stall = INT16_MIN;
|
||||
ctx.mv.block = block;
|
||||
ctx.mv.register_demand = live_vars.register_demand[block->index].data();
|
||||
ctx.mv.register_demand = program->live.register_demand[block->index].data();
|
||||
|
||||
/* go through all instructions and find memory loads */
|
||||
unsigned num_stores = 0;
|
||||
|
|
@ -1224,12 +1224,12 @@ schedule_block(sched_ctx& ctx, Program* program, Block* block, live& live_vars)
|
|||
/* resummarize the block's register demand */
|
||||
block->register_demand = RegisterDemand();
|
||||
for (unsigned idx = 0; idx < block->instructions.size(); idx++) {
|
||||
block->register_demand.update(live_vars.register_demand[block->index][idx]);
|
||||
block->register_demand.update(program->live.register_demand[block->index][idx]);
|
||||
}
|
||||
}
|
||||
|
||||
void
|
||||
schedule_program(Program* program, live& live_vars)
|
||||
schedule_program(Program* program)
|
||||
{
|
||||
/* don't use program->max_reg_demand because that is affected by max_waves_per_simd */
|
||||
RegisterDemand demand;
|
||||
|
|
@ -1278,7 +1278,7 @@ schedule_program(Program* program, live& live_vars)
|
|||
}
|
||||
|
||||
for (Block& block : program->blocks)
|
||||
schedule_block(ctx, program, &block, live_vars);
|
||||
schedule_block(ctx, program, &block);
|
||||
|
||||
/* update max_reg_demand and num_waves */
|
||||
RegisterDemand new_demand;
|
||||
|
|
|
|||
|
|
@ -69,7 +69,6 @@ struct spill_ctx {
|
|||
Program* program;
|
||||
aco::monotonic_buffer_resource memory;
|
||||
|
||||
live& live_vars;
|
||||
std::vector<aco::map<Temp, Temp>> renames;
|
||||
std::vector<aco::unordered_map<Temp, uint32_t>> spills_entry;
|
||||
std::vector<aco::unordered_map<Temp, uint32_t>> spills_exit;
|
||||
|
|
@ -89,8 +88,8 @@ struct spill_ctx {
|
|||
unsigned vgpr_spill_slots;
|
||||
Temp scratch_rsrc;
|
||||
|
||||
spill_ctx(const RegisterDemand target_pressure_, Program* program_, live& live_vars_)
|
||||
: target_pressure(target_pressure_), program(program_), memory(), live_vars(live_vars_),
|
||||
spill_ctx(const RegisterDemand target_pressure_, Program* program_)
|
||||
: target_pressure(target_pressure_), program(program_), memory(),
|
||||
renames(program->blocks.size(), aco::map<Temp, Temp>(memory)),
|
||||
spills_entry(program->blocks.size(), aco::unordered_map<Temp, uint32_t>(memory)),
|
||||
spills_exit(program->blocks.size(), aco::unordered_map<Temp, uint32_t>(memory)),
|
||||
|
|
@ -173,7 +172,7 @@ gather_ssa_use_info(spill_ctx& ctx)
|
|||
{
|
||||
unsigned instruction_idx = 0;
|
||||
for (Block& block : ctx.program->blocks) {
|
||||
IDSet& live_set = ctx.live_vars.live_out[block.index];
|
||||
IDSet& live_set = ctx.program->live.live_out[block.index];
|
||||
|
||||
for (int i = block.instructions.size() - 1; i >= 0; i--) {
|
||||
aco_ptr<Instruction>& instr = block.instructions[i];
|
||||
|
|
@ -300,12 +299,12 @@ RegisterDemand
|
|||
get_demand_before(spill_ctx& ctx, unsigned block_idx, unsigned idx)
|
||||
{
|
||||
if (idx == 0) {
|
||||
RegisterDemand demand = ctx.live_vars.register_demand[block_idx][idx];
|
||||
RegisterDemand demand = ctx.program->live.register_demand[block_idx][idx];
|
||||
aco_ptr<Instruction>& instr = ctx.program->blocks[block_idx].instructions[idx];
|
||||
aco_ptr<Instruction> instr_before(nullptr);
|
||||
return get_demand_before(demand, instr, instr_before);
|
||||
} else {
|
||||
return ctx.live_vars.register_demand[block_idx][idx - 1];
|
||||
return ctx.program->live.register_demand[block_idx][idx - 1];
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -335,7 +334,7 @@ get_live_in_demand(spill_ctx& ctx, unsigned block_idx)
|
|||
* reg_pressure if the branch instructions define sgprs. */
|
||||
for (unsigned pred : block.linear_preds)
|
||||
reg_pressure.sgpr =
|
||||
std::max<int16_t>(reg_pressure.sgpr, ctx.live_vars.register_demand[pred].back().sgpr);
|
||||
std::max<int16_t>(reg_pressure.sgpr, ctx.program->live.register_demand[pred].back().sgpr);
|
||||
|
||||
return reg_pressure;
|
||||
}
|
||||
|
|
@ -350,7 +349,7 @@ init_live_in_vars(spill_ctx& ctx, Block* block, unsigned block_idx)
|
|||
return {0, 0};
|
||||
|
||||
/* live-in variables at the beginning of the current block */
|
||||
const IDSet& live_in = ctx.live_vars.live_out[block_idx];
|
||||
const IDSet& live_in = ctx.program->live.live_out[block_idx];
|
||||
|
||||
/* loop header block */
|
||||
if (block->kind & block_kind_loop_header) {
|
||||
|
|
@ -404,7 +403,8 @@ init_live_in_vars(spill_ctx& ctx, Block* block, unsigned block_idx)
|
|||
for (unsigned t : live_in) {
|
||||
Temp var = Temp(t, ctx.program->temp_rc[t]);
|
||||
if (var.type() != type || ctx.spills_entry[block_idx].count(var) ||
|
||||
!ctx.live_vars.live_out[block_idx - 1].count(t) || var.regClass().is_linear_vgpr())
|
||||
!ctx.program->live.live_out[block_idx - 1].count(t) ||
|
||||
var.regClass().is_linear_vgpr())
|
||||
continue;
|
||||
|
||||
unsigned can_remat = ctx.remat.count(var);
|
||||
|
|
@ -514,7 +514,7 @@ init_live_in_vars(spill_ctx& ctx, Block* block, unsigned block_idx)
|
|||
uint32_t spill_id = 0;
|
||||
for (unsigned pred_idx : preds) {
|
||||
/* variable is not even live at the predecessor: probably from a phi */
|
||||
if (!ctx.live_vars.live_out[pred_idx].count(t)) {
|
||||
if (!ctx.program->live.live_out[pred_idx].count(t)) {
|
||||
spill = false;
|
||||
break;
|
||||
}
|
||||
|
|
@ -618,7 +618,7 @@ add_coupling_code(spill_ctx& ctx, Block* block, IDSet& live_in)
|
|||
if (block->linear_preds.size() == 1 &&
|
||||
!(block->kind & (block_kind_loop_exit | block_kind_loop_header))) {
|
||||
assert(ctx.processed[block->linear_preds[0]]);
|
||||
assert(ctx.live_vars.register_demand[block_idx].size() == block->instructions.size());
|
||||
assert(ctx.program->live.register_demand[block_idx].size() == block->instructions.size());
|
||||
|
||||
ctx.renames[block_idx] = ctx.renames[block->linear_preds[0]];
|
||||
if (!block->logical_preds.empty() && block->logical_preds[0] != block->linear_preds[0]) {
|
||||
|
|
@ -724,7 +724,7 @@ add_coupling_code(spill_ctx& ctx, Block* block, IDSet& live_in)
|
|||
|
||||
for (unsigned pred_idx : preds) {
|
||||
/* variable is dead at predecessor, it must be from a phi: this works because of CSSA form */
|
||||
if (!ctx.live_vars.live_out[pred_idx].count(pair.first.id()))
|
||||
if (!ctx.program->live.live_out[pred_idx].count(pair.first.id()))
|
||||
continue;
|
||||
|
||||
/* variable is already spilled at predecessor */
|
||||
|
|
@ -846,9 +846,9 @@ add_coupling_code(spill_ctx& ctx, Block* block, IDSet& live_in)
|
|||
|
||||
Block::edge_vec& preds = rc.is_linear() ? block->linear_preds : block->logical_preds;
|
||||
/* if a variable is dead at any predecessor, it must be from a phi */
|
||||
const bool is_dead =
|
||||
std::any_of(preds.begin(), preds.end(),
|
||||
[&](unsigned pred) { return !ctx.live_vars.live_out[pred].count(var.id()); });
|
||||
const bool is_dead = std::any_of(
|
||||
preds.begin(), preds.end(),
|
||||
[&](unsigned pred) { return !ctx.program->live.live_out[pred].count(var.id()); });
|
||||
if (is_dead)
|
||||
continue;
|
||||
|
||||
|
|
@ -933,11 +933,10 @@ add_coupling_code(spill_ctx& ctx, Block* block, IDSet& live_in)
|
|||
if (!ctx.processed[block_idx]) {
|
||||
assert(!(block->kind & block_kind_loop_header));
|
||||
RegisterDemand demand_before = get_demand_before(ctx, block_idx, idx);
|
||||
ctx.live_vars.register_demand[block->index].erase(
|
||||
ctx.live_vars.register_demand[block->index].begin(),
|
||||
ctx.live_vars.register_demand[block->index].begin() + idx);
|
||||
ctx.live_vars.register_demand[block->index].insert(
|
||||
ctx.live_vars.register_demand[block->index].begin(), instructions.size(), demand_before);
|
||||
std::vector<RegisterDemand>& register_demand =
|
||||
ctx.program->live.register_demand[block->index];
|
||||
register_demand.erase(register_demand.begin(), register_demand.begin() + idx);
|
||||
register_demand.insert(register_demand.begin(), instructions.size(), demand_before);
|
||||
}
|
||||
|
||||
std::vector<aco_ptr<Instruction>>::iterator start = std::next(block->instructions.begin(), idx);
|
||||
|
|
@ -983,7 +982,7 @@ process_block(spill_ctx& ctx, unsigned block_idx, Block* block, RegisterDemand s
|
|||
continue;
|
||||
|
||||
if (op.isFirstKill())
|
||||
ctx.live_vars.live_out[block_idx].erase(op.tempId());
|
||||
ctx.program->live.live_out[block_idx].erase(op.tempId());
|
||||
ctx.ssa_infos[op.tempId()].num_uses--;
|
||||
|
||||
if (!current_spills.count(op.getTemp()))
|
||||
|
|
@ -1000,7 +999,7 @@ process_block(spill_ctx& ctx, unsigned block_idx, Block* block, RegisterDemand s
|
|||
/* check if register demand is low enough before and after the current instruction */
|
||||
if (block->register_demand.exceeds(ctx.target_pressure)) {
|
||||
|
||||
RegisterDemand new_demand = ctx.live_vars.register_demand[block_idx][idx];
|
||||
RegisterDemand new_demand = ctx.program->live.register_demand[block_idx][idx];
|
||||
new_demand.update(get_demand_before(ctx, block_idx, idx));
|
||||
|
||||
/* if reg pressure is too high, spill variable with furthest next use */
|
||||
|
|
@ -1013,7 +1012,7 @@ process_block(spill_ctx& ctx, unsigned block_idx, Block* block, RegisterDemand s
|
|||
if (new_demand.vgpr - spilled_registers.vgpr > ctx.target_pressure.vgpr)
|
||||
type = RegType::vgpr;
|
||||
|
||||
for (unsigned t : ctx.live_vars.live_out[block_idx]) {
|
||||
for (unsigned t : ctx.program->live.live_out[block_idx]) {
|
||||
RegClass rc = ctx.program->temp_rc[t];
|
||||
Temp var = Temp(t, rc);
|
||||
if (rc.type() != type || current_spills.count(var) || rc.is_linear_vgpr())
|
||||
|
|
@ -1071,7 +1070,7 @@ process_block(spill_ctx& ctx, unsigned block_idx, Block* block, RegisterDemand s
|
|||
|
||||
for (const Definition& def : instr->definitions) {
|
||||
if (def.isTemp() && !def.isKill())
|
||||
ctx.live_vars.live_out[block_idx].insert(def.tempId());
|
||||
ctx.program->live.live_out[block_idx].insert(def.tempId());
|
||||
}
|
||||
/* rename operands */
|
||||
for (Operand& op : instr->operands) {
|
||||
|
|
@ -1112,7 +1111,7 @@ spill_block(spill_ctx& ctx, unsigned block_idx)
|
|||
|
||||
if (!(block->kind & block_kind_loop_header)) {
|
||||
/* add spill/reload code on incoming control flow edges */
|
||||
add_coupling_code(ctx, block, ctx.live_vars.live_out[block_idx]);
|
||||
add_coupling_code(ctx, block, ctx.program->live.live_out[block_idx]);
|
||||
}
|
||||
|
||||
assert(ctx.spills_exit[block_idx].empty());
|
||||
|
|
@ -1691,7 +1690,7 @@ assign_spill_slots(spill_ctx& ctx, unsigned spills_to_vgpr)
|
|||
} /* end namespace */
|
||||
|
||||
void
|
||||
spill(Program* program, live& live_vars)
|
||||
spill(Program* program)
|
||||
{
|
||||
program->config->spilled_vgprs = 0;
|
||||
program->config->spilled_sgprs = 0;
|
||||
|
|
@ -1703,7 +1702,7 @@ spill(Program* program, live& live_vars)
|
|||
return;
|
||||
|
||||
/* lower to CSSA before spilling to ensure correctness w.r.t. phis */
|
||||
lower_to_cssa(program, live_vars);
|
||||
lower_to_cssa(program);
|
||||
|
||||
/* calculate target register demand */
|
||||
const RegisterDemand demand = program->max_reg_demand; /* current max */
|
||||
|
|
@ -1733,7 +1732,7 @@ spill(Program* program, live& live_vars)
|
|||
const RegisterDemand target(vgpr_limit - extra_vgprs, sgpr_limit - extra_sgprs);
|
||||
|
||||
/* initialize ctx */
|
||||
spill_ctx ctx(target, program, live_vars);
|
||||
spill_ctx ctx(target, program);
|
||||
gather_ssa_use_info(ctx);
|
||||
get_rematerialize_info(ctx);
|
||||
|
||||
|
|
@ -1745,7 +1744,7 @@ spill(Program* program, live& live_vars)
|
|||
assign_spill_slots(ctx, extra_vgprs);
|
||||
|
||||
/* update live variable information */
|
||||
live_vars = live_var_analysis(program);
|
||||
live_var_analysis(program);
|
||||
|
||||
assert(program->num_waves > 0);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1198,7 +1198,7 @@ validate_ra(Program* program)
|
|||
return false;
|
||||
|
||||
bool err = false;
|
||||
aco::live live_vars = aco::live_var_analysis(program);
|
||||
aco::live_var_analysis(program);
|
||||
std::vector<std::vector<Temp>> phi_sgpr_ops(program->blocks.size());
|
||||
uint16_t sgpr_limit = get_addr_sgpr_from_waves(program, program->num_waves);
|
||||
|
||||
|
|
@ -1287,7 +1287,7 @@ validate_ra(Program* program)
|
|||
std::array<unsigned, 2048> regs; /* register file in bytes */
|
||||
regs.fill(0);
|
||||
|
||||
IDSet live = live_vars.live_out[block.index];
|
||||
IDSet live = program->live.live_out[block.index];
|
||||
/* remove killed p_phi sgpr operands */
|
||||
for (Temp tmp : phi_sgpr_ops[block.index])
|
||||
live.erase(tmp.id());
|
||||
|
|
|
|||
|
|
@ -247,8 +247,8 @@ finish_ra_test(ra_test_policy policy)
|
|||
}
|
||||
|
||||
program->workgroup_size = program->wave_size;
|
||||
aco::live live_vars = aco::live_var_analysis(program.get());
|
||||
aco::register_allocation(program.get(), live_vars, policy);
|
||||
aco::live_var_analysis(program.get());
|
||||
aco::register_allocation(program.get(), policy);
|
||||
|
||||
if (aco::validate_ra(program.get())) {
|
||||
fail_test("Validation after register allocation failed");
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue