Edit
IABSD.fr/xenocara/lib/mesa/src/amd/compiler/aco_live_var_analysis.cpp
jsg
- lib/mesa/src/amd/compiler/aco_live_var_analysis.cpp
-
/* * Copyright © 2018 Valve Corporation * Copyright © 2018 Google * * SPDX-License-Identifier: MIT */ #include "aco_ir.h" namespace aco { RegisterDemand get_live_changes(Instruction* instr) { RegisterDemand changes; for (const Definition& def : instr->definitions) { if (!def.isTemp() || def.isKill()) continue; changes += def.getTemp(); } for (const Operand& op : instr->operands) { if (!op.isTemp() || !op.isFirstKill()) continue; changes -= op.getTemp(); } return changes; } RegisterDemand get_temp_registers(Instruction* instr) { RegisterDemand demand_before; RegisterDemand demand_after; for (Definition def : instr->definitions) { if (def.isKill()) demand_after += def.getTemp(); else if (def.isTemp()) demand_before -= def.getTemp(); } for (Operand op : instr->operands) { if (op.isFirstKill() || op.isCopyKill()) { demand_before += op.getTemp(); if (op.isLateKill()) demand_after += op.getTemp(); } else if (op.isClobbered() && !op.isKill()) { demand_before += op.getTemp(); } } demand_after.update(demand_before); return demand_after; } namespace { struct live_ctx { monotonic_buffer_resource m; Program* program; int32_t worklist; uint32_t handled_once; }; bool instr_needs_vcc(Instruction* instr) { if (instr->isVOPC()) return true; if (instr->isVOP2() && !instr->isVOP3()) { if (instr->operands.size() == 3 && instr->operands[2].isTemp() && instr->operands[2].regClass().type() == RegType::sgpr) return true; if (instr->definitions.size() == 2) return true; } return false; } IDSet compute_live_out(live_ctx& ctx, Block* block) { IDSet live(ctx.m); if (block->logical_succs.empty()) { /* Linear blocks: * Directly insert the successor if it is a linear block as well. */ for (unsigned succ : block->linear_succs) { if (ctx.program->blocks[succ].logical_preds.empty()) { live.insert(ctx.program->live.live_in[succ]); } else { for (unsigned t : ctx.program->live.live_in[succ]) { if (ctx.program->temp_rc[t].is_linear()) live.insert(t); } } } } else { /* Logical blocks: * Linear successors are either linear blocks or logical targets. */ live = IDSet(ctx.program->live.live_in[block->linear_succs[0]], ctx.m); if (block->linear_succs.size() == 2) live.insert(ctx.program->live.live_in[block->linear_succs[1]]); /* At most one logical target needs a separate insertion. */ if (block->logical_succs.back() != block->linear_succs.back()) { for (unsigned t : ctx.program->live.live_in[block->logical_succs.back()]) { if (!ctx.program->temp_rc[t].is_linear()) live.insert(t); } } else { assert(block->logical_succs[0] == block->linear_succs[0]); } } /* Handle phi operands */ if (block->linear_succs.size() == 1 && block->linear_succs[0] >= ctx.handled_once) { Block& succ = ctx.program->blocks[block->linear_succs[0]]; auto it = std::find(succ.linear_preds.begin(), succ.linear_preds.end(), block->index); unsigned op_idx = std::distance(succ.linear_preds.begin(), it); for (aco_ptr<Instruction>& phi : succ.instructions) { if (!is_phi(phi)) break; if (phi->opcode == aco_opcode::p_phi || phi->definitions[0].isKill()) continue; if (phi->operands[op_idx].isTemp()) live.insert(phi->operands[op_idx].tempId()); } } if (block->logical_succs.size() == 1 && block->logical_succs[0] >= ctx.handled_once) { Block& succ = ctx.program->blocks[block->logical_succs[0]]; auto it = std::find(succ.logical_preds.begin(), succ.logical_preds.end(), block->index); unsigned op_idx = std::distance(succ.logical_preds.begin(), it); for (aco_ptr<Instruction>& phi : succ.instructions) { if (!is_phi(phi)) break; if (phi->opcode == aco_opcode::p_linear_phi || phi->definitions[0].isKill()) continue; if (phi->operands[op_idx].isTemp()) live.insert(phi->operands[op_idx].tempId()); } } return live; } void process_live_temps_per_block(live_ctx& ctx, Block* block) { RegisterDemand new_demand; block->register_demand = RegisterDemand(); IDSet live = compute_live_out(ctx, block); /* initialize register demand */ for (unsigned t : live) new_demand += Temp(t, ctx.program->temp_rc[t]); /* traverse the instructions backwards */ int idx; for (idx = block->instructions.size() - 1; idx >= 0; idx--) { Instruction* insn = block->instructions[idx].get(); if (is_phi(insn)) break; ctx.program->needs_vcc |= instr_needs_vcc(insn); insn->register_demand = RegisterDemand(new_demand.vgpr, new_demand.sgpr); bool has_vgpr_def = false; /* KILL */ for (Definition& definition : insn->definitions) { has_vgpr_def |= definition.regClass().type() == RegType::vgpr && !definition.regClass().is_linear_vgpr(); if (!definition.isTemp()) { continue; } if (definition.isFixed() && definition.physReg() == vcc) ctx.program->needs_vcc = true; const Temp temp = definition.getTemp(); const size_t n = live.erase(temp.id()); if (n) { new_demand -= temp; definition.setKill(false); } else { insn->register_demand += temp; definition.setKill(true); } } if (ctx.program->gfx_level >= GFX10 && insn->isVALU() && insn->definitions.back().regClass() == s2) { /* RDNA2 ISA doc, 6.2.4. Wave64 Destination Restrictions: * The first pass of a wave64 VALU instruction may not overwrite a scalar value used by * the second half. */ bool carry_in = insn->opcode == aco_opcode::v_addc_co_u32 || insn->opcode == aco_opcode::v_subb_co_u32 || insn->opcode == aco_opcode::v_subbrev_co_u32; for (unsigned op_idx = 0; op_idx < (carry_in ? 2 : insn->operands.size()); op_idx++) { if (insn->operands[op_idx].isOfType(RegType::sgpr)) insn->operands[op_idx].setLateKill(true); } } else if (insn->opcode == aco_opcode::p_bpermute_readlane || insn->opcode == aco_opcode::p_bpermute_permlane || insn->opcode == aco_opcode::p_bpermute_shared_vgpr || insn->opcode == aco_opcode::p_dual_src_export_gfx11 || insn->opcode == aco_opcode::v_mqsad_u32_u8) { for (Operand& op : insn->operands) op.setLateKill(true); } else if (insn->opcode == aco_opcode::p_interp_gfx11 && insn->operands.size() == 7) { insn->operands[5].setLateKill(true); /* we re-use the destination reg in the middle */ } else if (insn->opcode == aco_opcode::v_interp_p1_f32 && ctx.program->dev.has_16bank_lds) { insn->operands[0].setLateKill(true); } else if (insn->opcode == aco_opcode::p_init_scratch) { insn->operands.back().setLateKill(true); } else if (instr_info.classes[(int)insn->opcode] == instr_class::wmma) { insn->operands[0].setLateKill(true); insn->operands[1].setLateKill(true); } /* Check if a definition clobbers some operand */ int op_idx = get_op_fixed_to_def(insn); if (op_idx != -1) insn->operands[op_idx].setClobbered(true); /* we need to do this in a separate loop because the next one can * setKill() for several operands at once and we don't want to * overwrite that in a later iteration */ for (Operand& op : insn->operands) { op.setKill(false); /* Linear vgprs must be late kill: this is to ensure linear VGPR operands and * normal VGPR definitions don't try to use the same register, which is problematic * because of assignment restrictions. */ if (op.hasRegClass() && op.regClass().is_linear_vgpr() && !op.isUndefined() && has_vgpr_def) op.setLateKill(true); } /* GEN */ RegisterDemand operand_demand; for (unsigned i = 0; i < insn->operands.size(); ++i) { Operand& operand = insn->operands[i]; if (!operand.isTemp()) continue; const Temp temp = operand.getTemp(); if (operand.isPrecolored()) { assert(!operand.isLateKill()); ctx.program->needs_vcc |= operand.physReg() == vcc; /* Check if this operand gets overwritten by a precolored definition. */ if (std::any_of(insn->definitions.begin(), insn->definitions.end(), [=](Definition def) { return def.isFixed() && def.physReg() + def.size() > operand.physReg() && operand.physReg() + operand.size() > def.physReg(); })) operand.setClobbered(true); /* Check if another precolored operand uses the same temporary. * This assumes that operands of one instruction are not precolored twice to * the same register. In this case, register pressure might be overestimated. */ for (unsigned j = i + 1; !operand.isCopyKill() && j < insn->operands.size(); ++j) { if (insn->operands[j].isPrecolored() && insn->operands[j].getTemp() == temp) { operand_demand += temp; insn->operands[j].setCopyKill(true); } } } if (operand.isKill()) continue; if (live.insert(temp.id()).second) { operand.setFirstKill(true); for (unsigned j = i + 1; j < insn->operands.size(); ++j) { if (insn->operands[j].isTemp() && insn->operands[j].getTemp() == temp) insn->operands[j].setKill(true); } if (operand.isLateKill()) insn->register_demand += temp; new_demand += temp; } else if (operand.isClobbered()) { operand_demand += temp; } } operand_demand += new_demand; insn->register_demand.update(operand_demand); block->register_demand.update(insn->register_demand); } /* handle phi definitions */ for (int phi_idx = 0; phi_idx <= idx; phi_idx++) { Instruction* insn = block->instructions[phi_idx].get(); insn->register_demand = new_demand; assert(is_phi(insn) && insn->definitions.size() == 1); if (!insn->definitions[0].isTemp()) { assert(insn->definitions[0].isFixed() && insn->definitions[0].physReg() == exec); continue; } Definition& definition = insn->definitions[0]; ctx.program->needs_vcc |= definition.isFixed() && definition.physReg() == vcc; const size_t n = live.erase(definition.tempId()); if (n && (definition.isKill() || ctx.handled_once > block->index)) { Block::edge_vec& preds = insn->opcode == aco_opcode::p_phi ? block->logical_preds : block->linear_preds; for (unsigned i = 0; i < preds.size(); i++) { if (insn->operands[i].isTemp()) ctx.worklist = std::max<int>(ctx.worklist, preds[i]); } } definition.setKill(!n); } /* handle phi operands */ for (int phi_idx = 0; phi_idx <= idx; phi_idx++) { Instruction* insn = block->instructions[phi_idx].get(); assert(is_phi(insn)); /* Ignore dead phis. */ if (insn->definitions[0].isKill()) continue; for (Operand& operand : insn->operands) { if (!operand.isTemp()) continue; /* set if the operand is killed by this (or another) phi instruction */ operand.setKill(!live.count(operand.tempId())); } } if (ctx.program->live.live_in[block->index].insert(live)) { if (block->linear_preds.size()) { assert(block->logical_preds.empty() || block->logical_preds.back() <= block->linear_preds.back()); ctx.worklist = std::max<int>(ctx.worklist, block->linear_preds.back()); } else { ASSERTED bool is_valid = validate_ir(ctx.program); assert(!is_valid); } } block->live_in_demand = new_demand; block->register_demand.update(block->live_in_demand); ctx.program->max_reg_demand.update(block->register_demand); ctx.handled_once = std::min(ctx.handled_once, block->index); assert(!block->linear_preds.empty() || (new_demand == RegisterDemand() && live.empty())); } unsigned calc_waves_per_workgroup(Program* program) { /* When workgroup size is not known, just go with wave_size */ unsigned workgroup_size = program->workgroup_size == UINT_MAX ? program->wave_size : program->workgroup_size; return align(workgroup_size, program->wave_size) / program->wave_size; } } /* end namespace */ bool uses_scratch(Program* program) { /* RT uses scratch but we don't yet know how much. */ return program->config->scratch_bytes_per_wave || program->stage == raytracing_cs; } 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 = uses_scratch(program) && program->gfx_level == GFX9; if (program->gfx_level >= GFX10) { assert(!program->dev.xnack_enabled); return 0; } else if (program->gfx_level >= GFX8) { if (needs_flat_scr) return 6; else if (program->dev.xnack_enabled) return 4; else if (program->needs_vcc) return 2; else return 0; } else { assert(!program->dev.xnack_enabled); if (needs_flat_scr) return 4; else if (program->needs_vcc) return 2; else return 0; } } uint16_t get_sgpr_alloc(Program* program, uint16_t addressable_sgprs) { uint16_t sgprs = addressable_sgprs + get_extra_sgprs(program); uint16_t granule = program->dev.sgpr_alloc_granule; return ALIGN_NPOT(std::max(sgprs, granule), granule); } uint16_t get_vgpr_alloc(Program* program, uint16_t addressable_vgprs) { assert(addressable_vgprs <= program->dev.vgpr_limit); uint16_t granule = program->dev.vgpr_alloc_granule; return ALIGN_NPOT(std::max(addressable_vgprs, granule), granule); } unsigned round_down(unsigned a, unsigned b) { return a - (a % b); } uint16_t get_addr_sgpr_from_waves(Program* program, uint16_t waves) { /* it's not possible to allocate more than 128 SGPRs */ uint16_t sgprs = std::min(program->dev.physical_sgprs / waves, 128); sgprs = round_down(sgprs, program->dev.sgpr_alloc_granule); sgprs -= get_extra_sgprs(program); return std::min(sgprs, program->dev.sgpr_limit); } uint16_t get_addr_vgpr_from_waves(Program* program, uint16_t waves) { uint16_t vgprs = program->dev.physical_vgprs / waves; vgprs = vgprs / program->dev.vgpr_alloc_granule * program->dev.vgpr_alloc_granule; vgprs -= program->config->num_shared_vgprs / 2; return std::min(vgprs, program->dev.vgpr_limit); } void calc_min_waves(Program* program) { unsigned waves_per_workgroup = calc_waves_per_workgroup(program); unsigned simd_per_cu_wgp = program->dev.simd_per_cu * (program->wgp_mode ? 2 : 1); program->min_waves = DIV_ROUND_UP(waves_per_workgroup, simd_per_cu_wgp); } uint16_t max_suitable_waves(Program* program, uint16_t waves) { unsigned num_simd = program->dev.simd_per_cu * (program->wgp_mode ? 2 : 1); unsigned waves_per_workgroup = calc_waves_per_workgroup(program); unsigned num_workgroups = waves * num_simd / waves_per_workgroup; /* Adjust #workgroups for LDS */ unsigned lds_per_workgroup = align(program->config->lds_size * program->dev.lds_encoding_granule, program->dev.lds_alloc_granule); if (program->stage == fragment_fs) { /* PS inputs are moved from PC (parameter cache) to LDS before PS waves are launched. * Each PS input occupies 3x vec4 of LDS space. See Figure 10.3 in GCN3 ISA manual. * These limit occupancy the same way as other stages' LDS usage does. */ unsigned lds_bytes_per_interp = 3 * 16; unsigned lds_param_bytes = lds_bytes_per_interp * program->info.ps.num_inputs; lds_per_workgroup += align(lds_param_bytes, program->dev.lds_alloc_granule); } unsigned lds_limit = program->wgp_mode ? program->dev.lds_limit * 2 : program->dev.lds_limit; if (lds_per_workgroup) num_workgroups = std::min(num_workgroups, lds_limit / lds_per_workgroup); /* Hardware limitation */ if (waves_per_workgroup > 1) num_workgroups = std::min(num_workgroups, program->wgp_mode ? 32u : 16u); /* Adjust #waves for workgroup multiples: * In cases like waves_per_workgroup=3 or lds=65536 and * waves_per_workgroup=1, we want the maximum possible number of waves per * SIMD and not the minimum. so DIV_ROUND_UP is used */ unsigned workgroup_waves = num_workgroups * waves_per_workgroup; return DIV_ROUND_UP(workgroup_waves, num_simd); } void update_vgpr_sgpr_demand(Program* program, const RegisterDemand new_demand) { assert(program->min_waves >= 1); uint16_t sgpr_limit = get_addr_sgpr_from_waves(program, program->min_waves); uint16_t vgpr_limit = get_addr_vgpr_from_waves(program, program->min_waves); /* this won't compile, register pressure reduction necessary */ if (new_demand.vgpr > vgpr_limit || new_demand.sgpr > sgpr_limit) { program->num_waves = 0; program->max_reg_demand = new_demand; } else { program->num_waves = program->dev.physical_sgprs / get_sgpr_alloc(program, new_demand.sgpr); uint16_t vgpr_demand = get_vgpr_alloc(program, new_demand.vgpr) + program->config->num_shared_vgprs / 2; program->num_waves = std::min<uint16_t>(program->num_waves, program->dev.physical_vgprs / vgpr_demand); program->num_waves = std::min(program->num_waves, program->dev.max_waves_per_simd); /* Adjust for LDS and workgroup multiples and calculate max_reg_demand */ program->num_waves = max_suitable_waves(program, program->num_waves); program->max_reg_demand.vgpr = get_addr_vgpr_from_waves(program, program->num_waves); program->max_reg_demand.sgpr = get_addr_sgpr_from_waves(program, program->num_waves); } } void live_var_analysis(Program* program) { program->live.live_in.clear(); program->live.memory.release(); program->live.live_in.resize(program->blocks.size(), IDSet(program->live.memory)); program->max_reg_demand = RegisterDemand(); program->needs_vcc = program->gfx_level >= GFX10; live_ctx ctx; ctx.program = program; ctx.worklist = program->blocks.size() - 1; ctx.handled_once = program->blocks.size(); /* this implementation assumes that the block idx corresponds to the block's position in * program->blocks vector */ while (ctx.worklist >= 0) { process_live_temps_per_block(ctx, &program->blocks[ctx.worklist--]); } /* calculate the program's register demand and number of waves */ if (program->progress < CompilationProgress::after_ra) update_vgpr_sgpr_demand(program, program->max_reg_demand); } } // namespace aco