radeonsi: set threadgroup size to 0 for threadgroups with only 1 wave

This has no effect on Wave64.

Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Acked-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>

diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index 3cfb1207b4a94ef7bdb1504c319977d6a335585f..8d06859f1f2d262103ec7966e0fb6a51f4ad6f8d 100644 (file)
--- 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;