*/
#include "aco_ir.h"
+#include "util/u_math.h"
#include <set>
#include <vector>
#include "vulkan/radv_shader.h"
namespace aco {
-namespace {
+RegisterDemand get_live_changes(aco_ptr<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(aco_ptr<Instruction>& instr)
+{
+ RegisterDemand temp_registers;
+
+ for (Definition def : instr->definitions) {
+ if (!def.isTemp())
+ continue;
+ if (def.isKill())
+ temp_registers += def.getTemp();
+ }
+
+ for (Operand op : instr->operands) {
+ if (op.isTemp() && op.isLateKill() && op.isFirstKill())
+ temp_registers += op.getTemp();
+ }
+
+ return temp_registers;
+}
+RegisterDemand get_demand_before(RegisterDemand demand, aco_ptr<Instruction>& instr, aco_ptr<Instruction>& instr_before)
+{
+ demand -= get_live_changes(instr);
+ demand -= get_temp_registers(instr);
+ if (instr_before)
+ demand += get_temp_registers(instr_before);
+ return demand;
+}
+
+namespace {
void process_live_temps_per_block(Program *program, live& lives, Block* block,
std::set<unsigned>& worklist, std::vector<uint16_t>& phi_sgpr_ops)
{
register_demand.resize(block->instructions.size());
block->register_demand = RegisterDemand();
-
- std::set<Temp> live_sgprs;
- std::set<Temp> live_vgprs;
+ TempSet live = lives.live_out[block->index];
/* add the live_out_exec to live */
bool exec_live = false;
if (block->live_out_exec != Temp()) {
- live_sgprs.insert(block->live_out_exec);
- new_demand.sgpr += 2;
+ live.insert(block->live_out_exec);
exec_live = true;
}
- /* split the live-outs from this block into the temporary sets */
- std::vector<std::set<Temp>>& live_temps = lives.live_out;
- for (const Temp temp : live_temps[block->index]) {
- const bool inserted = temp.is_linear()
- ? live_sgprs.insert(temp).second
- : live_vgprs.insert(temp).second;
- if (inserted) {
- new_demand += temp;
- }
- }
+ /* initialize register demand */
+ for (Temp t : live)
+ new_demand += t;
new_demand.sgpr -= phi_sgpr_ops[block->index];
/* traverse the instructions backwards */
- for (int idx = block->instructions.size() -1; idx >= 0; idx--)
- {
- /* substract the 2 sgprs from exec */
+ int idx;
+ for (idx = block->instructions.size() -1; idx >= 0; idx--) {
+ Instruction *insn = block->instructions[idx].get();
+ if (is_phi(insn))
+ break;
+
+ /* substract the 1 or 2 sgprs from exec */
if (exec_live)
- assert(new_demand.sgpr >= 2);
- register_demand[idx] = RegisterDemand(new_demand.vgpr, new_demand.sgpr - (exec_live ? 2 : 0));
+ assert(new_demand.sgpr >= (int16_t) program->lane_mask.size());
+ register_demand[idx] = RegisterDemand(new_demand.vgpr, new_demand.sgpr - (exec_live ? program->lane_mask.size() : 0));
- Instruction *insn = block->instructions[idx].get();
/* KILL */
for (Definition& definition : insn->definitions) {
if (!definition.isTemp()) {
continue;
}
+ if ((definition.isFixed() || definition.hasHint()) && definition.physReg() == vcc)
+ program->needs_vcc = true;
const Temp temp = definition.getTemp();
- size_t n = 0;
- if (temp.is_linear())
- n = live_sgprs.erase(temp);
- else
- n = live_vgprs.erase(temp);
+ const size_t n = live.erase(temp);
if (n) {
new_demand -= temp;
}
/* GEN */
- if (insn->opcode == aco_opcode::p_phi ||
- insn->opcode == aco_opcode::p_linear_phi) {
- /* directly insert into the predecessors live-out set */
- std::vector<unsigned>& preds = insn->opcode == aco_opcode::p_phi
- ? block->logical_preds
- : block->linear_preds;
- for (unsigned i = 0; i < preds.size(); ++i)
- {
- Operand &operand = insn->operands[i];
- if (!operand.isTemp()) {
- continue;
- }
- /* check if we changed an already processed block */
- const bool inserted = live_temps[preds[i]].insert(operand.getTemp()).second;
- if (inserted) {
- operand.setFirstKill(true);
- worklist.insert(preds[i]);
- if (insn->opcode == aco_opcode::p_phi && operand.getTemp().type() == RegType::sgpr)
- phi_sgpr_ops[preds[i]] += operand.size();
- }
- }
- } else if (insn->opcode == aco_opcode::p_logical_end) {
+ if (insn->opcode == aco_opcode::p_logical_end) {
new_demand.sgpr += phi_sgpr_ops[block->index];
} else {
+ /* 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);
+
for (unsigned i = 0; i < insn->operands.size(); ++i)
{
Operand& operand = insn->operands[i];
- if (!operand.isTemp()) {
+ if (!operand.isTemp())
continue;
- }
+ if (operand.isFixed() && operand.physReg() == vcc)
+ program->needs_vcc = true;
const Temp temp = operand.getTemp();
- const bool inserted = temp.is_linear()
- ? live_sgprs.insert(temp).second
- : live_vgprs.insert(temp).second;
+ const bool inserted = live.insert(temp).second;
if (inserted) {
operand.setFirstKill(true);
for (unsigned j = i + 1; j < insn->operands.size(); ++j) {
insn->operands[j].setKill(true);
}
}
+ if (operand.isLateKill())
+ register_demand[idx] += temp;
new_demand += temp;
- } else {
- operand.setKill(false);
}
if (operand.isFixed() && operand.physReg() == exec)
block->register_demand.update(register_demand[idx]);
}
- /* now, we have the live-in sets and need to merge them into the live-out sets */
- for (unsigned pred_idx : block->logical_preds) {
- for (Temp vgpr : live_vgprs) {
- auto it = live_temps[pred_idx].insert(vgpr);
- if (it.second)
- worklist.insert(pred_idx);
- }
+ /* update block's register demand for a last time */
+ if (exec_live)
+ assert(new_demand.sgpr >= (int16_t) program->lane_mask.size());
+ new_demand.sgpr -= exec_live ? program->lane_mask.size() : 0;
+ block->register_demand.update(new_demand);
+
+ /* handle phi definitions */
+ int phi_idx = idx;
+ while (phi_idx >= 0) {
+ register_demand[phi_idx] = new_demand;
+ Instruction *insn = block->instructions[phi_idx].get();
+
+ assert(is_phi(insn));
+ assert(insn->definitions.size() == 1 && insn->definitions[0].isTemp());
+ Definition& definition = insn->definitions[0];
+ if ((definition.isFixed() || definition.hasHint()) && definition.physReg() == vcc)
+ program->needs_vcc = true;
+ const Temp temp = definition.getTemp();
+ const size_t n = live.erase(temp);
+
+ if (n)
+ definition.setKill(false);
+ else
+ definition.setKill(true);
+
+ phi_idx--;
}
- for (unsigned pred_idx : block->linear_preds) {
- for (Temp sgpr : live_sgprs) {
- auto it = live_temps[pred_idx].insert(sgpr);
+ /* now, we need to merge the live-ins into the live-out sets */
+ for (Temp t : live) {
+ std::vector<unsigned>& preds = t.is_linear() ? block->linear_preds : block->logical_preds;
+
+#ifndef NDEBUG
+ if (preds.empty())
+ aco_err(program, "Temporary never defined or are defined after use: %%%d in BB%d", t.id(), block->index);
+#endif
+
+ for (unsigned pred_idx : preds) {
+ auto it = lives.live_out[pred_idx].insert(t);
if (it.second)
worklist.insert(pred_idx);
}
}
- if (!(block->index != 0 || (live_vgprs.empty() && live_sgprs.empty()))) {
- aco_print_program(program, stderr);
- fprintf(stderr, "These temporaries are never defined or are defined after use:\n");
- for (Temp vgpr : live_vgprs)
- fprintf(stderr, "%%%d\n", vgpr.id());
- for (Temp sgpr : live_sgprs)
- fprintf(stderr, "%%%d\n", sgpr.id());
- abort();
+ /* handle phi operands */
+ phi_idx = idx;
+ while (phi_idx >= 0) {
+ Instruction *insn = block->instructions[phi_idx].get();
+ assert(is_phi(insn));
+ /* directly insert into the predecessors live-out set */
+ std::vector<unsigned>& preds = insn->opcode == aco_opcode::p_phi
+ ? block->logical_preds
+ : block->linear_preds;
+ for (unsigned i = 0; i < preds.size(); ++i) {
+ Operand &operand = insn->operands[i];
+ if (!operand.isTemp())
+ continue;
+ 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.getTemp()).second;
+ if (inserted) {
+ operand.setKill(true);
+ worklist.insert(preds[i]);
+ if (insn->opcode == aco_opcode::p_phi && operand.getTemp().type() == RegType::sgpr)
+ phi_sgpr_ops[preds[i]] += operand.size();
+ }
+ }
+ phi_idx--;
}
- assert(block->index != 0 || new_demand == RegisterDemand());
+ assert(block->index != 0 || (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 */
+uint16_t get_extra_sgprs(Program *program)
+{
+ if (program->chip_class >= GFX10) {
+ assert(!program->needs_flat_scr);
+ assert(!program->xnack_enabled);
+ return 2;
+ } else if (program->chip_class >= GFX8) {
+ if (program->needs_flat_scr)
+ return 6;
+ else if (program->xnack_enabled)
+ return 4;
+ else if (program->needs_vcc)
+ return 2;
+ else
+ return 0;
+ } else {
+ assert(!program->xnack_enabled);
+ if (program->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)
+{
+ assert(addressable_sgprs <= program->sgpr_limit);
+ uint16_t sgprs = addressable_sgprs + get_extra_sgprs(program);
+ uint16_t granule = program->sgpr_alloc_granule + 1;
+ return align(std::max(sgprs, granule), granule);
+}
+
+uint16_t get_vgpr_alloc(Program *program, uint16_t addressable_vgprs)
+{
+ assert(addressable_vgprs <= program->vgpr_limit);
+ uint16_t granule = program->vgpr_alloc_granule + 1;
+ return align(std::max(addressable_vgprs, granule), granule);
+}
+
+uint16_t get_addr_sgpr_from_waves(Program *program, uint16_t max_waves)
+{
+ uint16_t sgprs = program->physical_sgprs / max_waves & ~program->sgpr_alloc_granule;
+ sgprs -= get_extra_sgprs(program);
+ return std::min(sgprs, program->sgpr_limit);
+}
+
+uint16_t get_addr_vgpr_from_waves(Program *program, uint16_t max_waves)
+{
+ uint16_t vgprs = 256 / max_waves & ~program->vgpr_alloc_granule;
+ return std::min(vgprs, program->vgpr_limit);
+}
+
+void calc_min_waves(Program* program)
+{
+ unsigned waves_per_workgroup = calc_waves_per_workgroup(program);
+ /* currently min_waves is in wave64 waves */
+ if (program->wave_size == 32)
+ waves_per_workgroup = DIV_ROUND_UP(waves_per_workgroup, 2);
+
+ unsigned simd_per_cu = 4; /* TODO: different on Navi */
+ bool wgp = program->chip_class >= GFX10; /* assume WGP is used on Navi */
+ unsigned simd_per_cu_wgp = wgp ? simd_per_cu * 2 : simd_per_cu;
+
+ program->min_waves = DIV_ROUND_UP(waves_per_workgroup, simd_per_cu_wgp);
+}
+
void update_vgpr_sgpr_demand(Program* program, const RegisterDemand new_demand)
{
- // TODO: also take shared mem into account
- const int16_t total_sgpr_regs = program->chip_class >= GFX8 ? 800 : 512;
- const int16_t max_addressible_sgpr = program->sgpr_limit;
- /* VGPRs are allocated in chunks of 4 */
- const int16_t rounded_vgpr_demand = std::max<int16_t>(4, (new_demand.vgpr + 3) & ~3);
- /* SGPRs are allocated in chunks of 16 between 8 and 104. VCC occupies the last 2 registers */
- const int16_t rounded_sgpr_demand = std::min(std::max<int16_t>(8, (new_demand.sgpr + 2 + 7) & ~7), max_addressible_sgpr);
+ /* TODO: max_waves_per_simd, simd_per_cu and the number of physical vgprs for Navi */
+ unsigned max_waves_per_simd = 10;
+ if ((program->family >= CHIP_POLARIS10 && program->family <= CHIP_VEGAM) || program->chip_class >= GFX10_3)
+ max_waves_per_simd = 8;
+ unsigned simd_per_cu = 4;
+
+ bool wgp = program->chip_class >= GFX10; /* assume WGP is used on Navi */
+ unsigned simd_per_cu_wgp = wgp ? simd_per_cu * 2 : simd_per_cu;
+ unsigned lds_limit = wgp ? program->lds_limit * 2 : program->lds_limit;
+
/* this won't compile, register pressure reduction necessary */
- if (new_demand.vgpr > 256 || new_demand.sgpr > max_addressible_sgpr) {
+ if (new_demand.vgpr > program->vgpr_limit || new_demand.sgpr > program->sgpr_limit) {
program->num_waves = 0;
program->max_reg_demand = new_demand;
} else {
- program->num_waves = std::min<uint16_t>(10,
- std::min<uint16_t>(256 / rounded_vgpr_demand,
- total_sgpr_regs / rounded_sgpr_demand));
+ program->num_waves = program->physical_sgprs / get_sgpr_alloc(program, new_demand.sgpr);
+ program->num_waves = std::min<uint16_t>(program->num_waves, 256 / get_vgpr_alloc(program, new_demand.vgpr));
+ program->max_waves = max_waves_per_simd;
- program->max_reg_demand = { int16_t((256 / program->num_waves) & ~3), std::min<int16_t>(((total_sgpr_regs / program->num_waves) & ~7) - 2, max_addressible_sgpr)};
+ /* adjust max_waves for workgroup and LDS limits */
+ unsigned waves_per_workgroup = calc_waves_per_workgroup(program);
+ unsigned workgroups_per_cu_wgp = max_waves_per_simd * simd_per_cu_wgp / waves_per_workgroup;
+ if (program->config->lds_size) {
+ unsigned lds = program->config->lds_size * program->lds_alloc_granule;
+ workgroups_per_cu_wgp = std::min(workgroups_per_cu_wgp, lds_limit / lds);
+ }
+ if (waves_per_workgroup > 1 && program->chip_class < GFX10)
+ workgroups_per_cu_wgp = std::min(workgroups_per_cu_wgp, 16u); /* TODO: is this a SI-only limit? what about Navi? */
+
+ /* 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 */
+ program->max_waves = std::min<uint16_t>(program->max_waves, DIV_ROUND_UP(workgroups_per_cu_wgp * waves_per_workgroup, simd_per_cu_wgp));
+
+ /* incorporate max_waves and calculate max_reg_demand */
+ program->num_waves = std::min<uint16_t>(program->num_waves, program->max_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);
}
}
std::vector<uint16_t> phi_sgpr_ops(program->blocks.size());
RegisterDemand new_demand;
+ program->needs_vcc = false;
+
/* this implementation assumes that the block idx corresponds to the block's position in program->blocks vector */
for (Block& block : program->blocks)
worklist.insert(block.index);