From: Marek Olšák Date: Fri, 12 Jul 2019 21:22:30 +0000 (-0400) Subject: radeonsi: set threadgroup size to 0 for threadgroups with only 1 wave X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=e2c8ff009e5e077ae93ad19c63bd06d5c23d8376;p=mesa.git 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 --- 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;