diff options
| author | Timur Kristóf <timur.kristof@gmail.com> | 2019-11-22 17:07:34 +0100 |
|---|---|---|
| committer | Daniel Schürmann <daniel@schuermann.dev> | 2019-12-04 10:36:01 +0000 |
| commit | ed815d503efadc2ad0171b5ad296f2a84bac528b (patch) | |
| tree | 4d16d09f2f69b670ea65383ebff208bf7fa04ba9 | |
| parent | b8f2edb452e064422e3f32b89aae15dfd879399d (diff) | |
| download | mesa-ed815d503efadc2ad0171b5ad296f2a84bac528b.tar.gz | |
aco/wave32: Use wave_size for barrier intrinsic.
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
| -rw-r--r-- | src/amd/compiler/aco_insert_waitcnt.cpp | 4 | ||||
| -rw-r--r-- | 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; } |