aco: fix waitcnt insertion on GFX10.3
[mesa.git] / 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 wait_imm kill(Instruction* instr, wait_ctx& ctx, memory_sync_info sync_info)
477 {
478 wait_imm imm;
479 if (ctx.exp_cnt || ctx.vm_cnt || ctx.lgkm_cnt)
480 imm.combine(check_instr(instr, ctx));
481
482 imm.combine(parse_wait_instr(ctx, instr));
483
484
485 /* It's required to wait for scalar stores before "writing back" data.
486 * It shouldn't cost anything anyways since we're about to do s_endpgm.
487 */
488 if (ctx.lgkm_cnt && instr->opcode == aco_opcode::s_dcache_wb) {
489 assert(ctx.chip_class >= GFX8);
490 imm.lgkm = 0;
491 }
492
493 if (ctx.chip_class >= GFX10 && instr->format == Format::SMEM) {
494 /* GFX10: A store followed by a load at the same address causes a problem because
495 * the load doesn't load the correct values unless we wait for the store first.
496 * This is NOT mitigated by an s_nop.
497 *
498 * TODO: Refine this when we have proper alias analysis.
499 */
500 SMEM_instruction *smem = static_cast<SMEM_instruction *>(instr);
501 if (ctx.pending_s_buffer_store &&
502 !smem->definitions.empty() &&
503 !smem->sync.can_reorder()) {
504 imm.lgkm = 0;
505 }
506 }
507
508 if (instr->opcode == aco_opcode::p_barrier)
509 imm.combine(perform_barrier(ctx, static_cast<Pseudo_barrier_instruction *>(instr)->sync, semantic_acqrel));
510 else
511 imm.combine(perform_barrier(ctx, sync_info, semantic_release));
512
513 if (!imm.empty()) {
514 if (ctx.pending_flat_vm && imm.vm != wait_imm::unset_counter)
515 imm.vm = 0;
516 if (ctx.pending_flat_lgkm && imm.lgkm != wait_imm::unset_counter)
517 imm.lgkm = 0;
518
519 /* reset counters */
520 ctx.exp_cnt = std::min(ctx.exp_cnt, imm.exp);
521 ctx.vm_cnt = std::min(ctx.vm_cnt, imm.vm);
522 ctx.lgkm_cnt = std::min(ctx.lgkm_cnt, imm.lgkm);
523 ctx.vs_cnt = std::min(ctx.vs_cnt, imm.vs);
524
525 /* update barrier wait imms */
526 for (unsigned i = 0; i < storage_count; i++) {
527 wait_imm& bar = ctx.barrier_imm[i];
528 uint16_t& bar_ev = ctx.barrier_events[i];
529 if (bar.exp != wait_imm::unset_counter && imm.exp <= bar.exp) {
530 bar.exp = wait_imm::unset_counter;
531 bar_ev &= ~exp_events;
532 }
533 if (bar.vm != wait_imm::unset_counter && imm.vm <= bar.vm) {
534 bar.vm = wait_imm::unset_counter;
535 bar_ev &= ~(vm_events & ~event_flat);
536 }
537 if (bar.lgkm != wait_imm::unset_counter && imm.lgkm <= bar.lgkm) {
538 bar.lgkm = wait_imm::unset_counter;
539 bar_ev &= ~(lgkm_events & ~event_flat);
540 }
541 if (bar.vs != wait_imm::unset_counter && imm.vs <= bar.vs) {
542 bar.vs = wait_imm::unset_counter;
543 bar_ev &= ~vs_events;
544 }
545 if (bar.vm == wait_imm::unset_counter && bar.lgkm == wait_imm::unset_counter)
546 bar_ev &= ~event_flat;
547 }
548
549 /* remove all gprs with higher counter from map */
550 std::map<PhysReg,wait_entry>::iterator it = ctx.gpr_map.begin();
551 while (it != ctx.gpr_map.end())
552 {
553 if (imm.exp != wait_imm::unset_counter && imm.exp <= it->second.imm.exp)
554 ctx.wait_and_remove_from_entry(it->first, it->second, counter_exp);
555 if (imm.vm != wait_imm::unset_counter && imm.vm <= it->second.imm.vm)
556 ctx.wait_and_remove_from_entry(it->first, it->second, counter_vm);
557 if (imm.lgkm != wait_imm::unset_counter && imm.lgkm <= it->second.imm.lgkm)
558 ctx.wait_and_remove_from_entry(it->first, it->second, counter_lgkm);
559 if (imm.vs != wait_imm::unset_counter && imm.vs <= it->second.imm.vs)
560 ctx.wait_and_remove_from_entry(it->first, it->second, counter_vs);
561 if (!it->second.counters)
562 it = ctx.gpr_map.erase(it);
563 else
564 it++;
565 }
566 }
567
568 if (imm.vm == 0)
569 ctx.pending_flat_vm = false;
570 if (imm.lgkm == 0) {
571 ctx.pending_flat_lgkm = false;
572 ctx.pending_s_buffer_store = false;
573 }
574
575 return imm;
576 }
577
578 void update_barrier_counter(uint8_t *ctr, unsigned max)
579 {
580 if (*ctr != wait_imm::unset_counter && *ctr < max)
581 (*ctr)++;
582 }
583
584 void update_barrier_imm(wait_ctx& ctx, uint8_t counters, wait_event event, memory_sync_info sync)
585 {
586 for (unsigned i = 0; i < storage_count; i++) {
587 wait_imm& bar = ctx.barrier_imm[i];
588 uint16_t& bar_ev = ctx.barrier_events[i];
589 if (sync.storage & (1 << i) && !(sync.semantics & semantic_private)) {
590 bar_ev |= event;
591 if (counters & counter_lgkm)
592 bar.lgkm = 0;
593 if (counters & counter_vm)
594 bar.vm = 0;
595 if (counters & counter_exp)
596 bar.exp = 0;
597 if (counters & counter_vs)
598 bar.vs = 0;
599 } else if (!(bar_ev & ctx.unordered_events) && !(ctx.unordered_events & event)) {
600 if (counters & counter_lgkm && (bar_ev & lgkm_events) == event)
601 update_barrier_counter(&bar.lgkm, ctx.max_lgkm_cnt);
602 if (counters & counter_vm && (bar_ev & vm_events) == event)
603 update_barrier_counter(&bar.vm, ctx.max_vm_cnt);
604 if (counters & counter_exp && (bar_ev & exp_events) == event)
605 update_barrier_counter(&bar.exp, ctx.max_exp_cnt);
606 if (counters & counter_vs && (bar_ev & vs_events) == event)
607 update_barrier_counter(&bar.vs, ctx.max_vs_cnt);
608 }
609 }
610 }
611
612 void update_counters(wait_ctx& ctx, wait_event event, memory_sync_info sync=memory_sync_info())
613 {
614 uint8_t counters = get_counters_for_event(event);
615
616 if (counters & counter_lgkm && ctx.lgkm_cnt <= ctx.max_lgkm_cnt)
617 ctx.lgkm_cnt++;
618 if (counters & counter_vm && ctx.vm_cnt <= ctx.max_vm_cnt)
619 ctx.vm_cnt++;
620 if (counters & counter_exp && ctx.exp_cnt <= ctx.max_exp_cnt)
621 ctx.exp_cnt++;
622 if (counters & counter_vs && ctx.vs_cnt <= ctx.max_vs_cnt)
623 ctx.vs_cnt++;
624
625 update_barrier_imm(ctx, counters, event, sync);
626
627 if (ctx.unordered_events & event)
628 return;
629
630 if (ctx.pending_flat_lgkm)
631 counters &= ~counter_lgkm;
632 if (ctx.pending_flat_vm)
633 counters &= ~counter_vm;
634
635 for (std::pair<const PhysReg,wait_entry>& e : ctx.gpr_map) {
636 wait_entry& entry = e.second;
637
638 if (entry.events & ctx.unordered_events)
639 continue;
640
641 assert(entry.events);
642
643 if ((counters & counter_exp) && (entry.events & exp_events) == event && entry.imm.exp < ctx.max_exp_cnt)
644 entry.imm.exp++;
645 if ((counters & counter_lgkm) && (entry.events & lgkm_events) == event && entry.imm.lgkm < ctx.max_lgkm_cnt)
646 entry.imm.lgkm++;
647 if ((counters & counter_vm) && (entry.events & vm_events) == event && entry.imm.vm < ctx.max_vm_cnt)
648 entry.imm.vm++;
649 if ((counters & counter_vs) && (entry.events & vs_events) == event && entry.imm.vs < ctx.max_vs_cnt)
650 entry.imm.vs++;
651 }
652 }
653
654 void update_counters_for_flat_load(wait_ctx& ctx, memory_sync_info sync=memory_sync_info())
655 {
656 assert(ctx.chip_class < GFX10);
657
658 if (ctx.lgkm_cnt <= ctx.max_lgkm_cnt)
659 ctx.lgkm_cnt++;
660 if (ctx.vm_cnt <= ctx.max_vm_cnt)
661 ctx.vm_cnt++;
662
663 update_barrier_imm(ctx, counter_vm | counter_lgkm, event_flat, sync);
664
665 for (std::pair<PhysReg,wait_entry> e : ctx.gpr_map)
666 {
667 if (e.second.counters & counter_vm)
668 e.second.imm.vm = 0;
669 if (e.second.counters & counter_lgkm)
670 e.second.imm.lgkm = 0;
671 }
672 ctx.pending_flat_lgkm = true;
673 ctx.pending_flat_vm = true;
674 }
675
676 void insert_wait_entry(wait_ctx& ctx, PhysReg reg, RegClass rc, wait_event event, bool wait_on_read,
677 bool has_sampler=false)
678 {
679 uint16_t counters = get_counters_for_event(event);
680 wait_imm imm;
681 if (counters & counter_lgkm)
682 imm.lgkm = 0;
683 if (counters & counter_vm)
684 imm.vm = 0;
685 if (counters & counter_exp)
686 imm.exp = 0;
687 if (counters & counter_vs)
688 imm.vs = 0;
689
690 wait_entry new_entry(event, imm, !rc.is_linear(), wait_on_read);
691 new_entry.has_vmem_nosampler = (event & event_vmem) && !has_sampler;
692 new_entry.has_vmem_sampler = (event & event_vmem) && has_sampler;
693
694 for (unsigned i = 0; i < rc.size(); i++) {
695 auto it = ctx.gpr_map.emplace(PhysReg{reg.reg()+i}, new_entry);
696 if (!it.second)
697 it.first->second.join(new_entry);
698 }
699
700 if (ctx.collect_statistics) {
701 unsigned counters_todo = counters;
702 while (counters_todo) {
703 unsigned i = u_bit_scan(&counters_todo);
704 ctx.unwaited_instrs[i].insert(std::make_pair(ctx.gen_instr, 0u));
705 for (unsigned j = 0; j < rc.size(); j++)
706 ctx.reg_instrs[i][PhysReg{reg.reg()+j}].insert(ctx.gen_instr);
707 }
708 }
709 }
710
711 void insert_wait_entry(wait_ctx& ctx, Operand op, wait_event event, bool has_sampler=false)
712 {
713 if (!op.isConstant() && !op.isUndefined())
714 insert_wait_entry(ctx, op.physReg(), op.regClass(), event, false, has_sampler);
715 }
716
717 void insert_wait_entry(wait_ctx& ctx, Definition def, wait_event event, bool has_sampler=false)
718 {
719 insert_wait_entry(ctx, def.physReg(), def.regClass(), event, true, has_sampler);
720 }
721
722 void gen(Instruction* instr, wait_ctx& ctx)
723 {
724 switch (instr->format) {
725 case Format::EXP: {
726 Export_instruction* exp_instr = static_cast<Export_instruction*>(instr);
727
728 wait_event ev;
729 if (exp_instr->dest <= 9)
730 ev = event_exp_mrt_null;
731 else if (exp_instr->dest <= 15)
732 ev = event_exp_pos;
733 else
734 ev = event_exp_param;
735 update_counters(ctx, ev);
736
737 /* insert new entries for exported vgprs */
738 for (unsigned i = 0; i < 4; i++)
739 {
740 if (exp_instr->enabled_mask & (1 << i)) {
741 unsigned idx = exp_instr->compressed ? i >> 1 : i;
742 assert(idx < exp_instr->operands.size());
743 insert_wait_entry(ctx, exp_instr->operands[idx], ev);
744
745 }
746 }
747 insert_wait_entry(ctx, exec, s2, ev, false);
748 break;
749 }
750 case Format::FLAT: {
751 FLAT_instruction *flat = static_cast<FLAT_instruction*>(instr);
752 if (ctx.chip_class < GFX10 && !instr->definitions.empty())
753 update_counters_for_flat_load(ctx, flat->sync);
754 else
755 update_counters(ctx, event_flat, flat->sync);
756
757 if (!instr->definitions.empty())
758 insert_wait_entry(ctx, instr->definitions[0], event_flat);
759 break;
760 }
761 case Format::SMEM: {
762 SMEM_instruction *smem = static_cast<SMEM_instruction*>(instr);
763 update_counters(ctx, event_smem, smem->sync);
764
765 if (!instr->definitions.empty())
766 insert_wait_entry(ctx, instr->definitions[0], event_smem);
767 else if (ctx.chip_class >= GFX10 &&
768 !smem->sync.can_reorder())
769 ctx.pending_s_buffer_store = true;
770
771 break;
772 }
773 case Format::DS: {
774 DS_instruction *ds = static_cast<DS_instruction*>(instr);
775 update_counters(ctx, ds->gds ? event_gds : event_lds, ds->sync);
776 if (ds->gds)
777 update_counters(ctx, event_gds_gpr_lock);
778
779 if (!instr->definitions.empty())
780 insert_wait_entry(ctx, instr->definitions[0], ds->gds ? event_gds : event_lds);
781
782 if (ds->gds) {
783 for (const Operand& op : instr->operands)
784 insert_wait_entry(ctx, op, event_gds_gpr_lock);
785 insert_wait_entry(ctx, exec, s2, event_gds_gpr_lock, false);
786 }
787 break;
788 }
789 case Format::MUBUF:
790 case Format::MTBUF:
791 case Format::MIMG:
792 case Format::GLOBAL: {
793 wait_event ev = !instr->definitions.empty() || ctx.chip_class < GFX10 ? event_vmem : event_vmem_store;
794 update_counters(ctx, ev, get_sync_info(instr));
795
796 bool has_sampler = instr->format == Format::MIMG && !instr->operands[1].isUndefined() && instr->operands[1].regClass() == s4;
797
798 if (!instr->definitions.empty())
799 insert_wait_entry(ctx, instr->definitions[0], ev, has_sampler);
800
801 if (ctx.chip_class == GFX6 &&
802 instr->format != Format::MIMG &&
803 instr->operands.size() == 4) {
804 ctx.exp_cnt++;
805 update_counters(ctx, event_vmem_gpr_lock);
806 insert_wait_entry(ctx, instr->operands[3], event_vmem_gpr_lock);
807 } else if (ctx.chip_class == GFX6 &&
808 instr->format == Format::MIMG &&
809 instr->operands[1].regClass().type() == RegType::vgpr) {
810 ctx.exp_cnt++;
811 update_counters(ctx, event_vmem_gpr_lock);
812 insert_wait_entry(ctx, instr->operands[1], event_vmem_gpr_lock);
813 }
814
815 break;
816 }
817 case Format::SOPP: {
818 if (instr->opcode == aco_opcode::s_sendmsg ||
819 instr->opcode == aco_opcode::s_sendmsghalt)
820 update_counters(ctx, event_sendmsg);
821 }
822 default:
823 break;
824 }
825 }
826
827 void emit_waitcnt(wait_ctx& ctx, std::vector<aco_ptr<Instruction>>& instructions, wait_imm imm)
828 {
829 if (imm.vs != wait_imm::unset_counter) {
830 assert(ctx.chip_class >= GFX10);
831 SOPK_instruction* waitcnt_vs = create_instruction<SOPK_instruction>(aco_opcode::s_waitcnt_vscnt, Format::SOPK, 0, 1);
832 waitcnt_vs->definitions[0] = Definition(sgpr_null, s1);
833 waitcnt_vs->imm = imm.vs;
834 instructions.emplace_back(waitcnt_vs);
835 imm.vs = wait_imm::unset_counter;
836 }
837 if (!imm.empty()) {
838 SOPP_instruction* waitcnt = create_instruction<SOPP_instruction>(aco_opcode::s_waitcnt, Format::SOPP, 0, 0);
839 waitcnt->imm = imm.pack(ctx.chip_class);
840 waitcnt->block = -1;
841 instructions.emplace_back(waitcnt);
842 }
843 }
844
845 void handle_block(Program *program, Block& block, wait_ctx& ctx)
846 {
847 std::vector<aco_ptr<Instruction>> new_instructions;
848
849 wait_imm queued_imm;
850
851 for (aco_ptr<Instruction>& instr : block.instructions) {
852 bool is_wait = !parse_wait_instr(ctx, instr.get()).empty();
853
854 memory_sync_info sync_info = get_sync_info(instr.get());
855 queued_imm.combine(kill(instr.get(), ctx, sync_info));
856
857 ctx.gen_instr = instr.get();
858 gen(instr.get(), ctx);
859
860 if (instr->format != Format::PSEUDO_BARRIER && !is_wait) {
861 if (!queued_imm.empty()) {
862 emit_waitcnt(ctx, new_instructions, queued_imm);
863 queued_imm = wait_imm();
864 }
865 new_instructions.emplace_back(std::move(instr));
866
867 queued_imm.combine(perform_barrier(ctx, sync_info, semantic_acquire));
868
869 if (ctx.collect_statistics)
870 ctx.advance_unwaited_instrs();
871 }
872 }
873
874 if (!queued_imm.empty())
875 emit_waitcnt(ctx, new_instructions, queued_imm);
876
877 block.instructions.swap(new_instructions);
878 }
879
880 } /* end namespace */
881
882 static uint32_t calculate_score(std::vector<wait_ctx> &ctx_vec, uint32_t event_mask)
883 {
884 double result = 0.0;
885 unsigned num_waits = 0;
886 while (event_mask) {
887 unsigned event_index = u_bit_scan(&event_mask);
888 for (const wait_ctx &ctx : ctx_vec) {
889 for (unsigned dist : ctx.wait_distances[event_index]) {
890 double score = dist;
891 /* for many events, excessive distances provide little benefit, so
892 * decrease the score in that case. */
893 double threshold = INFINITY;
894 double inv_strength = 0.000001;
895 switch (1 << event_index) {
896 case event_smem:
897 threshold = 70.0;
898 inv_strength = 75.0;
899 break;
900 case event_vmem:
901 case event_vmem_store:
902 case event_flat:
903 threshold = 230.0;
904 inv_strength = 150.0;
905 break;
906 case event_lds:
907 threshold = 16.0;
908 break;
909 default:
910 break;
911 }
912 if (score > threshold) {
913 score -= threshold;
914 score = threshold + score / (1.0 + score / inv_strength);
915 }
916
917 /* we don't want increases in high scores to hide decreases in low scores,
918 * so raise to the power of 0.1 before averaging. */
919 result += pow(score, 0.1);
920 num_waits++;
921 }
922 }
923 }
924 return round(pow(result / num_waits, 10.0) * 10.0);
925 }
926
927 void insert_wait_states(Program* program)
928 {
929 /* per BB ctx */
930 std::vector<bool> done(program->blocks.size());
931 std::vector<wait_ctx> in_ctx(program->blocks.size(), wait_ctx(program));
932 std::vector<wait_ctx> out_ctx(program->blocks.size(), wait_ctx(program));
933
934 std::stack<unsigned> loop_header_indices;
935 unsigned loop_progress = 0;
936
937 for (unsigned i = 0; i < program->blocks.size();) {
938 Block& current = program->blocks[i++];
939 wait_ctx ctx = in_ctx[current.index];
940
941 if (current.kind & block_kind_loop_header) {
942 loop_header_indices.push(current.index);
943 } else if (current.kind & block_kind_loop_exit) {
944 bool repeat = false;
945 if (loop_progress == loop_header_indices.size()) {
946 i = loop_header_indices.top();
947 repeat = true;
948 }
949 loop_header_indices.pop();
950 loop_progress = std::min<unsigned>(loop_progress, loop_header_indices.size());
951 if (repeat)
952 continue;
953 }
954
955 bool changed = false;
956 for (unsigned b : current.linear_preds)
957 changed |= ctx.join(&out_ctx[b], false);
958 for (unsigned b : current.logical_preds)
959 changed |= ctx.join(&out_ctx[b], true);
960
961 if (done[current.index] && !changed) {
962 in_ctx[current.index] = std::move(ctx);
963 continue;
964 } else {
965 in_ctx[current.index] = ctx;
966 }
967
968 if (current.instructions.empty()) {
969 out_ctx[current.index] = std::move(ctx);
970 continue;
971 }
972
973 loop_progress = std::max<unsigned>(loop_progress, current.loop_nest_depth);
974 done[current.index] = true;
975
976 handle_block(program, current, ctx);
977
978 out_ctx[current.index] = std::move(ctx);
979 }
980
981 if (program->collect_statistics) {
982 program->statistics[statistic_vmem_score] =
983 calculate_score(out_ctx, event_vmem | event_flat | event_vmem_store);
984 program->statistics[statistic_smem_score] =
985 calculate_score(out_ctx, event_smem);
986 }
987 }
988
989 }
990