From 57d8799284fecbda30374206da31e87ae153056a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Timur=20Krist=C3=B3f?= Date: Tue, 29 Sep 2020 18:06:44 +0200 Subject: [PATCH] aco: Optimize thread_id_in_threadgroup when there is just one wave. MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Timur Kristóf Reviewed-by: Rhys Perry Part-of: --- src/amd/compiler/aco_instruction_selection.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index 0f3cc7d9b06..d6071e065b4 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -3883,9 +3883,12 @@ Temp thread_id_in_threadgroup(isel_context *ctx) /* tid_in_tg = wave_id * wave_size + tid_in_wave */ Builder bld(ctx->program, ctx->block); + Temp tid_in_wave = emit_mbcnt(ctx, bld.tmp(v1)); + + if (ctx->program->workgroup_size <= ctx->program->wave_size) + return tid_in_wave; Temp wave_id_in_tg = wave_id_in_threadgroup(ctx); - Temp tid_in_wave = emit_mbcnt(ctx, bld.tmp(v1)); Temp num_pre_threads = bld.sop2(aco_opcode::s_lshl_b32, bld.def(s1), bld.def(s1, scc), wave_id_in_tg, Operand(ctx->program->wave_size == 64 ? 6u : 5u)); return bld.vadd32(bld.def(v1), Operand(num_pre_threads), Operand(tid_in_wave));