2 * Copyright © 2018 Valve Corporation
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:
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
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
31 #include "vulkan/radv_shader.h"
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.
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.
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
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 {
62 event_vmem_store
= 1 << 4, /* GFX10+ */
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,
73 enum counter_type
: uint8_t {
75 counter_lgkm
= 1 << 1,
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
;
86 uint8_t get_counters_for_event(wait_event ev
)
96 case event_vmem_store
:
99 return counter_vm
| counter_lgkm
;
101 case event_exp_param
:
102 case event_exp_mrt_null
:
103 case event_gds_gpr_lock
:
104 case event_vmem_gpr_lock
:
111 uint16_t get_events_for_counter(counter_type ctr
)
127 static const uint8_t unset_counter
= 0xff;
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_
) {}
139 wait_imm(enum chip_class chip
, uint16_t packed
) : vs(unset_counter
)
143 vm
|= (packed
>> 10) & 0x30;
145 exp
= (packed
>> 4) & 0x7;
147 lgkm
= (packed
>> 8) & 0xf;
149 lgkm
|= (packed
>> 8) & 0x30;
152 uint16_t pack(enum chip_class chip
) const
155 assert(exp
== unset_counter
|| exp
<= 0x7);
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);
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);
169 assert(lgkm
== unset_counter
|| lgkm
<= 0xf);
170 assert(vm
== unset_counter
|| vm
<= 0xf);
171 imm
= ((lgkm
& 0xf) << 8) | ((exp
& 0x7) << 4) | (vm
& 0xf);
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 */
181 bool combine(const wait_imm
& other
)
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
);
193 return vm
== unset_counter
&& exp
== unset_counter
&&
194 lgkm
== unset_counter
&& vs
== unset_counter
;
200 uint16_t events
; /* use wait_event notion */
201 uint8_t counters
; /* use counter_type notion */
204 bool has_vmem_nosampler
:1;
205 bool has_vmem_sampler
:1;
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) {}
212 bool join(const wait_entry
& other
)
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
);
229 void remove_counter(counter_type counter
)
231 counters
&= ~counter
;
233 if (counter
== counter_lgkm
) {
234 imm
.lgkm
= wait_imm::unset_counter
;
235 events
&= ~(event_smem
| event_lds
| event_gds
| event_sendmsg
);
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;
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
);
250 if (counter
== counter_vs
) {
251 imm
.vs
= wait_imm::unset_counter
;
252 events
&= ~event_vmem_store
;
255 if (!(counters
& counter_lgkm
) && !(counters
& counter_vm
))
256 events
&= ~event_flat
;
262 enum chip_class chip_class
;
264 uint16_t max_exp_cnt
;
265 uint16_t max_lgkm_cnt
;
267 uint16_t unordered_events
= event_smem
| event_flat
;
271 uint8_t lgkm_cnt
= 0;
273 bool pending_flat_lgkm
= false;
274 bool pending_flat_vm
= false;
275 bool pending_s_buffer_store
= false; /* GFX10 workaround */
277 wait_imm barrier_imm
[storage_count
];
278 uint16_t barrier_events
[storage_count
] = {}; /* use wait_event notion */
280 std::map
<PhysReg
,wait_entry
> gpr_map
;
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
];
290 wait_ctx(Program
*program_
)
292 chip_class(program_
->chip_class
),
293 max_vm_cnt(program_
->chip_class
>= GFX9
? 62 : 14),
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
) {}
300 bool join(const wait_ctx
* other
, bool logical
)
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
);
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
;
317 for (std::pair
<PhysReg
,wait_entry
> entry
: other
->gpr_map
)
319 std::map
<PhysReg
,wait_entry
>::iterator it
= gpr_map
.find(entry
.first
);
320 if (entry
.second
.logical
!= logical
)
323 if (it
!= gpr_map
.end()) {
324 changed
|= it
->second
.join(entry
.second
);
326 gpr_map
.insert(entry
);
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
];
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
);
344 pos
->second
= std::min(pos
->second
, instr
.second
);
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());
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())
362 unsigned distance
= pos
->second
;
363 unsigned events
= entry
.events
& get_events_for_counter(counter
);
365 unsigned event_idx
= u_bit_scan(&events
);
366 wait_distances
[event_idx
].push_back(distance
);
369 unwaited_instrs
[counter_idx
].erase(instr
);
371 reg_instrs
[counter_idx
][reg
].clear();
374 entry
.remove_counter(counter
);
377 void advance_unwaited_instrs()
379 for (unsigned i
= 0; i
< num_counters
; i
++) {
380 for (auto it
= unwaited_instrs
[i
].begin(); it
!= unwaited_instrs
[i
].end(); ++it
)
386 wait_imm
check_instr(Instruction
* instr
, wait_ctx
& ctx
)
390 for (const Operand op
: instr
->operands
) {
391 if (op
.isConstant() || op
.isUndefined())
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
)
401 wait
.combine(it
->second
.imm
);
405 for (const Definition
& def
: instr
->definitions
) {
406 /* check consecutively written gprs */
407 for (unsigned j
= 0; j
< def
.getTemp().size(); j
++)
409 PhysReg reg
{def
.physReg() + j
};
411 std::map
<PhysReg
,wait_entry
>::iterator it
= ctx
.gpr_map
.find(reg
);
412 if (it
== ctx
.gpr_map
.end())
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
)
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
))
428 wait
.combine(it
->second
.imm
);
435 wait_imm
parse_wait_instr(wait_ctx
& ctx
, Instruction
*instr
)
437 if (instr
->opcode
== aco_opcode::s_waitcnt_vscnt
&&
438 instr
->definitions
[0].physReg() == sgpr_null
) {
440 imm
.vs
= std::min
<uint8_t>(imm
.vs
, static_cast<SOPK_instruction
*>(instr
)->imm
);
442 } else if (instr
->opcode
== aco_opcode::s_waitcnt
) {
443 return wait_imm(ctx
.chip_class
, static_cast<SOPP_instruction
*>(instr
)->imm
);
448 wait_imm
perform_barrier(wait_ctx
& ctx
, memory_sync_info sync
, unsigned semantics
)
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
;
455 unsigned idx
= u_bit_scan(&storage
);
457 /* LDS is private to the workgroup */
458 sync_scope bar_scope_lds
= MIN2(sync
.scope
, scope_workgroup
);
460 uint16_t events
= ctx
.barrier_events
[idx
];
461 if (bar_scope_lds
<= subgroup_scope
)
462 events
&= ~event_lds
;
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
);
469 imm
.combine(ctx
.barrier_imm
[idx
]);
476 void force_waitcnt(wait_ctx
& ctx
, wait_imm
& imm
)
485 if (ctx
.chip_class
>= GFX10
) {
491 wait_imm
kill(Instruction
* instr
, wait_ctx
& ctx
, memory_sync_info sync_info
)
495 if (debug_flags
& DEBUG_FORCE_WAITCNT
) {
496 /* Force emitting waitcnt states right after the instruction if there is
497 * something to wait for.
499 force_waitcnt(ctx
, imm
);
502 if (ctx
.exp_cnt
|| ctx
.vm_cnt
|| ctx
.lgkm_cnt
)
503 imm
.combine(check_instr(instr
, ctx
));
505 imm
.combine(parse_wait_instr(ctx
, instr
));
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.
511 if (ctx
.lgkm_cnt
&& instr
->opcode
== aco_opcode::s_dcache_wb
) {
512 assert(ctx
.chip_class
>= GFX8
);
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.
521 * TODO: Refine this when we have proper alias analysis.
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()) {
531 if (instr
->opcode
== aco_opcode::p_barrier
)
532 imm
.combine(perform_barrier(ctx
, static_cast<Pseudo_barrier_instruction
*>(instr
)->sync
, semantic_acqrel
));
534 imm
.combine(perform_barrier(ctx
, sync_info
, semantic_release
));
537 if (ctx
.pending_flat_vm
&& imm
.vm
!= wait_imm::unset_counter
)
539 if (ctx
.pending_flat_lgkm
&& imm
.lgkm
!= wait_imm::unset_counter
)
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
);
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
;
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
);
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
);
564 if (bar
.vs
!= wait_imm::unset_counter
&& imm
.vs
<= bar
.vs
) {
565 bar
.vs
= wait_imm::unset_counter
;
566 bar_ev
&= ~vs_events
;
568 if (bar
.vm
== wait_imm::unset_counter
&& bar
.lgkm
== wait_imm::unset_counter
)
569 bar_ev
&= ~event_flat
;
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())
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
);
592 ctx
.pending_flat_vm
= false;
594 ctx
.pending_flat_lgkm
= false;
595 ctx
.pending_s_buffer_store
= false;
601 void update_barrier_counter(uint8_t *ctr
, unsigned max
)
603 if (*ctr
!= wait_imm::unset_counter
&& *ctr
< max
)
607 void update_barrier_imm(wait_ctx
& ctx
, uint8_t counters
, wait_event event
, memory_sync_info sync
)
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
)) {
614 if (counters
& counter_lgkm
)
616 if (counters
& counter_vm
)
618 if (counters
& counter_exp
)
620 if (counters
& counter_vs
)
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
);
635 void update_counters(wait_ctx
& ctx
, wait_event event
, memory_sync_info sync
=memory_sync_info())
637 uint8_t counters
= get_counters_for_event(event
);
639 if (counters
& counter_lgkm
&& ctx
.lgkm_cnt
<= ctx
.max_lgkm_cnt
)
641 if (counters
& counter_vm
&& ctx
.vm_cnt
<= ctx
.max_vm_cnt
)
643 if (counters
& counter_exp
&& ctx
.exp_cnt
<= ctx
.max_exp_cnt
)
645 if (counters
& counter_vs
&& ctx
.vs_cnt
<= ctx
.max_vs_cnt
)
648 update_barrier_imm(ctx
, counters
, event
, sync
);
650 if (ctx
.unordered_events
& event
)
653 if (ctx
.pending_flat_lgkm
)
654 counters
&= ~counter_lgkm
;
655 if (ctx
.pending_flat_vm
)
656 counters
&= ~counter_vm
;
658 for (std::pair
<const PhysReg
,wait_entry
>& e
: ctx
.gpr_map
) {
659 wait_entry
& entry
= e
.second
;
661 if (entry
.events
& ctx
.unordered_events
)
664 assert(entry
.events
);
666 if ((counters
& counter_exp
) && (entry
.events
& exp_events
) == event
&& entry
.imm
.exp
< ctx
.max_exp_cnt
)
668 if ((counters
& counter_lgkm
) && (entry
.events
& lgkm_events
) == event
&& entry
.imm
.lgkm
< ctx
.max_lgkm_cnt
)
670 if ((counters
& counter_vm
) && (entry
.events
& vm_events
) == event
&& entry
.imm
.vm
< ctx
.max_vm_cnt
)
672 if ((counters
& counter_vs
) && (entry
.events
& vs_events
) == event
&& entry
.imm
.vs
< ctx
.max_vs_cnt
)
677 void update_counters_for_flat_load(wait_ctx
& ctx
, memory_sync_info sync
=memory_sync_info())
679 assert(ctx
.chip_class
< GFX10
);
681 if (ctx
.lgkm_cnt
<= ctx
.max_lgkm_cnt
)
683 if (ctx
.vm_cnt
<= ctx
.max_vm_cnt
)
686 update_barrier_imm(ctx
, counter_vm
| counter_lgkm
, event_flat
, sync
);
688 for (std::pair
<PhysReg
,wait_entry
> e
: ctx
.gpr_map
)
690 if (e
.second
.counters
& counter_vm
)
692 if (e
.second
.counters
& counter_lgkm
)
693 e
.second
.imm
.lgkm
= 0;
695 ctx
.pending_flat_lgkm
= true;
696 ctx
.pending_flat_vm
= true;
699 void insert_wait_entry(wait_ctx
& ctx
, PhysReg reg
, RegClass rc
, wait_event event
, bool wait_on_read
,
700 bool has_sampler
=false)
702 uint16_t counters
= get_counters_for_event(event
);
704 if (counters
& counter_lgkm
)
706 if (counters
& counter_vm
)
708 if (counters
& counter_exp
)
710 if (counters
& counter_vs
)
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
;
717 for (unsigned i
= 0; i
< rc
.size(); i
++) {
718 auto it
= ctx
.gpr_map
.emplace(PhysReg
{reg
.reg()+i
}, new_entry
);
720 it
.first
->second
.join(new_entry
);
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
);
734 void insert_wait_entry(wait_ctx
& ctx
, Operand op
, wait_event event
, bool has_sampler
=false)
736 if (!op
.isConstant() && !op
.isUndefined())
737 insert_wait_entry(ctx
, op
.physReg(), op
.regClass(), event
, false, has_sampler
);
740 void insert_wait_entry(wait_ctx
& ctx
, Definition def
, wait_event event
, bool has_sampler
=false)
742 insert_wait_entry(ctx
, def
.physReg(), def
.regClass(), event
, true, has_sampler
);
745 void gen(Instruction
* instr
, wait_ctx
& ctx
)
747 switch (instr
->format
) {
749 Export_instruction
* exp_instr
= static_cast<Export_instruction
*>(instr
);
752 if (exp_instr
->dest
<= 9)
753 ev
= event_exp_mrt_null
;
754 else if (exp_instr
->dest
<= 15)
757 ev
= event_exp_param
;
758 update_counters(ctx
, ev
);
760 /* insert new entries for exported vgprs */
761 for (unsigned i
= 0; i
< 4; i
++)
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
);
770 insert_wait_entry(ctx
, exec
, s2
, ev
, false);
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
);
778 update_counters(ctx
, event_flat
, flat
->sync
);
780 if (!instr
->definitions
.empty())
781 insert_wait_entry(ctx
, instr
->definitions
[0], event_flat
);
785 SMEM_instruction
*smem
= static_cast<SMEM_instruction
*>(instr
);
786 update_counters(ctx
, event_smem
, smem
->sync
);
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;
797 DS_instruction
*ds
= static_cast<DS_instruction
*>(instr
);
798 update_counters(ctx
, ds
->gds
? event_gds
: event_lds
, ds
->sync
);
800 update_counters(ctx
, event_gds_gpr_lock
);
802 if (!instr
->definitions
.empty())
803 insert_wait_entry(ctx
, instr
->definitions
[0], ds
->gds
? event_gds
: event_lds
);
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);
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
));
819 bool has_sampler
= instr
->format
== Format::MIMG
&& !instr
->operands
[1].isUndefined() && instr
->operands
[1].regClass() == s4
;
821 if (!instr
->definitions
.empty())
822 insert_wait_entry(ctx
, instr
->definitions
[0], ev
, has_sampler
);
824 if (ctx
.chip_class
== GFX6
&&
825 instr
->format
!= Format::MIMG
&&
826 instr
->operands
.size() == 4) {
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
) {
834 update_counters(ctx
, event_vmem_gpr_lock
);
835 insert_wait_entry(ctx
, instr
->operands
[1], event_vmem_gpr_lock
);
841 if (instr
->opcode
== aco_opcode::s_sendmsg
||
842 instr
->opcode
== aco_opcode::s_sendmsghalt
)
843 update_counters(ctx
, event_sendmsg
);
850 void emit_waitcnt(wait_ctx
& ctx
, std::vector
<aco_ptr
<Instruction
>>& instructions
, wait_imm imm
)
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
;
861 SOPP_instruction
* waitcnt
= create_instruction
<SOPP_instruction
>(aco_opcode::s_waitcnt
, Format::SOPP
, 0, 0);
862 waitcnt
->imm
= imm
.pack(ctx
.chip_class
);
864 instructions
.emplace_back(waitcnt
);
868 void handle_block(Program
*program
, Block
& block
, wait_ctx
& ctx
)
870 std::vector
<aco_ptr
<Instruction
>> new_instructions
;
874 for (aco_ptr
<Instruction
>& instr
: block
.instructions
) {
875 bool is_wait
= !parse_wait_instr(ctx
, instr
.get()).empty();
877 memory_sync_info sync_info
= get_sync_info(instr
.get());
878 queued_imm
.combine(kill(instr
.get(), ctx
, sync_info
));
880 ctx
.gen_instr
= instr
.get();
881 gen(instr
.get(), ctx
);
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();
888 new_instructions
.emplace_back(std::move(instr
));
890 queued_imm
.combine(perform_barrier(ctx
, sync_info
, semantic_acquire
));
892 if (ctx
.collect_statistics
)
893 ctx
.advance_unwaited_instrs();
897 if (!queued_imm
.empty())
898 emit_waitcnt(ctx
, new_instructions
, queued_imm
);
900 block
.instructions
.swap(new_instructions
);
903 } /* end namespace */
905 static uint32_t calculate_score(std::vector
<wait_ctx
> &ctx_vec
, uint32_t event_mask
)
908 unsigned num_waits
= 0;
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
]) {
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
) {
924 case event_vmem_store
:
927 inv_strength
= 150.0;
935 if (score
> threshold
) {
937 score
= threshold
+ score
/ (1.0 + score
/ inv_strength
);
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);
947 return round(pow(result
/ num_waits
, 10.0) * 10.0);
950 void insert_wait_states(Program
* program
)
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
));
957 std::stack
<unsigned> loop_header_indices
;
958 unsigned loop_progress
= 0;
960 for (unsigned i
= 0; i
< program
->blocks
.size();) {
961 Block
& current
= program
->blocks
[i
++];
962 wait_ctx ctx
= in_ctx
[current
.index
];
964 if (current
.kind
& block_kind_loop_header
) {
965 loop_header_indices
.push(current
.index
);
966 } else if (current
.kind
& block_kind_loop_exit
) {
968 if (loop_progress
== loop_header_indices
.size()) {
969 i
= loop_header_indices
.top();
972 loop_header_indices
.pop();
973 loop_progress
= std::min
<unsigned>(loop_progress
, loop_header_indices
.size());
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);
984 if (done
[current
.index
] && !changed
) {
985 in_ctx
[current
.index
] = std::move(ctx
);
988 in_ctx
[current
.index
] = ctx
;
991 if (current
.instructions
.empty()) {
992 out_ctx
[current
.index
] = std::move(ctx
);
996 loop_progress
= std::max
<unsigned>(loop_progress
, current
.loop_nest_depth
);
997 done
[current
.index
] = true;
999 handle_block(program
, current
, ctx
);
1001 out_ctx
[current
.index
] = std::move(ctx
);
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
);