src/amd/compiler/aco_insert_waitcnt.cpp

   1 /*
   2  * Copyright © 2018 Valve Corporation
   3  *
   4  * Permission is hereby granted, free of charge, to any person obtaining a
   5  * copy of this software and associated documentation files (the "Software"),
   6  * to deal in the Software without restriction, including without limitation
   7  * the rights to use, copy, modify, merge, publish, distribute, sublicense,
   8  * and/or sell copies of the Software, and to permit persons to whom the
   9  * Software is furnished to do so, subject to the following conditions:
  10  *
  11  * The above copyright notice and this permission notice (including the next
  12  * paragraph) shall be included in all copies or substantial portions of the
  13  * Software.
  14  *
  15  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
  16  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
  17  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
  18  * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
  19  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
  20  * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
  21  * IN THE SOFTWARE.
  22  *
  23  */
  24
  25 #include <algorithm>
  26 #include <map>
  27 #include <stack>
  28 #include <math.h>
  29
  30 #include "aco_ir.h"
  31 #include "vulkan/radv_shader.h"
  32
  33 namespace aco {
  34
  35 namespace {
  36
  37 /**
  38  * The general idea of this pass is:
  39  * The CFG is traversed in reverse postorder (forward) and loops are processed
  40  * several times until no progress is made.
  41  * Per BB two wait_ctx is maintained: an in-context and out-context.
  42  * The in-context is the joined out-contexts of the predecessors.
  43  * The context contains a map: gpr -> wait_entry
  44  * consisting of the information about the cnt values to be waited for.
  45  * Note: After merge-nodes, it might occur that for the same register
  46  *       multiple cnt values are to be waited for.
  47  *
  48  * The values are updated according to the encountered instructions:
  49  * - additional events increment the counter of waits of the same type
  50  * - or erase gprs with counters higher than to be waited for.
  51  */
  52
  53 // TODO: do a more clever insertion of wait_cnt (lgkm_cnt) when there is a load followed by a use of a previous load
  54
  55 /* Instructions of the same event will finish in-order except for smem
  56  * and maybe flat. Instructions of different events may not finish in-order. */
  57 enum wait_event : uint16_t {
  58    event_smem = 1 << 0,
  59    event_lds = 1 << 1,
  60    event_gds = 1 << 2,
  61    event_vmem = 1 << 3,
  62    event_vmem_store = 1 << 4, /* GFX10+ */
  63    event_flat = 1 << 5,
  64    event_exp_pos = 1 << 6,
  65    event_exp_param = 1 << 7,
  66    event_exp_mrt_null = 1 << 8,
  67    event_gds_gpr_lock = 1 << 9,
  68    event_vmem_gpr_lock = 1 << 10,
  69    event_sendmsg = 1 << 11,
  70    num_events = 12,
  71 };
  72
  73 enum counter_type : uint8_t {
  74    counter_exp = 1 << 0,
  75    counter_lgkm = 1 << 1,
  76    counter_vm = 1 << 2,
  77    counter_vs = 1 << 3,
  78    num_counters = 4,
  79 };
  80
  81 static const uint16_t exp_events = event_exp_pos | event_exp_param | event_exp_mrt_null | event_gds_gpr_lock | event_vmem_gpr_lock;
  82 static const uint16_t lgkm_events = event_smem | event_lds | event_gds | event_flat | event_sendmsg;
  83 static const uint16_t vm_events = event_vmem | event_flat;
  84 static const uint16_t vs_events = event_vmem_store;
  85
  86 uint8_t get_counters_for_event(wait_event ev)
  87 {
  88    switch (ev) {
  89    case event_smem:
  90    case event_lds:
  91    case event_gds:
  92    case event_sendmsg:
  93       return counter_lgkm;
  94    case event_vmem:
  95       return counter_vm;
  96    case event_vmem_store:
  97       return counter_vs;
  98    case event_flat:
  99       return counter_vm | counter_lgkm;
 100    case event_exp_pos:
 101    case event_exp_param:
 102    case event_exp_mrt_null:
 103    case event_gds_gpr_lock:
 104    case event_vmem_gpr_lock:
 105       return counter_exp;
 106    default:
 107       return 0;
 108    }
 109 }
 110
 111 uint16_t get_events_for_counter(counter_type ctr)
 112 {
 113    switch (ctr) {
 114    case counter_exp:
 115       return exp_events;
 116    case counter_lgkm:
 117       return lgkm_events;
 118    case counter_vm:
 119       return vm_events;
 120    case counter_vs:
 121       return vs_events;
 122    }
 123    return 0;
 124 }
 125
 126 struct wait_imm {
 127    static const uint8_t unset_counter = 0xff;
 128
 129    uint8_t vm;
 130    uint8_t exp;
 131    uint8_t lgkm;
 132    uint8_t vs;
 133
 134    wait_imm() :
 135       vm(unset_counter), exp(unset_counter), lgkm(unset_counter), vs(unset_counter) {}
 136    wait_imm(uint16_t vm_, uint16_t exp_, uint16_t lgkm_, uint16_t vs_) :
 137       vm(vm_), exp(exp_), lgkm(lgkm_), vs(vs_) {}
 138
 139    wait_imm(enum chip_class chip, uint16_t packed) : vs(unset_counter)
 140    {
 141       vm = packed & 0xf;
 142       if (chip >= GFX9)
 143          vm |= (packed >> 10) & 0x30;
 144
 145       exp = (packed >> 4) & 0x7;
 146
 147       lgkm = (packed >> 8) & 0xf;
 148       if (chip >= GFX10)
 149          lgkm |= (packed >> 8) & 0x30;
 150    }
 151
 152    uint16_t pack(enum chip_class chip) const
 153    {
 154       uint16_t imm = 0;
 155       assert(exp == unset_counter || exp <= 0x7);
 156       switch (chip) {
 157       case GFX10:
 158       case GFX10_3:
 159          assert(lgkm == unset_counter || lgkm <= 0x3f);
 160          assert(vm == unset_counter || vm <= 0x3f);
 161          imm = ((vm & 0x30) << 10) | ((lgkm & 0x3f) << 8) | ((exp & 0x7) << 4) | (vm & 0xf);
 162          break;
 163       case GFX9:
 164          assert(lgkm == unset_counter || lgkm <= 0xf);
 165          assert(vm == unset_counter || vm <= 0x3f);
 166          imm = ((vm & 0x30) << 10) | ((lgkm & 0xf) << 8) | ((exp & 0x7) << 4) | (vm & 0xf);
 167          break;
 168       default:
 169          assert(lgkm == unset_counter || lgkm <= 0xf);
 170          assert(vm == unset_counter || vm <= 0xf);
 171          imm = ((lgkm & 0xf) << 8) | ((exp & 0x7) << 4) | (vm & 0xf);
 172          break;
 173       }
 174       if (chip < GFX9 && vm == wait_imm::unset_counter)
 175          imm |= 0xc000; /* should have no effect on pre-GFX9 and now we won't have to worry about the architecture when interpreting the immediate */
 176       if (chip < GFX10 && lgkm == wait_imm::unset_counter)
 177          imm |= 0x3000; /* should have no effect on pre-GFX10 and now we won't have to worry about the architecture when interpreting the immediate */
 178       return imm;
 179    }
 180
 181    bool combine(const wait_imm& other)
 182    {
 183       bool changed = other.vm < vm || other.exp < exp || other.lgkm < lgkm || other.vs < vs;
 184       vm = std::min(vm, other.vm);
 185       exp = std::min(exp, other.exp);
 186       lgkm = std::min(lgkm, other.lgkm);
 187       vs = std::min(vs, other.vs);
 188       return changed;
 189    }
 190
 191    bool empty() const
 192    {
 193       return vm == unset_counter && exp == unset_counter &&
 194              lgkm == unset_counter && vs == unset_counter;
 195    }
 196 };
 197
 198 struct wait_entry {
 199    wait_imm imm;
 200    uint16_t events; /* use wait_event notion */
 201    uint8_t counters; /* use counter_type notion */
 202    bool wait_on_read:1;
 203    bool logical:1;
 204    bool has_vmem_nosampler:1;
 205    bool has_vmem_sampler:1;
 206
 207    wait_entry(wait_event event, wait_imm imm, bool logical, bool wait_on_read)
 208            : imm(imm), events(event), counters(get_counters_for_event(event)),
 209              wait_on_read(wait_on_read), logical(logical),
 210              has_vmem_nosampler(false), has_vmem_sampler(false) {}
 211
 212    bool join(const wait_entry& other)
 213    {
 214       bool changed = (other.events & ~events) ||
 215                      (other.counters & ~counters) ||
 216                      (other.wait_on_read && !wait_on_read) ||
 217                      (other.has_vmem_nosampler && !has_vmem_nosampler) ||
 218                      (other.has_vmem_sampler && !has_vmem_sampler);
 219       events |= other.events;
 220       counters |= other.counters;
 221       changed |= imm.combine(other.imm);
 222       wait_on_read |= other.wait_on_read;
 223       has_vmem_nosampler |= other.has_vmem_nosampler;
 224       has_vmem_sampler |= other.has_vmem_sampler;
 225       assert(logical == other.logical);
 226       return changed;
 227    }
 228
 229    void remove_counter(counter_type counter)
 230    {
 231       counters &= ~counter;
 232
 233       if (counter == counter_lgkm) {
 234          imm.lgkm = wait_imm::unset_counter;
 235          events &= ~(event_smem | event_lds | event_gds | event_sendmsg);
 236       }
 237
 238       if (counter == counter_vm) {
 239          imm.vm = wait_imm::unset_counter;
 240          events &= ~event_vmem;
 241          has_vmem_nosampler = false;
 242          has_vmem_sampler = false;
 243       }
 244
 245       if (counter == counter_exp) {
 246          imm.exp = wait_imm::unset_counter;
 247          events &= ~(event_exp_pos | event_exp_param | event_exp_mrt_null | event_gds_gpr_lock | event_vmem_gpr_lock);
 248       }
 249
 250       if (counter == counter_vs) {
 251          imm.vs = wait_imm::unset_counter;
 252          events &= ~event_vmem_store;
 253       }
 254
 255       if (!(counters & counter_lgkm) && !(counters & counter_vm))
 256          events &= ~event_flat;
 257    }
 258 };
 259
 260 struct wait_ctx {
 261    Program *program;
 262    enum chip_class chip_class;
 263    uint16_t max_vm_cnt;
 264    uint16_t max_exp_cnt;
 265    uint16_t max_lgkm_cnt;
 266    uint16_t max_vs_cnt;
 267    uint16_t unordered_events = event_smem | event_flat;
 268
 269    uint8_t vm_cnt = 0;
 270    uint8_t exp_cnt = 0;
 271    uint8_t lgkm_cnt = 0;
 272    uint8_t vs_cnt = 0;
 273    bool pending_flat_lgkm = false;
 274    bool pending_flat_vm = false;
 275    bool pending_s_buffer_store = false; /* GFX10 workaround */
 276
 277    wait_imm barrier_imm[storage_count];
 278    uint16_t barrier_events[storage_count] = {}; /* use wait_event notion */
 279
 280    std::map<PhysReg,wait_entry> gpr_map;
 281
 282    /* used for vmem/smem scores */
 283    bool collect_statistics;
 284    Instruction *gen_instr;
 285    std::map<Instruction *, unsigned> unwaited_instrs[num_counters];
 286    std::map<PhysReg,std::set<Instruction *>> reg_instrs[num_counters];
 287    std::vector<unsigned> wait_distances[num_events];
 288
 289    wait_ctx() {}
 290    wait_ctx(Program *program_)
 291            : program(program_),
 292              chip_class(program_->chip_class),
 293              max_vm_cnt(program_->chip_class >= GFX9 ? 62 : 14),
 294              max_exp_cnt(6),
 295              max_lgkm_cnt(program_->chip_class >= GFX10 ? 62 : 14),
 296              max_vs_cnt(program_->chip_class >= GFX10 ? 62 : 0),
 297              unordered_events(event_smem | (program_->chip_class < GFX10 ? event_flat : 0)),
 298              collect_statistics(program_->collect_statistics) {}
 299
 300    bool join(const wait_ctx* other, bool logical)
 301    {
 302       bool changed = other->exp_cnt > exp_cnt ||
 303                      other->vm_cnt > vm_cnt ||
 304                      other->lgkm_cnt > lgkm_cnt ||
 305                      other->vs_cnt > vs_cnt ||
 306                      (other->pending_flat_lgkm && !pending_flat_lgkm) ||
 307                      (other->pending_flat_vm && !pending_flat_vm);
 308
 309       exp_cnt = std::max(exp_cnt, other->exp_cnt);
 310       vm_cnt = std::max(vm_cnt, other->vm_cnt);
 311       lgkm_cnt = std::max(lgkm_cnt, other->lgkm_cnt);
 312       vs_cnt = std::max(vs_cnt, other->vs_cnt);
 313       pending_flat_lgkm |= other->pending_flat_lgkm;
 314       pending_flat_vm |= other->pending_flat_vm;
 315       pending_s_buffer_store |= other->pending_s_buffer_store;
 316
 317       for (std::pair<PhysReg,wait_entry> entry : other->gpr_map)
 318       {
 319          std::map<PhysReg,wait_entry>::iterator it = gpr_map.find(entry.first);
 320          if (entry.second.logical != logical)
 321             continue;
 322
 323          if (it != gpr_map.end()) {
 324             changed |= it->second.join(entry.second);
 325          } else {
 326             gpr_map.insert(entry);
 327             changed = true;
 328          }
 329       }
 330
 331       for (unsigned i = 0; i < storage_count; i++) {
 332          changed |= barrier_imm[i].combine(other->barrier_imm[i]);
 333          changed |= other->barrier_events[i] & ~barrier_events[i];
 334          barrier_events[i] |= other->barrier_events[i];
 335       }
 336
 337       /* these are used for statistics, so don't update "changed" */
 338       for (unsigned i = 0; i < num_counters; i++) {
 339          for (std::pair<Instruction *, unsigned> instr : other->unwaited_instrs[i]) {
 340             auto pos = unwaited_instrs[i].find(instr.first);
 341             if (pos == unwaited_instrs[i].end())
 342                unwaited_instrs[i].insert(instr);
 343             else
 344                pos->second = std::min(pos->second, instr.second);
 345          }
 346          /* don't use a foreach loop to avoid copies */
 347          for (auto it = other->reg_instrs[i].begin(); it != other->reg_instrs[i].end(); ++it)
 348             reg_instrs[i][it->first].insert(it->second.begin(), it->second.end());
 349       }
 350
 351       return changed;
 352    }
 353
 354    void wait_and_remove_from_entry(PhysReg reg, wait_entry& entry, counter_type counter) {
 355       if (collect_statistics && (entry.counters & counter)) {
 356          unsigned counter_idx = ffs(counter) - 1;
 357          for (Instruction *instr : reg_instrs[counter_idx][reg]) {
 358             auto pos = unwaited_instrs[counter_idx].find(instr);
 359             if (pos == unwaited_instrs[counter_idx].end())
 360                continue;
 361
 362             unsigned distance = pos->second;
 363             unsigned events = entry.events & get_events_for_counter(counter);
 364             while (events) {
 365                unsigned event_idx = u_bit_scan(&events);
 366                wait_distances[event_idx].push_back(distance);
 367             }
 368
 369             unwaited_instrs[counter_idx].erase(instr);
 370          }
 371          reg_instrs[counter_idx][reg].clear();
 372       }
 373
 374       entry.remove_counter(counter);
 375    }
 376
 377    void advance_unwaited_instrs()
 378    {
 379       for (unsigned i = 0; i < num_counters; i++) {
 380          for (auto it = unwaited_instrs[i].begin(); it != unwaited_instrs[i].end(); ++it)
 381             it->second++;
 382       }
 383    }
 384 };
 385
 386 wait_imm check_instr(Instruction* instr, wait_ctx& ctx)
 387 {
 388    wait_imm wait;
 389
 390    for (const Operand op : instr->operands) {
 391       if (op.isConstant() || op.isUndefined())
 392          continue;
 393
 394       /* check consecutively read gprs */
 395       for (unsigned j = 0; j < op.size(); j++) {
 396          PhysReg reg{op.physReg() + j};
 397          std::map<PhysReg,wait_entry>::iterator it = ctx.gpr_map.find(reg);
 398          if (it == ctx.gpr_map.end() || !it->second.wait_on_read)
 399             continue;
 400
 401          wait.combine(it->second.imm);
 402       }
 403    }
 404
 405    for (const Definition& def : instr->definitions) {
 406       /* check consecutively written gprs */
 407       for (unsigned j = 0; j < def.getTemp().size(); j++)
 408       {
 409          PhysReg reg{def.physReg() + j};
 410
 411          std::map<PhysReg,wait_entry>::iterator it = ctx.gpr_map.find(reg);
 412          if (it == ctx.gpr_map.end())
 413             continue;
 414
 415          /* Vector Memory reads and writes return in the order they were issued */
 416          bool has_sampler = instr->format == Format::MIMG && !instr->operands[1].isUndefined() && instr->operands[1].regClass() == s4;
 417          if (instr->isVMEM() && ((it->second.events & vm_events) == event_vmem) &&
 418              it->second.has_vmem_nosampler == !has_sampler && it->second.has_vmem_sampler == has_sampler)
 419             continue;
 420
 421          /* LDS reads and writes return in the order they were issued. same for GDS */
 422          if (instr->format == Format::DS) {
 423             bool gds = static_cast<DS_instruction*>(instr)->gds;
 424             if ((it->second.events & lgkm_events) == (gds ? event_gds : event_lds))
 425                continue;
 426          }
 427
 428          wait.combine(it->second.imm);
 429       }
 430    }
 431
 432    return wait;
 433 }
 434
 435 wait_imm parse_wait_instr(wait_ctx& ctx, Instruction *instr)
 436 {
 437    if (instr->opcode == aco_opcode::s_waitcnt_vscnt &&
 438        instr->definitions[0].physReg() == sgpr_null) {
 439       wait_imm imm;
 440       imm.vs = std::min<uint8_t>(imm.vs, static_cast<SOPK_instruction*>(instr)->imm);
 441       return imm;
 442    } else if (instr->opcode == aco_opcode::s_waitcnt) {
 443       return wait_imm(ctx.chip_class, static_cast<SOPP_instruction*>(instr)->imm);
 444    }
 445    return wait_imm();
 446 }
 447
 448 wait_imm perform_barrier(wait_ctx& ctx, memory_sync_info sync, unsigned semantics)
 449 {
 450    wait_imm imm;
 451    sync_scope subgroup_scope = ctx.program->workgroup_size <= ctx.program->wave_size ? scope_workgroup : scope_subgroup;
 452    if ((sync.semantics & semantics) && sync.scope > subgroup_scope) {
 453       unsigned storage = sync.storage;
 454       while (storage) {
 455          unsigned idx = u_bit_scan(&storage);
 456
 457          /* LDS is private to the workgroup */
 458          sync_scope bar_scope_lds = MIN2(sync.scope, scope_workgroup);
 459
 460          uint16_t events = ctx.barrier_events[idx];
 461          if (bar_scope_lds <= subgroup_scope)
 462             events &= ~event_lds;
 463
 464          /* in non-WGP, the L1/L0 cache keeps all memory operations in-order for the same workgroup */
 465          if (ctx.chip_class < GFX10 && sync.scope <= scope_workgroup)
 466             events &= ~(event_vmem | event_vmem_store | event_smem);
 467
 468          if (events)
 469             imm.combine(ctx.barrier_imm[idx]);
 470       }
 471    }
 472
 473    return imm;
 474 }
 475
 476 void force_waitcnt(wait_ctx& ctx, wait_imm& imm)
 477 {
 478    if (ctx.vm_cnt)
 479       imm.vm = 0;
 480    if (ctx.exp_cnt)
 481       imm.exp = 0;
 482    if (ctx.lgkm_cnt)
 483       imm.lgkm = 0;
 484
 485    if (ctx.chip_class >= GFX10) {
 486       if (ctx.vs_cnt)
 487          imm.vs = 0;
 488    }
 489 }
 490
 491 wait_imm kill(Instruction* instr, wait_ctx& ctx, memory_sync_info sync_info)
 492 {
 493    wait_imm imm;
 494
 495    if (debug_flags & DEBUG_FORCE_WAITCNT) {
 496       /* Force emitting waitcnt states right after the instruction if there is
 497        * something to wait for.
 498        */
 499       force_waitcnt(ctx, imm);
 500    }
 501
 502    if (ctx.exp_cnt || ctx.vm_cnt || ctx.lgkm_cnt)
 503       imm.combine(check_instr(instr, ctx));
 504
 505    imm.combine(parse_wait_instr(ctx, instr));
 506
 507
 508    /* It's required to wait for scalar stores before "writing back" data.
 509     * It shouldn't cost anything anyways since we're about to do s_endpgm.
 510     */
 511    if (ctx.lgkm_cnt && instr->opcode == aco_opcode::s_dcache_wb) {
 512       assert(ctx.chip_class >= GFX8);
 513       imm.lgkm = 0;
 514    }
 515
 516    if (ctx.chip_class >= GFX10 && instr->format == Format::SMEM) {
 517       /* GFX10: A store followed by a load at the same address causes a problem because
 518        * the load doesn't load the correct values unless we wait for the store first.
 519        * This is NOT mitigated by an s_nop.
 520        *
 521        * TODO: Refine this when we have proper alias analysis.
 522        */
 523       SMEM_instruction *smem = static_cast<SMEM_instruction *>(instr);
 524       if (ctx.pending_s_buffer_store &&
 525           !smem->definitions.empty() &&
 526           !smem->sync.can_reorder()) {
 527          imm.lgkm = 0;
 528       }
 529    }
 530
 531    if (instr->opcode == aco_opcode::p_barrier)
 532       imm.combine(perform_barrier(ctx, static_cast<Pseudo_barrier_instruction *>(instr)->sync, semantic_acqrel));
 533    else
 534       imm.combine(perform_barrier(ctx, sync_info, semantic_release));
 535
 536    if (!imm.empty()) {
 537       if (ctx.pending_flat_vm && imm.vm != wait_imm::unset_counter)
 538          imm.vm = 0;
 539       if (ctx.pending_flat_lgkm && imm.lgkm != wait_imm::unset_counter)
 540          imm.lgkm = 0;
 541
 542       /* reset counters */
 543       ctx.exp_cnt = std::min(ctx.exp_cnt, imm.exp);
 544       ctx.vm_cnt = std::min(ctx.vm_cnt, imm.vm);
 545       ctx.lgkm_cnt = std::min(ctx.lgkm_cnt, imm.lgkm);
 546       ctx.vs_cnt = std::min(ctx.vs_cnt, imm.vs);
 547
 548       /* update barrier wait imms */
 549       for (unsigned i = 0; i < storage_count; i++) {
 550          wait_imm& bar = ctx.barrier_imm[i];
 551          uint16_t& bar_ev = ctx.barrier_events[i];
 552          if (bar.exp != wait_imm::unset_counter && imm.exp <= bar.exp) {
 553             bar.exp = wait_imm::unset_counter;
 554             bar_ev &= ~exp_events;
 555          }
 556          if (bar.vm != wait_imm::unset_counter && imm.vm <= bar.vm) {
 557             bar.vm = wait_imm::unset_counter;
 558             bar_ev &= ~(vm_events & ~event_flat);
 559          }
 560          if (bar.lgkm != wait_imm::unset_counter && imm.lgkm <= bar.lgkm) {
 561             bar.lgkm = wait_imm::unset_counter;
 562             bar_ev &= ~(lgkm_events & ~event_flat);
 563          }
 564          if (bar.vs != wait_imm::unset_counter && imm.vs <= bar.vs) {
 565             bar.vs = wait_imm::unset_counter;
 566             bar_ev &= ~vs_events;
 567          }
 568          if (bar.vm == wait_imm::unset_counter && bar.lgkm == wait_imm::unset_counter)
 569             bar_ev &= ~event_flat;
 570       }
 571
 572       /* remove all gprs with higher counter from map */
 573       std::map<PhysReg,wait_entry>::iterator it = ctx.gpr_map.begin();
 574       while (it != ctx.gpr_map.end())
 575       {
 576          if (imm.exp != wait_imm::unset_counter && imm.exp <= it->second.imm.exp)
 577             ctx.wait_and_remove_from_entry(it->first, it->second, counter_exp);
 578          if (imm.vm != wait_imm::unset_counter && imm.vm <= it->second.imm.vm)
 579             ctx.wait_and_remove_from_entry(it->first, it->second, counter_vm);
 580          if (imm.lgkm != wait_imm::unset_counter && imm.lgkm <= it->second.imm.lgkm)
 581             ctx.wait_and_remove_from_entry(it->first, it->second, counter_lgkm);
 582          if (imm.vs != wait_imm::unset_counter && imm.vs <= it->second.imm.vs)
 583             ctx.wait_and_remove_from_entry(it->first, it->second, counter_vs);
 584          if (!it->second.counters)
 585             it = ctx.gpr_map.erase(it);
 586          else
 587             it++;
 588       }
 589    }
 590
 591    if (imm.vm == 0)
 592       ctx.pending_flat_vm = false;
 593    if (imm.lgkm == 0) {
 594       ctx.pending_flat_lgkm = false;
 595       ctx.pending_s_buffer_store = false;
 596    }
 597
 598    return imm;
 599 }
 600
 601 void update_barrier_counter(uint8_t *ctr, unsigned max)
 602 {
 603    if (*ctr != wait_imm::unset_counter && *ctr < max)
 604       (*ctr)++;
 605 }
 606
 607 void update_barrier_imm(wait_ctx& ctx, uint8_t counters, wait_event event, memory_sync_info sync)
 608 {
 609    for (unsigned i = 0; i < storage_count; i++) {
 610       wait_imm& bar = ctx.barrier_imm[i];
 611       uint16_t& bar_ev = ctx.barrier_events[i];
 612       if (sync.storage & (1 << i) && !(sync.semantics & semantic_private)) {
 613          bar_ev |= event;
 614          if (counters & counter_lgkm)
 615             bar.lgkm = 0;
 616          if (counters & counter_vm)
 617             bar.vm = 0;
 618          if (counters & counter_exp)
 619             bar.exp = 0;
 620          if (counters & counter_vs)
 621             bar.vs = 0;
 622       } else if (!(bar_ev & ctx.unordered_events) && !(ctx.unordered_events & event)) {
 623          if (counters & counter_lgkm && (bar_ev & lgkm_events) == event)
 624             update_barrier_counter(&bar.lgkm, ctx.max_lgkm_cnt);
 625          if (counters & counter_vm && (bar_ev & vm_events) == event)
 626             update_barrier_counter(&bar.vm, ctx.max_vm_cnt);
 627          if (counters & counter_exp && (bar_ev & exp_events) == event)
 628             update_barrier_counter(&bar.exp, ctx.max_exp_cnt);
 629          if (counters & counter_vs && (bar_ev & vs_events) == event)
 630             update_barrier_counter(&bar.vs, ctx.max_vs_cnt);
 631       }
 632    }
 633 }
 634
 635 void update_counters(wait_ctx& ctx, wait_event event, memory_sync_info sync=memory_sync_info())
 636 {
 637    uint8_t counters = get_counters_for_event(event);
 638
 639    if (counters & counter_lgkm && ctx.lgkm_cnt <= ctx.max_lgkm_cnt)
 640       ctx.lgkm_cnt++;
 641    if (counters & counter_vm && ctx.vm_cnt <= ctx.max_vm_cnt)
 642       ctx.vm_cnt++;
 643    if (counters & counter_exp && ctx.exp_cnt <= ctx.max_exp_cnt)
 644       ctx.exp_cnt++;
 645    if (counters & counter_vs && ctx.vs_cnt <= ctx.max_vs_cnt)
 646       ctx.vs_cnt++;
 647
 648    update_barrier_imm(ctx, counters, event, sync);
 649
 650    if (ctx.unordered_events & event)
 651       return;
 652
 653    if (ctx.pending_flat_lgkm)
 654       counters &= ~counter_lgkm;
 655    if (ctx.pending_flat_vm)
 656       counters &= ~counter_vm;
 657
 658    for (std::pair<const PhysReg,wait_entry>& e : ctx.gpr_map) {
 659       wait_entry& entry = e.second;
 660
 661       if (entry.events & ctx.unordered_events)
 662          continue;
 663
 664       assert(entry.events);
 665
 666       if ((counters & counter_exp) && (entry.events & exp_events) == event && entry.imm.exp < ctx.max_exp_cnt)
 667          entry.imm.exp++;
 668       if ((counters & counter_lgkm) && (entry.events & lgkm_events) == event && entry.imm.lgkm < ctx.max_lgkm_cnt)
 669          entry.imm.lgkm++;
 670       if ((counters & counter_vm) && (entry.events & vm_events) == event && entry.imm.vm < ctx.max_vm_cnt)
 671          entry.imm.vm++;
 672       if ((counters & counter_vs) && (entry.events & vs_events) == event && entry.imm.vs < ctx.max_vs_cnt)
 673          entry.imm.vs++;
 674    }
 675 }
 676
 677 void update_counters_for_flat_load(wait_ctx& ctx, memory_sync_info sync=memory_sync_info())
 678 {
 679    assert(ctx.chip_class < GFX10);
 680
 681    if (ctx.lgkm_cnt <= ctx.max_lgkm_cnt)
 682       ctx.lgkm_cnt++;
 683    if (ctx.vm_cnt <= ctx.max_vm_cnt)
 684       ctx.vm_cnt++;
 685
 686    update_barrier_imm(ctx, counter_vm | counter_lgkm, event_flat, sync);
 687
 688    for (std::pair<PhysReg,wait_entry> e : ctx.gpr_map)
 689    {
 690       if (e.second.counters & counter_vm)
 691          e.second.imm.vm = 0;
 692       if (e.second.counters & counter_lgkm)
 693          e.second.imm.lgkm = 0;
 694    }
 695    ctx.pending_flat_lgkm = true;
 696    ctx.pending_flat_vm = true;
 697 }
 698
 699 void insert_wait_entry(wait_ctx& ctx, PhysReg reg, RegClass rc, wait_event event, bool wait_on_read,
 700                        bool has_sampler=false)
 701 {
 702    uint16_t counters = get_counters_for_event(event);
 703    wait_imm imm;
 704    if (counters & counter_lgkm)
 705       imm.lgkm = 0;
 706    if (counters & counter_vm)
 707       imm.vm = 0;
 708    if (counters & counter_exp)
 709       imm.exp = 0;
 710    if (counters & counter_vs)
 711       imm.vs = 0;
 712
 713    wait_entry new_entry(event, imm, !rc.is_linear(), wait_on_read);
 714    new_entry.has_vmem_nosampler = (event & event_vmem) && !has_sampler;
 715    new_entry.has_vmem_sampler = (event & event_vmem) && has_sampler;
 716
 717    for (unsigned i = 0; i < rc.size(); i++) {
 718       auto it = ctx.gpr_map.emplace(PhysReg{reg.reg()+i}, new_entry);
 719       if (!it.second)
 720          it.first->second.join(new_entry);
 721    }
 722
 723    if (ctx.collect_statistics) {
 724       unsigned counters_todo = counters;
 725       while (counters_todo) {
 726          unsigned i = u_bit_scan(&counters_todo);
 727          ctx.unwaited_instrs[i].insert(std::make_pair(ctx.gen_instr, 0u));
 728          for (unsigned j = 0; j < rc.size(); j++)
 729             ctx.reg_instrs[i][PhysReg{reg.reg()+j}].insert(ctx.gen_instr);
 730       }
 731    }
 732 }
 733
 734 void insert_wait_entry(wait_ctx& ctx, Operand op, wait_event event, bool has_sampler=false)
 735 {
 736    if (!op.isConstant() && !op.isUndefined())
 737       insert_wait_entry(ctx, op.physReg(), op.regClass(), event, false, has_sampler);
 738 }
 739
 740 void insert_wait_entry(wait_ctx& ctx, Definition def, wait_event event, bool has_sampler=false)
 741 {
 742    insert_wait_entry(ctx, def.physReg(), def.regClass(), event, true, has_sampler);
 743 }
 744
 745 void gen(Instruction* instr, wait_ctx& ctx)
 746 {
 747    switch (instr->format) {
 748    case Format::EXP: {
 749       Export_instruction* exp_instr = static_cast<Export_instruction*>(instr);
 750
 751       wait_event ev;
 752       if (exp_instr->dest <= 9)
 753          ev = event_exp_mrt_null;
 754       else if (exp_instr->dest <= 15)
 755          ev = event_exp_pos;
 756       else
 757          ev = event_exp_param;
 758       update_counters(ctx, ev);
 759
 760       /* insert new entries for exported vgprs */
 761       for (unsigned i = 0; i < 4; i++)
 762       {
 763          if (exp_instr->enabled_mask & (1 << i)) {
 764             unsigned idx = exp_instr->compressed ? i >> 1 : i;
 765             assert(idx < exp_instr->operands.size());
 766             insert_wait_entry(ctx, exp_instr->operands[idx], ev);
 767
 768          }
 769       }
 770       insert_wait_entry(ctx, exec, s2, ev, false);
 771       break;
 772    }
 773    case Format::FLAT: {
 774       FLAT_instruction *flat = static_cast<FLAT_instruction*>(instr);
 775       if (ctx.chip_class < GFX10 && !instr->definitions.empty())
 776          update_counters_for_flat_load(ctx, flat->sync);
 777       else
 778          update_counters(ctx, event_flat, flat->sync);
 779
 780       if (!instr->definitions.empty())
 781          insert_wait_entry(ctx, instr->definitions[0], event_flat);
 782       break;
 783    }
 784    case Format::SMEM: {
 785       SMEM_instruction *smem = static_cast<SMEM_instruction*>(instr);
 786       update_counters(ctx, event_smem, smem->sync);
 787
 788       if (!instr->definitions.empty())
 789          insert_wait_entry(ctx, instr->definitions[0], event_smem);
 790       else if (ctx.chip_class >= GFX10 &&
 791                !smem->sync.can_reorder())
 792          ctx.pending_s_buffer_store = true;
 793
 794       break;
 795    }
 796    case Format::DS: {
 797       DS_instruction *ds = static_cast<DS_instruction*>(instr);
 798       update_counters(ctx, ds->gds ? event_gds : event_lds, ds->sync);
 799       if (ds->gds)
 800          update_counters(ctx, event_gds_gpr_lock);
 801
 802       if (!instr->definitions.empty())
 803          insert_wait_entry(ctx, instr->definitions[0], ds->gds ? event_gds : event_lds);
 804
 805       if (ds->gds) {
 806          for (const Operand& op : instr->operands)
 807             insert_wait_entry(ctx, op, event_gds_gpr_lock);
 808          insert_wait_entry(ctx, exec, s2, event_gds_gpr_lock, false);
 809       }
 810       break;
 811    }
 812    case Format::MUBUF:
 813    case Format::MTBUF:
 814    case Format::MIMG:
 815    case Format::GLOBAL: {
 816       wait_event ev = !instr->definitions.empty() || ctx.chip_class < GFX10 ? event_vmem : event_vmem_store;
 817       update_counters(ctx, ev, get_sync_info(instr));
 818
 819       bool has_sampler = instr->format == Format::MIMG && !instr->operands[1].isUndefined() && instr->operands[1].regClass() == s4;
 820
 821       if (!instr->definitions.empty())
 822          insert_wait_entry(ctx, instr->definitions[0], ev, has_sampler);
 823
 824       if (ctx.chip_class == GFX6 &&
 825           instr->format != Format::MIMG &&
 826           instr->operands.size() == 4) {
 827          ctx.exp_cnt++;
 828          update_counters(ctx, event_vmem_gpr_lock);
 829          insert_wait_entry(ctx, instr->operands[3], event_vmem_gpr_lock);
 830       } else if (ctx.chip_class == GFX6 &&
 831                  instr->format == Format::MIMG &&
 832                  instr->operands[1].regClass().type() == RegType::vgpr) {
 833          ctx.exp_cnt++;
 834          update_counters(ctx, event_vmem_gpr_lock);
 835          insert_wait_entry(ctx, instr->operands[1], event_vmem_gpr_lock);
 836       }
 837
 838       break;
 839    }
 840    case Format::SOPP: {
 841       if (instr->opcode == aco_opcode::s_sendmsg ||
 842           instr->opcode == aco_opcode::s_sendmsghalt)
 843          update_counters(ctx, event_sendmsg);
 844    }
 845    default:
 846       break;
 847    }
 848 }
 849
 850 void emit_waitcnt(wait_ctx& ctx, std::vector<aco_ptr<Instruction>>& instructions, wait_imm imm)
 851 {
 852    if (imm.vs != wait_imm::unset_counter) {
 853       assert(ctx.chip_class >= GFX10);
 854       SOPK_instruction* waitcnt_vs = create_instruction<SOPK_instruction>(aco_opcode::s_waitcnt_vscnt, Format::SOPK, 0, 1);
 855       waitcnt_vs->definitions[0] = Definition(sgpr_null, s1);
 856       waitcnt_vs->imm = imm.vs;
 857       instructions.emplace_back(waitcnt_vs);
 858       imm.vs = wait_imm::unset_counter;
 859    }
 860    if (!imm.empty()) {
 861       SOPP_instruction* waitcnt = create_instruction<SOPP_instruction>(aco_opcode::s_waitcnt, Format::SOPP, 0, 0);
 862       waitcnt->imm = imm.pack(ctx.chip_class);
 863       waitcnt->block = -1;
 864       instructions.emplace_back(waitcnt);
 865    }
 866 }
 867
 868 void handle_block(Program *program, Block& block, wait_ctx& ctx)
 869 {
 870    std::vector<aco_ptr<Instruction>> new_instructions;
 871
 872    wait_imm queued_imm;
 873
 874    for (aco_ptr<Instruction>& instr : block.instructions) {
 875       bool is_wait = !parse_wait_instr(ctx, instr.get()).empty();
 876
 877       memory_sync_info sync_info = get_sync_info(instr.get());
 878       queued_imm.combine(kill(instr.get(), ctx, sync_info));
 879
 880       ctx.gen_instr = instr.get();
 881       gen(instr.get(), ctx);
 882
 883       if (instr->format != Format::PSEUDO_BARRIER && !is_wait) {
 884          if (!queued_imm.empty()) {
 885             emit_waitcnt(ctx, new_instructions, queued_imm);
 886             queued_imm = wait_imm();
 887          }
 888          new_instructions.emplace_back(std::move(instr));
 889
 890          queued_imm.combine(perform_barrier(ctx, sync_info, semantic_acquire));
 891
 892          if (ctx.collect_statistics)
 893             ctx.advance_unwaited_instrs();
 894       }
 895    }
 896
 897    if (!queued_imm.empty())
 898       emit_waitcnt(ctx, new_instructions, queued_imm);
 899
 900    block.instructions.swap(new_instructions);
 901 }
 902
 903 } /* end namespace */
 904
 905 static uint32_t calculate_score(std::vector<wait_ctx> &ctx_vec, uint32_t event_mask)
 906 {
 907    double result = 0.0;
 908    unsigned num_waits = 0;
 909    while (event_mask) {
 910       unsigned event_index = u_bit_scan(&event_mask);
 911       for (const wait_ctx &ctx : ctx_vec) {
 912          for (unsigned dist : ctx.wait_distances[event_index]) {
 913             double score = dist;
 914             /* for many events, excessive distances provide little benefit, so
 915              * decrease the score in that case. */
 916             double threshold = INFINITY;
 917             double inv_strength = 0.000001;
 918             switch (1 << event_index) {
 919             case event_smem:
 920                threshold = 70.0;
 921                inv_strength = 75.0;
 922                break;
 923             case event_vmem:
 924             case event_vmem_store:
 925             case event_flat:
 926                threshold = 230.0;
 927                inv_strength = 150.0;
 928                break;
 929             case event_lds:
 930                threshold = 16.0;
 931                break;
 932             default:
 933                break;
 934             }
 935             if (score > threshold) {
 936                score -= threshold;
 937                score = threshold + score / (1.0 + score / inv_strength);
 938             }
 939
 940             /* we don't want increases in high scores to hide decreases in low scores,
 941              * so raise to the power of 0.1 before averaging. */
 942             result += pow(score, 0.1);
 943             num_waits++;
 944          }
 945       }
 946    }
 947    return round(pow(result / num_waits, 10.0) * 10.0);
 948 }
 949
 950 void insert_wait_states(Program* program)
 951 {
 952    /* per BB ctx */
 953    std::vector<bool> done(program->blocks.size());
 954    std::vector<wait_ctx> in_ctx(program->blocks.size(), wait_ctx(program));
 955    std::vector<wait_ctx> out_ctx(program->blocks.size(), wait_ctx(program));
 956
 957    std::stack<unsigned> loop_header_indices;
 958    unsigned loop_progress = 0;
 959
 960    for (unsigned i = 0; i < program->blocks.size();) {
 961       Block& current = program->blocks[i++];
 962       wait_ctx ctx = in_ctx[current.index];
 963
 964       if (current.kind & block_kind_loop_header) {
 965          loop_header_indices.push(current.index);
 966       } else if (current.kind & block_kind_loop_exit) {
 967          bool repeat = false;
 968          if (loop_progress == loop_header_indices.size()) {
 969             i = loop_header_indices.top();
 970             repeat = true;
 971          }
 972          loop_header_indices.pop();
 973          loop_progress = std::min<unsigned>(loop_progress, loop_header_indices.size());
 974          if (repeat)
 975             continue;
 976       }
 977
 978       bool changed = false;
 979       for (unsigned b : current.linear_preds)
 980          changed |= ctx.join(&out_ctx[b], false);
 981       for (unsigned b : current.logical_preds)
 982          changed |= ctx.join(&out_ctx[b], true);
 983
 984       if (done[current.index] && !changed) {
 985          in_ctx[current.index] = std::move(ctx);
 986          continue;
 987       } else {
 988          in_ctx[current.index] = ctx;
 989       }
 990
 991       if (current.instructions.empty()) {
 992          out_ctx[current.index] = std::move(ctx);
 993          continue;
 994       }
 995
 996       loop_progress = std::max<unsigned>(loop_progress, current.loop_nest_depth);
 997       done[current.index] = true;
 998
 999       handle_block(program, current, ctx);
1000
1001       out_ctx[current.index] = std::move(ctx);
1002    }
1003
1004    if (program->collect_statistics) {
1005       program->statistics[statistic_vmem_score] =
1006          calculate_score(out_ctx, event_vmem | event_flat | event_vmem_store);
1007       program->statistics[statistic_smem_score] =
1008          calculate_score(out_ctx, event_smem);
1009    }
1010 }
1011
1012 }
1013