2019-09-17 12:22:17 +01:00
|
|
|
/*
|
|
|
|
* Copyright © 2018 Valve Corporation
|
|
|
|
*
|
|
|
|
* Permission is hereby granted, free of charge, to any person obtaining a
|
|
|
|
* copy of this software and associated documentation files (the "Software"),
|
|
|
|
* to deal in the Software without restriction, including without limitation
|
|
|
|
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
|
|
|
|
* and/or sell copies of the Software, and to permit persons to whom the
|
|
|
|
* Software is furnished to do so, subject to the following conditions:
|
|
|
|
*
|
|
|
|
* The above copyright notice and this permission notice (including the next
|
|
|
|
* paragraph) shall be included in all copies or substantial portions of the
|
|
|
|
* Software.
|
|
|
|
*
|
|
|
|
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
|
|
|
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
|
|
|
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
|
|
|
|
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
|
|
|
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
|
|
|
|
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
|
|
|
|
* IN THE SOFTWARE.
|
|
|
|
*
|
|
|
|
*/
|
|
|
|
|
|
|
|
#include "aco_ir.h"
|
2019-10-14 17:46:02 +01:00
|
|
|
#include "aco_builder.h"
|
2019-09-17 12:22:17 +01:00
|
|
|
#include <unordered_set>
|
|
|
|
#include <algorithm>
|
|
|
|
|
|
|
|
#include "amdgfxregs.h"
|
|
|
|
|
|
|
|
#define SMEM_WINDOW_SIZE (350 - ctx.num_waves * 35)
|
|
|
|
#define VMEM_WINDOW_SIZE (1024 - ctx.num_waves * 64)
|
|
|
|
#define POS_EXP_WINDOW_SIZE 512
|
2019-08-29 16:17:32 +01:00
|
|
|
#define SMEM_MAX_MOVES (64 - ctx.num_waves * 4)
|
aco: fix def-use distance calculation when scheduling.
This change also increases the VMEM_MAX_MOVES
to mitigate some of the scheduling changes.
Totals from 34301 (24.61% of 139391) affected shaders:
SGPRs: 2515440 -> 2552304 (+1.47%); split: -1.25%, +2.71%
VGPRs: 1786676 -> 1794724 (+0.45%); split: -0.31%, +0.76%
CodeSize: 151079856 -> 151209828 (+0.09%); split: -0.06%, +0.15%
MaxWaves: 392454 -> 388966 (-0.89%); split: +0.39%, -1.28%
Instrs: 28870746 -> 28895907 (+0.09%); split: -0.09%, +0.17%
Cycles: 960450680 -> 961315796 (+0.09%); split: -0.09%, +0.18%
VMEM: 19027987 -> 19796223 (+4.04%); split: +7.49%, -3.45%
SMEM: 2434691 -> 2394829 (-1.64%); split: +2.80%, -4.43%
VClause: 551776 -> 543051 (-1.58%); split: -1.73%, +0.15%
SClause: 1230147 -> 1227637 (-0.20%); split: -1.40%, +1.20%
Copies: 1957640 -> 1963617 (+0.31%); split: -1.11%, +1.41%
Branches: 611747 -> 612504 (+0.12%); split: -0.11%, +0.23%
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-13 09:13:49 +01:00
|
|
|
#define VMEM_MAX_MOVES (256 - ctx.num_waves * 16)
|
2019-10-18 13:05:00 +01:00
|
|
|
/* creating clauses decreases def-use distances, so make it less aggressive the lower num_waves is */
|
aco: create VMEM clauses slightly more aggressive
Totals from 3325 (2.39% of 139391) affected shaders (NAVI10):
SGPRs: 331528 -> 331056 (-0.14%); split: -0.14%, +0.00%
VGPRs: 306164 -> 337764 (+10.32%); split: -0.02%, +10.34%
CodeSize: 38843180 -> 38865388 (+0.06%); split: -0.04%, +0.10%
MaxWaves: 18908 -> 17028 (-9.94%); split: +0.01%, -9.95%
Instrs: 7423908 -> 7427934 (+0.05%); split: -0.06%, +0.12%
Cycles: 527411756 -> 526388408 (-0.19%); split: -0.21%, +0.02%
VMEM: 1148421 -> 992660 (-13.56%); split: +0.10%, -13.67%
SMEM: 227337 -> 232380 (+2.22%); split: +2.26%, -0.04%
VClause: 146416 -> 111171 (-24.07%); split: -24.10%, +0.03%
SClause: 243674 -> 243689 (+0.01%); split: -0.00%, +0.01%
Copies: 663496 -> 660333 (-0.48%); split: -0.85%, +0.37%
Branches: 223725 -> 223721 (-0.00%)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-12 17:19:11 +01:00
|
|
|
#define VMEM_CLAUSE_MAX_GRAB_DIST (ctx.num_waves * 8)
|
2019-09-17 12:22:17 +01:00
|
|
|
#define POS_EXP_MAX_MOVES 512
|
|
|
|
|
|
|
|
namespace aco {
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
enum MoveResult {
|
|
|
|
move_success,
|
|
|
|
move_fail_ssa,
|
|
|
|
move_fail_rar,
|
|
|
|
move_fail_pressure,
|
|
|
|
};
|
|
|
|
|
|
|
|
struct MoveState {
|
|
|
|
RegisterDemand max_registers;
|
|
|
|
|
|
|
|
Block *block;
|
|
|
|
Instruction *current;
|
|
|
|
RegisterDemand *register_demand;
|
|
|
|
bool improved_rar;
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
std::vector<bool> depends_on;
|
2019-11-06 16:38:57 +00:00
|
|
|
/* 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. */
|
2019-09-17 12:22:17 +01:00
|
|
|
std::vector<bool> RAR_dependencies;
|
2019-11-06 16:38:57 +00:00
|
|
|
std::vector<bool> RAR_dependencies_clause;
|
|
|
|
|
|
|
|
int source_idx;
|
|
|
|
int insert_idx, insert_idx_clause;
|
|
|
|
RegisterDemand total_demand, total_demand_clause;
|
|
|
|
|
|
|
|
/* for moving instructions before the current instruction to after it */
|
|
|
|
void downwards_init(int current_idx, bool improved_rar, bool may_form_clauses);
|
|
|
|
MoveResult downwards_move(bool clause);
|
|
|
|
void downwards_skip();
|
|
|
|
|
|
|
|
/* for moving instructions after the first use of the current instruction upwards */
|
|
|
|
void upwards_init(int source_idx, bool improved_rar);
|
|
|
|
bool upwards_check_deps();
|
|
|
|
void upwards_set_insert_idx(int before);
|
|
|
|
MoveResult upwards_move();
|
|
|
|
void upwards_skip();
|
|
|
|
|
|
|
|
private:
|
|
|
|
void downwards_advance_helper();
|
|
|
|
};
|
2019-10-18 13:05:00 +01:00
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
struct sched_ctx {
|
2019-09-17 12:22:17 +01:00
|
|
|
int16_t num_waves;
|
|
|
|
int16_t last_SMEM_stall;
|
|
|
|
int last_SMEM_dep_idx;
|
2019-11-06 16:38:57 +00:00
|
|
|
MoveState mv;
|
2019-09-17 12:22:17 +01:00
|
|
|
};
|
|
|
|
|
|
|
|
/* 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>
|
2019-11-06 16:38:57 +00:00
|
|
|
void move_element(T begin_it, size_t idx, size_t before) {
|
2019-09-17 12:22:17 +01:00
|
|
|
if (idx < before) {
|
2019-11-06 16:38:57 +00:00
|
|
|
auto begin = std::next(begin_it, idx);
|
|
|
|
auto end = std::next(begin_it, before);
|
2019-09-17 12:22:17 +01:00
|
|
|
std::rotate(begin, begin + 1, end);
|
|
|
|
} else if (idx > before) {
|
2019-11-06 16:38:57 +00:00
|
|
|
auto begin = std::next(begin_it, before);
|
|
|
|
auto end = std::next(begin_it, idx + 1);
|
2019-09-17 12:22:17 +01:00
|
|
|
std::rotate(begin, end - 1, end);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
void MoveState::downwards_advance_helper()
|
|
|
|
{
|
|
|
|
source_idx--;
|
|
|
|
total_demand.update(register_demand[source_idx]);
|
|
|
|
}
|
|
|
|
|
|
|
|
void MoveState::downwards_init(int current_idx, bool improved_rar_, bool may_form_clauses)
|
|
|
|
{
|
|
|
|
improved_rar = improved_rar_;
|
|
|
|
source_idx = current_idx;
|
|
|
|
|
|
|
|
insert_idx = current_idx + 1;
|
|
|
|
insert_idx_clause = current_idx;
|
|
|
|
|
|
|
|
total_demand = total_demand_clause = register_demand[current_idx];
|
|
|
|
|
|
|
|
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;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
/* update total_demand/source_idx */
|
|
|
|
downwards_advance_helper();
|
|
|
|
}
|
|
|
|
|
|
|
|
MoveResult MoveState::downwards_move(bool clause)
|
|
|
|
{
|
|
|
|
aco_ptr<Instruction>& instr = block->instructions[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 ? (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 (clause) {
|
|
|
|
for (const Operand& op : instr->operands) {
|
|
|
|
if (op.isTemp()) {
|
|
|
|
depends_on[op.tempId()] = true;
|
|
|
|
if (op.isFirstKill())
|
|
|
|
RAR_dependencies[op.tempId()] = true;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
int dest_insert_idx = clause ? insert_idx_clause : insert_idx;
|
|
|
|
RegisterDemand register_pressure = clause ? total_demand_clause : total_demand;
|
|
|
|
|
2020-02-21 20:14:03 +00:00
|
|
|
const RegisterDemand candidate_diff = get_live_changes(instr);
|
|
|
|
const RegisterDemand temp = get_temp_registers(instr);
|
2019-11-06 16:38:57 +00:00
|
|
|
if (RegisterDemand(register_pressure - candidate_diff).exceeds(max_registers))
|
|
|
|
return move_fail_pressure;
|
2020-02-21 20:14:03 +00:00
|
|
|
const RegisterDemand temp2 = get_temp_registers(block->instructions[dest_insert_idx - 1]);
|
2019-11-06 16:38:57 +00:00
|
|
|
const RegisterDemand new_demand = register_demand[dest_insert_idx - 1] - temp2 + temp;
|
|
|
|
if (new_demand.exceeds(max_registers))
|
|
|
|
return move_fail_pressure;
|
|
|
|
|
|
|
|
/* move the candidate below the memory load */
|
|
|
|
move_element(block->instructions.begin(), source_idx, dest_insert_idx);
|
|
|
|
|
|
|
|
/* update register pressure */
|
|
|
|
move_element(register_demand, source_idx, dest_insert_idx);
|
|
|
|
for (int i = source_idx; i < dest_insert_idx - 1; i++)
|
|
|
|
register_demand[i] -= candidate_diff;
|
|
|
|
register_demand[dest_insert_idx - 1] = new_demand;
|
|
|
|
total_demand_clause -= candidate_diff;
|
|
|
|
insert_idx_clause--;
|
|
|
|
if (!clause) {
|
|
|
|
total_demand -= candidate_diff;
|
|
|
|
insert_idx--;
|
|
|
|
}
|
|
|
|
|
|
|
|
downwards_advance_helper();
|
|
|
|
return move_success;
|
|
|
|
}
|
|
|
|
|
|
|
|
void MoveState::downwards_skip()
|
|
|
|
{
|
|
|
|
aco_ptr<Instruction>& instr = block->instructions[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;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
total_demand_clause.update(register_demand[source_idx]);
|
|
|
|
|
|
|
|
downwards_advance_helper();
|
|
|
|
}
|
|
|
|
|
|
|
|
void MoveState::upwards_init(int source_idx_, bool improved_rar_)
|
|
|
|
{
|
|
|
|
source_idx = source_idx_;
|
|
|
|
improved_rar = improved_rar_;
|
|
|
|
|
|
|
|
insert_idx = -1;
|
|
|
|
|
|
|
|
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;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
bool MoveState::upwards_check_deps()
|
|
|
|
{
|
|
|
|
aco_ptr<Instruction>& instr = block->instructions[source_idx];
|
|
|
|
for (const Operand& op : instr->operands) {
|
|
|
|
if (op.isTemp() && depends_on[op.tempId()])
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
void MoveState::upwards_set_insert_idx(int before)
|
|
|
|
{
|
|
|
|
insert_idx = before;
|
|
|
|
total_demand = register_demand[before - 1];
|
|
|
|
}
|
|
|
|
|
|
|
|
MoveResult MoveState::upwards_move()
|
|
|
|
{
|
|
|
|
assert(insert_idx >= 0);
|
|
|
|
|
|
|
|
aco_ptr<Instruction>& instr = block->instructions[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 */
|
2020-02-21 20:14:03 +00:00
|
|
|
const RegisterDemand candidate_diff = get_live_changes(instr);
|
|
|
|
const RegisterDemand temp = get_temp_registers(instr);
|
2019-11-06 16:38:57 +00:00
|
|
|
if (RegisterDemand(total_demand + candidate_diff).exceeds(max_registers))
|
|
|
|
return move_fail_pressure;
|
2020-02-21 20:14:03 +00:00
|
|
|
const RegisterDemand temp2 = get_temp_registers(block->instructions[insert_idx - 1]);
|
2019-11-06 16:38:57 +00:00
|
|
|
const RegisterDemand new_demand = register_demand[insert_idx - 1] - 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(), source_idx, insert_idx);
|
|
|
|
|
|
|
|
/* update register pressure */
|
|
|
|
move_element(register_demand, source_idx, insert_idx);
|
|
|
|
for (int i = insert_idx + 1; i <= source_idx; i++)
|
|
|
|
register_demand[i] += candidate_diff;
|
|
|
|
register_demand[insert_idx] = new_demand;
|
|
|
|
total_demand += candidate_diff;
|
|
|
|
|
|
|
|
insert_idx++;
|
|
|
|
|
|
|
|
total_demand.update(register_demand[source_idx]);
|
|
|
|
source_idx++;
|
|
|
|
|
|
|
|
return move_success;
|
|
|
|
}
|
|
|
|
|
|
|
|
void MoveState::upwards_skip()
|
|
|
|
{
|
|
|
|
if (insert_idx >= 0) {
|
|
|
|
aco_ptr<Instruction>& instr = block->instructions[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;
|
|
|
|
}
|
|
|
|
total_demand.update(register_demand[source_idx]);
|
|
|
|
}
|
|
|
|
|
|
|
|
source_idx++;
|
|
|
|
}
|
|
|
|
|
2020-03-18 09:40:06 +00:00
|
|
|
bool is_gs_or_done_sendmsg(const Instruction *instr)
|
2019-10-14 17:46:02 +01:00
|
|
|
{
|
|
|
|
if (instr->opcode == aco_opcode::s_sendmsg) {
|
2021-01-21 16:13:34 +00:00
|
|
|
uint16_t imm = instr->sopp().imm;
|
2019-10-14 17:46:02 +01:00
|
|
|
return (imm & sendmsg_id_mask) == _sendmsg_gs ||
|
|
|
|
(imm & sendmsg_id_mask) == _sendmsg_gs_done;
|
|
|
|
}
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2020-03-18 09:40:06 +00:00
|
|
|
bool is_done_sendmsg(const Instruction *instr)
|
2019-10-14 17:46:02 +01:00
|
|
|
{
|
2021-01-20 14:49:08 +00:00
|
|
|
if (instr->opcode == aco_opcode::s_sendmsg)
|
2021-01-21 16:13:34 +00:00
|
|
|
return (instr->sopp().imm & sendmsg_id_mask) == _sendmsg_gs_done;
|
2019-10-14 17:46:02 +01:00
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
memory_sync_info get_sync_info_with_hack(const Instruction* instr)
|
2019-10-14 17:46:02 +01:00
|
|
|
{
|
2020-06-26 15:54:22 +01:00
|
|
|
memory_sync_info sync = get_sync_info(instr);
|
2021-01-20 15:27:16 +00:00
|
|
|
if (instr->isSMEM() && !instr->operands.empty() && instr->operands[0].bytes() == 16) {
|
2020-06-26 15:54:22 +01:00
|
|
|
// FIXME: currently, it doesn't seem beneficial to omit this due to how our scheduler works
|
|
|
|
sync.storage = (storage_class)(sync.storage | storage_buffer);
|
2020-08-27 12:11:46 +01:00
|
|
|
sync.semantics = (memory_semantics)((sync.semantics | semantic_private) & ~semantic_can_reorder);
|
2019-10-14 17:46:02 +01:00
|
|
|
}
|
2020-06-26 15:54:22 +01:00
|
|
|
return sync;
|
2019-10-14 17:46:02 +01:00
|
|
|
}
|
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
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;
|
|
|
|
};
|
2019-11-07 14:48:51 +00:00
|
|
|
|
|
|
|
struct hazard_query {
|
|
|
|
bool contains_spill;
|
2020-06-26 15:54:22 +01:00
|
|
|
bool contains_sendmsg;
|
aco: allow to schedule SALU/SMEM through exec changes
Totals from 16794 (12.05% of 139391) affected shaders (NAVI10):
SGPRs: 757760 -> 762048 (+0.57%); split: -0.39%, +0.95%
VGPRs: 402844 -> 402744 (-0.02%); split: -0.04%, +0.02%
CodeSize: 22290900 -> 22285068 (-0.03%); split: -0.06%, +0.04%
MaxWaves: 294163 -> 294222 (+0.02%); split: +0.03%, -0.01%
Instrs: 4190074 -> 4188513 (-0.04%); split: -0.08%, +0.04%
Cycles: 40685028 -> 40678640 (-0.02%); split: -0.03%, +0.02%
VMEM: 7711867 -> 7704315 (-0.10%); split: +0.28%, -0.38%
SMEM: 942472 -> 1007052 (+6.85%); split: +7.15%, -0.30%
VClause: 92990 -> 92974 (-0.02%); split: -0.03%, +0.01%
SClause: 263700 -> 263810 (+0.04%); split: -0.38%, +0.42%
Copies: 277467 -> 276988 (-0.17%); split: -0.37%, +0.20%
Branches: 45899 -> 45896 (-0.01%)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-12 15:16:46 +01:00
|
|
|
bool uses_exec;
|
2020-06-26 15:54:22 +01:00
|
|
|
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) */
|
2019-11-07 14:48:51 +00:00
|
|
|
};
|
|
|
|
|
|
|
|
void init_hazard_query(hazard_query *query) {
|
|
|
|
query->contains_spill = false;
|
2020-06-26 15:54:22 +01:00
|
|
|
query->contains_sendmsg = false;
|
aco: allow to schedule SALU/SMEM through exec changes
Totals from 16794 (12.05% of 139391) affected shaders (NAVI10):
SGPRs: 757760 -> 762048 (+0.57%); split: -0.39%, +0.95%
VGPRs: 402844 -> 402744 (-0.02%); split: -0.04%, +0.02%
CodeSize: 22290900 -> 22285068 (-0.03%); split: -0.06%, +0.04%
MaxWaves: 294163 -> 294222 (+0.02%); split: +0.03%, -0.01%
Instrs: 4190074 -> 4188513 (-0.04%); split: -0.08%, +0.04%
Cycles: 40685028 -> 40678640 (-0.02%); split: -0.03%, +0.02%
VMEM: 7711867 -> 7704315 (-0.10%); split: +0.28%, -0.38%
SMEM: 942472 -> 1007052 (+6.85%); split: +7.15%, -0.30%
VClause: 92990 -> 92974 (-0.02%); split: -0.03%, +0.01%
SClause: 263700 -> 263810 (+0.04%); split: -0.38%, +0.42%
Copies: 277467 -> 276988 (-0.17%); split: -0.37%, +0.20%
Branches: 45899 -> 45896 (-0.01%)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-12 15:16:46 +01:00
|
|
|
query->uses_exec = false;
|
2020-06-26 15:54:22 +01:00
|
|
|
memset(&query->mem_events, 0, sizeof(query->mem_events));
|
|
|
|
query->aliasing_storage = 0;
|
|
|
|
query->aliasing_storage_smem = 0;
|
|
|
|
}
|
|
|
|
|
|
|
|
void add_memory_event(memory_event_set *set, Instruction *instr, memory_sync_info *sync)
|
|
|
|
{
|
|
|
|
set->has_control_barrier |= is_done_sendmsg(instr);
|
|
|
|
if (instr->opcode == aco_opcode::p_barrier) {
|
2021-01-21 16:13:34 +00:00
|
|
|
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;
|
2020-06-26 15:54:22 +01:00
|
|
|
}
|
|
|
|
|
|
|
|
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;
|
|
|
|
}
|
2019-11-07 14:48:51 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
void add_to_hazard_query(hazard_query *query, Instruction *instr)
|
2019-09-17 12:22:17 +01:00
|
|
|
{
|
2019-11-07 14:48:51 +00:00
|
|
|
if (instr->opcode == aco_opcode::p_spill || instr->opcode == aco_opcode::p_reload)
|
|
|
|
query->contains_spill = true;
|
2020-06-26 15:54:22 +01:00
|
|
|
query->contains_sendmsg |= instr->opcode == aco_opcode::s_sendmsg;
|
aco: allow to schedule SALU/SMEM through exec changes
Totals from 16794 (12.05% of 139391) affected shaders (NAVI10):
SGPRs: 757760 -> 762048 (+0.57%); split: -0.39%, +0.95%
VGPRs: 402844 -> 402744 (-0.02%); split: -0.04%, +0.02%
CodeSize: 22290900 -> 22285068 (-0.03%); split: -0.06%, +0.04%
MaxWaves: 294163 -> 294222 (+0.02%); split: +0.03%, -0.01%
Instrs: 4190074 -> 4188513 (-0.04%); split: -0.08%, +0.04%
Cycles: 40685028 -> 40678640 (-0.02%); split: -0.03%, +0.02%
VMEM: 7711867 -> 7704315 (-0.10%); split: +0.28%, -0.38%
SMEM: 942472 -> 1007052 (+6.85%); split: +7.15%, -0.30%
VClause: 92990 -> 92974 (-0.02%); split: -0.03%, +0.01%
SClause: 263700 -> 263810 (+0.04%); split: -0.38%, +0.42%
Copies: 277467 -> 276988 (-0.17%); split: -0.37%, +0.20%
Branches: 45899 -> 45896 (-0.01%)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-12 15:16:46 +01:00
|
|
|
query->uses_exec |= needs_exec_mask(instr);
|
2020-06-26 15:54:22 +01:00
|
|
|
|
|
|
|
memory_sync_info sync = get_sync_info_with_hack(instr);
|
2019-11-07 14:48:51 +00:00
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
add_memory_event(&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;
|
2021-01-20 15:27:16 +00:00
|
|
|
if (instr->isSMEM())
|
2020-06-26 15:54:22 +01:00
|
|
|
query->aliasing_storage_smem |= storage;
|
|
|
|
else
|
|
|
|
query->aliasing_storage |= storage;
|
|
|
|
}
|
2019-11-07 14:48:51 +00:00
|
|
|
}
|
|
|
|
|
|
|
|
enum HazardResult {
|
|
|
|
hazard_success,
|
|
|
|
hazard_fail_reorder_vmem_smem,
|
|
|
|
hazard_fail_reorder_ds,
|
|
|
|
hazard_fail_reorder_sendmsg,
|
|
|
|
hazard_fail_spill,
|
2020-02-11 14:15:32 +00:00
|
|
|
hazard_fail_export,
|
|
|
|
hazard_fail_barrier,
|
2019-11-07 14:48:51 +00:00
|
|
|
/* Must stop at these failures. The hazard query code doesn't consider them
|
|
|
|
* when added. */
|
|
|
|
hazard_fail_exec,
|
2020-04-01 14:38:43 +01:00
|
|
|
hazard_fail_unreorderable,
|
2019-11-07 14:48:51 +00:00
|
|
|
};
|
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
HazardResult perform_hazard_query(hazard_query *query, Instruction *instr, bool upwards)
|
2019-11-07 14:48:51 +00:00
|
|
|
{
|
aco: allow to schedule SALU/SMEM through exec changes
Totals from 16794 (12.05% of 139391) affected shaders (NAVI10):
SGPRs: 757760 -> 762048 (+0.57%); split: -0.39%, +0.95%
VGPRs: 402844 -> 402744 (-0.02%); split: -0.04%, +0.02%
CodeSize: 22290900 -> 22285068 (-0.03%); split: -0.06%, +0.04%
MaxWaves: 294163 -> 294222 (+0.02%); split: +0.03%, -0.01%
Instrs: 4190074 -> 4188513 (-0.04%); split: -0.08%, +0.04%
Cycles: 40685028 -> 40678640 (-0.02%); split: -0.03%, +0.02%
VMEM: 7711867 -> 7704315 (-0.10%); split: +0.28%, -0.38%
SMEM: 942472 -> 1007052 (+6.85%); split: +7.15%, -0.30%
VClause: 92990 -> 92974 (-0.02%); split: -0.03%, +0.01%
SClause: 263700 -> 263810 (+0.04%); split: -0.38%, +0.42%
Copies: 277467 -> 276988 (-0.17%); split: -0.37%, +0.20%
Branches: 45899 -> 45896 (-0.01%)
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-12 15:16:46 +01:00
|
|
|
/* don't schedule discards downwards */
|
|
|
|
if (!upwards && instr->opcode == aco_opcode::p_exit_early_if)
|
|
|
|
return hazard_fail_unreorderable;
|
|
|
|
|
|
|
|
if (query->uses_exec) {
|
|
|
|
for (const Definition& def : instr->definitions) {
|
|
|
|
if (def.isFixed() && def.physReg() == exec)
|
|
|
|
return hazard_fail_exec;
|
|
|
|
}
|
2019-11-07 14:48:51 +00:00
|
|
|
}
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
/* don't move exports so that they stay closer together */
|
2021-01-20 15:27:16 +00:00
|
|
|
if (instr->isEXP())
|
2020-02-11 14:15:32 +00:00
|
|
|
return hazard_fail_export;
|
2019-09-17 12:22:17 +01:00
|
|
|
|
2020-04-01 14:38:43 +01:00
|
|
|
/* don't move non-reorderable instructions */
|
|
|
|
if (instr->opcode == aco_opcode::s_memtime ||
|
|
|
|
instr->opcode == aco_opcode::s_memrealtime ||
|
2020-06-18 14:45:31 +01:00
|
|
|
instr->opcode == aco_opcode::s_setprio ||
|
|
|
|
instr->opcode == aco_opcode::s_getreg_b32)
|
2020-04-01 14:38:43 +01:00
|
|
|
return hazard_fail_unreorderable;
|
2019-09-17 12:22:17 +01:00
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
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(&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))
|
2020-05-01 11:04:27 +01:00
|
|
|
return hazard_fail_barrier;
|
2020-06-26 15:54:22 +01:00
|
|
|
if ((first->bar_classes && (second->bar_release || second->access_release)) ||
|
|
|
|
((first->access_relaxed | first->access_atomic) & (second->bar_release | second->access_release)))
|
2020-02-11 14:15:32 +00:00
|
|
|
return hazard_fail_barrier;
|
2020-06-26 15:54:22 +01:00
|
|
|
|
|
|
|
/* don't move memory barriers around other memory barriers */
|
|
|
|
if (first->bar_classes && second->bar_classes)
|
2019-11-07 14:48:51 +00:00
|
|
|
return hazard_fail_barrier;
|
|
|
|
|
2020-08-04 19:20:21 +01:00
|
|
|
/* 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_atomic_counter | storage_image | storage_shared;
|
|
|
|
if (first->has_control_barrier && ((second->access_atomic | second->access_relaxed) & control_classes))
|
|
|
|
return hazard_fail_barrier;
|
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
/* don't move memory loads/stores past potentially aliasing loads/stores */
|
2021-01-20 15:27:16 +00:00
|
|
|
unsigned aliasing_storage = instr->isSMEM() ?
|
2020-06-26 15:54:22 +01:00
|
|
|
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;
|
2019-11-07 14:48:51 +00:00
|
|
|
return hazard_fail_reorder_vmem_smem;
|
2020-06-26 15:54:22 +01:00
|
|
|
}
|
2019-11-07 14:48:51 +00:00
|
|
|
|
|
|
|
if ((instr->opcode == aco_opcode::p_spill || instr->opcode == aco_opcode::p_reload) &&
|
|
|
|
query->contains_spill)
|
|
|
|
return hazard_fail_spill;
|
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
if (instr->opcode == aco_opcode::s_sendmsg && query->contains_sendmsg)
|
|
|
|
return hazard_fail_reorder_sendmsg;
|
|
|
|
|
2019-11-07 14:48:51 +00:00
|
|
|
return hazard_success;
|
2019-09-17 12:22:17 +01:00
|
|
|
}
|
|
|
|
|
|
|
|
void schedule_SMEM(sched_ctx& ctx, Block* block,
|
|
|
|
std::vector<RegisterDemand>& register_demand,
|
|
|
|
Instruction* current, int idx)
|
|
|
|
{
|
|
|
|
assert(idx != 0);
|
|
|
|
int window_size = SMEM_WINDOW_SIZE;
|
|
|
|
int max_moves = SMEM_MAX_MOVES;
|
|
|
|
int16_t k = 0;
|
|
|
|
|
2019-10-10 17:04:06 +01:00
|
|
|
/* don't move s_memtime/s_memrealtime */
|
|
|
|
if (current->opcode == aco_opcode::s_memtime || current->opcode == aco_opcode::s_memrealtime)
|
|
|
|
return;
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
/* first, check if we have instructions before current to move down */
|
2019-11-07 14:48:51 +00:00
|
|
|
hazard_query hq;
|
|
|
|
init_hazard_query(&hq);
|
|
|
|
add_to_hazard_query(&hq, current);
|
2019-09-17 12:22:17 +01:00
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_init(idx, false, false);
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
for (int candidate_idx = idx - 1; k < max_moves && candidate_idx > (int) idx - window_size; candidate_idx--) {
|
|
|
|
assert(candidate_idx >= 0);
|
2019-11-06 16:38:57 +00:00
|
|
|
assert(candidate_idx == ctx.mv.source_idx);
|
2019-09-17 12:22:17 +01:00
|
|
|
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;
|
2019-10-31 16:33:35 +00:00
|
|
|
if (candidate->isVMEM())
|
|
|
|
break;
|
2019-09-17 12:22:17 +01:00
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
bool can_move_down = true;
|
2019-11-07 14:48:51 +00:00
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
HazardResult haz = perform_hazard_query(&hq, candidate.get(), false);
|
2020-02-11 16:55:39 +00:00
|
|
|
if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill || haz == hazard_fail_reorder_sendmsg || haz == hazard_fail_barrier || haz == hazard_fail_export)
|
2019-09-17 12:22:17 +01:00
|
|
|
can_move_down = false;
|
2020-02-11 14:15:32 +00:00
|
|
|
else if (haz != hazard_success)
|
2019-11-07 14:48:51 +00:00
|
|
|
break;
|
|
|
|
|
|
|
|
/* don't use LDS/GDS instructions to hide latency since it can
|
|
|
|
* significanly worsen LDS scheduling */
|
2021-01-20 15:27:16 +00:00
|
|
|
if (candidate->isDS() || !can_move_down) {
|
2019-11-07 14:48:51 +00:00
|
|
|
add_to_hazard_query(&hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
|
|
|
}
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
MoveResult res = ctx.mv.downwards_move(false);
|
|
|
|
if (res == move_fail_ssa || res == move_fail_rar) {
|
2019-11-07 14:48:51 +00:00
|
|
|
add_to_hazard_query(&hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
2019-11-06 16:38:57 +00:00
|
|
|
} else if (res == move_fail_pressure) {
|
2019-09-17 12:22:17 +01:00
|
|
|
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 */
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_init(idx + 1, false);
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
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++) {
|
2019-11-06 16:38:57 +00:00
|
|
|
assert(candidate_idx == ctx.mv.source_idx);
|
2019-09-17 12:22:17 +01:00
|
|
|
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 */
|
2019-11-06 16:38:57 +00:00
|
|
|
bool is_dependency = !found_dependency && !ctx.mv.upwards_check_deps();
|
2019-10-10 15:31:40 +01:00
|
|
|
/* no need to steal from following VMEM instructions */
|
|
|
|
if (is_dependency && candidate->isVMEM())
|
|
|
|
break;
|
2019-11-07 14:48:51 +00:00
|
|
|
|
|
|
|
if (found_dependency) {
|
2020-06-26 15:54:22 +01:00
|
|
|
HazardResult haz = perform_hazard_query(&hq, candidate.get(), true);
|
2020-02-11 14:15:32 +00:00
|
|
|
if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill ||
|
2020-02-11 16:55:39 +00:00
|
|
|
haz == hazard_fail_reorder_sendmsg || haz == hazard_fail_barrier ||
|
|
|
|
haz == hazard_fail_export)
|
2019-11-07 14:48:51 +00:00
|
|
|
is_dependency = true;
|
2020-02-11 14:15:32 +00:00
|
|
|
else if (haz != hazard_success)
|
2019-11-07 14:48:51 +00:00
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
if (is_dependency) {
|
|
|
|
if (!found_dependency) {
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_set_insert_idx(candidate_idx);
|
2019-11-07 14:48:51 +00:00
|
|
|
init_hazard_query(&hq);
|
2019-09-17 12:22:17 +01:00
|
|
|
found_dependency = true;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2019-11-07 14:48:51 +00:00
|
|
|
if (is_dependency || !found_dependency) {
|
|
|
|
if (found_dependency)
|
|
|
|
add_to_hazard_query(&hq, candidate.get());
|
|
|
|
else
|
|
|
|
k++;
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
2019-10-31 16:33:35 +00:00
|
|
|
}
|
2019-11-06 16:38:57 +00:00
|
|
|
|
|
|
|
MoveResult res = ctx.mv.upwards_move();
|
|
|
|
if (res == move_fail_ssa || res == move_fail_rar) {
|
|
|
|
/* no need to steal from following VMEM instructions */
|
|
|
|
if (res == move_fail_ssa && candidate->isVMEM())
|
2019-10-31 16:33:35 +00:00
|
|
|
break;
|
2019-11-07 14:48:51 +00:00
|
|
|
add_to_hazard_query(&hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
2019-11-06 16:38:57 +00:00
|
|
|
} else if (res == move_fail_pressure) {
|
2019-09-17 12:22:17 +01:00
|
|
|
break;
|
|
|
|
}
|
|
|
|
k++;
|
|
|
|
}
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.last_SMEM_dep_idx = found_dependency ? ctx.mv.insert_idx : 0;
|
2019-09-17 12:22:17 +01:00
|
|
|
ctx.last_SMEM_stall = 10 - ctx.num_waves - k;
|
|
|
|
}
|
|
|
|
|
|
|
|
void schedule_VMEM(sched_ctx& ctx, Block* block,
|
|
|
|
std::vector<RegisterDemand>& register_demand,
|
|
|
|
Instruction* current, int idx)
|
|
|
|
{
|
|
|
|
assert(idx != 0);
|
|
|
|
int window_size = VMEM_WINDOW_SIZE;
|
|
|
|
int max_moves = VMEM_MAX_MOVES;
|
2019-10-18 13:05:00 +01:00
|
|
|
int clause_max_grab_dist = VMEM_CLAUSE_MAX_GRAB_DIST;
|
2019-09-17 12:22:17 +01:00
|
|
|
int16_t k = 0;
|
|
|
|
|
|
|
|
/* first, check if we have instructions before current to move down */
|
2019-11-07 14:48:51 +00:00
|
|
|
hazard_query indep_hq;
|
|
|
|
hazard_query clause_hq;
|
|
|
|
init_hazard_query(&indep_hq);
|
|
|
|
init_hazard_query(&clause_hq);
|
|
|
|
add_to_hazard_query(&indep_hq, current);
|
2019-09-17 12:22:17 +01:00
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_init(idx, true, true);
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
for (int candidate_idx = idx - 1; k < max_moves && candidate_idx > (int) idx - window_size; candidate_idx--) {
|
2019-11-06 16:38:57 +00:00
|
|
|
assert(candidate_idx == ctx.mv.source_idx);
|
2019-09-17 12:22:17 +01:00
|
|
|
assert(candidate_idx >= 0);
|
|
|
|
aco_ptr<Instruction>& candidate = block->instructions[candidate_idx];
|
2021-01-20 15:27:16 +00:00
|
|
|
bool is_vmem = candidate->isVMEM() || candidate->isFlatLike();
|
2019-09-17 12:22:17 +01:00
|
|
|
|
|
|
|
/* 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;
|
2019-10-18 13:05:00 +01:00
|
|
|
|
|
|
|
bool part_of_clause = false;
|
2019-11-27 17:27:36 +00:00
|
|
|
if (current->isVMEM() == candidate->isVMEM()) {
|
2019-11-06 16:38:57 +00:00
|
|
|
int grab_dist = ctx.mv.insert_idx_clause - candidate_idx;
|
2019-10-18 13:05:00 +01:00
|
|
|
/* 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. */
|
2020-11-30 17:53:23 +00:00
|
|
|
part_of_clause = grab_dist < clause_max_grab_dist &&
|
|
|
|
should_form_clause(current, candidate.get());
|
2019-10-18 13:05:00 +01:00
|
|
|
}
|
2019-09-17 12:22:17 +01:00
|
|
|
|
|
|
|
/* if current depends on candidate, add additional dependencies and continue */
|
2019-11-27 17:27:36 +00:00
|
|
|
bool can_move_down = !is_vmem || part_of_clause;
|
2019-09-17 12:22:17 +01:00
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
HazardResult haz = perform_hazard_query(part_of_clause ? &clause_hq : &indep_hq, candidate.get(), false);
|
2020-02-11 14:15:32 +00:00
|
|
|
if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill ||
|
2020-02-11 16:55:39 +00:00
|
|
|
haz == hazard_fail_reorder_sendmsg || haz == hazard_fail_barrier ||
|
|
|
|
haz == hazard_fail_export)
|
2019-09-17 12:22:17 +01:00
|
|
|
can_move_down = false;
|
2019-11-07 14:48:51 +00:00
|
|
|
else if (haz != hazard_success)
|
|
|
|
break;
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
if (!can_move_down) {
|
2019-11-07 14:48:51 +00:00
|
|
|
add_to_hazard_query(&indep_hq, candidate.get());
|
|
|
|
add_to_hazard_query(&clause_hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
|
|
|
}
|
|
|
|
|
2020-07-24 15:49:43 +01:00
|
|
|
Instruction *candidate_ptr = candidate.get();
|
2019-11-06 16:38:57 +00:00
|
|
|
MoveResult res = ctx.mv.downwards_move(part_of_clause);
|
|
|
|
if (res == move_fail_ssa || res == move_fail_rar) {
|
2019-11-07 14:48:51 +00:00
|
|
|
add_to_hazard_query(&indep_hq, candidate.get());
|
|
|
|
add_to_hazard_query(&clause_hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
2019-11-06 16:38:57 +00:00
|
|
|
} else if (res == move_fail_pressure) {
|
2019-09-17 12:22:17 +01:00
|
|
|
break;
|
2019-10-18 13:05:00 +01:00
|
|
|
}
|
2020-07-24 15:49:43 +01:00
|
|
|
if (part_of_clause)
|
|
|
|
add_to_hazard_query(&indep_hq, candidate_ptr);
|
aco: fix def-use distance calculation when scheduling.
This change also increases the VMEM_MAX_MOVES
to mitigate some of the scheduling changes.
Totals from 34301 (24.61% of 139391) affected shaders:
SGPRs: 2515440 -> 2552304 (+1.47%); split: -1.25%, +2.71%
VGPRs: 1786676 -> 1794724 (+0.45%); split: -0.31%, +0.76%
CodeSize: 151079856 -> 151209828 (+0.09%); split: -0.06%, +0.15%
MaxWaves: 392454 -> 388966 (-0.89%); split: +0.39%, -1.28%
Instrs: 28870746 -> 28895907 (+0.09%); split: -0.09%, +0.17%
Cycles: 960450680 -> 961315796 (+0.09%); split: -0.09%, +0.18%
VMEM: 19027987 -> 19796223 (+4.04%); split: +7.49%, -3.45%
SMEM: 2434691 -> 2394829 (-1.64%); split: +2.80%, -4.43%
VClause: 551776 -> 543051 (-1.58%); split: -1.73%, +0.15%
SClause: 1230147 -> 1227637 (-0.20%); split: -1.40%, +1.20%
Copies: 1957640 -> 1963617 (+0.31%); split: -1.11%, +1.41%
Branches: 611747 -> 612504 (+0.12%); split: -0.11%, +0.23%
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-13 09:13:49 +01:00
|
|
|
else
|
|
|
|
k++;
|
2019-09-17 12:22:17 +01:00
|
|
|
if (candidate_idx < ctx.last_SMEM_dep_idx)
|
|
|
|
ctx.last_SMEM_stall++;
|
|
|
|
}
|
|
|
|
|
|
|
|
/* find the first instruction depending on current or find another VMEM */
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_init(idx + 1, true);
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
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++) {
|
2019-11-06 16:38:57 +00:00
|
|
|
assert(candidate_idx == ctx.mv.source_idx);
|
2019-09-17 12:22:17 +01:00
|
|
|
assert(candidate_idx < (int) block->instructions.size());
|
|
|
|
aco_ptr<Instruction>& candidate = block->instructions[candidate_idx];
|
2021-01-20 15:27:16 +00:00
|
|
|
bool is_vmem = candidate->isVMEM() || candidate->isFlatLike();
|
2019-09-17 12:22:17 +01:00
|
|
|
|
|
|
|
if (candidate->opcode == aco_opcode::p_logical_end)
|
|
|
|
break;
|
|
|
|
|
|
|
|
/* check if candidate depends on current */
|
2019-10-31 16:33:35 +00:00
|
|
|
bool is_dependency = false;
|
2019-11-07 14:48:51 +00:00
|
|
|
if (found_dependency) {
|
2020-06-26 15:54:22 +01:00
|
|
|
HazardResult haz = perform_hazard_query(&indep_hq, candidate.get(), true);
|
2019-11-07 14:48:51 +00:00
|
|
|
if (haz == hazard_fail_reorder_ds || haz == hazard_fail_spill ||
|
2020-02-11 14:15:32 +00:00
|
|
|
haz == hazard_fail_reorder_vmem_smem || haz == hazard_fail_reorder_sendmsg ||
|
2020-02-11 16:55:39 +00:00
|
|
|
haz == hazard_fail_barrier || haz == hazard_fail_export)
|
2019-11-07 14:48:51 +00:00
|
|
|
is_dependency = true;
|
2020-02-11 14:15:32 +00:00
|
|
|
else if (haz != hazard_success)
|
2019-11-07 14:48:51 +00:00
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
is_dependency |= !found_dependency && !ctx.mv.upwards_check_deps();
|
2019-09-17 12:22:17 +01:00
|
|
|
if (is_dependency) {
|
|
|
|
if (!found_dependency) {
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_set_insert_idx(candidate_idx);
|
2019-11-07 14:48:51 +00:00
|
|
|
init_hazard_query(&indep_hq);
|
2019-09-17 12:22:17 +01:00
|
|
|
found_dependency = true;
|
|
|
|
}
|
2019-11-27 17:27:36 +00:00
|
|
|
} else if (is_vmem) {
|
2019-10-31 16:33:35 +00:00
|
|
|
/* don't move up dependencies of other VMEM instructions */
|
2019-08-28 11:08:12 +01:00
|
|
|
for (const Definition& def : candidate->definitions) {
|
|
|
|
if (def.isTemp())
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.depends_on[def.tempId()] = true;
|
2019-08-28 11:08:12 +01:00
|
|
|
}
|
2019-09-17 12:22:17 +01:00
|
|
|
}
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
if (is_dependency || !found_dependency) {
|
2019-11-07 14:48:51 +00:00
|
|
|
if (found_dependency)
|
|
|
|
add_to_hazard_query(&indep_hq, candidate.get());
|
aco: fix def-use distance calculation when scheduling.
This change also increases the VMEM_MAX_MOVES
to mitigate some of the scheduling changes.
Totals from 34301 (24.61% of 139391) affected shaders:
SGPRs: 2515440 -> 2552304 (+1.47%); split: -1.25%, +2.71%
VGPRs: 1786676 -> 1794724 (+0.45%); split: -0.31%, +0.76%
CodeSize: 151079856 -> 151209828 (+0.09%); split: -0.06%, +0.15%
MaxWaves: 392454 -> 388966 (-0.89%); split: +0.39%, -1.28%
Instrs: 28870746 -> 28895907 (+0.09%); split: -0.09%, +0.17%
Cycles: 960450680 -> 961315796 (+0.09%); split: -0.09%, +0.18%
VMEM: 19027987 -> 19796223 (+4.04%); split: +7.49%, -3.45%
SMEM: 2434691 -> 2394829 (-1.64%); split: +2.80%, -4.43%
VClause: 551776 -> 543051 (-1.58%); split: -1.73%, +0.15%
SClause: 1230147 -> 1227637 (-0.20%); split: -1.40%, +1.20%
Copies: 1957640 -> 1963617 (+0.31%); split: -1.11%, +1.41%
Branches: 611747 -> 612504 (+0.12%); split: -0.11%, +0.23%
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/7903>
2020-08-13 09:13:49 +01:00
|
|
|
else
|
|
|
|
k++;
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
|
|
|
}
|
2019-11-06 16:38:57 +00:00
|
|
|
|
|
|
|
MoveResult res = ctx.mv.upwards_move();
|
|
|
|
if (res == move_fail_ssa || res == move_fail_rar) {
|
2019-11-07 14:48:51 +00:00
|
|
|
add_to_hazard_query(&indep_hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.upwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
2019-11-06 16:38:57 +00:00
|
|
|
} else if (res == move_fail_pressure) {
|
2019-09-17 12:22:17 +01:00
|
|
|
break;
|
|
|
|
}
|
|
|
|
k++;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
void schedule_position_export(sched_ctx& ctx, Block* block,
|
|
|
|
std::vector<RegisterDemand>& register_demand,
|
|
|
|
Instruction* current, int idx)
|
|
|
|
{
|
|
|
|
assert(idx != 0);
|
|
|
|
int window_size = POS_EXP_WINDOW_SIZE;
|
|
|
|
int max_moves = POS_EXP_MAX_MOVES;
|
|
|
|
int16_t k = 0;
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_init(idx, true, false);
|
2019-09-17 12:22:17 +01:00
|
|
|
|
2019-11-07 14:48:51 +00:00
|
|
|
hazard_query hq;
|
|
|
|
init_hazard_query(&hq);
|
|
|
|
add_to_hazard_query(&hq, current);
|
2019-09-17 12:22:17 +01:00
|
|
|
|
|
|
|
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;
|
2021-01-20 15:27:16 +00:00
|
|
|
if (candidate->isVMEM() || candidate->isSMEM() || candidate->isFlatLike())
|
2019-09-17 12:22:17 +01:00
|
|
|
break;
|
|
|
|
|
2020-06-26 15:54:22 +01:00
|
|
|
HazardResult haz = perform_hazard_query(&hq, candidate.get(), false);
|
2020-04-01 14:38:43 +01:00
|
|
|
if (haz == hazard_fail_exec || haz == hazard_fail_unreorderable)
|
2019-09-17 12:22:17 +01:00
|
|
|
break;
|
|
|
|
|
2019-11-07 14:48:51 +00:00
|
|
|
if (haz != hazard_success) {
|
|
|
|
add_to_hazard_query(&hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
|
|
|
}
|
|
|
|
|
2019-11-06 16:38:57 +00:00
|
|
|
MoveResult res = ctx.mv.downwards_move(false);
|
|
|
|
if (res == move_fail_ssa || res == move_fail_rar) {
|
2019-11-07 14:48:51 +00:00
|
|
|
add_to_hazard_query(&hq, candidate.get());
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.downwards_skip();
|
2019-09-17 12:22:17 +01:00
|
|
|
continue;
|
2019-11-06 16:38:57 +00:00
|
|
|
} else if (res == move_fail_pressure) {
|
2019-09-17 12:22:17 +01:00
|
|
|
break;
|
|
|
|
}
|
|
|
|
k++;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
void schedule_block(sched_ctx& ctx, Program *program, Block* block, live& live_vars)
|
|
|
|
{
|
|
|
|
ctx.last_SMEM_dep_idx = 0;
|
|
|
|
ctx.last_SMEM_stall = INT16_MIN;
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.block = block;
|
|
|
|
ctx.mv.register_demand = live_vars.register_demand[block->index].data();
|
2019-09-17 12:22:17 +01:00
|
|
|
|
|
|
|
/* go through all instructions and find memory loads */
|
|
|
|
for (unsigned idx = 0; idx < block->instructions.size(); idx++) {
|
|
|
|
Instruction* current = block->instructions[idx].get();
|
|
|
|
|
2021-01-20 15:27:16 +00:00
|
|
|
if (block->kind & block_kind_export_end && current->isEXP()) {
|
2021-01-21 16:13:34 +00:00
|
|
|
unsigned target = current->exp().dest;
|
2020-08-13 15:54:57 +01:00
|
|
|
if (target >= V_008DFC_SQ_EXP_POS && target < V_008DFC_SQ_EXP_PRIM) {
|
|
|
|
ctx.mv.current = current;
|
|
|
|
schedule_position_export(ctx, block, live_vars.register_demand[block->index], current, idx);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
if (current->definitions.empty())
|
|
|
|
continue;
|
|
|
|
|
2021-01-20 15:27:16 +00:00
|
|
|
if (current->isVMEM() || current->isFlatLike()) {
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.current = current;
|
2019-09-17 12:22:17 +01:00
|
|
|
schedule_VMEM(ctx, block, live_vars.register_demand[block->index], current, idx);
|
2019-11-06 16:38:57 +00:00
|
|
|
}
|
|
|
|
|
2021-01-20 15:27:16 +00:00
|
|
|
if (current->isSMEM()) {
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.current = current;
|
2019-09-17 12:22:17 +01:00
|
|
|
schedule_SMEM(ctx, block, live_vars.register_demand[block->index], current, idx);
|
2019-11-06 16:38:57 +00:00
|
|
|
}
|
2019-09-17 12:22:17 +01:00
|
|
|
}
|
|
|
|
|
|
|
|
/* resummarize the block's register demand */
|
|
|
|
block->register_demand = RegisterDemand();
|
|
|
|
for (unsigned idx = 0; idx < block->instructions.size(); idx++) {
|
|
|
|
block->register_demand.update(live_vars.register_demand[block->index][idx]);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
|
|
|
|
void schedule_program(Program *program, live& live_vars)
|
|
|
|
{
|
aco: fix max_waves_per_simd on Polaris, VegaM and GFX10.3
fossil-db (Polaris):
Totals from 20263 (14.75% of 137414) affected shaders:
SGPRs: 871407 -> 871679 (+0.03%); split: -0.00%, +0.03%
VGPRs: 513828 -> 550028 (+7.05%); split: -1.68%, +8.72%
CodeSize: 18869680 -> 18828148 (-0.22%); split: -0.23%, +0.01%
MaxWaves: 162012 -> 162030 (+0.01%); split: +0.01%, -0.00%
Instrs: 3629172 -> 3618817 (-0.29%); split: -0.30%, +0.02%
Cycles: 15682244 -> 15638244 (-0.28%); split: -0.30%, +0.02%
VMEM: 10675942 -> 10673344 (-0.02%); split: +0.18%, -0.21%
SMEM: 1209717 -> 1206088 (-0.30%); split: +0.03%, -0.33%
VClause: 81780 -> 81227 (-0.68%); split: -0.73%, +0.06%
SClause: 231724 -> 231561 (-0.07%); split: -0.07%, +0.00%
Copies: 187126 -> 180831 (-3.36%); split: -3.62%, +0.26%
Branches: 26841 -> 26837 (-0.01%); split: -0.03%, +0.01%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5546>
2020-06-18 14:30:51 +01:00
|
|
|
/* 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);
|
2021-02-05 13:38:08 +00:00
|
|
|
demand.vgpr += program->config->num_shared_vgprs / 2;
|
aco: fix max_waves_per_simd on Polaris, VegaM and GFX10.3
fossil-db (Polaris):
Totals from 20263 (14.75% of 137414) affected shaders:
SGPRs: 871407 -> 871679 (+0.03%); split: -0.00%, +0.03%
VGPRs: 513828 -> 550028 (+7.05%); split: -1.68%, +8.72%
CodeSize: 18869680 -> 18828148 (-0.22%); split: -0.23%, +0.01%
MaxWaves: 162012 -> 162030 (+0.01%); split: +0.01%, -0.00%
Instrs: 3629172 -> 3618817 (-0.29%); split: -0.30%, +0.02%
Cycles: 15682244 -> 15638244 (-0.28%); split: -0.30%, +0.02%
VMEM: 10675942 -> 10673344 (-0.02%); split: +0.18%, -0.21%
SMEM: 1209717 -> 1206088 (-0.30%); split: +0.03%, -0.33%
VClause: 81780 -> 81227 (-0.68%); split: -0.73%, +0.06%
SClause: 231724 -> 231561 (-0.07%); split: -0.07%, +0.00%
Copies: 187126 -> 180831 (-3.36%); split: -3.62%, +0.26%
Branches: 26841 -> 26837 (-0.01%); split: -0.03%, +0.01%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5546>
2020-06-18 14:30:51 +01:00
|
|
|
|
2019-09-17 12:22:17 +01:00
|
|
|
sched_ctx ctx;
|
2019-11-06 16:38:57 +00:00
|
|
|
ctx.mv.depends_on.resize(program->peekAllocationId());
|
|
|
|
ctx.mv.RAR_dependencies.resize(program->peekAllocationId());
|
|
|
|
ctx.mv.RAR_dependencies_clause.resize(program->peekAllocationId());
|
2019-09-17 12:22:17 +01:00
|
|
|
/* 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. */
|
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
2021-01-19 11:37:52 +00:00
|
|
|
//TODO: account for possible uneven num_waves on GFX10+
|
2021-01-28 13:07:11 +00:00
|
|
|
unsigned wave_fac = program->dev.physical_vgprs / 256;
|
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
2021-01-19 11:37:52 +00:00
|
|
|
if (program->num_waves <= 5 * wave_fac)
|
2019-08-29 16:17:32 +01:00
|
|
|
ctx.num_waves = program->num_waves;
|
aco: fix max_waves_per_simd on Polaris, VegaM and GFX10.3
fossil-db (Polaris):
Totals from 20263 (14.75% of 137414) affected shaders:
SGPRs: 871407 -> 871679 (+0.03%); split: -0.00%, +0.03%
VGPRs: 513828 -> 550028 (+7.05%); split: -1.68%, +8.72%
CodeSize: 18869680 -> 18828148 (-0.22%); split: -0.23%, +0.01%
MaxWaves: 162012 -> 162030 (+0.01%); split: +0.01%, -0.00%
Instrs: 3629172 -> 3618817 (-0.29%); split: -0.30%, +0.02%
Cycles: 15682244 -> 15638244 (-0.28%); split: -0.30%, +0.02%
VMEM: 10675942 -> 10673344 (-0.02%); split: +0.18%, -0.21%
SMEM: 1209717 -> 1206088 (-0.30%); split: +0.03%, -0.33%
VClause: 81780 -> 81227 (-0.68%); split: -0.73%, +0.06%
SClause: 231724 -> 231561 (-0.07%); split: -0.07%, +0.00%
Copies: 187126 -> 180831 (-3.36%); split: -3.62%, +0.26%
Branches: 26841 -> 26837 (-0.01%); split: -0.03%, +0.01%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5546>
2020-06-18 14:30:51 +01:00
|
|
|
else if (demand.vgpr >= 29)
|
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
2021-01-19 11:37:52 +00:00
|
|
|
ctx.num_waves = 5 * wave_fac;
|
aco: fix max_waves_per_simd on Polaris, VegaM and GFX10.3
fossil-db (Polaris):
Totals from 20263 (14.75% of 137414) affected shaders:
SGPRs: 871407 -> 871679 (+0.03%); split: -0.00%, +0.03%
VGPRs: 513828 -> 550028 (+7.05%); split: -1.68%, +8.72%
CodeSize: 18869680 -> 18828148 (-0.22%); split: -0.23%, +0.01%
MaxWaves: 162012 -> 162030 (+0.01%); split: +0.01%, -0.00%
Instrs: 3629172 -> 3618817 (-0.29%); split: -0.30%, +0.02%
Cycles: 15682244 -> 15638244 (-0.28%); split: -0.30%, +0.02%
VMEM: 10675942 -> 10673344 (-0.02%); split: +0.18%, -0.21%
SMEM: 1209717 -> 1206088 (-0.30%); split: +0.03%, -0.33%
VClause: 81780 -> 81227 (-0.68%); split: -0.73%, +0.06%
SClause: 231724 -> 231561 (-0.07%); split: -0.07%, +0.00%
Copies: 187126 -> 180831 (-3.36%); split: -3.62%, +0.26%
Branches: 26841 -> 26837 (-0.01%); split: -0.03%, +0.01%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/5546>
2020-06-18 14:30:51 +01:00
|
|
|
else if (demand.vgpr >= 25)
|
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
2021-01-19 11:37:52 +00:00
|
|
|
ctx.num_waves = 6 * wave_fac;
|
2019-08-29 16:17:32 +01:00
|
|
|
else
|
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
2021-01-19 11:37:52 +00:00
|
|
|
ctx.num_waves = 7 * wave_fac;
|
2019-12-18 16:18:35 +00:00
|
|
|
ctx.num_waves = std::max<uint16_t>(ctx.num_waves, program->min_waves);
|
2020-12-10 17:54:43 +00:00
|
|
|
ctx.num_waves = std::min<uint16_t>(ctx.num_waves, program->num_waves);
|
2019-08-29 16:17:32 +01:00
|
|
|
|
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
2021-01-19 11:37:52 +00:00
|
|
|
/* VMEM_MAX_MOVES and such assume pre-GFX10 wave count */
|
|
|
|
ctx.num_waves = std::max<uint16_t>(ctx.num_waves / wave_fac, 1);
|
|
|
|
|
2020-12-10 17:54:43 +00:00
|
|
|
assert(ctx.num_waves > 0);
|
aco: fix num_waves on GFX10+
There are half the SIMDs per CU and physical_vgprs should be 512 instead
of 256.
fossil-db (GFX10.3):
Totals from 3622 (2.60% of 139391) affected shaders:
VGPRs: 298192 -> 289732 (-2.84%); split: -3.43%, +0.59%
CodeSize: 29443432 -> 29458388 (+0.05%); split: -0.00%, +0.06%
MaxWaves: 21703 -> 23395 (+7.80%); split: +7.84%, -0.05%
Instrs: 5677920 -> 5681438 (+0.06%); split: -0.01%, +0.07%
Cycles: 280715524 -> 280895676 (+0.06%); split: -0.00%, +0.07%
VMEM: 981142 -> 981894 (+0.08%); split: +0.18%, -0.10%
SMEM: 243315 -> 243454 (+0.06%); split: +0.07%, -0.02%
VClause: 88991 -> 89767 (+0.87%); split: -0.02%, +0.89%
SClause: 200660 -> 200659 (-0.00%); split: -0.00%, +0.00%
Copies: 430729 -> 434160 (+0.80%); split: -0.07%, +0.86%
Branches: 158004 -> 158021 (+0.01%); split: -0.01%, +0.02%
Signed-off-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/8523>
2021-01-19 11:37:52 +00:00
|
|
|
ctx.mv.max_registers = { int16_t(get_addr_vgpr_from_waves(program, ctx.num_waves * wave_fac) - 2),
|
|
|
|
int16_t(get_addr_sgpr_from_waves(program, ctx.num_waves * wave_fac))};
|
2019-09-17 12:22:17 +01:00
|
|
|
|
|
|
|
for (Block& block : program->blocks)
|
|
|
|
schedule_block(ctx, program, &block, live_vars);
|
|
|
|
|
|
|
|
/* 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);
|
|
|
|
|
|
|
|
/* if enabled, this code asserts that register_demand is updated correctly */
|
|
|
|
#if 0
|
|
|
|
int prev_num_waves = program->num_waves;
|
|
|
|
const RegisterDemand prev_max_demand = program->max_reg_demand;
|
|
|
|
|
|
|
|
std::vector<RegisterDemand> demands(program->blocks.size());
|
|
|
|
for (unsigned j = 0; j < program->blocks.size(); j++) {
|
|
|
|
demands[j] = program->blocks[j].register_demand;
|
|
|
|
}
|
|
|
|
|
2020-10-08 09:12:58 +01:00
|
|
|
live live_vars2 = aco::live_var_analysis(program);
|
2019-09-17 12:22:17 +01:00
|
|
|
|
|
|
|
for (unsigned j = 0; j < program->blocks.size(); j++) {
|
|
|
|
Block &b = program->blocks[j];
|
|
|
|
for (unsigned i = 0; i < b.instructions.size(); i++)
|
|
|
|
assert(live_vars.register_demand[b.index][i] == live_vars2.register_demand[b.index][i]);
|
|
|
|
assert(b.register_demand == demands[j]);
|
|
|
|
}
|
|
|
|
|
|
|
|
assert(program->max_reg_demand == prev_max_demand);
|
|
|
|
assert(program->num_waves == prev_num_waves);
|
|
|
|
#endif
|
|
|
|
}
|
|
|
|
|
|
|
|
}
|