From ed815d503efadc2ad0171b5ad296f2a84bac528b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Timur=20Krist=C3=B3f?= Date: Fri, 22 Nov 2019 17:07:34 +0100 Subject: [PATCH] aco/wave32: Use wave_size for barrier intrinsic. 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 Reviewed-by: Daniel Schürmann --- src/amd/compiler/aco_insert_waitcnt.cpp | 4 ++-- src/amd/compiler/aco_instruction_selection.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/src/amd/compiler/aco_insert_waitcnt.cpp b/src/amd/compiler/aco_insert_waitcnt.cpp index 7960902d690..a8343d18894 100644 --- a/src/amd/compiler/aco_insert_waitcnt.cpp +++ b/src/amd/compiler/aco_insert_waitcnt.cpp @@ -399,7 +399,7 @@ wait_imm kill(Instruction* instr, wait_ctx& ctx) switch (instr->opcode) { case aco_opcode::p_memory_barrier_all: for (unsigned i = 0; i < barrier_count; i++) { - if ((1 << i) == barrier_shared && workgroup_size <= 64) + if ((1 << i) == barrier_shared && workgroup_size <= ctx.program->wave_size) continue; imm.combine(ctx.barrier_imm[i]); } @@ -414,7 +414,7 @@ wait_imm kill(Instruction* instr, wait_ctx& ctx) imm.combine(ctx.barrier_imm[ffs(barrier_image) - 1]); break; case aco_opcode::p_memory_barrier_shared: - if (workgroup_size > 64) + if (workgroup_size > ctx.program->wave_size) imm.combine(ctx.barrier_imm[ffs(barrier_shared) - 1]); break; default: diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index 3d061bbb448..70cee225670 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -5631,7 +5631,7 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr) case nir_intrinsic_barrier: { unsigned* bsize = ctx->program->info->cs.block_size; unsigned workgroup_size = bsize[0] * bsize[1] * bsize[2]; - if (workgroup_size > 64) + if (workgroup_size > ctx->program->wave_size) bld.sopp(aco_opcode::s_barrier); break; }