From: Timur Kristóf Date: Fri, 22 Nov 2019 16:07:34 +0000 (+0100) Subject: aco/wave32: Use wave_size for barrier intrinsic. X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=ed815d503efadc2ad0171b5ad296f2a84bac528b;p=mesa.git aco/wave32: Use wave_size for barrier intrinsic. Signed-off-by: Timur Kristóf Reviewed-by: Rhys Perry Reviewed-by: Daniel Schürmann --- 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; }