Edit
IABSD.fr/xenocara/lib/mesa/src/amd/compiler/aco_scheduler.cpp
jsg
- lib/mesa/src/amd/compiler/aco_scheduler.cpp
-
/* * Copyright © 2018 Valve Corporation * * SPDX-License-Identifier: MIT */ #include "aco_builder.h" #include "aco_ir.h" #include "common/amdgfxregs.h" #include <algorithm> #include <unordered_set> #include <vector> #define SMEM_WINDOW_SIZE (350 - ctx.num_waves * 35) #define VMEM_WINDOW_SIZE (1024 - ctx.num_waves * 64) #define LDS_WINDOW_SIZE 64 #define POS_EXP_WINDOW_SIZE 512 #define SMEM_MAX_MOVES (64 - ctx.num_waves * 4) #define VMEM_MAX_MOVES (256 - ctx.num_waves * 16) #define LDSDIR_MAX_MOVES 10 #define LDS_MAX_MOVES 32 /* creating clauses decreases def-use distances, so make it less aggressive the lower num_waves is */ #define VMEM_CLAUSE_MAX_GRAB_DIST (ctx.num_waves * 2) #define VMEM_STORE_CLAUSE_MAX_GRAB_DIST (ctx.num_waves * 4) #define POS_EXP_MAX_MOVES 512 namespace aco { namespace { enum MoveResult { move_success, move_fail_ssa, move_fail_rar, move_fail_pressure, }; /** * Cursor for downwards moves, where a single instruction is moved towards * or below a group of instruction that hardware can execute as a clause. */ struct DownwardsCursor { int source_idx; /* Current instruction to consider for moving */ int insert_idx_clause; /* First clause instruction */ int insert_idx; /* First instruction *after* the clause */ /* Maximum demand of all clause instructions, * i.e. from insert_idx_clause (inclusive) to insert_idx (exclusive) */ RegisterDemand clause_demand; /* Maximum demand of instructions from source_idx to insert_idx_clause (both exclusive) */ RegisterDemand total_demand; DownwardsCursor(int current_idx, RegisterDemand initial_clause_demand) : source_idx(current_idx - 1), insert_idx_clause(current_idx), insert_idx(current_idx + 1), clause_demand(initial_clause_demand) {} void verify_invariants(const Block* block); }; /** * Cursor for upwards moves, where a single instruction is moved below * another instruction. */ struct UpwardsCursor { int source_idx; /* Current instruction to consider for moving */ int insert_idx; /* Instruction to move in front of */ /* Maximum demand of instructions from insert_idx (inclusive) to source_idx (exclusive) */ RegisterDemand total_demand; UpwardsCursor(int source_idx_) : source_idx(source_idx_) { insert_idx = -1; /* to be initialized later */ } bool has_insert_idx() const { return insert_idx != -1; } void verify_invariants(const Block* block); }; struct MoveState { RegisterDemand max_registers; Block* block; Instruction* current; bool improved_rar; std::vector<bool> depends_on; /* Two are needed because, for downwards VMEM scheduling, one needs to * exclude the instructions in the clause, since new instructions in the * clause are not moved past any other instructions in the clause. */ std::vector<bool> RAR_dependencies; std::vector<bool> RAR_dependencies_clause; /* for moving instructions before the current instruction to after it */ DownwardsCursor downwards_init(int current_idx, bool improved_rar, bool may_form_clauses); MoveResult downwards_move(DownwardsCursor&, bool clause); void downwards_skip(DownwardsCursor&); /* for moving instructions after the first use of the current instruction upwards */ UpwardsCursor upwards_init(int source_idx, bool improved_rar); bool upwards_check_deps(UpwardsCursor&); void upwards_update_insert_idx(UpwardsCursor&); MoveResult upwards_move(UpwardsCursor&); void upwards_skip(UpwardsCursor&); }; struct sched_ctx { amd_gfx_level gfx_level; int16_t num_waves; int16_t last_SMEM_stall; int last_SMEM_dep_idx; MoveState mv; bool schedule_pos_exports = true; unsigned schedule_pos_export_div = 1; }; /* This scheduler is a simple bottom-up pass based on ideas from * "A Novel Lightweight Instruction Scheduling Algorithm for Just-In-Time Compiler" * from Xiaohua Shi and Peng Guo. * The basic approach is to iterate over all instructions. When a memory instruction * is encountered it tries to move independent instructions from above and below * between the memory instruction and it's first user. * The novelty is that this scheduler cares for the current register pressure: * Instructions will only be moved if the register pressure won't exceed a certain bound. */ template <typename T> void move_element(T begin_it, size_t idx, size_t before) { if (idx < before) { auto begin = std::next(begin_it, idx); auto end = std::next(begin_it, before); std::rotate(begin, begin + 1, end); } else if (idx > before) { auto begin = std::next(begin_it, before); auto end = std::next(begin_it, idx + 1); std::rotate(begin, end - 1, end); } } void DownwardsCursor::verify_invariants(const Block* block) { assert(source_idx < insert_idx_clause); assert(insert_idx_clause < insert_idx); #ifndef NDEBUG RegisterDemand reference_demand; for (int i = source_idx + 1; i < insert_idx_clause; ++i) { reference_demand.update(block->instructions[i]->register_demand); } assert(total_demand == reference_demand); reference_demand = {}; for (int i = insert_idx_clause; i < insert_idx; ++i) { reference_demand.update(block->instructions[i]->register_demand); } assert(clause_demand == reference_demand); #endif } DownwardsCursor MoveState::downwards_init(int current_idx, bool improved_rar_, bool may_form_clauses) { improved_rar = improved_rar_; std::fill(depends_on.begin(), depends_on.end(), false); if (improved_rar) { std::fill(RAR_dependencies.begin(), RAR_dependencies.end(), false); if (may_form_clauses) std::fill(RAR_dependencies_clause.begin(), RAR_dependencies_clause.end(), false); } for (const Operand& op : current->operands) { if (op.isTemp()) { depends_on[op.tempId()] = true; if (improved_rar && op.isFirstKill()) RAR_dependencies[op.tempId()] = true; } } DownwardsCursor cursor(current_idx, block->instructions[current_idx]->register_demand); cursor.verify_invariants(block); return cursor; } /* If add_to_clause is true, the current clause is extended by moving the * instruction at source_idx in front of the clause. Otherwise, the instruction * is moved past the end of the clause without extending it */ MoveResult MoveState::downwards_move(DownwardsCursor& cursor, bool add_to_clause) { aco_ptr<Instruction>& instr = block->instructions[cursor.source_idx]; for (const Definition& def : instr->definitions) if (def.isTemp() && depends_on[def.tempId()]) return move_fail_ssa; /* check if one of candidate's operands is killed by depending instruction */ std::vector<bool>& RAR_deps = improved_rar ? (add_to_clause ? RAR_dependencies_clause : RAR_dependencies) : depends_on; for (const Operand& op : instr->operands) { if (op.isTemp() && RAR_deps[op.tempId()]) { // FIXME: account for difference in register pressure return move_fail_rar; } } if (add_to_clause) { for (const Operand& op : instr->operands) { if (op.isTemp()) { depends_on[op.tempId()] = true; if (op.isFirstKill()) RAR_dependencies[op.tempId()] = true; } } } const int dest_insert_idx = add_to_clause ? cursor.insert_idx_clause : cursor.insert_idx; RegisterDemand register_pressure = cursor.total_demand; if (!add_to_clause) { register_pressure.update(cursor.clause_demand); } /* Check the new demand of the instructions being moved over */ const RegisterDemand candidate_diff = get_live_changes(instr.get()); if (RegisterDemand(register_pressure - candidate_diff).exceeds(max_registers)) return move_fail_pressure; /* New demand for the moved instruction */ const RegisterDemand temp = get_temp_registers(instr.get()); const RegisterDemand temp2 = get_temp_registers(block->instructions[dest_insert_idx - 1].get()); const RegisterDemand new_demand = block->instructions[dest_insert_idx - 1]->register_demand - temp2 + temp; if (new_demand.exceeds(max_registers)) return move_fail_pressure; /* move the candidate below the memory load */ move_element(block->instructions.begin(), cursor.source_idx, dest_insert_idx); /* update register pressure */ for (int i = cursor.source_idx; i < dest_insert_idx - 1; i++) block->instructions[i]->register_demand -= candidate_diff; block->instructions[dest_insert_idx - 1]->register_demand = new_demand; cursor.insert_idx_clause--; if (cursor.source_idx != cursor.insert_idx_clause) { /* Update demand if we moved over any instructions before the clause */ cursor.total_demand -= candidate_diff; } else { assert(cursor.total_demand == RegisterDemand{}); } if (add_to_clause) { cursor.clause_demand.update(new_demand); } else { cursor.clause_demand -= candidate_diff; cursor.insert_idx--; } cursor.source_idx--; cursor.verify_invariants(block); return move_success; } void MoveState::downwards_skip(DownwardsCursor& cursor) { aco_ptr<Instruction>& instr = block->instructions[cursor.source_idx]; for (const Operand& op : instr->operands) { if (op.isTemp()) { depends_on[op.tempId()] = true; if (improved_rar && op.isFirstKill()) { RAR_dependencies[op.tempId()] = true; RAR_dependencies_clause[op.tempId()] = true; } } } cursor.total_demand.update(instr->register_demand); cursor.source_idx--; cursor.verify_invariants(block); } void UpwardsCursor::verify_invariants(const Block* block) { #ifndef NDEBUG if (!has_insert_idx()) { return; } assert(insert_idx < source_idx); RegisterDemand reference_demand; for (int i = insert_idx; i < source_idx; ++i) { reference_demand.update(block->instructions[i]->register_demand); } assert(total_demand == reference_demand); #endif } UpwardsCursor MoveState::upwards_init(int source_idx, bool improved_rar_) { improved_rar = improved_rar_; std::fill(depends_on.begin(), depends_on.end(), false); std::fill(RAR_dependencies.begin(), RAR_dependencies.end(), false); for (const Definition& def : current->definitions) { if (def.isTemp()) depends_on[def.tempId()] = true; } return UpwardsCursor(source_idx); } bool MoveState::upwards_check_deps(UpwardsCursor& cursor) { aco_ptr<Instruction>& instr = block->instructions[cursor.source_idx]; for (const Operand& op : instr->operands) { if (op.isTemp() && depends_on[op.tempId()]) return false; } return true; } void MoveState::upwards_update_insert_idx(UpwardsCursor& cursor) { cursor.insert_idx = cursor.source_idx; cursor.total_demand = block->instructions[cursor.insert_idx]->register_demand; } MoveResult MoveState::upwards_move(UpwardsCursor& cursor) { assert(cursor.has_insert_idx()); aco_ptr<Instruction>& instr = block->instructions[cursor.source_idx]; for (const Operand& op : instr->operands) { if (op.isTemp() && depends_on[op.tempId()]) return move_fail_ssa; } /* check if candidate uses/kills an operand which is used by a dependency */ for (const Operand& op : instr->operands) { if (op.isTemp() && (!improved_rar || op.isFirstKill()) && RAR_dependencies[op.tempId()]) return move_fail_rar; } /* check if register pressure is low enough: the diff is negative if register pressure is * decreased */ const RegisterDemand candidate_diff = get_live_changes(instr.get()); const RegisterDemand temp = get_temp_registers(instr.get()); if (RegisterDemand(cursor.total_demand + candidate_diff).exceeds(max_registers)) return move_fail_pressure; const RegisterDemand temp2 = get_temp_registers(block->instructions[cursor.insert_idx - 1].get()); const RegisterDemand new_demand = block->instructions[cursor.insert_idx - 1]->register_demand - temp2 + candidate_diff + temp; if (new_demand.exceeds(max_registers)) return move_fail_pressure; /* move the candidate above the insert_idx */ move_element(block->instructions.begin(), cursor.source_idx, cursor.insert_idx); /* update register pressure */ block->instructions[cursor.insert_idx]->register_demand = new_demand; for (int i = cursor.insert_idx + 1; i <= cursor.source_idx; i++) block->instructions[i]->register_demand += candidate_diff; cursor.total_demand += candidate_diff; cursor.total_demand.update(block->instructions[cursor.source_idx]->register_demand); cursor.insert_idx++; cursor.source_idx++; cursor.verify_invariants(block); return move_success; } void MoveState::upwards_skip(UpwardsCursor& cursor) { if (cursor.has_insert_idx()) { aco_ptr<Instruction>& instr = block->instructions[cursor.source_idx]; for (const Definition& def : instr->definitions) { if (def.isTemp()) depends_on[def.tempId()] = true; } for (const Operand& op : instr->operands) { if (op.isTemp()) RAR_dependencies[op.tempId()] = true; } cursor.total_demand.update(instr->register_demand); } cursor.source_idx++; cursor.verify_invariants(block); } bool is_done_sendmsg(amd_gfx_level gfx_level, const Instruction* instr) { if (gfx_level <= GFX10_3 && instr->opcode == aco_opcode::s_sendmsg) return (instr->salu().imm & sendmsg_id_mask) == sendmsg_gs_done; return false; } bool is_pos_prim_export(amd_gfx_level gfx_level, const Instruction* instr) { /* Because of NO_PC_EXPORT=1, a done=1 position or primitive export can launch PS waves before * the NGG/VS wave finishes if there are no parameter exports. */ return instr->opcode == aco_opcode::exp && instr->exp().dest >= V_008DFC_SQ_EXP_POS && instr->exp().dest <= V_008DFC_SQ_EXP_PRIM && gfx_level >= GFX10; } memory_sync_info get_sync_info_with_hack(const Instruction* instr) { memory_sync_info sync = get_sync_info(instr); if (instr->isSMEM() && !instr->operands.empty() && instr->operands[0].bytes() == 16) { // FIXME: currently, it doesn't seem beneficial to omit this due to how our scheduler works sync.storage = (storage_class)(sync.storage | storage_buffer); sync.semantics = (memory_semantics)((sync.semantics | semantic_private) & ~semantic_can_reorder); } return sync; } struct memory_event_set { bool has_control_barrier; unsigned bar_acquire; unsigned bar_release; unsigned bar_classes; unsigned access_acquire; unsigned access_release; unsigned access_relaxed; unsigned access_atomic; }; struct hazard_query { amd_gfx_level gfx_level; bool contains_spill; bool contains_sendmsg; bool uses_exec; bool writes_exec; memory_event_set mem_events; unsigned aliasing_storage; /* storage classes which are accessed (non-SMEM) */ unsigned aliasing_storage_smem; /* storage classes which are accessed (SMEM) */ }; void init_hazard_query(const sched_ctx& ctx, hazard_query* query) { query->gfx_level = ctx.gfx_level; query->contains_spill = false; query->contains_sendmsg = false; query->uses_exec = false; query->writes_exec = false; memset(&query->mem_events, 0, sizeof(query->mem_events)); query->aliasing_storage = 0; query->aliasing_storage_smem = 0; } void add_memory_event(amd_gfx_level gfx_level, memory_event_set* set, Instruction* instr, memory_sync_info* sync) { set->has_control_barrier |= is_done_sendmsg(gfx_level, instr); set->has_control_barrier |= is_pos_prim_export(gfx_level, instr); if (instr->opcode == aco_opcode::p_barrier) { Pseudo_barrier_instruction& bar = instr->barrier(); if (bar.sync.semantics & semantic_acquire) set->bar_acquire |= bar.sync.storage; if (bar.sync.semantics & semantic_release) set->bar_release |= bar.sync.storage; set->bar_classes |= bar.sync.storage; set->has_control_barrier |= bar.exec_scope > scope_invocation; } if (!sync->storage) return; if (sync->semantics & semantic_acquire) set->access_acquire |= sync->storage; if (sync->semantics & semantic_release) set->access_release |= sync->storage; if (!(sync->semantics & semantic_private)) { if (sync->semantics & semantic_atomic) set->access_atomic |= sync->storage; else set->access_relaxed |= sync->storage; } } void add_to_hazard_query(hazard_query* query, Instruction* instr) { if (instr->opcode == aco_opcode::p_spill || instr->opcode == aco_opcode::p_reload) query->contains_spill = true; query->contains_sendmsg |= instr->opcode == aco_opcode::s_sendmsg; query->uses_exec |= needs_exec_mask(instr); for (const Definition& def : instr->definitions) { if (def.isFixed() && def.physReg() == exec) query->writes_exec = true; } memory_sync_info sync = get_sync_info_with_hack(instr); add_memory_event(query->gfx_level, &query->mem_events, instr, &sync); if (!(sync.semantics & semantic_can_reorder)) { unsigned storage = sync.storage; /* images and buffer/global memory can alias */ // TODO: more precisely, buffer images and // buffer/global memory can alias if (storage & (storage_buffer | storage_image)) storage |= storage_buffer | storage_image; if (instr->isSMEM()) query->aliasing_storage_smem |= storage; else query->aliasing_storage |= storage; } } enum HazardResult { hazard_success, hazard_fail_reorder_vmem_smem, hazard_fail_reorder_ds, hazard_fail_reorder_sendmsg, hazard_fail_spill, hazard_fail_export, hazard_fail_barrier, /* Must stop at these failures. The hazard query code doesn't consider them * when added. */ hazard_fail_exec, hazard_fail_unreorderable, }; HazardResult perform_hazard_query(hazard_query* query, Instruction* instr, bool upwards) { /* don't schedule discards downwards */ if (!upwards && instr->opcode == aco_opcode::p_exit_early_if_not) return hazard_fail_unreorderable; /* In Primitive Ordered Pixel Shading, await overlapped waves as late as possible, and notify * overlapping waves that they can continue execution as early as possible. */ if (upwards) { if (instr->opcode == aco_opcode::p_pops_gfx9_add_exiting_wave_id || is_wait_export_ready(query->gfx_level, instr)) { return hazard_fail_unreorderable; } } else { if (instr->opcode == aco_opcode::p_pops_gfx9_ordered_section_done) { return hazard_fail_unreorderable; } } if (query->uses_exec || query->writes_exec) { for (const Definition& def : instr->definitions) { if (def.isFixed() && def.physReg() == exec) return hazard_fail_exec; } } if (query->writes_exec && needs_exec_mask(instr)) return hazard_fail_exec; /* Don't move exports so that they stay closer together. * Since GFX11, export order matters. MRTZ must come first, * then color exports sorted from first to last. * Also, with Primitive Ordered Pixel Shading on GFX11+, the `done` export must not be moved * above the memory accesses before the queue family scope (more precisely, fragment interlock * scope, but it's not available in ACO) release barrier that is expected to be inserted before * the export, as well as before any `s_wait_event export_ready` which enters the ordered * section, because the `done` export exits the ordered section. */ if (instr->isEXP() || instr->opcode == aco_opcode::p_dual_src_export_gfx11) return hazard_fail_export; /* don't move non-reorderable instructions */ if (instr->opcode == aco_opcode::s_memtime || instr->opcode == aco_opcode::s_memrealtime || instr->opcode == aco_opcode::s_setprio || instr->opcode == aco_opcode::s_getreg_b32 || instr->opcode == aco_opcode::p_shader_cycles_hi_lo_hi || instr->opcode == aco_opcode::p_init_scratch || instr->opcode == aco_opcode::p_jump_to_epilog || instr->opcode == aco_opcode::s_sendmsg_rtn_b32 || instr->opcode == aco_opcode::s_sendmsg_rtn_b64 || instr->opcode == aco_opcode::p_end_with_regs || instr->opcode == aco_opcode::s_nop || instr->opcode == aco_opcode::s_sleep || instr->opcode == aco_opcode::s_trap) return hazard_fail_unreorderable; memory_event_set instr_set; memset(&instr_set, 0, sizeof(instr_set)); memory_sync_info sync = get_sync_info_with_hack(instr); add_memory_event(query->gfx_level, &instr_set, instr, &sync); memory_event_set* first = &instr_set; memory_event_set* second = &query->mem_events; if (upwards) std::swap(first, second); /* everything after barrier(acquire) happens after the atomics/control_barriers before * everything after load(acquire) happens after the load */ if ((first->has_control_barrier || first->access_atomic) && second->bar_acquire) return hazard_fail_barrier; if (((first->access_acquire || first->bar_acquire) && second->bar_classes) || ((first->access_acquire | first->bar_acquire) & (second->access_relaxed | second->access_atomic))) return hazard_fail_barrier; /* everything before barrier(release) happens before the atomics/control_barriers after * * everything before store(release) happens before the store */ if (first->bar_release && (second->has_control_barrier || second->access_atomic)) return hazard_fail_barrier; if ((first->bar_classes && (second->bar_release || second->access_release)) || ((first->access_relaxed | first->access_atomic) & (second->bar_release | second->access_release))) return hazard_fail_barrier; /* don't move memory barriers around other memory barriers */ if (first->bar_classes && second->bar_classes) return hazard_fail_barrier; /* Don't move memory accesses to before control barriers. I don't think * this is necessary for the Vulkan memory model, but it might be for GLSL450. */ unsigned control_classes = storage_buffer | storage_image | storage_shared | storage_task_payload; if (first->has_control_barrier && ((second->access_atomic | second->access_relaxed) & control_classes)) return hazard_fail_barrier; /* don't move memory loads/stores past potentially aliasing loads/stores */ unsigned aliasing_storage = instr->isSMEM() ? query->aliasing_storage_smem : query->aliasing_storage; if ((sync.storage & aliasing_storage) && !(sync.semantics & semantic_can_reorder)) { unsigned intersect = sync.storage & aliasing_storage; if (intersect & storage_shared) return hazard_fail_reorder_ds; return hazard_fail_reorder_vmem_smem; } if ((instr->opcode == aco_opcode::p_spill || instr->opcode == aco_opcode::p_reload) && query->contains_spill) return hazard_fail_spill; if (instr->opcode == aco_opcode::s_sendmsg && query->contains_sendmsg) return hazard_fail_reorder_sendmsg; return hazard_success; } unsigned get_likely_cost(Instruction* instr) { if (instr->opcode == aco_opcode::p_split_vector || instr->opcode == aco_opcode::p_extract_vector) { unsigned cost = 0; for (Definition def : instr->definitions) { if (instr->operands[0].isKill() && def.regClass().type() == instr->operands[0].regClass().type()) continue; cost += def.size(); } return cost; } else if (instr->opcode == aco_opcode::p_create_vector) { unsigned cost = 0; for (Operand op : instr->operands) { if (op.isTemp() && op.isFirstKill() && op.regClass().type() == instr->definitions[0].regClass().type()) continue; cost += op.size(); } return cost; } else { /* For the moment, just assume the same cost for all other instructions. */ return 1; } } void schedule_SMEM(sched_ctx& ctx, Block* block, Instruction* current, int idx) { assert(idx != 0); int window_size = SMEM_WINDOW_SIZE; int max_moves = SMEM_MAX_MOVES; int16_t k = 0; /* don't move s_memtime/s_memrealtime */ if (current->opcode == aco_opcode::s_memtime || current->opcode == aco_opcode::s_memrealtime || current->opcode == aco_opcode::s_sendmsg_rtn_b32 || current->opcode == aco_opcode::s_sendmsg_rtn_b64) return; /* first, check if we have instructions before current to move down */ hazard_query hq; init_hazard_query(ctx, &hq); add_to_hazard_query(&hq, current); DownwardsCursor cursor = ctx.mv.downwards_init(idx, false, false); for (int candidate_idx = idx - 1; k < max_moves && candidate_idx > (int)idx - window_size; candidate_idx--) { assert(candidate_idx >= 0); assert(candidate_idx == cursor.source_idx); aco_ptr<Instruction>& candidate = block->instructions[candidate_idx]; /* break if we'd make the previous SMEM instruction stall */ bool can_stall_prev_smem = idx <= ctx.last_SMEM_dep_idx && candidate_idx < ctx.last_SMEM_dep_idx; if (can_stall_prev_smem && ctx.last_SMEM_stall >= 0) break; /* break when encountering another MEM instruction, logical_start or barriers */ if (candidate->opcode == aco_opcode::p_logical_start) break; /* only move VMEM instructions below descriptor loads. be more aggressive at higher num_waves * to help create more vmem clauses */ if ((candidate->isVMEM() || candidate->isFlatLike()) && (cursor.insert_idx - cursor.source_idx > (ctx.num_waves * 4) || current->operands[0].size() == 4)) break; /* don't move descriptor loads below buffer loads */ if (candidate->isSMEM() && !candidate->operands.empty() && current->operands[0].size() == 4 && candidate->operands[0].size() == 2) break; bool can_move_down = true; HazardResult haz = perform_hazard_query(&hq, candidate.get(), false); if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill || haz == hazard_fail_reorder_sendmsg || haz == hazard_fail_barrier || haz == hazard_fail_export) can_move_down = false; else if (haz != hazard_success) break; /* don't use LDS/GDS instructions to hide latency since it can * significantly worsen LDS scheduling */ if (candidate->isDS() || !can_move_down) { add_to_hazard_query(&hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } MoveResult res = ctx.mv.downwards_move(cursor, false); if (res == move_fail_ssa || res == move_fail_rar) { add_to_hazard_query(&hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } else if (res == move_fail_pressure) { break; } if (candidate_idx < ctx.last_SMEM_dep_idx) ctx.last_SMEM_stall++; k++; } /* find the first instruction depending on current or find another MEM */ UpwardsCursor up_cursor = ctx.mv.upwards_init(idx + 1, false); bool found_dependency = false; /* second, check if we have instructions after current to move up */ for (int candidate_idx = idx + 1; k < max_moves && candidate_idx < (int)idx + window_size; candidate_idx++) { assert(candidate_idx == up_cursor.source_idx); assert(candidate_idx < (int)block->instructions.size()); aco_ptr<Instruction>& candidate = block->instructions[candidate_idx]; if (candidate->opcode == aco_opcode::p_logical_end) break; /* check if candidate depends on current */ bool is_dependency = !found_dependency && !ctx.mv.upwards_check_deps(up_cursor); /* no need to steal from following VMEM instructions */ if (is_dependency && (candidate->isVMEM() || candidate->isFlatLike())) break; if (found_dependency) { HazardResult haz = perform_hazard_query(&hq, candidate.get(), true); if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill || haz == hazard_fail_reorder_sendmsg || haz == hazard_fail_barrier || haz == hazard_fail_export) is_dependency = true; else if (haz != hazard_success) break; } if (is_dependency) { if (!found_dependency) { ctx.mv.upwards_update_insert_idx(up_cursor); init_hazard_query(ctx, &hq); found_dependency = true; } } if (is_dependency || !found_dependency) { if (found_dependency) add_to_hazard_query(&hq, candidate.get()); else k++; ctx.mv.upwards_skip(up_cursor); continue; } MoveResult res = ctx.mv.upwards_move(up_cursor); if (res == move_fail_ssa || res == move_fail_rar) { /* no need to steal from following VMEM instructions */ if (res == move_fail_ssa && (candidate->isVMEM() || candidate->isFlatLike())) break; add_to_hazard_query(&hq, candidate.get()); ctx.mv.upwards_skip(up_cursor); continue; } else if (res == move_fail_pressure) { break; } k++; } ctx.last_SMEM_dep_idx = found_dependency ? up_cursor.insert_idx : 0; ctx.last_SMEM_stall = 10 - ctx.num_waves - k; } void schedule_VMEM(sched_ctx& ctx, Block* block, Instruction* current, int idx) { assert(idx != 0); int window_size = VMEM_WINDOW_SIZE; int max_moves = VMEM_MAX_MOVES; int clause_max_grab_dist = VMEM_CLAUSE_MAX_GRAB_DIST; bool only_clauses = false; int16_t k = 0; /* first, check if we have instructions before current to move down */ hazard_query indep_hq; hazard_query clause_hq; init_hazard_query(ctx, &indep_hq); init_hazard_query(ctx, &clause_hq); add_to_hazard_query(&indep_hq, current); DownwardsCursor cursor = ctx.mv.downwards_init(idx, true, true); for (int candidate_idx = idx - 1; k < max_moves && candidate_idx > (int)idx - window_size; candidate_idx--) { assert(candidate_idx == cursor.source_idx); assert(candidate_idx >= 0); aco_ptr<Instruction>& candidate = block->instructions[candidate_idx]; bool is_vmem = candidate->isVMEM() || candidate->isFlatLike(); /* break when encountering another VMEM instruction, logical_start or barriers */ if (candidate->opcode == aco_opcode::p_logical_start) break; /* break if we'd make the previous SMEM instruction stall */ bool can_stall_prev_smem = idx <= ctx.last_SMEM_dep_idx && candidate_idx < ctx.last_SMEM_dep_idx; if (can_stall_prev_smem && ctx.last_SMEM_stall >= 0) break; bool part_of_clause = false; if (current->isVMEM() == candidate->isVMEM()) { int grab_dist = cursor.insert_idx_clause - candidate_idx; /* We can't easily tell how much this will decrease the def-to-use * distances, so just use how far it will be moved as a heuristic. */ part_of_clause = grab_dist < clause_max_grab_dist + k && should_form_clause(current, candidate.get()); } /* if current depends on candidate, add additional dependencies and continue */ bool can_move_down = !is_vmem || part_of_clause || candidate->definitions.empty(); if (only_clauses) { /* In case of high register pressure, only try to form clauses, * and only if the previous clause is not larger * than the current one will be. */ if (part_of_clause) { int clause_size = cursor.insert_idx - cursor.insert_idx_clause; int prev_clause_size = 1; while (should_form_clause(current, block->instructions[candidate_idx - prev_clause_size].get())) prev_clause_size++; if (prev_clause_size > clause_size + 1) break; } else { can_move_down = false; } } HazardResult haz = perform_hazard_query(part_of_clause ? &clause_hq : &indep_hq, candidate.get(), false); if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill || haz == hazard_fail_reorder_sendmsg || haz == hazard_fail_barrier || haz == hazard_fail_export) can_move_down = false; else if (haz != hazard_success) break; if (!can_move_down) { if (part_of_clause) break; add_to_hazard_query(&indep_hq, candidate.get()); add_to_hazard_query(&clause_hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } Instruction* candidate_ptr = candidate.get(); MoveResult res = ctx.mv.downwards_move(cursor, part_of_clause); if (res == move_fail_ssa || res == move_fail_rar) { if (part_of_clause) break; add_to_hazard_query(&indep_hq, candidate.get()); add_to_hazard_query(&clause_hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } else if (res == move_fail_pressure) { only_clauses = true; if (part_of_clause) break; add_to_hazard_query(&indep_hq, candidate.get()); add_to_hazard_query(&clause_hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } if (part_of_clause) add_to_hazard_query(&indep_hq, candidate_ptr); else k++; if (candidate_idx < ctx.last_SMEM_dep_idx) ctx.last_SMEM_stall++; } /* find the first instruction depending on current or find another VMEM */ UpwardsCursor up_cursor = ctx.mv.upwards_init(idx + 1, true); bool found_dependency = false; /* second, check if we have instructions after current to move up */ for (int candidate_idx = idx + 1; k < max_moves && candidate_idx < (int)idx + window_size; candidate_idx++) { assert(candidate_idx == up_cursor.source_idx); assert(candidate_idx < (int)block->instructions.size()); aco_ptr<Instruction>& candidate = block->instructions[candidate_idx]; bool is_vmem = candidate->isVMEM() || candidate->isFlatLike(); if (candidate->opcode == aco_opcode::p_logical_end) break; /* check if candidate depends on current */ bool is_dependency = false; if (found_dependency) { HazardResult haz = perform_hazard_query(&indep_hq, candidate.get(), true); if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill || haz == hazard_fail_reorder_vmem_smem || haz == hazard_fail_reorder_sendmsg || haz == hazard_fail_barrier || haz == hazard_fail_export) is_dependency = true; else if (haz != hazard_success) break; } is_dependency |= !found_dependency && !ctx.mv.upwards_check_deps(up_cursor); if (is_dependency) { if (!found_dependency) { ctx.mv.upwards_update_insert_idx(up_cursor); init_hazard_query(ctx, &indep_hq); found_dependency = true; } } else if (is_vmem) { /* don't move up dependencies of other VMEM instructions */ for (const Definition& def : candidate->definitions) { if (def.isTemp()) ctx.mv.depends_on[def.tempId()] = true; } } if (is_dependency || !found_dependency) { if (found_dependency) add_to_hazard_query(&indep_hq, candidate.get()); else k++; ctx.mv.upwards_skip(up_cursor); continue; } MoveResult res = ctx.mv.upwards_move(up_cursor); if (res == move_fail_ssa || res == move_fail_rar) { add_to_hazard_query(&indep_hq, candidate.get()); ctx.mv.upwards_skip(up_cursor); continue; } else if (res == move_fail_pressure) { break; } k++; } } void schedule_LDS(sched_ctx& ctx, Block* block, Instruction* current, int idx) { assert(idx != 0); int window_size = LDS_WINDOW_SIZE; int max_moves = current->isLDSDIR() ? LDSDIR_MAX_MOVES : LDS_MAX_MOVES; int16_t k = 0; /* first, check if we have instructions before current to move down */ hazard_query hq; init_hazard_query(ctx, &hq); add_to_hazard_query(&hq, current); DownwardsCursor cursor = ctx.mv.downwards_init(idx, true, false); for (int i = 0; k < max_moves && i < window_size; i++) { aco_ptr<Instruction>& candidate = block->instructions[cursor.source_idx]; bool is_mem = candidate->isVMEM() || candidate->isFlatLike() || candidate->isSMEM(); if (candidate->opcode == aco_opcode::p_logical_start || is_mem) break; if (candidate->isDS() || candidate->isLDSDIR()) { add_to_hazard_query(&hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } if (perform_hazard_query(&hq, candidate.get(), false) != hazard_success || ctx.mv.downwards_move(cursor, false) != move_success) break; k++; } /* second, check if we have instructions after current to move up */ bool found_dependency = false; int i = 0; UpwardsCursor up_cursor = ctx.mv.upwards_init(idx + 1, true); /* find the first instruction depending on current */ for (; k < max_moves && i < window_size; i++) { aco_ptr<Instruction>& candidate = block->instructions[up_cursor.source_idx]; bool is_mem = candidate->isVMEM() || candidate->isFlatLike() || candidate->isSMEM(); if (candidate->opcode == aco_opcode::p_logical_end || is_mem) break; /* check if candidate depends on current */ if (!ctx.mv.upwards_check_deps(up_cursor)) { init_hazard_query(ctx, &hq); add_to_hazard_query(&hq, candidate.get()); ctx.mv.upwards_update_insert_idx(up_cursor); ctx.mv.upwards_skip(up_cursor); found_dependency = true; i++; break; } ctx.mv.upwards_skip(up_cursor); } for (; found_dependency && k < max_moves && i < window_size; i++) { aco_ptr<Instruction>& candidate = block->instructions[up_cursor.source_idx]; bool is_mem = candidate->isVMEM() || candidate->isFlatLike() || candidate->isSMEM(); if (candidate->opcode == aco_opcode::p_logical_end || is_mem) break; HazardResult haz = perform_hazard_query(&hq, candidate.get(), true); if (haz == hazard_fail_exec || haz == hazard_fail_unreorderable) break; if (haz != hazard_success || ctx.mv.upwards_move(up_cursor) != move_success) { add_to_hazard_query(&hq, candidate.get()); ctx.mv.upwards_skip(up_cursor); } else { k++; } } } void schedule_position_export(sched_ctx& ctx, Block* block, Instruction* current, int idx) { assert(idx != 0); int window_size = POS_EXP_WINDOW_SIZE / ctx.schedule_pos_export_div; int max_moves = POS_EXP_MAX_MOVES / ctx.schedule_pos_export_div; int16_t k = 0; DownwardsCursor cursor = ctx.mv.downwards_init(idx, true, false); hazard_query hq; init_hazard_query(ctx, &hq); add_to_hazard_query(&hq, current); for (int candidate_idx = idx - 1; k < max_moves && candidate_idx > (int)idx - window_size; candidate_idx--) { assert(candidate_idx >= 0); aco_ptr<Instruction>& candidate = block->instructions[candidate_idx]; if (candidate->opcode == aco_opcode::p_logical_start) break; if (candidate->isVMEM() || candidate->isSMEM() || candidate->isFlatLike()) break; HazardResult haz = perform_hazard_query(&hq, candidate.get(), false); if (haz == hazard_fail_exec || haz == hazard_fail_unreorderable) break; if (haz != hazard_success) { add_to_hazard_query(&hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } MoveResult res = ctx.mv.downwards_move(cursor, false); if (res == move_fail_ssa || res == move_fail_rar) { add_to_hazard_query(&hq, candidate.get()); ctx.mv.downwards_skip(cursor); continue; } else if (res == move_fail_pressure) { break; } k++; } } unsigned schedule_VMEM_store(sched_ctx& ctx, Block* block, Instruction* current, int idx) { hazard_query hq; init_hazard_query(ctx, &hq); DownwardsCursor cursor = ctx.mv.downwards_init(idx, true, true); int skip = 0; for (int16_t k = 0; k < VMEM_STORE_CLAUSE_MAX_GRAB_DIST;) { aco_ptr<Instruction>& candidate = block->instructions[cursor.source_idx]; if (candidate->opcode == aco_opcode::p_logical_start) break; if (!should_form_clause(current, candidate.get())) { add_to_hazard_query(&hq, candidate.get()); ctx.mv.downwards_skip(cursor); k += get_likely_cost(candidate.get()); continue; } if (perform_hazard_query(&hq, candidate.get(), false) != hazard_success || ctx.mv.downwards_move(cursor, true) != move_success) break; skip++; } return skip; } void schedule_block(sched_ctx& ctx, Program* program, Block* block) { ctx.last_SMEM_dep_idx = 0; ctx.last_SMEM_stall = INT16_MIN; ctx.mv.block = block; /* go through all instructions and find memory loads */ unsigned num_stores = 0; for (unsigned idx = 0; idx < block->instructions.size(); idx++) { Instruction* current = block->instructions[idx].get(); if (current->opcode == aco_opcode::p_logical_end) break; if (block->kind & block_kind_export_end && current->isEXP() && ctx.schedule_pos_exports) { unsigned target = current->exp().dest; if (target >= V_008DFC_SQ_EXP_POS && target < V_008DFC_SQ_EXP_PRIM) { ctx.mv.current = current; schedule_position_export(ctx, block, current, idx); } } if (current->definitions.empty()) { num_stores += current->isVMEM() || current->isFlatLike() ? 1 : 0; continue; } if (current->isVMEM() || current->isFlatLike()) { ctx.mv.current = current; schedule_VMEM(ctx, block, current, idx); } if (current->isSMEM()) { ctx.mv.current = current; schedule_SMEM(ctx, block, current, idx); } if (current->isLDSDIR() || (current->isDS() && !current->ds().gds)) { ctx.mv.current = current; schedule_LDS(ctx, block, current, idx); } } /* GFX11 benefits from creating VMEM store clauses. */ if (num_stores > 1 && program->gfx_level >= GFX11) { for (int idx = block->instructions.size() - 1; idx >= 0; idx--) { Instruction* current = block->instructions[idx].get(); if (!current->definitions.empty() || !(current->isVMEM() || current->isFlatLike())) continue; ctx.mv.current = current; idx -= schedule_VMEM_store(ctx, block, current, idx); } } /* resummarize the block's register demand */ block->register_demand = block->live_in_demand; for (const aco_ptr<Instruction>& instr : block->instructions) block->register_demand.update(instr->register_demand); } } /* end namespace */ void schedule_program(Program* program) { /* don't use program->max_reg_demand because that is affected by max_waves_per_simd */ RegisterDemand demand; for (Block& block : program->blocks) demand.update(block.register_demand); demand.vgpr += program->config->num_shared_vgprs / 2; sched_ctx ctx; ctx.gfx_level = program->gfx_level; ctx.mv.depends_on.resize(program->peekAllocationId()); ctx.mv.RAR_dependencies.resize(program->peekAllocationId()); ctx.mv.RAR_dependencies_clause.resize(program->peekAllocationId()); /* Allowing the scheduler to reduce the number of waves to as low as 5 * improves performance of Thrones of Britannia significantly and doesn't * seem to hurt anything else. */ // TODO: account for possible uneven num_waves on GFX10+ unsigned wave_fac = program->dev.physical_vgprs / 256; if (program->num_waves <= 5 * wave_fac) ctx.num_waves = program->num_waves; else if (demand.vgpr >= 29) ctx.num_waves = 5 * wave_fac; else if (demand.vgpr >= 25) ctx.num_waves = 6 * wave_fac; else ctx.num_waves = 7 * wave_fac; ctx.num_waves = std::max<uint16_t>(ctx.num_waves, program->min_waves); ctx.num_waves = std::min<uint16_t>(ctx.num_waves, program->num_waves); ctx.num_waves = max_suitable_waves(program, ctx.num_waves); /* VMEM_MAX_MOVES and such assume pre-GFX10 wave count */ ctx.num_waves = std::max<uint16_t>(ctx.num_waves / wave_fac, 1); assert(ctx.num_waves > 0); ctx.mv.max_registers = { int16_t(get_addr_vgpr_from_waves( program, std::max<uint16_t>(ctx.num_waves * wave_fac, program->min_waves)) - 2), int16_t(get_addr_sgpr_from_waves( program, std::max<uint16_t>(ctx.num_waves * wave_fac, program->min_waves)))}; /* NGG culling shaders are very sensitive to position export scheduling. * Schedule less aggressively when early primitive export is used, and * keep the position export at the very bottom when late primitive export is used. */ if (program->info.hw_stage == AC_HW_NEXT_GEN_GEOMETRY_SHADER) { ctx.schedule_pos_exports = program->info.schedule_ngg_pos_exports; ctx.schedule_pos_export_div = 4; } for (Block& block : program->blocks) schedule_block(ctx, program, &block); /* update max_reg_demand and num_waves */ RegisterDemand new_demand; for (Block& block : program->blocks) { new_demand.update(block.register_demand); } update_vgpr_sgpr_demand(program, new_demand); /* Validate live variable information */ if (!validate_live_vars(program)) abort(); } } // namespace aco