KonstantinSeurer
/
mesa
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Timur Kristóf 2019-11-22 17:07:34 +01:00 committed by Daniel Schürmann
parent b8f2edb452
commit ed815d503e
2 changed files with 3 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -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:

2
src/amd/compiler/aco_instruction_selection.cpp
View File

@ -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;
}