OpenHarmony/third_party_mesa3d
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

aco/gfx11: don't use v_bfrev_b32 with wave64

Browse Source 
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 <pendingchaos02@gmail.com>
Reviewed-by: Georg Lehmann <dadschoorse@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/29912>
This commit is contained in:
Rhys Perry
2024-06-26 14:53:07 +01:00
committed by Marge Bot
parent 52e9370c13
commit 1643c933ef
1 changed files with 3 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
src/amd/compiler/aco_lower_to_hw_instr.cpp
View File

@@ -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()) {