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

[mesa.git] / src / gallium / drivers / radeonsi / si_shader.c

diff --git a/src/gallium/drivers/radeonsi/si_shader.c b/src/gallium/drivers/radeonsi/si_shader.c
index a3578ebf720efda6f44d437108b56e1a73d10421..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 */
@@ -7023,6 +7023,7 @@ int si_compile_tgsi_shader(struct si_screen *sscreen,
                        memset(&gs_prolog_key, 0, sizeof(gs_prolog_key));
                        gs_prolog_key.gs_prolog.states = shader->key.part.gs.prolog;
                        gs_prolog_key.gs_prolog.is_monolithic = true;
+                       gs_prolog_key.gs_prolog.as_ngg = shader->key.as_ngg;
                        si_build_gs_prolog_function(&ctx, &gs_prolog_key);
                        gs_prolog = ctx.main_fn;
 
@@ -7260,11 +7261,6 @@ si_get_shader_part(struct si_screen *sscreen,
        result->key = *key;
 
        struct si_shader shader = {};
-       struct si_shader_context ctx;
-
-       si_init_shader_ctx(&ctx, sscreen, compiler);
-       ctx.shader = &shader;
-       ctx.type = type;
 
        switch (type) {
        case PIPE_SHADER_VERTEX:
@@ -7278,6 +7274,7 @@ si_get_shader_part(struct si_screen *sscreen,
                break;
        case PIPE_SHADER_GEOMETRY:
                assert(prolog);
+               shader.key.as_ngg = key->gs_prolog.as_ngg;
                break;
        case PIPE_SHADER_FRAGMENT:
                if (prolog)
@@ -7289,6 +7286,11 @@ si_get_shader_part(struct si_screen *sscreen,
                unreachable("bad shader part");
        }
 
+       struct si_shader_context ctx;
+       si_init_shader_ctx(&ctx, sscreen, compiler);
+       ctx.shader = &shader;
+       ctx.type = type;
+
        build(&ctx, key);
 
        /* Compile. */
@@ -7624,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;
 
@@ -7699,6 +7701,7 @@ static bool si_shader_select_gs_parts(struct si_screen *sscreen,
        union si_shader_part_key prolog_key;
        memset(&prolog_key, 0, sizeof(prolog_key));
        prolog_key.gs_prolog.states = shader->key.part.gs.prolog;
+       prolog_key.gs_prolog.as_ngg = shader->key.as_ngg;
 
        shader->prolog2 = si_get_shader_part(sscreen, &sscreen->gs_prologs,
                                            PIPE_SHADER_GEOMETRY, true,