diff options
| author | Timur Kristóf <[email protected]> | 2020-02-26 17:41:04 +0100 |
|---|---|---|
| committer | Marge Bot <[email protected]> | 2020-03-11 08:34:10 +0000 |
| commit | a8d15ab6daf0a07476e9dfabe513c0f1e0f3bf82 (patch) | |
| tree | a5de53b5e58f7acb436704b285065e421f8a4184 /src | |
| parent | 2489e4dfd183919028d5a346c2dffc6138c7269f (diff) | |
aco: Implement control_barrier for tessellation control shaders.
Signed-off-by: Timur Kristóf <[email protected]>
Reviewed-by: Rhys Perry <[email protected]>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3964>
Diffstat (limited to 'src')
| -rw-r--r-- | src/amd/compiler/aco_instruction_selection.cpp | 27 |
1 files changed, 24 insertions, 3 deletions
diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index 0be69aa47e8..069f3d24f62 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -6315,10 +6315,31 @@ void visit_intrinsic(isel_context *ctx, nir_intrinsic_instr *instr) visit_get_buffer_size(ctx, instr); break; case nir_intrinsic_control_barrier: { - unsigned* bsize = ctx->program->info->cs.block_size; - unsigned workgroup_size = bsize[0] * bsize[1] * bsize[2]; - if (workgroup_size > ctx->program->wave_size) + if (ctx->program->chip_class == GFX6 && ctx->shader->info.stage == MESA_SHADER_TESS_CTRL) { + /* GFX6 only (thanks to a hw bug workaround): + * The real barrier instruction isn’t needed, because an entire patch + * always fits into a single wave. + */ + break; + } + + if (ctx->shader->info.stage == MESA_SHADER_COMPUTE) { + unsigned* bsize = ctx->program->info->cs.block_size; + unsigned workgroup_size = bsize[0] * bsize[1] * bsize[2]; + if (workgroup_size > ctx->program->wave_size) + bld.sopp(aco_opcode::s_barrier); + } else if (ctx->shader->info.stage == MESA_SHADER_TESS_CTRL) { + /* For each patch provided during rendering, n TCS shader invocations will be processed, + * where n is the number of vertices in the output patch. + */ + unsigned workgroup_size = ctx->tcs_num_patches * ctx->shader->info.tess.tcs_vertices_out; + if (workgroup_size > ctx->program->wave_size) + bld.sopp(aco_opcode::s_barrier); + } else { + /* We don't know the workgroup size, so always emit the s_barrier. */ bld.sopp(aco_opcode::s_barrier); + } + break; } case nir_intrinsic_group_memory_barrier: |