fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-01-04 02:40:11 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

aco: Implement control_barrier for tessellation control shaders.

Browse source 
Signed-off-by: Timur Kristóf <timur.kristof@gmail.com>
Reviewed-by: Rhys Perry <pendingchaos02@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/3964>
This commit is contained in:
Timur Kristóf 2020-02-26 17:41:04 +01:00 • committed by Marge Bot
parent 2489e4dfd1
commit a8d15ab6da
1 changed files with 24 additions and 3 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

27
src/amd/compiler/aco_instruction_selection.cpp
View file

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