#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 += program->lane_mask.size();
+ 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 */
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;
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;
}
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)
definition.setKill(false);
phi_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);
- }
- }
+ /* 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 : block->linear_preds) {
- for (Temp sgpr : live_sgprs) {
- auto it = live_temps[pred_idx].insert(sgpr);
+ for (unsigned pred_idx : preds) {
+ auto it = lives.live_out[pred_idx].insert(t);
if (it.second)
worklist.insert(pred_idx);
}
if (operand.isFixed() && operand.physReg() == vcc)
program->needs_vcc = true;
/* check if we changed an already processed block */
- const bool inserted = live_temps[preds[i]].insert(operand.getTemp()).second;
+ const bool inserted = lives.live_out[preds[i]].insert(operand.getTemp()).second;
if (inserted) {
operand.setKill(true);
worklist.insert(preds[i]);
phi_idx--;
}
- if ((block->logical_preds.empty() && !live_vgprs.empty()) ||
- (block->linear_preds.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();
- }
-
- assert(block->index != 0 || new_demand == RegisterDemand());
+ assert(block->index != 0 || (new_demand == RegisterDemand() && live.empty()));
}
unsigned calc_waves_per_workgroup(Program *program)
{
- unsigned workgroup_size = program->wave_size;
- if (program->stage == compute_cs) {
- unsigned* bsize = program->info->cs.block_size;
- workgroup_size = bsize[0] * bsize[1] * bsize[2];
- }
+ /* 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 */
{
if (program->chip_class >= GFX10) {
assert(!program->needs_flat_scr);
- assert(!program->needs_xnack_mask);
+ assert(!program->xnack_enabled);
return 2;
} else if (program->chip_class >= GFX8) {
if (program->needs_flat_scr)
return 6;
- else if (program->needs_xnack_mask)
+ else if (program->xnack_enabled)
return 4;
else if (program->needs_vcc)
return 2;
else
return 0;
} else {
- assert(!program->needs_xnack_mask);
+ assert(!program->xnack_enabled);
if (program->needs_flat_scr)
return 4;
else if (program->needs_vcc)
{
/* 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 */
projects / mesa.git / blobdiff