From e2c8ff009e5e077ae93ad19c63bd06d5c23d8376 Mon Sep 17 00:00:00 2001 From: =?utf8?q?Marek=20Ol=C5=A1=C3=A1k?= Date: Fri, 12 Jul 2019 17:22:30 -0400 Subject: [PATCH] radeonsi: set threadgroup size to 0 for threadgroups with only 1 wave This has no effect on Wave64. Reviewed-by: Pierre-Eric Pelloux-Prayer Acked-by: Samuel Pitoiset --- src/gallium/drivers/radeonsi/si_shader.c | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c index 3cfb1207b4a..8d06859f1f2 100644 --- a/src/gallium/drivers/radeonsi/si_shader.c +++ b/src/gallium/drivers/radeonsi/si_shader.c @@ -4484,10 +4484,10 @@ static unsigned si_get_max_workgroup_size(const struct si_shader *shader) case PIPE_SHADER_TESS_CTRL: /* Return this so that LLVM doesn't remove s_barrier * instructions on chips where we use s_barrier. */ - return shader->selector->screen->info.chip_class >= GFX7 ? 128 : 64; + return shader->selector->screen->info.chip_class >= GFX7 ? 128 : 0; case PIPE_SHADER_GEOMETRY: - return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 64; + return shader->selector->screen->info.chip_class >= GFX9 ? 128 : 0; case PIPE_SHADER_COMPUTE: break; /* see below */ @@ -7626,7 +7626,7 @@ static void si_build_tcs_epilog_function(struct si_shader_context *ctx, /* Create the function. */ si_create_function(ctx, "tcs_epilog", NULL, 0, &fninfo, - ctx->screen->info.chip_class >= GFX7 ? 128 : 64); + ctx->screen->info.chip_class >= GFX7 ? 128 : 0); ac_declare_lds_as_pointer(&ctx->ac); func = ctx->main_fn; -- 2.30.2