mirror of
https://gitlab.freedesktop.org/mesa/mesa.git
synced 2026-01-08 15:00:11 +01:00
aco: move statistics enum to aco_shader_info.h
to make it accessible from the driver. Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19721>
This commit is contained in:
Daniel Schürmann
Marge Bot
parent
4306897979
commit
efc0835787
6 changed files with 41 additions and 42 deletions
|
|
@ -34,32 +34,31 @@
|
|||
#include <iostream>
|
||||
#include <vector>
|
||||
|
||||
static const std::array<aco_compiler_statistic_info, aco::num_statistics> statistic_infos = []()
|
||||
static const std::array<aco_compiler_statistic_info, aco_num_statistics> statistic_infos = []()
|
||||
{
|
||||
std::array<aco_compiler_statistic_info, aco::num_statistics> ret{};
|
||||
ret[aco::statistic_hash] =
|
||||
std::array<aco_compiler_statistic_info, aco_num_statistics> ret{};
|
||||
ret[aco_statistic_hash] =
|
||||
aco_compiler_statistic_info{"Hash", "CRC32 hash of code and constant data"};
|
||||
ret[aco::statistic_instructions] =
|
||||
ret[aco_statistic_instructions] =
|
||||
aco_compiler_statistic_info{"Instructions", "Instruction count"};
|
||||
ret[aco::statistic_copies] =
|
||||
ret[aco_statistic_copies] =
|
||||
aco_compiler_statistic_info{"Copies", "Copy instructions created for pseudo-instructions"};
|
||||
ret[aco::statistic_branches] = aco_compiler_statistic_info{"Branches", "Branch instructions"};
|
||||
ret[aco::statistic_latency] =
|
||||
ret[aco_statistic_branches] = aco_compiler_statistic_info{"Branches", "Branch instructions"};
|
||||
ret[aco_statistic_latency] =
|
||||
aco_compiler_statistic_info{"Latency", "Issue cycles plus stall cycles"};
|
||||
ret[aco::statistic_inv_throughput] = aco_compiler_statistic_info{
|
||||
ret[aco_statistic_inv_throughput] = aco_compiler_statistic_info{
|
||||
"Inverse Throughput", "Estimated busy cycles to execute one wave"};
|
||||
ret[aco::statistic_vmem_clauses] = aco_compiler_statistic_info{
|
||||
ret[aco_statistic_vmem_clauses] = aco_compiler_statistic_info{
|
||||
"VMEM Clause", "Number of VMEM clauses (includes 1-sized clauses)"};
|
||||
ret[aco::statistic_smem_clauses] = aco_compiler_statistic_info{
|
||||
ret[aco_statistic_smem_clauses] = aco_compiler_statistic_info{
|
||||
"SMEM Clause", "Number of SMEM clauses (includes 1-sized clauses)"};
|
||||
ret[aco::statistic_sgpr_presched] =
|
||||
ret[aco_statistic_sgpr_presched] =
|
||||
aco_compiler_statistic_info{"Pre-Sched SGPRs", "SGPR usage before scheduling"};
|
||||
ret[aco::statistic_vgpr_presched] =
|
||||
ret[aco_statistic_vgpr_presched] =
|
||||
aco_compiler_statistic_info{"Pre-Sched VGPRs", "VGPR usage before scheduling"};
|
||||
return ret;
|
||||
}();
|
||||
|
||||
const unsigned aco_num_statistics = aco::num_statistics;
|
||||
const aco_compiler_statistic_info* aco_statistic_infos = statistic_infos.data();
|
||||
|
||||
uint64_t
|
||||
|
|
@ -255,7 +254,7 @@ aco_compile_shader(const struct aco_compiler_options* options,
|
|||
|
||||
size_t stats_size = 0;
|
||||
if (program->collect_statistics)
|
||||
stats_size = aco::num_statistics * sizeof(uint32_t);
|
||||
stats_size = aco_num_statistics * sizeof(uint32_t);
|
||||
|
||||
(*build_binary)(binary,
|
||||
shaders[shader_count - 1]->info.stage,
|
||||
|
|
|
|||
|
|
@ -67,7 +67,6 @@ typedef void (aco_shader_part_callback)(void **priv_ptr,
|
|||
const char *disasm_str,
|
||||
uint32_t disasm_size);
|
||||
|
||||
extern const unsigned aco_num_statistics;
|
||||
extern const struct aco_compiler_statistic_info* aco_statistic_infos;
|
||||
|
||||
void aco_compile_shader(const struct aco_compiler_options* options,
|
||||
|
|
|
|||
|
|
@ -2110,20 +2110,6 @@ static constexpr Stage tess_eval_es(HWStage::ES,
|
|||
SWStage::TES); /* tesselation evaluation before geometry */
|
||||
static constexpr Stage geometry_gs(HWStage::GS, SWStage::GS);
|
||||
|
||||
enum statistic {
|
||||
statistic_hash,
|
||||
statistic_instructions,
|
||||
statistic_copies,
|
||||
statistic_branches,
|
||||
statistic_latency,
|
||||
statistic_inv_throughput,
|
||||
statistic_vmem_clauses,
|
||||
statistic_smem_clauses,
|
||||
statistic_sgpr_presched,
|
||||
statistic_vgpr_presched,
|
||||
num_statistics
|
||||
};
|
||||
|
||||
struct DeviceInfo {
|
||||
uint16_t lds_encoding_granule;
|
||||
uint16_t lds_alloc_granule;
|
||||
|
|
@ -2185,7 +2171,7 @@ public:
|
|||
CompilationProgress progress;
|
||||
|
||||
bool collect_statistics = false;
|
||||
uint32_t statistics[num_statistics];
|
||||
uint32_t statistics[aco_num_statistics];
|
||||
|
||||
float_mode next_fp_mode;
|
||||
unsigned next_loop_depth = 0;
|
||||
|
|
|
|||
|
|
@ -1653,7 +1653,7 @@ handle_operands(std::map<PhysReg, copy_operation>& copy_map, lower_context* ctx,
|
|||
bool skip_partial_copies = true;
|
||||
for (auto it = copy_map.begin();;) {
|
||||
if (copy_map.empty()) {
|
||||
ctx->program->statistics[statistic_copies] +=
|
||||
ctx->program->statistics[aco_statistic_copies] +=
|
||||
ctx->instructions.size() - num_instructions_before;
|
||||
return;
|
||||
}
|
||||
|
|
@ -1961,7 +1961,8 @@ handle_operands(std::map<PhysReg, copy_operation>& copy_map, lower_context* ctx,
|
|||
break;
|
||||
}
|
||||
}
|
||||
ctx->program->statistics[statistic_copies] += ctx->instructions.size() - num_instructions_before;
|
||||
ctx->program->statistics[aco_statistic_copies] +=
|
||||
ctx->instructions.size() - num_instructions_before;
|
||||
}
|
||||
|
||||
void
|
||||
|
|
|
|||
|
|
@ -194,6 +194,20 @@ struct aco_compiler_options {
|
|||
} debug;
|
||||
};
|
||||
|
||||
enum aco_statistic {
|
||||
aco_statistic_hash,
|
||||
aco_statistic_instructions,
|
||||
aco_statistic_copies,
|
||||
aco_statistic_branches,
|
||||
aco_statistic_latency,
|
||||
aco_statistic_inv_throughput,
|
||||
aco_statistic_vmem_clauses,
|
||||
aco_statistic_smem_clauses,
|
||||
aco_statistic_sgpr_presched,
|
||||
aco_statistic_vgpr_presched,
|
||||
aco_num_statistics
|
||||
};
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
|
|
|
|||
|
|
@ -40,8 +40,8 @@ collect_presched_stats(Program* program)
|
|||
RegisterDemand presched_demand;
|
||||
for (Block& block : program->blocks)
|
||||
presched_demand.update(block.register_demand);
|
||||
program->statistics[statistic_sgpr_presched] = presched_demand.sgpr;
|
||||
program->statistics[statistic_vgpr_presched] = presched_demand.vgpr;
|
||||
program->statistics[aco_statistic_sgpr_presched] = presched_demand.sgpr;
|
||||
program->statistics[aco_statistic_vgpr_presched] = presched_demand.vgpr;
|
||||
}
|
||||
|
||||
class BlockCycleEstimator {
|
||||
|
|
@ -438,21 +438,21 @@ collect_preasm_stats(Program* program)
|
|||
std::set<Instruction*> vmem_clause;
|
||||
std::set<Instruction*> smem_clause;
|
||||
|
||||
program->statistics[statistic_instructions] += block.instructions.size();
|
||||
program->statistics[aco_statistic_instructions] += block.instructions.size();
|
||||
|
||||
for (aco_ptr<Instruction>& instr : block.instructions) {
|
||||
if (instr->isSOPP() && instr->sopp().block != -1)
|
||||
program->statistics[statistic_branches]++;
|
||||
program->statistics[aco_statistic_branches]++;
|
||||
|
||||
if (instr->opcode == aco_opcode::p_constaddr)
|
||||
program->statistics[statistic_instructions] += 2;
|
||||
program->statistics[aco_statistic_instructions] += 2;
|
||||
|
||||
if ((instr->isVMEM() || instr->isScratch() || instr->isGlobal()) &&
|
||||
!instr->operands.empty()) {
|
||||
if (std::none_of(vmem_clause.begin(), vmem_clause.end(),
|
||||
[&](Instruction* other)
|
||||
{ return should_form_clause(instr.get(), other); }))
|
||||
program->statistics[statistic_vmem_clauses]++;
|
||||
program->statistics[aco_statistic_vmem_clauses]++;
|
||||
vmem_clause.insert(instr.get());
|
||||
} else {
|
||||
vmem_clause.clear();
|
||||
|
|
@ -462,7 +462,7 @@ collect_preasm_stats(Program* program)
|
|||
if (std::none_of(smem_clause.begin(), smem_clause.end(),
|
||||
[&](Instruction* other)
|
||||
{ return should_form_clause(instr.get(), other); }))
|
||||
program->statistics[statistic_smem_clauses]++;
|
||||
program->statistics[aco_statistic_smem_clauses]++;
|
||||
smem_clause.insert(instr.get());
|
||||
} else {
|
||||
smem_clause.clear();
|
||||
|
|
@ -545,8 +545,8 @@ collect_preasm_stats(Program* program)
|
|||
program->workgroup_size / (double)align(program->workgroup_size, program->wave_size);
|
||||
wave64_per_cycle *= max_utilization;
|
||||
|
||||
program->statistics[statistic_latency] = round(latency);
|
||||
program->statistics[statistic_inv_throughput] = round(1.0 / wave64_per_cycle);
|
||||
program->statistics[aco_statistic_latency] = round(latency);
|
||||
program->statistics[aco_statistic_inv_throughput] = round(1.0 / wave64_per_cycle);
|
||||
|
||||
if (debug_flags & DEBUG_PERF_INFO) {
|
||||
aco_print_program(program, stderr, print_no_ssa | print_perf_info);
|
||||
|
|
@ -571,7 +571,7 @@ collect_preasm_stats(Program* program)
|
|||
void
|
||||
collect_postasm_stats(Program* program, const std::vector<uint32_t>& code)
|
||||
{
|
||||
program->statistics[aco::statistic_hash] = util_hash_crc32(code.data(), code.size() * 4);
|
||||
program->statistics[aco_statistic_hash] = util_hash_crc32(code.data(), code.size() * 4);
|
||||
}
|
||||
|
||||
} // namespace aco
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue