src/amd/compiler/aco_live_var_analysis.cpp

   1 /*
   2  * Copyright © 2018 Valve Corporation
   3  * Copyright © 2018 Google
   4  *
   5  * Permission is hereby granted, free of charge, to any person obtaining a
   6  * copy of this software and associated documentation files (the "Software"),
   7  * to deal in the Software without restriction, including without limitation
   8  * the rights to use, copy, modify, merge, publish, distribute, sublicense,
   9  * and/or sell copies of the Software, and to permit persons to whom the
  10  * Software is furnished to do so, subject to the following conditions:
  11  *
  12  * The above copyright notice and this permission notice (including the next
  13  * paragraph) shall be included in all copies or substantial portions of the
  14  * Software.
  15  *
  16  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  17  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  18  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
  19  * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  20  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
  21  * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
  22  * IN THE SOFTWARE.
  23  *
  24  * Authors:
  25  *    Daniel Schürmann (daniel.schuermann@campus.tu-berlin.de)
  26  *    Bas Nieuwenhuizen (bas@basnieuwenhuizen.nl)
  27  *
  28  */
  29
  30 #include "aco_ir.h"
  31 #include "util/u_math.h"
  32
  33 #include <set>
  34 #include <vector>
  35
  36 #include "vulkan/radv_shader.h"
  37
  38 namespace aco {
  39 namespace {
  40
  41 void process_live_temps_per_block(Program *program, live& lives, Block* block,
  42                                   std::set<unsigned>& worklist, std::vector<uint16_t>& phi_sgpr_ops)
  43 {
  44    std::vector<RegisterDemand>& register_demand = lives.register_demand[block->index];
  45    RegisterDemand new_demand;
  46
  47    register_demand.resize(block->instructions.size());
  48    block->register_demand = RegisterDemand();
  49
  50    std::set<Temp> live_sgprs;
  51    std::set<Temp> live_vgprs;
  52
  53    /* add the live_out_exec to live */
  54    bool exec_live = false;
  55    if (block->live_out_exec != Temp()) {
  56       live_sgprs.insert(block->live_out_exec);
  57       new_demand.sgpr += program->lane_mask.size();
  58       exec_live = true;
  59    }
  60
  61    /* split the live-outs from this block into the temporary sets */
  62    std::vector<std::set<Temp>>& live_temps = lives.live_out;
  63    for (const Temp temp : live_temps[block->index]) {
  64       const bool inserted = temp.is_linear()
  65                           ? live_sgprs.insert(temp).second
  66                           : live_vgprs.insert(temp).second;
  67       if (inserted) {
  68          new_demand += temp;
  69       }
  70    }
  71    new_demand.sgpr -= phi_sgpr_ops[block->index];
  72
  73    /* traverse the instructions backwards */
  74    int idx;
  75    for (idx = block->instructions.size() -1; idx >= 0; idx--) {
  76       Instruction *insn = block->instructions[idx].get();
  77       if (is_phi(insn))
  78          break;
  79
  80       /* substract the 1 or 2 sgprs from exec */
  81       if (exec_live)
  82          assert(new_demand.sgpr >= (int16_t) program->lane_mask.size());
  83       register_demand[idx] = RegisterDemand(new_demand.vgpr, new_demand.sgpr - (exec_live ? program->lane_mask.size() : 0));
  84
  85       /* KILL */
  86       for (Definition& definition : insn->definitions) {
  87          if (!definition.isTemp()) {
  88             continue;
  89          }
  90
  91          const Temp temp = definition.getTemp();
  92          size_t n = 0;
  93          if (temp.is_linear())
  94             n = live_sgprs.erase(temp);
  95          else
  96             n = live_vgprs.erase(temp);
  97
  98          if (n) {
  99             new_demand -= temp;
 100             definition.setKill(false);
 101          } else {
 102             register_demand[idx] += temp;
 103             definition.setKill(true);
 104          }
 105
 106          if (definition.isFixed() && definition.physReg() == exec)
 107             exec_live = false;
 108       }
 109
 110       /* GEN */
 111       if (insn->opcode == aco_opcode::p_logical_end) {
 112          new_demand.sgpr += phi_sgpr_ops[block->index];
 113       } else {
 114          for (unsigned i = 0; i < insn->operands.size(); ++i)
 115          {
 116             Operand& operand = insn->operands[i];
 117             if (!operand.isTemp()) {
 118                continue;
 119             }
 120             const Temp temp = operand.getTemp();
 121             const bool inserted = temp.is_linear()
 122                                 ? live_sgprs.insert(temp).second
 123                                 : live_vgprs.insert(temp).second;
 124             if (inserted) {
 125                operand.setFirstKill(true);
 126                for (unsigned j = i + 1; j < insn->operands.size(); ++j) {
 127                   if (insn->operands[j].isTemp() && insn->operands[j].tempId() == operand.tempId()) {
 128                      insn->operands[j].setFirstKill(false);
 129                      insn->operands[j].setKill(true);
 130                   }
 131                }
 132                new_demand += temp;
 133             } else {
 134                operand.setKill(false);
 135             }
 136
 137             if (operand.isFixed() && operand.physReg() == exec)
 138                exec_live = true;
 139          }
 140       }
 141
 142       block->register_demand.update(register_demand[idx]);
 143    }
 144
 145    /* update block's register demand for a last time */
 146    if (exec_live)
 147       assert(new_demand.sgpr >= (int16_t) program->lane_mask.size());
 148    new_demand.sgpr -= exec_live ? program->lane_mask.size() : 0;
 149    block->register_demand.update(new_demand);
 150
 151    /* handle phi definitions */
 152    int phi_idx = idx;
 153    while (phi_idx >= 0) {
 154       register_demand[phi_idx] = new_demand;
 155       Instruction *insn = block->instructions[phi_idx].get();
 156
 157       assert(is_phi(insn));
 158       assert(insn->definitions.size() == 1 && insn->definitions[0].isTemp());
 159       Definition& definition = insn->definitions[0];
 160       const Temp temp = definition.getTemp();
 161       size_t n = 0;
 162
 163       if (temp.is_linear())
 164          n = live_sgprs.erase(temp);
 165       else
 166          n = live_vgprs.erase(temp);
 167
 168       if (n)
 169          definition.setKill(false);
 170       else
 171          definition.setKill(true);
 172
 173       phi_idx--;
 174    }
 175
 176    /* now, we have the live-in sets and need to merge them into the live-out sets */
 177    for (unsigned pred_idx : block->logical_preds) {
 178       for (Temp vgpr : live_vgprs) {
 179          auto it = live_temps[pred_idx].insert(vgpr);
 180          if (it.second)
 181             worklist.insert(pred_idx);
 182       }
 183    }
 184
 185    for (unsigned pred_idx : block->linear_preds) {
 186       for (Temp sgpr : live_sgprs) {
 187          auto it = live_temps[pred_idx].insert(sgpr);
 188          if (it.second)
 189             worklist.insert(pred_idx);
 190       }
 191    }
 192
 193    /* handle phi operands */
 194    phi_idx = idx;
 195    while (phi_idx >= 0) {
 196       Instruction *insn = block->instructions[phi_idx].get();
 197       assert(is_phi(insn));
 198       /* directly insert into the predecessors live-out set */
 199       std::vector<unsigned>& preds = insn->opcode == aco_opcode::p_phi
 200                                    ? block->logical_preds
 201                                    : block->linear_preds;
 202       for (unsigned i = 0; i < preds.size(); ++i) {
 203          Operand &operand = insn->operands[i];
 204          if (!operand.isTemp()) {
 205             continue;
 206          }
 207          /* check if we changed an already processed block */
 208          const bool inserted = live_temps[preds[i]].insert(operand.getTemp()).second;
 209          if (inserted) {
 210             operand.setKill(true);
 211             worklist.insert(preds[i]);
 212             if (insn->opcode == aco_opcode::p_phi && operand.getTemp().type() == RegType::sgpr)
 213                phi_sgpr_ops[preds[i]] += operand.size();
 214          }
 215       }
 216       phi_idx--;
 217    }
 218
 219    if (!(block->index != 0 || (live_vgprs.empty() && live_sgprs.empty()))) {
 220       aco_print_program(program, stderr);
 221       fprintf(stderr, "These temporaries are never defined or are defined after use:\n");
 222       for (Temp vgpr : live_vgprs)
 223          fprintf(stderr, "%%%d\n", vgpr.id());
 224       for (Temp sgpr : live_sgprs)
 225          fprintf(stderr, "%%%d\n", sgpr.id());
 226       abort();
 227    }
 228
 229    assert(block->index != 0 || new_demand == RegisterDemand());
 230 }
 231
 232 unsigned calc_waves_per_workgroup(Program *program)
 233 {
 234    unsigned workgroup_size = program->wave_size;
 235    if (program->stage == compute_cs) {
 236       unsigned* bsize = program->info->cs.block_size;
 237       workgroup_size = bsize[0] * bsize[1] * bsize[2];
 238    }
 239    return align(workgroup_size, program->wave_size) / program->wave_size;
 240 }
 241 } /* end namespace */
 242
 243 uint16_t get_extra_sgprs(Program *program)
 244 {
 245    if (program->chip_class >= GFX10) {
 246       assert(!program->needs_flat_scr);
 247       assert(!program->needs_xnack_mask);
 248       return 2;
 249    } else if (program->chip_class >= GFX8) {
 250       if (program->needs_flat_scr)
 251          return 6;
 252       else if (program->needs_xnack_mask)
 253          return 4;
 254       else if (program->needs_vcc)
 255          return 2;
 256       else
 257          return 0;
 258    } else {
 259       assert(!program->needs_xnack_mask);
 260       if (program->needs_flat_scr)
 261          return 4;
 262       else if (program->needs_vcc)
 263          return 2;
 264       else
 265          return 0;
 266    }
 267 }
 268
 269 uint16_t get_sgpr_alloc(Program *program, uint16_t addressable_sgprs)
 270 {
 271    assert(addressable_sgprs <= program->sgpr_limit);
 272    uint16_t sgprs = addressable_sgprs + get_extra_sgprs(program);
 273    uint16_t granule = program->sgpr_alloc_granule + 1;
 274    return align(std::max(sgprs, granule), granule);
 275 }
 276
 277 uint16_t get_vgpr_alloc(Program *program, uint16_t addressable_vgprs)
 278 {
 279    assert(addressable_vgprs <= program->vgpr_limit);
 280    uint16_t granule = program->vgpr_alloc_granule + 1;
 281    return align(std::max(addressable_vgprs, granule), granule);
 282 }
 283
 284 uint16_t get_addr_sgpr_from_waves(Program *program, uint16_t max_waves)
 285 {
 286     uint16_t sgprs = program->physical_sgprs / max_waves & ~program->sgpr_alloc_granule;
 287     sgprs -= get_extra_sgprs(program);
 288     return std::min(sgprs, program->sgpr_limit);
 289 }
 290
 291 uint16_t get_addr_vgpr_from_waves(Program *program, uint16_t max_waves)
 292 {
 293     uint16_t vgprs = 256 / max_waves & ~program->vgpr_alloc_granule;
 294     return std::min(vgprs, program->vgpr_limit);
 295 }
 296
 297 void calc_min_waves(Program* program)
 298 {
 299    unsigned waves_per_workgroup = calc_waves_per_workgroup(program);
 300    /* currently min_waves is in wave64 waves */
 301    if (program->wave_size == 32)
 302       waves_per_workgroup = DIV_ROUND_UP(waves_per_workgroup, 2);
 303
 304    unsigned simd_per_cu = 4; /* TODO: different on Navi */
 305    bool wgp = program->chip_class >= GFX10; /* assume WGP is used on Navi */
 306    unsigned simd_per_cu_wgp = wgp ? simd_per_cu * 2 : simd_per_cu;
 307
 308    program->min_waves = DIV_ROUND_UP(waves_per_workgroup, simd_per_cu_wgp);
 309 }
 310
 311 void update_vgpr_sgpr_demand(Program* program, const RegisterDemand new_demand)
 312 {
 313    /* TODO: max_waves_per_simd, simd_per_cu and the number of physical vgprs for Navi */
 314    unsigned max_waves_per_simd = 10;
 315    unsigned simd_per_cu = 4;
 316
 317    bool wgp = program->chip_class >= GFX10; /* assume WGP is used on Navi */
 318    unsigned simd_per_cu_wgp = wgp ? simd_per_cu * 2 : simd_per_cu;
 319    unsigned lds_limit = wgp ? program->lds_limit * 2 : program->lds_limit;
 320
 321    /* this won't compile, register pressure reduction necessary */
 322    if (new_demand.vgpr > program->vgpr_limit || new_demand.sgpr > program->sgpr_limit) {
 323       program->num_waves = 0;
 324       program->max_reg_demand = new_demand;
 325    } else {
 326       program->num_waves = program->physical_sgprs / get_sgpr_alloc(program, new_demand.sgpr);
 327       program->num_waves = std::min<uint16_t>(program->num_waves, 256 / get_vgpr_alloc(program, new_demand.vgpr));
 328       program->max_waves = max_waves_per_simd;
 329
 330       /* adjust max_waves for workgroup and LDS limits */
 331       unsigned waves_per_workgroup = calc_waves_per_workgroup(program);
 332       unsigned workgroups_per_cu_wgp = max_waves_per_simd * simd_per_cu_wgp / waves_per_workgroup;
 333       if (program->config->lds_size) {
 334          unsigned lds = program->config->lds_size * program->lds_alloc_granule;
 335          workgroups_per_cu_wgp = std::min(workgroups_per_cu_wgp, lds_limit / lds);
 336       }
 337       if (waves_per_workgroup > 1 && program->chip_class < GFX10)
 338          workgroups_per_cu_wgp = std::min(workgroups_per_cu_wgp, 16u); /* TODO: is this a SI-only limit? what about Navi? */
 339
 340       /* in cases like waves_per_workgroup=3 or lds=65536 and
 341        * waves_per_workgroup=1, we want the maximum possible number of waves per
 342        * SIMD and not the minimum. so DIV_ROUND_UP is used */
 343       program->max_waves = std::min<uint16_t>(program->max_waves, DIV_ROUND_UP(workgroups_per_cu_wgp * waves_per_workgroup, simd_per_cu_wgp));
 344
 345       /* incorporate max_waves and calculate max_reg_demand */
 346       program->num_waves = std::min<uint16_t>(program->num_waves, program->max_waves);
 347       program->max_reg_demand.vgpr = get_addr_vgpr_from_waves(program, program->num_waves);
 348       program->max_reg_demand.sgpr = get_addr_sgpr_from_waves(program, program->num_waves);
 349    }
 350 }
 351
 352 live live_var_analysis(Program* program,
 353                        const struct radv_nir_compiler_options *options)
 354 {
 355    live result;
 356    result.live_out.resize(program->blocks.size());
 357    result.register_demand.resize(program->blocks.size());
 358    std::set<unsigned> worklist;
 359    std::vector<uint16_t> phi_sgpr_ops(program->blocks.size());
 360    RegisterDemand new_demand;
 361
 362    /* this implementation assumes that the block idx corresponds to the block's position in program->blocks vector */
 363    for (Block& block : program->blocks)
 364       worklist.insert(block.index);
 365    while (!worklist.empty()) {
 366       std::set<unsigned>::reverse_iterator b_it = worklist.rbegin();
 367       unsigned block_idx = *b_it;
 368       worklist.erase(block_idx);
 369       process_live_temps_per_block(program, result, &program->blocks[block_idx], worklist, phi_sgpr_ops);
 370       new_demand.update(program->blocks[block_idx].register_demand);
 371    }
 372
 373    /* calculate the program's register demand and number of waves */
 374    update_vgpr_sgpr_demand(program, new_demand);
 375
 376    return result;
 377 }
 378
 379 }
 380