diff --git a/src/amd/compiler/aco_interface.cpp b/src/amd/compiler/aco_interface.cpp index 03544424c2e..57271308f32 100644 --- a/src/amd/compiler/aco_interface.cpp +++ b/src/amd/compiler/aco_interface.cpp @@ -34,32 +34,31 @@ #include #include -static const std::array statistic_infos = []() +static const std::array statistic_infos = []() { - std::array ret{}; - ret[aco::statistic_hash] = + std::array 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, diff --git a/src/amd/compiler/aco_interface.h b/src/amd/compiler/aco_interface.h index c404038c42d..7bef7727cf5 100644 --- a/src/amd/compiler/aco_interface.h +++ b/src/amd/compiler/aco_interface.h @@ -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, diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h index 9a10c67a49c..393527574a9 100644 --- a/src/amd/compiler/aco_ir.h +++ b/src/amd/compiler/aco_ir.h @@ -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; diff --git a/src/amd/compiler/aco_lower_to_hw_instr.cpp b/src/amd/compiler/aco_lower_to_hw_instr.cpp index 7984237c154..ac18c044589 100644 --- a/src/amd/compiler/aco_lower_to_hw_instr.cpp +++ b/src/amd/compiler/aco_lower_to_hw_instr.cpp @@ -1653,7 +1653,7 @@ handle_operands(std::map& 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& 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 diff --git a/src/amd/compiler/aco_shader_info.h b/src/amd/compiler/aco_shader_info.h index b09e25c1a8d..8fba17aedc9 100644 --- a/src/amd/compiler/aco_shader_info.h +++ b/src/amd/compiler/aco_shader_info.h @@ -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 diff --git a/src/amd/compiler/aco_statistics.cpp b/src/amd/compiler/aco_statistics.cpp index 7f150033dc7..ae790b10d25 100644 --- a/src/amd/compiler/aco_statistics.cpp +++ b/src/amd/compiler/aco_statistics.cpp @@ -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 vmem_clause; std::set smem_clause; - program->statistics[statistic_instructions] += block.instructions.size(); + program->statistics[aco_statistic_instructions] += block.instructions.size(); for (aco_ptr& 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& 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