From 1643c933efaee9908ba3a36694b0088069a8163c Mon Sep 17 00:00:00 2001 From: Rhys Perry Date: Wed, 26 Jun 2024 14:53:07 +0100 Subject: [PATCH] aco/gfx11: don't use v_bfrev_b32 with wave64 v_mov_b32 can be dual issued. fossil-db (navi31): Totals from 1792 (2.26% of 79395) affected shaders: CodeSize: 27462476 -> 27470308 (+0.03%); split: -0.00%, +0.03% Latency: 29403214 -> 29402713 (-0.00%); split: -0.00%, +0.00% InvThroughput: 5005863 -> 5004702 (-0.02%); split: -0.02%, +0.00% Signed-off-by: Rhys Perry Reviewed-by: Georg Lehmann Part-of: --- src/amd/compiler/aco_lower_to_hw_instr.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/amd/compiler/aco_lower_to_hw_instr.cpp b/src/amd/compiler/aco_lower_to_hw_instr.cpp index ad98f52cfcb..566361cfb5d 100644 --- a/src/amd/compiler/aco_lower_to_hw_instr.cpp +++ b/src/amd/compiler/aco_lower_to_hw_instr.cpp @@ -1275,7 +1275,9 @@ copy_constant(lower_context* ctx, Builder& bld, Definition dst, Operand op) if (dst.regClass().type() == RegType::sgpr) return copy_constant_sgpr(bld, dst, op.constantValue64()); - if (dst.bytes() == 4 && op.isLiteral()) { + bool dual_issue_mov = ctx->program->gfx_level >= GFX11 && ctx->program->wave_size == 64 && + ctx->program->workgroup_size > 32; + if (dst.bytes() == 4 && op.isLiteral() && !dual_issue_mov) { uint32_t imm = op.constantValue(); Operand rev_op = Operand::get_const(ctx->program->gfx_level, util_bitreverse(imm), 4); if (!rev_op.isLiteral()) {