From 91fe2a23617de09efe91038d7f836c6a051553a2 Mon Sep 17 00:00:00 2001 From: Bas Nieuwenhuizen Date: Mon, 7 Nov 2022 00:04:38 +0100 Subject: [PATCH] aco: Use more detailed wave64 timing for GFX10+. Also nabbed some dual issue stuff for GFX11 from LLVM. Part-of: --- src/amd/compiler/aco_statistics.cpp | 30 ++++++++++++++++++++++++++++- 1 file changed, 29 insertions(+), 1 deletion(-) diff --git a/src/amd/compiler/aco_statistics.cpp b/src/amd/compiler/aco_statistics.cpp index ae790b10d25..00d0edfe52c 100644 --- a/src/amd/compiler/aco_statistics.cpp +++ b/src/amd/compiler/aco_statistics.cpp @@ -104,6 +104,33 @@ struct perf_info { unsigned cost1; }; +static bool +is_dual_issue_capable(const Program& program, const Instruction& instruction) +{ + if (program.gfx_level < GFX11 || !instruction.isVALU()) + return false; + + /* Currently assumed to be just the instructions that are allowed as both + * VOPD X and VOPD Y operation. + */ + switch (instruction.opcode) { + case aco_opcode::v_fmac_f32: + case aco_opcode::v_fmaak_f32: + case aco_opcode::v_fmamk_f32: + case aco_opcode::v_mul_f32: + case aco_opcode::v_add_f32: + case aco_opcode::v_sub_f32: + case aco_opcode::v_subrev_f32: + case aco_opcode::v_mul_legacy_f32: + case aco_opcode::v_mov_b32: + case aco_opcode::v_cndmask_b32: + case aco_opcode::v_max_f32: + case aco_opcode::v_min_f32: + case aco_opcode::v_dot2c_f32_f16: return true; + default: return false; + } +} + static perf_info get_perf_info(Program* program, aco_ptr& instr) { @@ -358,7 +385,8 @@ BlockCycleEstimator::add(aco_ptr& instr) unsigned start; bool dual_issue = program->gfx_level >= GFX10 && program->wave_size == 64 && - is_vector(instr->opcode) && program->workgroup_size > 32; + is_vector(instr->opcode) && !is_dual_issue_capable(*program, *instr) && + program->workgroup_size > 32; for (unsigned i = 0; i < (dual_issue ? 2 : 1); i++) { cur_cycle += cycles_until_res_available(instr);